首先明确概念
- CPU上叫主机端，host
- GPU上叫设备端，device
函数执行空间指定符
函数执行空间指定符表示一个函数是在主机上执行还是在设备上执行，以及它是否可以从主机或设备调用。
global
__global__ 执行空间指定符声明一个函数为kernel。这样的函数是：
- 在设备上执行，
- 可以从主机调用，
- 对于计算能力为 5.0 或更高的设备可从设备调用（详见 CUDA 动态并行）。
__global__ 函数必须具有 void 返回类型，并且不能是类的成员。
对 __global__ 函数的任何调用都必须指定其执行配置，如 执行配置 中所述。
调用 __global__ 函数是异步的，意味着它在设备完成执行之前返回。
device
__device__ 执行空间指定符声明一个函数为：
- 在设备上执行，
- 仅可从设备调用。
__global__ 和 __device__ 执行空间指定符不能一起使用。
host
__host__ 执行空间指定符声明一个函数为：
- 在主机上执行，
- 仅可从主机调用。
平常写函数的时候不带这个指定符，默认在cpu上跑
仅声明一个函数具有 __host__ 执行空间指定符或不带任何 __host__、__device__ 或 __global__ 执行空间指定符，这俩效果一样；在任一情况下，函数只为主机编译。
__global__ 和 __host__ 执行空间指定符不能一起使用。
__device__ 和 __host__ 执行空间指定符可以一起使用，此时函数为主机和设备编译。在 应用程序兼容性 中引入的 __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
}
未定义行为
当出现以下情况时，跨执行空间的调用具有未定义行为：
__CUDA_ARCH__定义时，从__global__、__device__或__host__ __device__函数内部调用__host__函数。__CUDA_ARCH__未定义时，从__host__函数内部调用__device__函数。
noinlin 和 forceinline
编译器在适当时将任何 __device__ 函数内联。
__noinline__ 函数限定符可用作提示，让编译器尽可能不内联该函数。
__forceinline__ 函数限定符可用于强制编译器内联该函数。
__noinline__ 和 __forceinline__ 函数限定符不能一起使用，也不能应用于内联函数。
inline_hint
__inline_hint__ 限定符启用编译器中更积极的内联。与 __forceinline__ 不同，它并不暗示该函数是内联的。它可以用于在使用 LTO 时提高跨模块的内联。
__noinline__ 和 __forceinline__ 函数限定符均不能与 __inline_hint__ 函数限定符一起使用。
变量内存空间指定符
变量内存空间指定符表示device上变量的内存位置。
在设备代码中声明的自动变量，如果没有本节中描述的任何 __device__、__shared__ 和 __constant__ 内存空间指定符，通常位于寄存器中。但在某些情况下，编译器可能选择将其放在local memory中，这可能会导致性能下降，如 设备内存访问 中详述。
device
__device__ 内存空间指定符声明一个变量位于设备上。
最多只能与本节接下来的三个内存空间指定符中的一个一起使用，以进一步指明变量属于哪个内存空间。如果没有任何一个出现，则变量：
- 位于全局内存空间，
- 拥有其创建的 CUDA 上下文的生命周期，
- 每个设备有一个独立的对象，
- 可通过运行时库从网格内的所有线程和主机访问（
cudaGetSymbolAddress()/cudaGetSymbolSize()/cudaMemcpyToSymbol()/cudaMemcpyFromSymbol()）。
constant
__constant__ 内存空间指定符，可选地与 __device__ 一起使用，声明一个变量：
- 位于常量内存空间，
- 拥有其创建的 CUDA 上下文的生命周期，
- 每个设备有一个独立的对象，
- 可通过运行时库从网格内的所有线程和主机访问（
cudaGetSymbolAddress()/cudaGetSymbolSize()/cudaMemcpyToSymbol()/cudaMemcpyFromSymbol()）。
shared
__shared__ 内存空间指定符，可选地与 __device__ 一起使用，声明一个变量：
- 位于线程块的共享内存空间，
- 拥有块的生命周期，
- 每个块有一个独立的对象，
- 仅可由块内的所有线程访问，
- 没有固定的地址。
如下声明共享内存中的变量作为外部数组时:
extern shared float shared[];
数组的大小在启动时确定（见 执行配置）。以这种方式声明的所有变量，都从内存中的相同地址开始，因此必须通过偏移显式管理变量在数组中的布局。例如，如果想要
short array0[128];
float array1[64];
int array2[256];
在动态分配的共享内存中，可以以下列方式声明和初始化数组：
extern __shared__ float array[];
__device__ void func() // __device__ or __global__ function
{
short* array0 = (short*)array;
float* array1 = (float*)&array0[128];
int* array2 = (int*)&array1[64];
}
注意，指针需要与它们指向的类型对齐，因此以下代码无效，因为 array1 没有对齐到 4 字节:
extern __shared__ float array[];
__device__ void func() // __device__ or __global__ function
{
short* array0 = (short*)array;
float* array1 = (float*)&array0[127];
}
内置向量类型的对齐要求列在表 5
grid_constant
对于计算架构 7.0 及以上，__grid_constant__ 注解用于修饰一个 const-qualified 的非引用类型的 __global__ 函数参数，该参数：
- 拥有网格的生命周期，
- 是私有于网格的，即，该对象对主机线程和来自其他网格（包括子网格）的线程不可访问，
- 每个网格有一个独立的对象，即，网格中的所有线程看到相同的地址，
- 是只读的，即，修改
__grid_constant__对象或其任何子对象的行为是未定义的，包括可变成员。
要求：
- 用
__grid_constant__注解的内核参数必须是 const-qualified 的非引用类型。 - 所有函数声明必须在任何
__grid_constant__参数方面与之匹配。 - 函数模板专门化必须在任何
__grid_constant__参数方面与主模板声明匹配。 - 函数模板实例化指令必须在任何
__grid_constant__参数方面与主模板声明匹配。
如果取了 __global__ 函数参数的地址，编译器通常会在线程局部内存中复制内核参数，并使用复制的地址，部分支持 C++ 语义，允许每个线程修改其自己的局部副本的函数参数。用 __grid_constant__ 注解 __global__ 函数参数可确保编译器不会在线程局部内存中创建内核参数的副本，而是使用参数本身的通用地址。避免局部副本可能会提高性能。
__device__ void unknown_function(S const&);
__global__ void kernel(const __grid_constant__ S s) {
s.x += threadIdx.x; // 未定义行为：尝试修改只读内存
// 编译器不会为“s”创建每线程局部副本：
unknown_function(s);
}
managed
__managed__ 内存空间指定符，可选地与 __device__ 一起使用，声明一个变量：
- 可以从设备和主机代码引用，例如，可以取其地址，或可以直接从设备或主机函数读取或写入。
- 拥有应用程序的生命周期。
有关更多详情，请参见 __managed__ 内存空间指定符。
restrict
nvcc 通过 __restrict__ 关键字支持受限指针。
受限指针在 C99 中引入，用于减轻存在于 C 类型语言中的别名问题，该问题阻碍了从代码重排到常见子表达式消除的所有类型的优化。
以下是一个受别名问题影响的例子，使用受限指针可以帮助编译器减少指令数量：
void foo(const float* a,
const float* b,
float* c)
{
c[0] = a[0] * b[0];
c[1] = a[0] * b[0];
c[2] = a[0] * b[0] * a[1];
c[3] = a[0] * a[1];
c[4] = a[0] * b[0];
c[5] = b[0];
...
}
在 C 类型语言中，指针 a、b 和 c 可能被别名，所以通过 c 的任何写入都可能修改 a 或 b 的元素。这意味着为了保证功能正确性，编译器不能将 a[0] 和 b[0] 加载到寄存器中，相乘它们，并将结果存储到
c[0] 和 c[1] 中，因为如果比如 a[0] 实际上是 c[0] 的同一位置，则结果将与抽象执行模型不同。因此，编译器不能利用常见的子表达式。同样，编译器不能仅仅将 c[4] 的计算重新排序到 c[0] 和 c[1] 的计算附近，因为前面对 c[3] 的写入可能改变了 c[4] 的计算输入。
通过将 a、b 和 c 设为受限指针，程序员向编译器断言这些指针实际上不是别名，这在本例中意味着通过 c 的写入永远不会覆盖 a 或 b 的元素。这将函数原型更改如下：
void foo(const float* __restrict__ a,
const float* __restrict__ b,
float* __restrict__ 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;
...
}
这里的效果是减少了内存访问次数和计算次数。这被由于“缓存”加载和常见子表达式导致的寄存器压力增加所平衡。
由于寄存器压力在许多 CUDA 代码中是一个关键问题，restricted pointers的使用可能会对 CUDA 代码的性能产生负面影响，因为降低了占用率。
Built-in Vector Types 内建的向量类型
char, short, int, long, longlong, float, double
这些是从基本整数和浮点类型派生的向量类型。它们是结构体，第 1、2、3 和 4 个组件分别可以通过字段 x、y、z 和 w 访问。它们都带有形式为 make_<type name> 的构造函数；例如，
int2 make_int2(int x, int y);
用于创建一个值为 (x, y) 的 int2 类型的向量。
向量类型的对齐要求详细说明在以下表格中。
| Type | Alignment |
|---|---|
| char1, uchar1 | 1 |
| char2, uchar2 | 2 |
| char3, uchar3 | 1 |
| char4, uchar4 | 4 |
| short1, ushort1 | 2 |
| short2, ushort2 | 4 |
| short3, ushort3 | 2 |
| short4, ushort4 | 8 |
| int1, uint1 | 4 |
| int2, uint2 | 8 |
| int3, uint3 | 4 |
| int4, uint4 | 16 |
| long1, ulong1 | 4 if sizeof(long) is equal to sizeof(int) 8, otherwise |
| long2, ulong2 | 8 if sizeof(long) is equal to sizeof(int), 16, otherwise |
| long3, ulong3 | 4 if sizeof(long) is equal to sizeof(int), 8, otherwise |
| long4, ulong4 | 16 |
| longlong1, ulonglong1 | 8 |
| longlong2, ulonglong2 | 16 |
| longlong3, ulonglong3 | 8 |
| longlong4, ulonglong4 | 16 |
| float1 | 4 |
| float2 | 8 |
| float3 | 4 |
| float4 | 16 |
| double1 | 8 |
| double2 | 16 |
| double3 | 8 |
| double4 | 16 |
dim3
此类型是基于 uint3 的整数向量类型，用于指定尺寸。当定义 dim3 类型的变量时，任何未指定的组件都将初始化为 1。
内置变量
内置变量指定了网格和块的维度以及块和线程索引。它们仅在在设备上执行的函数中有效。
gridDim
此变量为 dim3 类型（见 dim3），包含网格的维度。
blockIdx
此变量为 uint3 类型（见 char, short, int, long, longlong, float, double），包含网格内的块索引。
blockDim
此变量为 dim3 类型（见 dim3），包含块的维度。
threadIdx
此变量为 uint3 类型（见 char, short, int, long, longlong, float, double），包含块内的线程索引。
warpSize
此变量为 int 类型，包含线程中的 warp 大小（关于 warp 的定义，请见 SIMT 架构）。
内存屏障函数
CUDA 编程模型假设设备具有弱序内存模型，即 CUDA 线程将数据写入共享内存、全局内存、页锁定的主机内存或对等设备的内存的顺序不一定是另一个 CUDA 或主机线程观察到的数据写入顺序。如果两个线程在没有同步的情况下从同一内存位置读取或写入，则其行为是未定义的。
在以下示例中，线程 1 执行 writeXY()，而线程 2 执行 readXY()。
__device__ int X = 1, Y = 2;
__device__ void writeXY()
{
X = 10;
Y = 20;
}
__device__ void readXY()
{
int B = Y;
int A = X;
}
两个线程同时从相同的内存位置 X 和 Y 读写。任何数据竞争都是未定义行为，并且没有定义的语义。A 和 B 的结果可以是任何值。
内存屏障函数可以用于在内存访问上强制执行顺序一致的排序。内存屏障函数在强制执行排序的范围上有所不同，但
它们独立于访问的内存空间（共享内存、全局内存、页锁定的主机内存和对等设备的内存）。
void __threadfence_block();
等同于 cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_block)，确保：
- 调用线程在调用
__threadfence_block()之前对所有内存的所有写入被该线程块中的所有线程观察到，发生在调用线程在调用__threadfence_block()之后对所有内存的所有写入之前； - 调用线程在调用
__threadfence_block()之前对所有内存的所有读取在调用__threadfence_block()之后对所有内存的所有读取之前排序。
void __threadfence();
等同于 cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_device)，确保调用线程在调用 __threadfence() 之后对所有内存的所有写入不会被设备中的任何线程观察到，发生在调用线程在调用 __threadfence() 之前对所有内存的任何写入之前。
void __threadfence_system();
等同于 cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_system)，确保调用线程在调用 __threadfence_system() 之前对所有内存的所有写入被设备中的所有线程、主机线程和对等设备中的所有线程观察到，发生在调用线程在调用 __threadfence_system() 之后对所有内存的所有写入之前。
__threadfence_system() 仅由计算能力为 2.x 及更高的设备支持。
在上述代码示例中，我们可以在代码中插入屏障，如下所示：
__device__ int X = 1, Y = 2;
__device__ void writeXY()
{
X = 10;
__threadfence();
Y = 20;
}
__device__ void readXY()
{
int B = Y;
__threadfence();
int A = X;
}
对于此代码，可以观察到以下结果：
A等于 1 和B等于 2，A等于 10 和B等于 2，A等于 10 和B等于 20。
第四种情况是不可能发生的，因为第一次写入操作必须在第二次写入之前被看到。如果线程 1 和线程 2 在同一个块中，那么只需要使用 __threadfence_block() 即可。但如果它们不在同一个块中，那么对于同一设备上的 CUDA 线程需要使用 __threadfence()，而对于来自不同设备的 CUDA 线程则必须使用 __threadfence_system()。
一个典型的应用场景是，一些线程需要使用由其他线程生成的数据。例如，在下面的代码示例中，展示了一个内核程序，它能在一次调用中计算一个含 N 个数字的数组的总和。每个块首先计算数组的一个子集的和，并将结果保存在全局内存中。当所有块都完成后，最后完成的块会从全局内存中读取这些部分和，并将它们加起来以得到最终结果。为了判定哪个块是最后完成的，每个块都会原子性地递增一个计数器，以此表示它已完成计算并存储了它的部分和（详见关于原子函数的部分）。最后一个块是接收到等于 gridDim.x-1 的计数器值的块。如果在存储部分和与递增计数器之间没有设置内存栅栏，那么计数器可能在部分和被存储之前就已递增，从而可能达到 gridDim.x-1 的值，导致最后一个块在这些部分和实际上还未在内存中更新之前就开始读取它们。
内存栅栏函数（Memory fence functions）仅影响线程的内存操作顺序，并不能保证这些操作对其他线程可见（像 __syncthreads() 那样仅对同一块内的线程可见，详见同步函数部分）。在下面的代码示例中，结果变量被声明为易失性（参见易失性限定符部分），以此确保其内存操作对其他线程可见。
__device__ unsigned int count = 0;
__shared__ bool isLastBlockDone;
__global__ void sum(const float* array, unsigned int N,
volatile float* result)
{
// Each block sums a subset of the input array.
float partialSum = calculatePartialSum(array, N);
if (threadIdx.x == 0) {
// Thread 0 of each block stores the partial sum
// to global memory. The compiler will use
// a store operation that bypasses the L1 cache
// since the "result" variable is declared as
// volatile. This ensures that the threads of
// the last block will read the correct partial
// sums computed by all other blocks.
result[blockIdx.x] = partialSum;
// Thread 0 makes sure that the incrementation
// of the "count" variable is only performed after
// the partial sum has been written to global memory.
__threadfence();
// Thread 0 signals that it is done.
unsigned int value = atomicInc(&count, gridDim.x);
// Thread 0 determines if its block is the last
// block to be done.
isLastBlockDone = (value == (gridDim.x - 1));
}
// Synchronize to make sure that each thread reads
// the correct value of isLastBlockDone.
__syncthreads();
if (isLastBlockDone) {
// The last block sums the partial sums
// stored in result[0 .. gridDim.x-1]
float totalSum = calculateTotalSum(result);
if (threadIdx.x == 0) {
// Thread 0 of last block stores the total sum
// to global memory and resets the count
// varialble, so that the next kernel call
// works properly.
result[0] = totalSum;
count = 0;
}
}
}
同步函数
void __syncthreads();
该函数会等待线程块中的所有线程都达到同一点，并确保这些线程在 __syncthreads() 之前对全局内存和共享内存的所有访问对块中的所有线程可见。
__syncthreads() 主要用于协调同一块中的线程之间的通信。当块内的线程访问共享或全局内存中的同一地址时，可能会出现读后写、写后读或写后写的风险。通过在这些内存访问之间同步线程，可以有效避免这类数据风险。
在条件代码中可以使用 __syncthreads()，但前提是整个线程块中的条件判断结果必须相同，否则代码可能会挂起或产生意外的副作用。
计算能力为 2.x 及以上的设备支持 __syncthreads() 的三种变体：
int __syncthreads_count(int predicate);
这个函数与 __syncthreads() 相同，但增加了一个功能：它会评估块中所有线程的 predicate（谓词）并返回 predicate 非零评估的线程数。
int __syncthreads_and(int predicate);
这个函数与 __syncthreads() 相同，但增加了一个功能：它会评估块中所有线程的 predicate 并仅在所有线程的 predicate 都为非零时返回非零值。
int __syncthreads_or(int predicate);
这个函数与 __syncthreads() 相同，但增加了一个功能：它会评估块中所有线程的 predicate 并仅在任一线程的 predicate 为非零时返回非零值。
void __syncwarp(unsigned mask=0xffffffff);
这个函数使得执行线程等待，直到掩码中指定的所有 warp 路径执行了 __syncwarp()（具有相同的掩码）后才继续执行。每个调用线程必须在掩码中有自己的位设置，且掩码中指定的所有未退出线程必须执行具有相同掩码的对应 __syncwarp()，否则结果是不确定的。
执行 __syncwarp() 可以保证参与屏障的线程之间的内存顺序。因此，如果 warp 内的线程想要通过内存进行通信，它们可以先向内存中存储数据，然后执行 __syncwarp()，之后再安全地读取其他线程在 warp 中存储的值。
注意
对于 .target sm_6x 或更低版本，掩码中的所有线程必须在收敛过程中执行相同的__syncwarp()，并且所有掩码中的值的并集必须等于活动掩码。否则，其行为是不确定的。