记录 | CUDA编程中声明内联函数的方法

下面将 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__ 来禁止内联优化；