【CUDA编程】系列博客参考NVIDIA官方文档“CUDA C++ Programming Guide（v12.6）”。
本文为原创文章，未经本人允许，禁止转载。转载请注明出处。

1.Function Execution Space Specifiers

函数执行空间限定符用于定义一个函数是运行在host上还是device上，以及是否可以从host或device调用。

1.1.`global`

__global__执行空间限定符将函数声明为kernel函数。这样的函数具有以下特性：

在device上执行。
从host端调用。
对于计算能力5.0及以上的device，可从device端调用。

一个__global__函数的返回类型必须是void，且不能是类的成员。

对__global__函数的任何调用都必须按照Execution Configuration章节中描述的指定其执行配置。

对__global__的调用是异步的，这意味着在device完成其执行之前，该调用就会返回。

1.2.`device`

__device__执行空间限定符声明的函数具有以下特性：

在device上执行。
只能从device端调用。

__global__和__device__不能一起使用。

1.3.`host`

__host__执行空间限定符声明的函数具有以下特性：

在host上执行。
只能从host端调用。

声明一个函数时，如果仅使用__host__限定符，或者不使用任何限定符（即不使用__global__、__device__或__host__），则该函数会被编译为仅供host使用。

__global__和__host__不能一起使用。

__device__和__host__可以一起使用，此时函数会同时为host和device生成代码。

__CUDA_ARCH__宏可以用于区分host和device之间的代码路径：

__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

当出现以下情况时，跨执行空间（cross-execution space）调用会表现为未定义行为：

如果__CUDA_ARCH__已定义，从__global__、__device__或__host__ __device__函数调用__host__函数。
如果__CUDA_ARCH__未定义，从__host__函数调用__device__函数。

1.5.`noinline`and`forceinline`

编译器会在适当时内联任何__device__函数。

__noinline__函数限定符可作为提示，告知编译器如果可能，不要内联该函数。

__forceinline__函数限定符可用于强制编译器内联该函数。

__noinline__和__forceinline__不能一起使用，也不能将它们应用于一个已经是内联的函数。

1.6.`__inline_hint__`

__inline_hint__限定符在编译器中启用了更积极的内联优化。与__forceinline__不同，它并不意味着函数必须被内联。它可以在使用LTO（Link Time Optimization，链接时优化）时提高跨模块内联的效果。

__noinline__和__forceinline__不能与__inline_hint__同时使用。

2.Variable Memory Space Specifiers

变量内存空间限定符用于表示变量在device上的内存位置。

在device代码中声明的自动变量，如果没使用__device__、__shared__或__constant__内存空间限定符，则通常存储在寄存器中。然而，在某些情况下，编译器也可能会选择将其放置在local memory中，这可能会对性能产生不利影响（详见：Device Memory Accesses）。

2.1.`device`

__device__内存空间限定符声明一个驻留在device上的变量。

在接下来的三个部分定义的内存空间限定符中，最多只能使用其中一个，并且可以与__device__一起使用以进一步指定变量所属的内存空间。如果没有指定其他限定符，变量将具有以下特性：

驻留在全局内存空间中。
拥有其创建时所在CUDA context的生命周期。
每个device上有个独立的对象。也就是说，使用__device__创建的对象，会在每个device上都有一个独立的副本。
可以从grid内的所有线程以及通过运行时库从host进行访问（例如cudaGetSymbolAddress()、cudaGetSymbolSize()、cudaMemcpyToSymbol()、cudaMemcpyFromSymbol()）。

2.2.`constant`

__constant__内存空间限定符，可以选择性的和__device__一起使用。__constant__声明的变量具有以下特性：

驻留在常量内存空间中。
拥有其创建时所在CUDA context的生命周期。
每个device上有个独立的对象。
可以从grid内的所有线程以及通过运行时库从host进行访问（例如cudaGetSymbolAddress()、cudaGetSymbolSize()、cudaMemcpyToSymbol()、cudaMemcpyFromSymbol()）。

如果host在有并发grid访问该常量时修改该常量，则该行为在grid生命周期内是未定义的。

2.3.`shared`

__shared__内存空间限定符，可以选择性的和__device__一起使用。__shared__声明的变量具有以下特性：

驻留在线程块的共享内存空间中。
具有线程块的生命周期。
每个线程块有一个独立的对象。
只能被线程块中的所有线程访问。
不具有常量地址（constant address）。

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

extern __shared__ float shared[];

数组的大小在启动时确定。以extern __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];
}

2.4.`__grid_constant__`

对于计算能力大于等于7.0的device，__grid_constant__可用于修饰__global__函数中的const参数（非引用类型），这种参数具有以下特性：

具有grid的生命周期。
对grid私有，即host线程或其他grid（包括sub-grid）中的线程不可访问该对象。
每个grid有一个独立的对象，即grid中的所有线程看到的是相同的地址。
只读，即修改__grid_constant__对象或其任何子对象（包括mutable成员）是未定义行为。

要求：

使用__grid_constant__修饰的kernel参数必须是具有const限定符的非引用类型。
所有函数声明必须与任何__grid_constant__参数的声明保持一致。
函数模板的特例化必须与主模板声明中关于__grid_constant__参数的声明保持一致。
函数模板的实例化指令必须与主模板声明中关于__grid_constant__参数的声明保持一致。

如果获取了一个__global__函数参数的地址，编译器通常会在线程local memory中创建该kernel函数参数的副本，并使用该副本的地址。这是为了部分支持C++的语义，即允许每个线程修改其函数参数的本地副本。然而，如果使用__grid_constant__修饰__global__函数的参数，就能确保编译器不会在线程local memory中创建该kernel函数参数的副本，而是直接使用参数本身的通用地址。通过避免本地副本的创建，可以提高性能。

__device__ void unknown_function(S const&);
__global__ void kernel(const __grid_constant__ S s) {
   s.x += threadIdx.x;  // Undefined Behavior: tried to modify read-only memory

   // Compiler will _not_ create a per-thread thread local copy of "s":
   unknown_function(s);
}

2.5.`managed`

__managed__内存空间限定符，可以选择性的和__device__一起使用。__managed__声明的变量具有以下特性：

可以被device代码和host代码引用，例如，可以获取其地址，或者直接从device或host函数中读取或写入。
拥有应用程序的生命周期。

2.6.`restrict`

nvcc通过__restrict__关键字支持受限指针（restricted pointers）。

受限指针是C99标准中引入的一种指针类型修饰符，通过关键字restrict实现，用于告诉编译器指针指向的内存区域不会与其他指针指向的内存区域重叠（即没有别名关系）。也就是说，在C/C++中，别名问题指的是不同的指针指向同一块内存。受限指针用于解决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];
    ...
}

在上述代码中，指针a、b、c都是外部传入的指针，编译器会默认假设这些指针可能有别名关系，比如，a[0]、b[0]、c[0]可能是同一块内存地址。这种假设会导致编译器选择保守的优化方式，以确保功能的正确性。比如c[0] = a[0] * b[0]，写入c[0]后，编译器不能确信a[0]或b[0]的值是否被修改，因此在后续的c[1] = a[0] * b[0]中，编译器会重新加载a[0]和b[0]。如果没有别名问题，a[0]和b[0]可以一次性加载到寄存器中，后续的计算可以复用寄存器中的值。但由于编译器无法确定别名情况，每次都需要从内存中重新加载，增加了内存访问的次数。同理，因为可能的别名问题，为了保证结果正确，编译器不能随意重排序这些操作，限制了优化空间。

通过将指针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代码中的关键问题，使用受限指针可能会对CUDA代码的性能产生负面影响，因为寄存器压力的增加可能会降低线程块的占用率。

3.Built-in Vector Types

3.1.char, short, int, long, longlong, float, double

这些是从基本的整型和浮点型派生的向量类型。其第一个、第二个、第三个和第四个分量可以分别通过字段x、y、z和w来访问。所有这些向量类型都配有形式为make_<type name>的构造函数。比如：

int2 make_int2(int x, int y);

该函数创建了一个类型为int2的向量，其值为(x, y)。

向量类型的对齐要求详见下表。

3.2.dim3

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

4.Built-in Variables

内置变量用于指定grid和线程块的维度，以及线程块和线程的索引。它们仅在device上执行的函数内有效。

4.1.gridDim

变量类型为dim3（见第3.2部分），表示grid的维度。

4.2.blockIdx

变量类型为uint3（见第3.1部分），表示线程块在grid中的索引。

4.3.blockDim

变量类型为dim3（见第3.2部分），表示线程块的维度。

4.4.threadIdx

变量类型为uint3（见第3.1部分），表示线程在线程块中的索引。

4.5.warpSize

变量类型为int，表示warp的大小（详见：SIMT Architecture）。

5.Memory Fence Functions

memory fence译为内存栅栏，和内存屏障（memory barrier）是一个意思。下文直接使用内存屏障这个名词。

CUDA编程模型假设device采用一个弱序内存模型（a weakly-ordered memory model），也就是说，内存操作可能不会按照代码的实际顺序发生，这种设计旨在提高性能。如果两个线程在没有同步的情况下对同一内存位置进行读取或写入，则行为未定义。

在下面的代码示例中，线程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进行读取和写入。任何数据竞争（data-race）都会导致未定义行为，并且没有明确的语义。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(); //作用范围为整个device

上述代码等价于cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_device)。

void __threadfence_system(); //作用范围为整个系统，包括device内所有线程、host线程和多device中的所有线程

上述代码等价于cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_system)。

只有计算能力在2.x及以上的device才支持__threadfence_system()。

在前面的代码示例中，我们可以通过如下方式插入内存屏障：

__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。

因为__threadfence()的作用，A等于1且B等于20的情况不会出现，因为B等于20就意味着X = 10肯定已经执行完了（因为__threadfence()保证了X = 10的执行先于Y = 20），并且A = X又必须在B = Y之后执行，此时A不可能等于1。

如果线程1和线程2属于同一个线程块，则使用__threadfence_block()就足够了。如果线程1和线程2不属于同一个线程块，但属于同一个device，则必须使用__threadfence()。如果线程1和线程2不属于同一个device，则必须使用__threadfence_system()。

注意，内存屏障函数仅能影响线程的内存操作顺序，并不能确保这些内存操作对其他线程可见，也不能确保线程同步。

一个常见的使用场景是线程需要使用其他线程生成的数据。在下面的代码示例中，kernel函数的输入是一个长度为N的一维数组，该函数的功能是计算该数组所有元素的和。让每个线程块负责计算一个数组子集的部分和，并将结果存储在全局内存中。当所有线程块完成后，最后一个线程块读取全局内存中的所有部分和，并将它们相加得到最终结果。为了确定哪个线程块最后完成计算，每个线程块通过原子操作递增一个计数器来表示它完成了计算并存储了部分和。最后一个线程块的计数器值应该是gridDim.x-1。如果在存储部分和与递增计数器之间没有设置内存屏障，则计数器可能在部分和写入到内存之前递增，从而导致最后一个线程块开始读取部分和时，部分和尚未被正确更新。通过将result变量声明为volatile来确保内存操作对其他线程可见。

__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) { //每个线程块中的线程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 incrementing
        // 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.
        //确保所有线程块都计算完毕
        //线程块的数量为gridDim.x - 1
        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
            // variable, so that the next kernel call
            // works properly.
            result[0] = totalSum;
            count = 0;
        }
    }
}

【CUDA编程】【26】【6.C++ Language Extensions】【Part1】

Function Execution Space Specifiers，Variable Memory Space Specifiers，Built-in Vector Types，Built-in Variables，Memory Fence Functions

1.Function Execution Space Specifiers

1.1.`global`

1.2.`device`

1.3.`host`

1.4.Undefined behavior

1.5.`noinline`and`forceinline`

1.6.`__inline_hint__`

2.Variable Memory Space Specifiers

2.1.`device`

2.2.`constant`

2.3.`shared`

2.4.`__grid_constant__`

2.5.`managed`

2.6.`restrict`

3.Built-in Vector Types

3.1.char, short, int, long, longlong, float, double

3.2.dim3

4.Built-in Variables

4.1.gridDim

4.2.blockIdx

4.3.blockDim

4.4.threadIdx

4.5.warpSize

5.Memory Fence Functions

CATALOG

FEATURED TAGS

1.Function Execution Space Specifiers

1.1.__global__

1.2.__device__

1.3.__host__

1.4.Undefined behavior

1.5.__noinline__and__forceinline__

1.6.__inline_hint__

2.Variable Memory Space Specifiers

2.1.__device__

2.2.__constant__

2.3.__shared__

2.4.__grid_constant__

2.5.__managed__

2.6.__restrict__

3.Built-in Vector Types

3.1.char, short, int, long, longlong, float, double

3.2.dim3

4.Built-in Variables

4.1.gridDim

4.2.blockIdx

4.3.blockDim

4.4.threadIdx

4.5.warpSize

5.Memory Fence Functions

CATALOG

FEATURED TAGS

1.1.`global`

1.2.`device`

1.3.`host`

1.5.`noinline`and`forceinline`

1.6.`__inline_hint__`

2.1.`device`

2.2.`constant`

2.3.`shared`

2.4.`__grid_constant__`

2.5.`managed`

2.6.`restrict`