下面将 say_hello()
声明为内联函数:
#include <cstdio>
#include <cuda_runtime.h>__device__ __inline__ void say_hello(){printf("Hello, world!\n");
}__global__ void kernel(){say_hello();
}int main(){kernel<<<1, 1>>>();cudaDeviceSynchronize();return 0;
}
需要注意的点:
● inline
在 C++ 中的效果是声明一个函数为 weak
符号(弱符号),和性能优化意义上的内联无关;
● 优化意义上的内联指把函数体直接放到调用者那里去;
● 因此 CUDA 编译器提供了一个 “私货” 关键字:__inline__
来声明一个函数为内联。不论是 CPU 函数还是 GPU 都可以使用,只要你用的 CUDA 编译器。GCC 编译器相应的私货则是 __attribute__(("inline"))
;
● 注意声明为 __inline__
不一定就保证内联了,如果函数太大,编译器可能会放弃内联化。因此 CUDA 还提供了 __forceinline__
这个关键字来强制一个函数为内联。GCC 也有相应的 __attribute__(("always_inline"))
;
● 此外,还有 __noinline__
来禁止内联优化;