CUDA从入门到放弃（十三）：C++语言扩展 C++ Language Extensions

1 Function Execution Space Specifiers 函数执行空间指定符

这些指定符定义函数是在主机还是设备上执行，以及它们是否可以跨平台调用。

1-1 __global__

__global__指定函数为内核，可在设备上执行，也可从主机或（计算能力5.0以上设备）设备上调用。必须有void返回类型，调用是异步的。

1-2 device

device 指定函数只在设备上执行，只能从设备调用。不能与__global__ 一起使用。

1-3 host

host__指定函数在主机上执行，只能从主机调用。等价于不使用任何执行空间指定符，函数仅编译为主机代码。不能与__global 一起使用，但可以与__device__ 一起使用，此时函数将同时编译为主机和设备代码。使用__CUDA_ARCH__ 宏可以区分主机和设备代码路径。

__host__ __device__ func()
{
	#if __CUDA_ARCH__ >= 800
   	// Device code path for compute capability 8.x
	#elif __CUDA_ARCH__ >= 700
   	// Device code path for compute capability 7.x
	#elif __CUDA_ARCH__ >= 600
   	// Device code path for compute capability 6.x
	#elif __CUDA_ARCH__ >= 500
   	// Device code path for compute capability 5.x
	#elif !defined(__CUDA_ARCH__)
   	// Host code path
	#endif
}

1-4 Undefined behavior

跨执行空间调用时，若__CUDA_ARCH__已定义，从__global__、device__或__host __device__函数中调用__host__函数，或__CUDA_ARCH__未定义时，从__host__函数中调用__device__函数，行为是未定义的。

1-5 noinline and forceinline

编译器会适当内联__device__函数。__noinline__修饰符提示编译器尽量避免内联，__forceinline__修饰符强制内联。两者不可同时使用，且不适用于内联函数。

1-6 inline_hint

inline_hint__修饰符促进编译器更积极内联，不同于__forceinline，它不保证函数一定会内联。使用LTO时，有助于改善跨模块内联。它不与__noinline__或__forceinline__修饰符共用。

2 变量内存空间指定符 Variable Memory Space Specifiers

变量内存空间指定符决定了变量在设备上的存储位置。

变量若未使用__device__、__shared__和__constant__指定符，通常放在寄存器中。但编译器可能将其置于本地内存，这可能影响性能。

2-1 device

__device__指定符表示变量驻留在设备上。它可以与其他指定符结合，进一步定义变量的内存空间。若无其他指定符，则变量位于全局内存，具有CUDA上下文的生命周期，每设备一个独立对象，可从网格内所有线程和主机访问。

2-2 constant

constant（可与__device__结合）指定变量驻留在常量内存，同样具有CUDA上下文的生命周期，每设备一个独立对象，可从网格内所有线程和主机访问。

2-3 shared

__shared__内存空间指定符（可与__device__一起使用）声明一个变量，该变量位于线程块的共享内存中，其生命周期与线程块相同，每个线程块有独立对象，仅可从该线程块内的线程访问，地址不固定。

在声明共享内存中的变量为外部数组时，如：

extern __shared__ float shared[];

数组的大小在启动时确定.

2-4 grid_constant

在计算架构7.0及以上的环境中，__grid_constant__注解用于标记const限定的非引用类型__global__函数参数。这些参数具有网格的生命周期，私属于网格，每个网格有一个独立的对象，且是只读的。编译器不会为这些参数创建线程本地内存中的副本，而是使用参数本身的地址，这有助于提升性能。

__device__ void unknown_function(S const&);  
__global__ void kernel(const __grid_constant__ S s) {  
   s.x += threadIdx.x;  // 未定义行为：尝试修改只读内存  
  
   // 编译器将不会为每个线程创建"s"的线程本地副本：  
   unknown_function(s);  
}

2-5 managed

__managed__是一个内存空间指定符，用于声明在主机和设备上均可访问的变量。这些变量与CUDA上下文同生命周期，主机访问时通过页锁定内存管理，设备访问时驻留在设备内存中。使用__managed__可以简化数据复制的代码，但程序员需注意潜在的同步和数据竞争问题。

2-6 restrict

nvcc支持使用__restrict__关键字来声明受限指针，这是C99中引入的特性，旨在解决C类语言中的别名问题，从而优化代码。在C中，指针可能指向相同的内存位置，这限制了编译器进行诸如重排序和公共子表达式消除等优化。通过声明受限指针，程序员告诉编译器这些指针不会指向相同位置，从而允许编译器进行更多优化。例如，在函数foo中，通过为参数a、b和c添加__restrict__修饰符，编译器可以更安全地优化指令，减少计算冗余，同时保持程序的正确性。

void foo(const float* __restrict__ a,  
         const float* __restrict__ b,  
         float* __restrict__ c)  
{  
    float t0 = a[0];  
    float t1 = b[0];  
    float t2 = t0 * t1;  
    float t3 = a[1];  
    c[0] = t2;  
    c[1] = t2;  
    c[4] = t2;  
    c[2] = t2 * t3;  
    c[3] = t0 * t3;  
    c[5] = t1;  
    ...  
}

为了让编译器优化器获益，所有指针参数必须声明为受限。使用__restrict__后，编译器能自由优化指令，减少内存访问和计算，但可能增加寄存器压力。在CUDA代码中，寄存器压力是关键问题，因此使用受限指针可能因减少占用率而降低性能。

3. 内置向量类型 Built-in Vector Types

3-1. char, short, int, long, longlong, float, double

这些是基于基本整数和浮点类型派生的向量类型。它们是结构体，其第1、第2、第3和第4个组件分别可以通过字段x、y、z和w来访问。它们都带有形式为make_<类型名>的构造函数函数；例如，

int2 make_int2(int x, int y);

这个函数会创建一个值为(x, y)的int2类型的向量。

3-2. dim3

这种类型是基于uint3的整数向量类型，用于指定维度。在定义dim3类型的变量时，未指定的任何组件都会初始化为1。

4 内置变量 Built-in Variables

内置变量用于指定网格和块的维度以及块和线程的索引。它们仅在设备上执行的函数内部有效。

4-1 gridDim

这个变量是dim3类型（参见dim3），包含网格的维度。

4-2 blockIdx

这个变量是uint3类型（参见char, short, int, long, longlong, float, double），包含网格内的块索引。

4-3 blockDim

这个变量是dim3类型（参见dim3），包含块的维度。

4-4 threadIdx

这个变量是uint3类型（参见char, short, int, long, longlong, float, double），包含块内的线程索引。

4-5 warpSize

这个变量是int类型，包含warp中的线程数。

5 内存屏障函数 Memory Fence Functions

CUDA编程模型使用弱序内存模型，即CUDA线程写入数据的顺序并不一定是其他线程观察到数据写入的顺序。两个线程无同步地读写同一内存位置会导致未定义行为。

内存屏障函数用于确保内存访问的顺序一致性，不受访问的内存空间类型影响。

void __threadfence_block() 

确保同一块内的线程能观察到调用前后的内存操作顺序。

void __threadfence() 

确保设备内所有线程观察到调用前后的内存操作顺序。

void __threadfence_system() 

确保设备、主机以及对等设备的所有线程观察到调用前后的内存操作顺序。

内存屏障函数仅控制线程的内存操作顺序，不保证对其他线程可见。

6 同步函数 Synchronization Functions

void __syncthreads(); 

同步同一线程块内所有线程，确保所有线程到达该点前对全局和共享内存的访问对所有线程可见。

int __syncthreads_count(int predicate); 

同步线程并返回谓词为真的线程数。

int __syncthreads_and(int predicate); 

同步线程并仅当所有线程谓词为真时返回非零。

int __syncthreads_or(int predicate); 

同步线程并仅当任意线程谓词为真时返回非零。

void __syncwarp(unsigned mask=0xffffffff); 

使线程等待，直到指定掩码中的所有warp通道执行该函数后继续执行，确保内存顺序。

参考资料
1 CUDA编程入门
2 CUDA编程入门极简教程
3 CUDA C++ Programming Guide
4 CUDA C++ Best Practices Guide
5 NVIDIA CUDA初级教程视频
6 CUDA专家手册 [GPU编程权威指南]
7 CUDA并行程序设计：GPU编程指南
8 CUDA C编程权威指南