5.4 C/C++ 语言扩展

5.4.1 函数与变量注解

5.4.1.1 执行空间说明符

执行空间说明符 __host__、__device__、__tile__、__global__ 和 __tile_global__ 用于指示函数是在主机端、SIMT（单指令多线程）还是 tile 上下文中执行。

表 39：执行空间说明符

执行空间说明符	在其中执行	可从中调用
	主机端	SIMT
`__host__`（无说明符）	✓	×
`__device__`	×	✓
`__global__`	×	✓
`__tile__`	×	×
`__tile_global__`	×	×
`__host__ __device__ __tile__`	✓	✓

__global__ 和 __tile_global__ 函数的约束：

必须返回 void。
不能是类、结构体或联合体的成员。
需要如核函数配置所述的执行配置。
不支持递归。
有关其他限制，请参阅 __global__ 函数参数。

对 __global__ 和 __tile_global__ 函数的调用是异步的。它们在设备完成执行之前就将控制权返回给主机线程。

声明了多个执行空间的函数（例如 __host__ __device__）会为每个上下文分别编译。可以使用 __CUDA_ARCH__ 宏来区分主机端和设备端代码路径：

__host__ __device__ void func() {
#if defined(__CUDA_ARCH__)
    // 设备代码路径
#else
    // 主机代码路径
#endif
}

5.4.1.2 内存空间说明符

内存空间说明符 __device__、__tile__、__managed__、__constant__ 和 __shared__ 用于指示变量在设备上的存储位置。下表总结了各内存空间的属性：

表 40：内存空间说明符

内存空间说明符	位置	可访问方	生命周期	唯一实例
`__device__`	设备全局内存	设备线程（网格）/ CUDA 运行时 API	程序 / CUDA 上下文	每个设备
`__tile__`	设备全局内存	Tile 块 / CUDA 运行时 API	程序 / CUDA 上下文	每个设备
`__constant__`	设备常量内存	设备线程（网格）/ CUDA 运行时 API	程序 / CUDA 上下文	每个设备
`__managed__`	主机和设备（自动）	主机 / 设备线程	程序	每个程序
`__shared__`	设备端（流多处理器，SM）	线程块内线程	线程块	线程块
无说明符	设备（寄存器）	单个线程	单个线程	单个线程

主机端可以使用 CUDA 运行时 API 函数 cudaGetSymbolAddress()、cudaGetSymbolSize()、cudaMemcpyToSymbol() 和 cudaMemcpyFromSymbol() 来访问 __device__、__tile__ 和 __constant__ 变量。
__constant__ 变量在设备代码中是只读的，只能通过 CUDA 运行时 API 从主机端修改。

以下示例演示了如何使用这些 API：

__device__    float device_var        = 4.0f; // 设备内存中的变量
__constant__ float constant_mem_var = 4.0f; // 常量内存中的变量
                                              // 为便于阅读，以下示例聚焦于设备变量。

int main() {
     float* device_ptr;
     cudaGetSymbolAddress((void**) &device_ptr, device_var); // 获取 device_var 的地址

     size_t symbol_size;
     cudaGetSymbolSize(&symbol_size, device_var); // 获取该符号的大小（4 字节）。

     float host_var;
     cudaMemcpyFromSymbol(&host_var, device_var, sizeof(host_var)); // 从设备复制到主机。

     host_var = 3.0f;
     cudaMemcpyToSymbol(device_var, &host_var, sizeof(host_var)); // 从主机复制到设备。
}

参见 Compiler Explorer 上的示例。

5.4.1.2.1 `shared` 内存

__shared__ 内存变量可以是静态大小（在编译时确定）或动态大小（在核函数启动时确定）。有关在运行时指定共享内存大小的详细信息，请参见核函数配置一节。

共享内存的约束：

动态大小的变量必须声明为外部数组或指针。
静态大小的变量不能在其声明中初始化。

以下示例演示了如何声明 __shared__ 变量并指定其大小：

extern __shared__ char dynamic_smem_pointer[];
// extern __shared__ char* dynamic_smem_pointer; // 另一种写法

__global__ void kernel() { // 或 __device__ 函数
    __shared__ int smem_var1[4];                   // 静态大小
    auto smem_var2 = (int*) dynamic_smem_pointer; // 动态大小
}

int main() {
    size_t shared_memory_size = 16;
    kernel<<<1, 1, shared_memory_size>>>();
    cudaDeviceSynchronize();
}

参见 Compiler Explorer 上的示例。

5.4.1.2.2 `managed` 内存

__managed__ 变量有以下限制：

__managed__ 变量的地址不是常量表达式。
__managed__ 变量不得具有引用类型 T&。
在 CUDA 运行时可能处于无效状态时，不得使用 __managed__ 变量的地址或值，包括以下情况：
在具有静态或 thread_local 存储期的对象的静态 / 动态初始化或析构过程中。
在 exit() 被调用后执行的代码中。例如，带有 __attribute__((destructor)) 的函数。
在 CUDA 运行时可能尚未初始化时执行的代码中。例如，带有 __attribute__((constructor)) 的函数。
__managed__ 变量不能用作 decltype() 表达式的无括号 id-expression 参数。
__managed__ 变量具有与动态分配的托管内存相同的连贯性和一致性行为。
另请参见局部变量的限制。

以下是 __managed__ 变量的合法与非法用法示例：

#include <cassert>

__device__ __managed__ int global_var = 10; // OK

int* ptr = &global_var; // ERROR：在静态初始化中使用了托管变量

struct MyStruct1 {
    int field;
    MyStruct1() : field(global_var) {};
};

struct MyStruct2 {
    ~MyStruct2() { global_var = 10; }
};

MyStruct1 temp1; // ERROR：在动态初始化中使用了托管变量

MyStruct2 temp2; // ERROR：在具有静态存储期的
                 //        对象的析构函数中使用了托管变量

__device__ __managed__ const int const_var = 10; // ERROR：const 限定类型

__device__ __managed__ int&                 reference = global_var; // ERROR：引用类型

template <int* Addr>
struct MyStruct3 {};

MyStruct3<&global_var> temp; // ERROR：托管变量的地址不是常量表达式

__global__ void kernel(int* ptr) {
    assert(ptr == &global_var); // OK
    global_var = 20;             // OK
}

int main() {
    int* ptr = &global_var;                 // OK
    kernel<<<1, 1>>>(ptr);
    cudaDeviceSynchronize();
    global_var++;                 // OK
    decltype(global_var) var1;    // ERROR：将托管变量用作 decltype 的无括号参数
      decltype((global_var)) var2; // OK
}

5.4.1.2.3 `tile` 变量

__tile__ 内存空间说明符与 __device__ 类似。标记为 __tile__ 的变量在设备全局内存中分配，可以被 CUDA 运行时 API 函数访问。与 __device__ 变量不同，__tile__ 变量可以从 tile 代码中直接访问。

一般而言，__device__ 变量不能按名称从 __tile__ 或 __tile_global__ 代码中直接访问，__tile__ 变量也不能按名称从 __device__ 或 __global__ 代码中直接访问。不过，这两种变量都分配在设备全局内存中，指向该内存的指针可以在两种上下文中访问。

以下示例展示了 tile 变量的合法与非法用法：

__device__ int x_device = 0;
__tile__   int x_tile = 0;

// ERROR：变量不能同时是 device 和 tile
__device__ __tile__ int x_device_tile = 0;

__global__ void simt_function(int* in) {
  x_device = 1; // OK：从 SIMT 代码访问 __device__ 变量
  x_tile = 1;   // ERROR：从 SIMT 代码访问 __tile__ 变量

    *in = 1; // OK：可以指向任何设备全局内存，包括 x_tile。
}

__tile_global__ void tile_function(int* in) {
  x_device = 1; // ERROR：从 tile 代码访问 __device__ 变量
  x_tile = 1;   // OK：从 tile 代码访问 __tile__ 变量

    *in = 1; // OK：可以指向任何设备全局内存，包括 x_device。
}

int main() {
   int* x_device_ptr;
   cudaGetSymbolAddress((void**) &x_device_ptr, x_device); // 获取 x_device 的地址

  int* x_tile_ptr;
  cudaGetSymbolAddress((void**) &x_tile_ptr, x_tile); // 获取 x_tile 的地址

    // OK：将指向 tile 变量的指针传入 SIMT 核函数
    simt_function<<<1,1>>>(x_tile_ptr);
    cudaDeviceSynchronize();

    // OK：将指向 device 变量的指针传入 tile 核函数
    tile_function<<<1,1>>>(x_device_ptr);
    cudaDeviceSynchronize();
}

__tile__ 变量不能包含指针或引用类型的子对象。例如：

__tile__ int* ptr; // ERROR

struct S1 { int* ptr; };
__tile__ S1 val; // ERROR

5.4.1.3 内联说明符

以下说明符可用于控制 __host__ 和 __device__ 函数的内联：

__noinline__：指示 nvcc 不内联该函数。
__forceinline__：强制 nvcc 在单个翻译单元内内联该函数。
__inline_hint__：在使用链接时优化（Link-Time Optimization）时，在翻译单元之间启用激进的内联。

这些说明符互斥。将其应用于 __tile__ 函数时会被忽略。

5.4.1.4 `restrict` 指针

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

指针别名（pointer aliasing）是指两个或多个指针引用重叠的内存区域。这会抑制诸如代码重排和公共子表达式消除等优化。

限定为 restrict 的指针是程序员作出的一项承诺：在该指针的生命周期内，其指向的内存只会通过该指针访问。这使得编译器能够执行更激进的优化。

访问该设备函数的所有线程仅从其中读取；或者
至多一个线程对其进行写入，且没有其他线程对其读取。

以下示例展示了一个别名问题，并演示了使用限定指针如何帮助编译器减少指令数量：

__device__
void device_function(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 可能存在别名，通过 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 的元素。函数原型变为：

__device__
void device_function(const float* __restrict__ a, const float* __restrict__ b,
                     float* __restrict__ c);

注意，为了让编译器优化器发挥作用，所有指针参数都必须是限定的。加上 __restrict__ 关键字后，编译器可以在保持与抽象执行模型完全相同的功能的同时，随意进行重排和公共子表达式消除：

__device__
void device_function(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;
     ...
}

参见 Compiler Explorer 上的示例。

结果是内存访问和计算次数减少，代价是寄存器压力增加——因为加载和公共子表达式被缓存在寄存器中。

由于寄存器压力在许多 CUDA 代码中是关键问题，使用限定指针可能因降低占用率而对性能产生负面影响。

对标记为 __restrict__ 的 __global__ 函数 const 指针的访问会被编译为只读缓存加载，类似于 PTX ld.global.nc 或 __ldg() 底层加载与存储函数指令：

__global__
void kernel1(const float* in, float* out) {
    *out = *in; // PTX：ld.global
}

__global__
void kernel2(const float* __restrict__ in, float* out) {
    *out = *in; // PTX：ld.global.nc
}

参见 Compiler Explorer 上的示例。

5.4.1.5 `__grid_constant__` 参数

用 __grid_constant__ 注解 __global__ 函数的参数，可以阻止编译器为该参数创建每线程副本。相反，网格中的所有线程都通过单个地址访问该参数，这可以提高性能。

__grid_constant__ 参数具有以下属性：

其生命周期为核函数执行期间。
它对单个核函数私有，即其他网格（包括子网格）的线程无法访问该对象。
核函数中的所有线程看到的地址相同。
它是只读的。修改 __grid_constant__ 对象或其任何子对象（包括可变成员）的行为是未定义的。

要求：

用 __grid_constant__ 注解的核函数参数必须是 const 限定的非引用类型。
所有函数声明必须与 __grid_constant__ 参数保持一致。
函数模板特化在 __grid_constant__ 参数方面必须与主模板声明匹配。
函数模板实例化在 __grid_constant__ 参数方面也必须与主模板声明匹配。
在 __tile_global__ 函数参数中，__grid_constant__ 注解会被忽略。

示例：

struct MyStruct {
    int         x;
    mutable int y;
};

__device__ void external_function(const MyStruct&);

__global__ void kernel(const __grid_constant__ MyStruct s) {
    // s.x++; // 编译错误：试图修改只读内存
    // s.y++; // 未定义行为：试图修改只读内存

      // 编译器不会为 "s" 创建每线程局部副本：
      external_function(s);
}

参见 Compiler Explorer 上的示例。

5.4.1.6 注解汇总

下表总结了 CUDA 注解，并报告了每个注解适用于哪个执行空间以及在哪里有效。

表 41：注解汇总

注解	`__host__` / `__device__` / `__host__ __device__`	`__global__`
`__noinline__`、`__forceinline__`、`__inline_hint__`	函数	×
`__restrict__`	指针参数	指针参数
`__grid_constant__`	×	参数
`__launch_bounds__`	×	函数
`__maxnreg__`	×	函数
`__cluster_dims__`	×	函数

5.4.2 内建类型与变量

5.4.2.1 主机编译器类型扩展

只要主机编译器支持，CUDA 允许使用非标准的算术类型。支持以下类型：

128 位整数类型 __int128。
在主机编译器定义了 __SIZEOF_INT128__ 宏时，Linux 上支持。
128 位浮点类型 __float128 和 _Float128 适用于计算能力 10.0 及更高版本的 GPU 设备。__float128 类型的常量表达式可能被编译器以更低精度的浮点表示来处理。
在主机编译器定义了 __SIZEOF_FLOAT128__ 或 __FLOAT128__ 宏时，Linux x86 上支持。
_Complex 类型仅在主机代码中支持。

128 位整数和浮点类型在 tile 代码中不受支持。

5.4.2.2 内建变量

用于在 x、y、z 维度上指定和获取网格与线程块的核函数配置的值，类型为 dim3。用于获取线程块和线程索引的变量，类型为 uint3。dim3 和 uint3 都是由名为 x、y、z 的三个无符号值组成的平凡结构体。在 C++11 及更高版本中，dim3 所有分量的默认值为 1。

内建的仅设备端变量：

dim3 gridDim：包含网格在 x、y、z 维度上的维度，即线程块的数量。
dim3 blockDim：包含线程块在 x、y、z 维度上的维度，即线程的数量。
uint3 blockIdx：包含线程块在网格内于 x、y、z 维度上的索引。
uint3 threadIdx：包含线程在线程块内于 x、y、z 维度上的索引。
int warpSize：定义为一个线程束（warp）中线程数量的运行时值，通常为 32。有关线程束的定义，另请参见线程束与 SIMT。

这些变量在 tile 代码中不受支持。要在 tile 代码中获取线程块 ID 或网格大小，请使用 cuda::tiles::bid() 和 cuda::tiles::num_blocks() API。

5.4.2.3 内建类型

CUDA 提供了从基本整数和浮点类型派生的向量类型，主机端和设备端都支持。下表列出了可用的向量类型。

表 42：向量类型

C++ 类型	基本向量 X1	向量 X2	向量 X3	向量 X4
signed char	char1	char2	char3	char4
unsigned char	uchar1	uchar2	uchar3	uchar4
signed short	short1	short2	short3	short4
unsigned short	ushort1	ushort2	ushort3	ushort4
signed int	int1	int2	int3	int4
unsigned	uint1	uint2	uint3	uint4
signed long	long1	long2	long3	long4_16a / long4_32a
unsigned long	ulong1	ulong2	ulong3	ulong4_16a / ulong4_32a
signed long long	longlong1	longlong2	longlong3	longlong4_16a / longlong4_32a
unsigned long long	ulonglong1	ulonglong2	ulonglong3	ulonglong4_16a / ulonglong4_32a
float	float1	float2	float3	float4
double	double1	double2	double3	double4_16a / double4_32a

注意，long4、ulong4、longlong4、ulonglong4 和 double4 已在 CUDA 13 中被弃用，将来可能被移除。

下表详细列出了向量类型的字节大小和对齐要求：

表 43：对齐要求

类型	大小	对齐
char1、uchar1	1	1
char2、uchar2	2	2
char3、uchar3	3	1
char4、uchar4	4	4
short1、ushort1	2	2
short2、ushort2	4	4
short3、ushort3	6	2
short4、ushort4	8	8
int1、uint1	4	4
int2、uint2	8	8
int3、uint3	12	4
int4、uint4	16	16
long1、ulong1	4/8 *	4/8 *
long2、ulong2	8/16 *	8/16 *
long3、ulong3	12/24 *	4/8 *
long4、ulong4（已弃用）	16/32 *	16 *
long4_16a、ulong4_16a	16/32 *	16
long4_32a、ulong4_32a	16/32 *	32
longlong1、ulonglong1	8	8
longlong2、ulonglong2	16	16
longlong3、ulonglong3	24	8
longlong4、ulonglong4（已弃用）	32	16
longlong4_16a、ulonglong4_16a	32	16
longlong4_32a、ulonglong4_32a	32	32
float1	4	4
float2	8	8
float3	12	4
float4	16	16
double1	8	8
double2	16	16
double3	24	8
double4（已弃用）	32	16
double4_16a	32	16
double4_32a	32	32

* 在 C++ LLP64 数据模型（Windows 64 位）中，long 为 4 字节，而在 C++ LP64 数据模型（Linux 64 位）中为 8 字节。

向量类型是结构体。它们的第 1、第 2、第 3 和第 4 个分量分别可通过 x、y、z 和 w 字段访问。

int sum(int4 value) {
    return value.x + value.y + value.z + value.w;
}

它们都有一个 make_<type_name>() 形式的工厂函数；例如：

int4 add_one(int x, int y, int z, int w) {
    return make_int4(x + 1, y + 1, z + 1, w + 1);
}

如果主机代码不用 nvcc 编译，可以通过包含 CUDA 工具包提供的 cuda_runtime.h 头文件来引入向量类型和相关函数。

5.4.3 核函数配置

对 __global__ 或 __tile_global__ 函数的任何调用都必须为该调用指定执行配置。该执行配置定义了将在设备上执行该函数的网格和线程块的维度，以及关联的流。

执行配置通过在函数名和括号括起来的参数列表之间插入形如 <<<grid_dim, block_dim, dynamic_smem_bytes, stream>>> 的表达式来指定，其中：

grid_dim 的类型为 dim3，指定网格的维度和大小，使得 grid_dim.x * grid_dim.y * grid_dim.z 等于所启动的线程块数量。
block_dim 的类型为 dim3，指定每个线程块的维度和大小，使得 block_dim.x * block_dim.y * block_dim.z 等于每个线程块的线程数量。对于 __tile_global__ 核函数，block_dim 必须为 1，因为编译器会在启动 tile 核函数时决定调度多少线程。
dynamic_smem_bytes 是可选的 size_t 参数，默认为 0。它指定本次调用中每个线程块除静态分配的内存外，动态分配的共享内存字节数。该内存供 extern __shared__ 数组使用（见 __shared__ 内存）。
stream 的类型为 cudaStream_t（指针），指定关联的流。stream 是可选参数，默认为 NULL。

以下示例展示了核函数的声明和调用：

__global__ void kernel(float* parameter);

kernel<<<grid_dim, block_dim, dynamic_smem_bytes>>>(parameter);

执行配置的参数先于实际函数的参数求值。

如果 grid_dim 或 block_dim 超过设备允许的最大尺寸（如计算能力一节所述），或者 dynamic_smem_bytes 在计入静态分配的内存后超过可用的共享内存，则该函数调用失败。

5.4.3.1 线程块簇

__global__ void __cluster_dims__(2, 1, 1) kernel(float* parameter);

默认形式的 __cluster_dims__() 指定核函数将作为网格簇启动。如果未指定某个簇维度，用户可以在启动时指定。若在启动时未指定某个维度，将导致启动时错误。__cluster_dims__ 属性不能用于 __tile_global__ 核函数，因为编译器决定如何在簇之间调度 tile 核函数。

线程块簇的维度也可以在运行时指定，具有簇的核函数可以使用 cudaLaunchKernelEx API 启动。该 API 接受类型为 cudaLaunchConfig_t 的配置参数、核函数指针和核函数参数。下例展示了运行时核函数配置。

__global__ void kernel(float parameter1, int parameter2) {}

int main() {
     cudaLaunchConfig_t config = {0};
     // 网格维度不受簇启动影响，仍按线程块数量枚举。
     // 网格维度应为簇大小的倍数。
     config.gridDim           = dim3{4}; // 4 个线程块
     config.blockDim          = dim3{32}; // 每个线程块 32 个线程
     config.dynamicSmemBytes = 1024;      // 1 KB

      cudaLaunchAttribute attribute[1];
      attribute[0].id               = cudaLaunchAttributeClusterDimension;
      attribute[0].val.clusterDim.x = 2; // X 维度的簇大小
      attribute[0].val.clusterDim.y = 1;
      attribute[0].val.clusterDim.z = 1;
      config.attrs    = attribute;
      config.numAttrs = 1;

      float parameter1 = 3.0f;
      int   parameter2 = 4;
      cudaLaunchKernelEx(&config, kernel, parameter1, parameter2);
}

参见 Compiler Explorer 上的示例。

5.4.3.2 启动界限

如核函数启动与占用率一节所述，使用更少的寄存器可以让更多线程和线程块驻留在多处理器上，从而提高性能。

因此，编译器使用启发式算法在保持寄存器溢出和指令数量最小的同时最小化寄存器使用。应用程序可以选择性地以启动界限（launch bounds）的形式向编译器提供额外信息来辅助这些启发式算法，启动界限通过在 __global__ 函数定义中使用 __launch_bounds__() 限定符指定：

__global__ void
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor,
                  maxBlocksPerCluster)
MyKernel(...) {
     ...
}

maxThreadsPerBlock 指定应用程序启动 MyKernel() 时每个线程块的最大线程数；它编译为 .maxntid PTX 指令。
minBlocksPerMultiprocessor 可选，指定每个多处理器期望的最小驻留线程块数；它编译为 .minnctapersm PTX 指令。
maxBlocksPerCluster 可选，指定应用程序启动 MyKernel() 时每个簇的最大线程块数；它编译为 .maxclusterrank PTX 指令。

如果指定了启动界限，编译器首先推导核函数应使用的寄存器数量的上限 L。这确保了 maxThreadsPerBlock 个线程的 minBlocksPerMultiprocessor 个线程块（如果未指定 minBlocksPerMultiprocessor，则为单个线程块）可以驻留在多处理器上。有关核函数使用的寄存器数量与每个线程块分配的寄存器数量之间的关系，请参见占用率一节。然后编译器按如下方式优化寄存器使用：

如果初始寄存器使用超过 L，编译器将其减少到小于或等于 L。这通常会导致本地内存使用增加和 / 或指令数量增加。
如果初始寄存器使用低于 L
如果指定了 maxThreadsPerBlock 但未指定 minBlocksPerMultiprocessor，编译器使用 maxThreadsPerBlock 来确定 n 个与 n+1 个驻留线程块之间转换的寄存器使用阈值。当少用一个寄存器就能容纳一个额外的驻留线程块时，就会发生这种情况。然后，编译器应用与未指定启动界限时类似的启发式算法。
如果同时指定了 minBlocksPerMultiprocessor 和 maxThreadsPerBlock，编译器可以将寄存器使用增加到 L，以减少指令数量并更好地隐藏单线程指令的延迟。

如果核函数在以下条件下执行，将启动失败：

每个线程块的线程数多于其启动界限 maxThreadsPerBlock。
每个簇的线程块数多于其启动界限 maxBlocksPerCluster。

CUDA 核函数所需的每线程资源可能以不理想的方式限制最大线程块大小。为了保持与未来硬件和工具包的前向兼容性，并确保至少一个线程块能在流多处理器（Streaming Multiprocessor，SM）上运行，开发者应包含单参数的 __launch_bounds__(maxThreadsPerBlock)，它指定核函数启动时的最大线程块大小。不这样做可能导致 "too many resources requested for launch"（为启动请求的资源过多）错误。在某些情况下，提供 __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor) 的双参数版本可以提高性能。minBlocksPerMultiprocessor 的最佳值应通过对每个核函数的详细分析来确定。

核函数的最优启动界限通常在不同主架构版本之间有所不同。以下代码示例说明了如何在设备代码中用 __CUDA_ARCH__ 宏管理这一点。

#define THREADS_PER_BLOCK          256

#if __CUDA_ARCH__ >= 900
    #define MY_KERNEL_MAX_THREADS            (2 * THREADS_PER_BLOCK)
    #define MY_KERNEL_MIN_BLOCKS             3
#else
    #define MY_KERNEL_MAX_THREADS            THREADS_PER_BLOCK
    #define MY_KERNEL_MIN_BLOCKS             2
#endif

__global__ void
__launch_bounds__(MY_KERNEL_MAX_THREADS, MY_KERNEL_MIN_BLOCKS)
MyKernel(...) {
    ...
}

当用每个线程块的最大线程数（即 __launch_bounds__() 的第一个参数）调用 MyKernel 时，人们很想把 MY_KERNEL_MAX_THREADS 用作执行配置中每个线程块的线程数：

// 主机代码
MyKernel<<<blocksPerGrid, MY_KERNEL_MAX_THREADS>>>(...);

然而这行不通，因为如执行空间说明符一节所述，__CUDA_ARCH__ 在主机代码中是未定义的。因此，MyKernel 将以每个线程块 256 个线程启动。应改为按以下方式确定每个线程块的线程数：

要么在编译时使用不依赖 __CUDA_ARCH__ 的宏或常量来确定，例如

c // 主机代码 MyKernel<<<blocksPerGrid, THREADS_PER_BLOCK>>>(...);

要么在运行时基于计算能力确定

```c // 主机代码 cudaGetDeviceProperties(&deviceProp, device); int threadsPerBlock = (deviceProp.major >= 9) ? 2 * THREADS_PER_BLOCK : THREADS_PER_BLOCK;

MyKernel<<>>(...); ```

--resource-usage 编译器选项报告寄存器使用情况。CUDA 分析器报告占用率，可用于推导驻留线程块的数量。

5.4.3.3 每个线程的最大寄存器数

为了支持底层性能调优，CUDA C++ 提供了 __maxnreg__() 函数限定符，它将性能调优信息传递给后端优化编译器。__maxnreg__() 限定符指定可分配给线程块中单个线程的最大寄存器数。在 __global__ 函数的定义中：

__global__ void
__maxnreg__(maxNumberRegistersPerThread)
MyKernel(...) {
    ...
}

maxNumberRegistersPerThread 变量指定要分配给核函数 MyKernel() 的线程块中单个线程的最大寄存器数；它编译为 .maxnreg PTX 指令。

__launch_bounds__() 和 __maxnreg__() 限定符不能同时应用于同一个核函数。

--maxrregcount <N> 编译器选项可用于控制文件中所有 __global__ 函数的寄存器使用。对于带有 __maxnreg__ 限定符的核函数，该选项被忽略。

5.4.4 同步原语

5.4.4.1 线程块同步函数

void __syncthreads();
int __syncthreads_count(int predicate);
int __syncthreads_and(int predicate);
int __syncthreads_or(int predicate);

这些内联函数（intrinsic）协调同一线程块内线程之间的通信。当线程块中的线程访问共享内存或全局内存中的相同地址时，可能发生读后写、写后读或写后写冒险。通过在这些访问之间同步线程，可以避免这些冒险。

这些内联函数具有以下语义：

__syncthreads*() 等待，直到线程块中的所有未退出线程同时到达程序中同一个 __syncthreads*() 内联函数调用或退出。
__syncthreads*() 在参与线程之间提供内存序：对 __syncthreads*() 内联函数的调用强先于（见 C++ 规范 [intro.races]）任何参与线程从等待中解除阻塞或退出。

以下示例展示了如何使用 __syncthreads() 同步线程块内的线程，并安全地对线程间共享的数组元素求和：

#include <cuda_runtime_api.h>
#include <memory.h>
#include <cstdlib>
#include <ctime>
#include <stdio.h>

// 假设 blockDim.x 为 128
__global__ void example_syncthreads(int* input_data, int* output_data)
{
     __shared__ int shared_data[128];
     shared_data[threadIdx.x] = input_data[blockDim.x*blockIdx.x + threadIdx.x];

       // 所有线程同步，保证对 'shared_data' 的所有写入在任何线程
       // 从 '__syncthreads()' 解除阻塞之前有序：
       __syncthreads();

       // 单个线程安全地读取 'shared_data'：
       if (threadIdx.x == 0) {
           float sum = 0;
           for (int i = 0; i < blockDim.x; ++i) {
               sum += shared_data[i];
           }
           output_data[blockIdx.x] = sum;
       }
}

void initArray(int* A, int length)
{
      std::srand(std::time({}));
    for(int i=0; i<length; i++)
    {
         A[i] = int(10.f * (rand() / (float)RAND_MAX));
    }
}


int main(int argc, char** arg)
{
    constexpr int block_size = 128;
    constexpr int input_length = 1024;
    constexpr int output_length = input_length / block_size;

       // 指向内存向量的指针
       int* input = nullptr;
       int* output = nullptr;
       int* comparisonResult = (int*)malloc(output_length*sizeof(int));

       // 使用统一内存分配缓冲区
       cudaMallocManaged(&input, input_length*sizeof(int));
       cudaMallocManaged(&output, output_length*sizeof(int));

       initArray(input, input_length);
       int grid_size = input_length / block_size;

       example_syncthreads<<<grid_size, block_size>>>(input, output);
     cudaDeviceSynchronize();

     for(int i=0; i<output_length; i++)
     {
         comparisonResult[i] = 0;
         for(int j=0; j < block_size; j++)
         {
             comparisonResult[i] += input[i*block_size + j];
         }
     }

    for(int i=0; i< output_length; i++)
    {
        if(output[i] != comparisonResult[i])
        {
            printf("Results do not match at index %d: %d != %d\n", i,
                   output[i], comparisonResult[i]);

            exit(-1);
        }
    }
    printf("Test passed\n");

     return 0;
}

__syncthreads*() 内联函数允许出现在条件代码中，但前提是该条件在整个线程块内求值为一致（uniform）。否则，执行可能挂起或产生意外的副作用。

以下示例演示了有效的行为：

// 假设 blockDim.x 为 128
__global__ void syncthreads_valid_behavior(int* input_data, int* output_data) {

     __shared__ int shared_data[128];
     shared_data[threadIdx.x] = input_data[threadIdx.x];
     if (blockIdx.x > 0) { // 正确，所有块线程的条件一致
         __syncthreads();
         output_data[threadIdx.x] = shared_data[128 - threadIdx.x];
     }
}

而以下示例表现出无效行为，例如核函数挂起或未定义行为：

// 假设 blockDim.x 为 128
__global__ void syncthreads_invalid_behavior1(int* input_data, int* output_data) {

     __shared__ int shared_data[256];
     shared_data[threadIdx.x] = input_data[threadIdx.x];
     if (threadIdx.x > 0) { // 错误，非一致条件
         __syncthreads();   // 未定义行为
         output_data[threadIdx.x] = shared_data[128 - threadIdx.x];
      }
}

// 假设 blockDim.x 为 128
__global__ void syncthreads_invalid_behavior2(int* input_data, int* output_data) {

     __shared__ int shared_data[256];
     shared_data[threadIdx.x] = input_data[threadIdx.x];
     for (int i = 0; i < blockDim.x; ++i) {
         if (i == threadIdx.x) { // 错误，非一致条件
             __syncthreads();    // 未定义行为
         }
     }
     output_data[threadIdx.x] = shared_data[128 - threadIdx.x];
}

带谓词的 __syncthreads() 变体：

int __syncthreads_count(int predicate);

除了对块中所有未退出线程求值一个谓词，并返回该谓词求值为非零值的线程数量之外，与 __syncthreads() 完全相同。

int __syncthreads_and(int predicate);

除了对块中所有未退出线程求值谓词之外，与 __syncthreads() 完全相同。返回非零值当且仅当所有线程的谓词求值均为非零。

int __syncthreads_or(int predicate);

除了对块中所有未退出线程求值谓词之外，与 __syncthreads() 完全相同。返回非零值当且仅当其中一个或多个线程的谓词求值为非零。

5.4.4.2 线程束同步函数

void __syncwarp(unsigned mask = 0xFFFFFFFF);

内联函数 __syncwarp() 协调同一线程束内线程之间的通信。当线程束中的某些线程访问共享内存或全局内存中的相同地址时，可能发生潜在的读后写、写后读或写后写冒险。通过在这些访问之间同步线程，可以避免这些数据冒险。

调用 __syncwarp(mask) 在 mask 指定的线程束内参与线程之间提供内存序：对 __syncwarp(mask) 的调用强先于（见 C++ 规范 [intro.races]）mask 指定的任一线程束线程从等待中解除阻塞或退出。

这些函数受线程束 __sync 内联函数约束的约束。

以下示例演示了如何使用 __syncwarp() 同步线程束内的线程，以安全地访问共享内存数组：

__global__ void example_syncwarp(int* input_data, int* output_data) {
    if (threadIdx.x < warpSize) {
        __shared__ int shared_data[warpSize];
        shared_data[threadIdx.x] = input_data[threadIdx.x];

            __syncwarp(); // 等价于 __syncwarp(0xFFFFFFFF)
            if (threadIdx.x == 0)
                output_data[0] = shared_data[1];
       }
}

5.4.4.3 内存栅栏函数

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 的结果值可以是任意值。

内存栅栏和同步函数强制内存访问的顺序一致性（sequentially consistent）序。这些函数在强制序的线程作用域上有所不同，但与所访问的内存空间无关，包括共享内存、全局内存、页锁定主机内存和对端设备内存。

提示

出于安全性和可移植性的考虑，建议尽可能使用 libcu++ 提供的 cuda::atomic_thread_fence。

线程块级内存栅栏

CUDA C++

// <cuda/atomic> 头文件
cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_block);

确保：

调用线程在调用 cuda::atomic_thread_fence() 之前对所有内存的所有写入，被调用线程所在线程块中的所有线程观察到发生在调用线程在调用 cuda::atomic_thread_fence() 之后对所有内存的所有写入之前；
调用线程在调用 cuda::atomic_thread_fence() 之前对所有内存的所有读取，序在调用线程在调用 cuda::atomic_thread_fence() 之后对所有内存的所有读取之前。

内联函数

void __threadfence_block();

确保： - 调用线程在调用 __threadfence_block() 之前对所有内存的所有写入，被调用线程所在线程块中的所有线程观察到发生在调用线程在调用 __threadfence_block() 之后对所有内存的所有写入之前； - 调用线程在调用 __threadfence_block() 之前对所有内存的所有读取，序在调用线程在调用 __threadfence_block() 之后对所有内存的所有读取之前。

设备级内存栅栏

CUDA C++

cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_device);

确保：

调用线程在调用 cuda::atomic_thread_fence() 之后对所有内存的写入，不会被设备中的任何线程观察到发生在调用线程在调用 cuda::atomic_thread_fence() 之前对所有内存的任何写入之前。

内联函数

void __threadfence();

确保：

调用线程在调用 __threadfence() 之后对所有内存的写入，不会被设备中的任何线程观察到发生在调用线程在调用 __threadfence() 之前对所有内存的任何写入之前。

系统级内存栅栏

CUDA C++

cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_system);

确保：

调用线程在调用 cuda::atomic_thread_fence() 之前对所有内存的所有写入，被设备中的所有线程、主机线程以及所有对端设备中的线程观察到发生在调用线程在调用 cuda::atomic_thread_fence() 之后对所有内存的所有写入之前。

内联函数

void __threadfence_system();

确保：

调用线程在调用 __threadfence_system() 之前对所有内存的所有写入，被设备中的所有线程、主机线程以及所有对端设备中的线程观察到发生在调用线程在调用 __threadfence_system() 之后对所有内存的所有写入之前。

在之前的代码示例中，可以如下在代码中插入内存栅栏：

CUDA C++

#include <cuda/atomic>

__device__ int X = 1, Y = 2;

__device__ void writeXY() {
     X = 10;
     cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_device);
     Y = 20;
}

__device__ void readXY() {
     int B = Y;
     cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_device);
     int A = 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，即 readXY() 在 writeXY() 之前执行；
A 等于 10 且 B 等于 20，即 writeXY() 在 readXY() 之前执行；
A 等于 10 且 B 等于 2；
A 为 1 且 B 为 20 的情况不可能出现，因为内存栅栏确保了对 X 的写入在对 Y 的写入之前可见。

如果线程 1 和 2 属于同一线程块，使用线程块级栅栏就足够了。如果线程 1 和 2 不属于同一线程块：若它们是同一设备的 CUDA 线程，必须使用设备级栅栏；若它们是两个不同设备的 CUDA 线程，必须使用系统级栅栏。

以下代码示例展示了一个常见用例，其中线程消费其他线程产生的数据。该核函数在单次调用中计算 N 个数字数组的和。

每个线程块先对数组的一个子集求和，并将结果存储到全局内存中。
当所有线程块完成后，最后一个线程块从全局内存中读取每个部分和并将它们相加以获得最终结果。
为了确定哪个线程块最后完成，每个线程块以原子方式递增计数器，以示其部分和的计算和存储已完成（进一步细节见原子函数一节）。最后一个线程块收到的计数器值等于 gridDim.x - 1。

如果在存储部分和与递增计数器之间没有栅栏，计数器可能在部分和存储之前递增。这可能导致计数器达到 gridDim.x - 1，并使最后一个线程块在部分和在内存中更新之前就开始读取它们。

注意

内存栅栏只影响内存操作的执行顺序；它不保证这些操作对其他线程的可见性。

在下面的代码示例中，通过将 result 变量声明为 volatile 来确保对其内存操作的可见性。更多细节请参见 volatile 限定变量一节。

#include <cuda/atomic>

__device__ int count = 0;

__global__ void sum(const float*     array,
                    int              N,
                    volatile float* result) {
    __shared__ bool isLastBlockDone;
    // 每个线程块对输入数组的一个子集求和。
     float partialSum = calculatePartialSum(array, N);

     if (threadIdx.x == 0) {
         // 每个线程块的线程 0 将部分和存储到全局内存。
         // 由于 "result" 变量声明为 volatile，编译器将使用绕过 L1 缓存的存储操作。
         // 这确保最后一个线程块的线程能读取所有其他线程块计算的正确部分和。
         result[blockIdx.x] = partialSum;

         // 线程 0 确保仅在部分和已写入全局内存之后才执行对 "count" 变量的递增。
         cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_device);

          // 线程 0 发出其已完成的信号。
          int count_old = atomicInc(&count, gridDim.x);

          // 线程 0 判断其线程块是否为最后完成的线程块。
          isLastBlockDone = (count_old == (gridDim.x - 1));
     }
     // 同步以确保每个线程读取 isLastBlockDone 的正确值。
     __syncthreads();

     if (isLastBlockDone) {
         // 最后一个线程块对存储在 result[0 .. gridDim.x-1] 中的部分和求和
         float totalSum = calculateTotalSum(result);

          if (threadIdx.x == 0) {
              // 最后一个线程块的线程 0 将总和存储到全局内存并重置 count 变量，
              // 以便下一次核函数调用正常工作。
              result[0] = totalSum;
              count     = 0;
          }
     }
}

5.4.5 原子函数

原子函数对共享数据执行读-改-写操作，使它们看起来像单步执行。原子性确保每个操作要么完全完成，要么根本不执行，为所有参与线程提供一致的数据视图。

CUDA 通过五种方式提供原子函数：

扩展的 CUDA C++ 原子函数，cuda::atomic 和 cuda::atomic_ref。

它们在主机代码和设备代码中都允许使用。
它们遵循 C++ 标准原子操作语义。
它们允许指定原子操作的线程作用域。

标准 C++ 原子函数，cuda::std::atomic 和 cuda::std::atomic_ref。

它们在主机代码和设备代码中都允许使用。
它们遵循 C++ 标准原子操作语义。
它们不允许指定原子操作的线程作用域。

编译器内建原子函数，__nv_atomic_<op>()。

它们自 CUDA 12.8 起可用。
它们仅允许在设备代码中使用。
它们遵循 C++ 标准原子内存序语义。
它们允许指定原子操作的线程作用域。
它们与 C++ 标准原子操作具有相同的内存序语义。
它们支持 cuda::std::atomic 和它们支持 cuda::std::atomic 和 cuda::std::atomic_ref 允许的数据类型的子集，128 位数据类型除外。
它们在 tile 代码中不受支持。

CUDA Tile C++ 原子函数（例如 cuda::tiles::atomic_load）：

它们只允许在 tile 代码中使用。
它们遵循 C++ 标准原子内存序语义。
它们允许通过 cuda::tiles::thread_scope 指定线程作用域。

传统原子函数，atomic<Op>()。

它们仅允许在设备代码中使用。
它们仅支持 memory_order_relaxed 的 C++ 原子内存语义。
它们允许将原子操作的线程作用域作为函数名的一部分来指定。
与内建原子函数不同，传统原子函数仅保证原子性，不引入同步点（栅栏）。
它们支持内建原子函数允许的数据类型的子集。原子加操作支持额外的数据类型。

提示

出于效率、安全性和可移植性的考虑，建议使用 libcu++ 提供的扩展 CUDA C++ 原子函数。

5.4.5.1 传统原子函数

传统原子函数对存储在全局内存或共享内存中的 32、64 或 128 位字执行原子读-改-写操作。例如，atomicAdd() 函数读取全局内存或共享内存中特定地址处的一个字，加上一个数，并将结果写回同一地址。

原子函数只能在设备函数中使用。
对于 __half2、__nv_bfloat162、float2 和 float4 等向量类型，读-改-写操作对向量的每个元素执行；不保证整个向量在单次访问中是原子的。

本节所述的原子函数的内存序为 cuda::std::memory_order_relaxed，并且只在特定的线程作用域下是原子的：

不带后缀的原子 API（例如 atomicAdd）在 cuda::thread_scope_device 作用域下是原子的。
带 _block 后缀的原子 API（例如 atomicAdd_block）在 cuda::thread_scope_block 作用域下是原子的。
带 _system 后缀的原子 API（例如 atomicAdd_system），如果满足特定条件，在 cuda::thread_scope_system 作用域下是原子的。

以下示例展示了 CPU 和 GPU 以原子方式更新地址 addr 处的整数值：

#include <cuda_runtime.h>

__global__ void atomicAdd_kernel(int* addr) {
    atomicAdd_system(addr, 10);
}

void test_atomicAdd(int device_id) {
    int* addr;
    cudaMallocManaged(&addr, 4);
    *addr = 0;

    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, device_id);
    if (deviceProp.concurrentManagedAccess != 1) {
        return; // 该设备不能与 CPU 并发地一致访问托管内存
    }

      atomicAdd_kernel<<<...>>>(addr);
      __sync_fetch_and_add(addr, 10); // CPU 原子操作
}

注意，任何原子操作都可以基于 atomicCAS()（比较并交换，Compare and Swap）实现。例如，单精度浮点数的 atomicAdd() 可以如下实现：

#include <cuda/memory>
#include <cuda/std/bit>

__device__ float customAtomicAdd(float* d_ptr, float value) {
     volatile unsigned* d_ptr_unsigned = reinterpret_cast<unsigned*>(d_ptr);
     unsigned old_value         = *d_ptr_unsigned;
     unsigned assumed;
     do {
          assumed                           = old_value;
          float     assumed_float           = cuda::std::bit_cast<float>(assumed);
          float     expected_value          = assumed_float + value;
          unsigned expected_value_unsigned = cuda::std::bit_cast<unsigned>(expected_value);
          old_value                         = atomicCAS(d_ptr_unsigned, assumed,
                                                        expected_value_unsigned);
     // 注意：使用整数比较以避免 NaN 时挂起（因为 NaN != NaN）
     } while (assumed != old_value);
     return cuda::std::bit_cast<float>(old_value);
}

参见 Compiler Explorer 上的示例。

5.4.5.1.1 `atomicAdd()`

T atomicAdd(T* address, T val);

该函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old + val。
将结果存回同一地址的内存。

该函数返回旧值。

atomicAdd() 支持以下数据类型：

int、unsigned、unsigned long long、float、double、__half2、__half。
__nv_bfloat16、__nv_bfloat162，适用于计算能力 8.x 及更高版本的设备。
float2、float4，适用于计算能力 9.x 及更高版本的设备，且仅支持全局内存地址。

对向量类型（如 __half2 或 float4）应用 atomicAdd() 的原子性是针对每个分量分别保证的；不保证整个向量作为单次访问是原子的。

5.4.5.1.2 `atomicSub()`

T atomicSub(T* address, T val);

该函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old - val。
将结果存回同一地址的内存。

该函数返回旧值。

atomicSub() 支持以下数据类型：

int、unsigned。

5.4.5.1.3 `atomicInc()`

unsigned atomicInc(unsigned* address, unsigned val);

该函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old >= val ? 0 : (old + 1)。
将结果存回同一地址的内存。

该函数返回旧值。

5.4.5.1.4 `atomicDec()`

unsigned atomicDec(unsigned* address, unsigned val);

该函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 (old == 0 || old > val) ? val : (old - 1)。
将结果存回同一地址的内存。

该函数返回旧值。

5.4.5.1.5 `atomicAnd()`

T atomicAnd(T* address, T val);

该函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old & val。
将结果存回同一地址的内存。

该函数返回旧值。

atomicAnd() 支持以下数据类型：

int、unsigned、unsigned long long。

5.4.5.1.6 `atomicOr()`

T atomicOr(T* address, T val);

该函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old | val。
将结果存回同一地址的内存。

该函数返回旧值。

atomicOr() 支持以下数据类型：

int、unsigned、unsigned long long。

5.4.5.1.7 `atomicXor()`

T atomicXor(T* address, T val);

该函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old ^ val。
将结果存回同一地址的内存。

该函数返回旧值。

atomicXor() 支持以下数据类型：

int、unsigned、unsigned long long。

5.4.5.1.8 `atomicMin()`

T atomicMin(T* address, T val);

该函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old 和 val 中的最小值。
将结果存回同一地址的内存。

该函数返回旧值。

atomicMin() 支持以下数据类型：

int、unsigned、unsigned long long、long long。

5.4.5.1.9 `atomicMax()`

T atomicMax(T* address, T val);

该函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old 和 val 中的最大值。
将结果存回同一地址的内存。

该函数返回旧值。

atomicMax() 支持以下数据类型：

int、unsigned、unsigned long long、long long。

5.4.5.1.10 `atomicExch()`

T atomicExch(T* address, T val);

template<typename T>
T atomicExch(T* address, T val); // 仅 128 位类型，计算能力 9.x 及更高版本

该函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
将 val 存回同一地址的内存。

该函数返回旧值。

atomicExch() 支持以下数据类型：

int、unsigned、unsigned long long、float。

C++ 模板函数 atomicExch() 支持 128 位类型，要求如下：

计算能力 9.x 及更高版本。
T 必须按 16 字节对齐，即 alignof(T) >= 16。
T 必须可平凡复制，即 std::is_trivially_copyable_v<T>。
对于 C++03 及更早版本：T 必须可平凡构造，即 std::is_default_constructible_v<T>。

5.4.5.1.11 `atomicCAS()`

T atomicCAS(T* address, T compare, T val);

template<typename T>
T atomicCAS(T* address, T compare, T val); // 仅 128 位类型，计算能力 9.x 及更高版本

该函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old == compare ? val : old。
将结果存回同一地址的内存。

该函数返回旧值。

atomicCAS() 支持以下数据类型：

int、unsigned、unsigned long long、unsigned short。

C++ 模板函数 atomicCAS() 支持 128 位类型，要求如下：

计算能力 9.x 及更高版本。
T 必须按 16 字节对齐，即 alignof(T) >= 16。
T 必须可平凡复制，即 std::is_trivially_copyable_v<T>。
对于 C++03 及更早版本：T 必须可平凡构造，即 std::is_default_constructible_v<T>。

5.4.5.2 内建原子函数

CUDA 12.8 及更高版本支持用于原子操作的 CUDA 编译器内建函数，遵循与 C++ 标准原子操作相同的内存序语义和 CUDA 线程作用域。这些函数遵循 GNU 原子内建函数的签名，额外带有一个线程作用域参数。

当支持内建原子函数时，nvcc 定义宏 __CUDACC_DEVICE_ATOMIC_BUILTINS__。

下面列出了内存序和线程作用域的原始枚举值，用作内建原子函数的 order 和 scope 参数：

// 原子内存序
enum {
   __NV_ATOMIC_RELAXED,
   __NV_ATOMIC_CONSUME,
   __NV_ATOMIC_ACQUIRE,
   __NV_ATOMIC_RELEASE,
   __NV_ATOMIC_ACQ_REL,
   __NV_ATOMIC_SEQ_CST
};

// 线程作用域
enum {
   __NV_THREAD_SCOPE_THREAD,
   __NV_THREAD_SCOPE_BLOCK,
   __NV_THREAD_SCOPE_CLUSTER,
   __NV_THREAD_SCOPE_DEVICE,
   __NV_THREAD_SCOPE_SYSTEM
};

内存序对应于 C++ 标准原子操作的内存序。
线程作用域遵循 cuda::thread_scope 定义。
__NV_ATOMIC_CONSUME 内存序目前使用更强的 __NV_ATOMIC_ACQUIRE 内存序实现。
__NV_THREAD_SCOPE_THREAD 线程作用域目前使用更宽的 __NV_THREAD_SCOPE_BLOCK 线程作用域实现。

示例：

__device__ T __nv_atomic_load_n(T* pointer,
                                int memory_order,
                                int thread_scope = __NV_THREAD_SCOPE_SYSTEM);

原子内建函数有以下限制：

它们只能在设备函数中使用。
它们不能操作本地内存。
不能取这些函数的地址。
order 和 scope 参数必须是整数常量；不能是变量。
线程作用域 __NV_THREAD_SCOPE_CLUSTER 在 sm_90 及更高架构上支持。

不支持用法的示例：

 // 不允许在主机函数中使用
 __host__ void bar() {
     unsigned u1 = 1, u2 = 2;
     __nv_atomic_load(&u1, &u2, __NV_ATOMIC_RELAXED, __NV_THREAD_SCOPE_SYSTEM);
 }

  // 不允许应用于本地内存
__device__ void foo() {
    unsigned a = 1, b;
    __nv_atomic_load(&a, &b, __NV_ATOMIC_RELAXED, __NV_THREAD_SCOPE_SYSTEM);
}

 // 不允许用作模板默认参数。
 // 不能取函数地址。
 template<void *F = __nv_atomic_load_n>
 class X {
     void *f = F; // 不能取函数地址。
 };

 // 不允许在构造函数初始化列表中调用。
 class Y {
     int a;
 public:
     __device__ Y(int *b): a(__nv_atomic_load_n(b, __NV_ATOMIC_RELAXED)) {}
 };

5.4.5.2.1 `__nv_atomic_fetch_add()`、`__nv_atomic_add()`

__device__ T     __nv_atomic_fetch_add(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

__device__ void __nv_atomic_add       (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

这些函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old + val。
将结果存回同一地址的内存。
__nv_atomic_fetch_add 返回旧值。
__nv_atomic_add 无返回值。

这些函数支持以下数据类型：

int、unsigned、unsigned long long、float、double。

5.4.5.2.2 `__nv_atomic_fetch_sub()`、`__nv_atomic_sub()`

__device__ T     __nv_atomic_fetch_sub(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

__device__ void __nv_atomic_sub       (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

这些函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old - val。
将结果存回同一地址的内存。
__nv_atomic_fetch_sub 返回旧值。
__nv_atomic_sub 无返回值。

这些函数支持以下数据类型：

int、unsigned、unsigned long long、float、double。

5.4.5.2.3 `__nv_atomic_fetch_and()`、`__nv_atomic_and()`

__device__ T     __nv_atomic_fetch_and(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

__device__ void __nv_atomic_and       (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

这些函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old & val。
将结果存回同一地址的内存。
__nv_atomic_fetch_and 返回旧值。
__nv_atomic_and 无返回值。

这些函数支持以下数据类型：

任何 4 或 8 字节的整型。

5.4.5.2.4 `__nv_atomic_fetch_or()`、`__nv_atomic_or()`

__device__ T     __nv_atomic_fetch_or(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

__device__ void __nv_atomic_or       (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

这些函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old | val。
将结果存回同一地址的内存。
__nv_atomic_fetch_or 返回旧值。
__nv_atomic_or 无返回值。

这些函数支持以下数据类型：

任何 4 或 8 字节的整型。

5.4.5.2.5 `__nv_atomic_fetch_xor()`、`__nv_atomic_xor()`

__device__ T     __nv_atomic_fetch_xor(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

__device__ void __nv_atomic_xor       (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

这些函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old ^ val。
将结果存回同一地址的内存。
__nv_atomic_fetch_xor 返回旧值。
__nv_atomic_xor 无返回值。

这些函数支持以下数据类型：

任何 4 或 8 字节的整型。

5.4.5.2.6 `__nv_atomic_fetch_min()`、`__nv_atomic_min()`

__device__ T     __nv_atomic_fetch_min(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

__device__ void __nv_atomic_min       (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

这些函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old 和 val 中的最小值。
将结果存回同一地址的内存。
__nv_atomic_fetch_min 返回旧值。
__nv_atomic_min 无返回值。

这些函数支持以下数据类型：

unsigned、int、unsigned long long、long long。

5.4.5.2.7 `__nv_atomic_fetch_max()`、`__nv_atomic_max()`

__device__ T     __nv_atomic_fetch_max(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

__device__ void __nv_atomic_max       (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

这些函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
计算 old 和 val 中的最大值。
将结果存回同一地址的内存。
__nv_atomic_fetch_max 返回旧值。
__nv_atomic_max 无返回值。

这些函数支持以下数据类型：

unsigned、int、unsigned long long、long long。

5.4.5.2.8 `__nv_atomic_exchange()`、`__nv_atomic_exchange_n()`

__device__ T     __nv_atomic_exchange_n(T* address, T val,        int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

__device__ void __nv_atomic_exchange (T* address, T* val, T* ret, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

这些函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
__nv_atomic_exchange_n 将 val 存储到 address 指向的位置。__nv_atomic_exchange 将 old 存储到 ret 指向的位置，并将 val 地址处的值存储到 address 指向的位置。
__nv_atomic_exchange_n 返回旧值。
__nv_atomic_exchange 无返回值。

这些函数支持以下数据类型：

任何 4、8 或 16 字节的数据类型。
16 字节数据类型在计算能力 9.x 及更高版本的设备上支持。

5.4.5.2.9 `__nv_atomic_compare_exchange()`、`__nv_atomic_compare_exchange_n()`

__device__ bool __nv_atomic_compare_exchange (T* address, T* expected, T* desired, bool weak, int success_order, int failure_order,
                                              int scope = __NV_THREAD_SCOPE_SYSTEM);

__device__ bool __nv_atomic_compare_exchange_n(T* address, T* expected, T desired, bool weak, int success_order, int failure_order,
                                               int scope = __NV_THREAD_SCOPE_SYSTEM);

这些函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
将 old 与 expected 指向的值比较。
如果相等，返回值为 true，并将 desired 存储到 address 指向的位置。否则返回 false，并将 old 存储到 expected 指向的位置。

参数 weak 被忽略，它选择 success_order 和 failure_order 中更强的内存序来执行比较并交换操作。

这些函数支持以下数据类型：

任何 2、4、8 或 16 字节的数据类型。
16 字节数据类型在计算能力 9.x 及更高版本的设备上支持。

5.4.5.2.10 `__nv_atomic_load()`、`__nv_atomic_load_n()`

__device__ void __nv_atomic_load (T* address, T* ret, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

__device__ T     __nv_atomic_load_n(T* address,       int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

这些函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
__nv_atomic_load 将 old 存储到 ret 指向的位置。__nv_atomic_load_n 返回 old。

这些函数支持以下数据类型：

任何 1、2、4、8 或 16 字节的数据类型。

order 不能是 __NV_ATOMIC_RELEASE 或 __NV_ATOMIC_ACQ_REL。

5.4.5.2.11 `__nv_atomic_store()`、`__nv_atomic_store_n()`

__device__ void __nv_atomic_store (T* address, T* val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

__device__ void __nv_atomic_store_n(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

这些函数在一个原子事务中执行以下操作：

读取位于全局内存或共享内存中地址 address 处的旧值。
__nv_atomic_store 读取 val 指向的值并存储到 address 指向的位置。__nv_atomic_store_n 将 val 存储到 address 指向的位置。

order 不能是 __NV_ATOMIC_CONSUME、__NV_ATOMIC_ACQUIRE 或 __NV_ATOMIC_ACQ_REL。

5.4.5.2.12 `__nv_atomic_thread_fence()`

__device__ void __nv_atomic_thread_fence(int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

该原子函数根据指定的内存序，在该线程请求的内存访问之间建立序。线程作用域参数指定了可能观察到该操作序效果的线程集合。

5.4.6 线程束函数

本节介绍允许线程束（warp）内的线程相互通信并执行计算的线程束函数。

提示

出于效率、安全性与可移植性的考虑，建议尽可能使用 CUB 线程束级"集合"原语（CUB Warp-Wide "Collective" Primitives）来执行线程束操作。

5.4.6.1 线程束活动掩码

unsigned __activemask();

该函数返回一个 32 位整数掩码，表示调用线程束中所有当前活动的线程。当调用 __activemask() 时，若线程束中第 N 个通道（lane）处于活动状态，则第 N 位置位。返回的掩码中用 0 位表示非活动线程。已退出程序的线程始终标记为非活动。

警告

__activemask() 不能用于确定哪些线程束通道执行了给定的分支。该函数面向机会主义的线程束级编程，仅提供线程束内活动线程的瞬时快照。

c // 检查是否至少有一个线程的谓词求值为真 if (pred) { // 无效：'at_least_one' 的值是非确定性的， // 在不同执行之间可能发生变化。 at_least_one = __activemask() > 0; }

注意，在 __activemask() 调用处汇聚的线程，不保证在后续指令处仍保持汇聚，除非这些指令是线程束同步内联函数（__sync）。例如，编译器可能会重排指令，活动线程的集合可能无法保留：

unsigned mask       = __activemask();              // 假设 mask == 0xFFFFFFFF（所有位都置位，所有线程都活动）

int       predicate = threadIdx.x % 2 == 0;        // 偶数线程为 1，奇数线程为 0

int       result    = __any_sync(mask, predicate); // 活动线程可能无法保留

5.4.6.2 线程束表决函数

int      __all_sync   (unsigned mask, int predicate);
int      __any_sync   (unsigned mask, int predicate);
unsigned __ballot_sync(unsigned mask, int predicate);

线程束表决（vote）函数使给定线程束的线程能够执行归约-广播（reduction-and-broadcast）操作。这些函数从线程束中每个未退出的线程接收一个整数谓词作为输入，并将其与零比较。然后比较结果按以下方式之一在线程束的活动线程之间组合（归约），并向每个参与线程广播单个返回值：

__all_sync(unsigned mask, predicate)：对 mask 中所有未退出的线程求值 predicate；若对全部线程 predicate 求值都为非零，则返回非零值。
__any_sync(unsigned mask, predicate)：对 mask 中所有未退出的线程求值 predicate；若对其中一个或多个线程 predicate 求值都为非零，则返回非零值。
__ballot_sync(unsigned mask, predicate)：对 mask 中所有未退出的线程求值 predicate，返回一个整数：若对线程束第 N 个线程 predicate 求值非零且第 N 个线程处于活动状态，则该整数的第 N 位置位；否则第 N 位为零。

这些函数受限于线程束 __sync 内联函数约束。

警告

这些内联函数不提供任何内存序。

5.4.6.3 线程束匹配函数

提示

建议使用 libcu++ 的 cuda::device::warp_match_all() 函数，作为 __match_all_sync 函数的通用且更安全的替代方案。

unsigned __match_any_sync(unsigned mask, T value);
unsigned __match_all_sync(unsigned mask, T value, int *pred);

线程束匹配（match）函数在线程束内的未退出线程之间对变量执行广播-比较（broadcast-and-compare）操作。

__match_any_sync：返回 mask 中具有相同按位值（bitwise value）的未退出线程的掩码。
__match_all_sync：若 mask 中所有未退出线程具有相同的按位值，则返回 mask；否则返回 0。若 mask 中所有未退出线程具有相同的按位值，则谓词 pred 置为 true；否则置为 false。

T 可以是 int、unsigned、long、unsigned long、long long、unsigned long long、float 或 double。

这些函数受限于线程束 __sync 内联函数约束。

警告

这些内联函数不提供任何内存序。

5.4.6.4 线程束归约函数

提示

出于效率、安全性与可移植性的考虑，建议尽可能使用 CUB 线程束级"集合"原语来执行线程束归约。

需要计算能力 8.x 或更高的设备支持。

T             __reduce_add_sync(unsigned mask, T value);
T             __reduce_min_sync(unsigned mask, T value);
T             __reduce_max_sync(unsigned mask, T value);

unsigned __reduce_and_sync(unsigned mask, unsigned value);
unsigned __reduce_or_sync (unsigned mask, unsigned value);
unsigned __reduce_xor_sync(unsigned mask, unsigned value);

__reduce_<op>_sync 内联函数先同步 mask 中指定的所有未退出线程，然后对 value 中提供的数据执行归约操作。

__reduce_add_sync、__reduce_min_sync、__reduce_max_sync：返回对 mask 中指定的每个未退出线程在 value 中提供的值分别执行算术加、min、max 归约操作的结果。T 可以是无符号或有符号整数。
__reduce_and_sync、__reduce_or_sync、__reduce_xor_sync：返回对 mask 中指定的每个未退出线程在 value 中提供的值执行按位 AND、OR、XOR 归约操作的结果。

这些函数受限于线程束 __sync 内联函数约束。

警告

这些内联函数不提供任何内存序。

5.4.6.5 线程束洗牌（Shuffle）函数

提示

建议使用 libcu++ 的 cuda::device::warp_shuffle() 函数，作为 __shfl_sync() 与 __shfl_<op>_sync() 内联函数的通用且更安全的替代方案。

T __shfl_sync     (unsigned mask, T value, int      srcLane,  int width=warpSize);
T __shfl_up_sync  (unsigned mask, T value, unsigned delta,    int width=warpSize);
T __shfl_down_sync(unsigned mask, T value, unsigned delta,    int width=warpSize);
T __shfl_xor_sync (unsigned mask, T value, int      laneMask, int width=warpSize);

线程束洗牌（shuffle）函数在线程束内的未退出线程之间交换值，而无需使用共享内存。

__shfl_sync()：从指定通道直接复制。 该内联函数返回 ID 由 srcLane 给定的线程所持有的 value 的值。

如果 width 小于 warpSize，则线程束的每个子部分（subsection）表现为一个独立实体，起始逻辑通道 ID 为 0。
如果 srcLane 超出 [0, width − 1] 范围，则结果对应于 srcLane % width 所表示的线程（位于同一子部分内）所持有的值。

__shfl_up_sync()：从 ID 比调用者小的通道复制。 该内联函数用调用者的通道 ID 减去 delta 计算出源通道 ID，返回该通道 ID 对应的线程所持有的 value 的值：其效果是将 value 沿线程束向上移动 delta 个通道。

如果 width 小于 warpSize，则线程束的每个子部分表现为一个独立实体，起始逻辑通道 ID 为 0。
源通道索引不会绕过（wrap around）width 的值，因此较低的 delta 个通道将保持不变。

__shfl_down_sync()：从 ID 比调用者大的通道复制。 该内联函数用调用者的通道 ID 加上 delta 计算出源通道 ID，返回该通道 ID 对应的线程所持有的 value 的值：其效果是将 value 沿线程束向下移动 delta 个通道。

如果 width 小于 warpSize，则线程束的每个子部分表现为一个独立实体，起始逻辑通道 ID 为 0。
与 __shfl_up_sync() 一样，源通道的 ID 编号不会绕过 width 的值，因此较高的 delta 个通道实际上保持不变。

__shfl_xor_sync()：根据自身通道 ID 的按位异或从某通道复制。 该内联函数将调用者的通道 ID 与 laneMask 执行按位异或，计算出源通道 ID，返回该通道 ID 对应的线程所持有的 value 的值。这种模式实现了蝴蝶形（butterfly）寻址模式，用于树形归约和广播。

如果 width 小于 warpSize，则每组 width 个连续的线程可以访问前面组的元素。但是，如果它们试图访问后面组的线程中的元素，则会返回它们自己的 value 值。

T 可以是：

int、unsigned、long、unsigned long、long long、unsigned long long、float 或 double。
__half 和 __half2，需包含 cuda_fp16.h 头文件。
__nv_bfloat16 和 __nv_bfloat162，需包含 cuda_bf16.h 头文件。

线程只能从正在积极参与该内联函数的另一个线程读取数据。如果目标线程处于非活动状态，则获取的值是未定义的。

width 必须是 [1, warpSize] 范围内 2 的幂，即 1、2、4、8、16 或 32。其他值将产生未定义的结果。

这些函数受限于线程束 __sync 内联函数约束。

线程束洗牌正确用法的示例：

int laneId = threadIdx.x % warpSize;
int data   = ...

// 线程束的所有线程都从通道 0 获取 'data'
int result1 = __shfl_sync(0xFFFFFFFF, data, 0);

if (laneId < 4) {
    // 通道 0、1、2、3 从通道 1 获取 'data'
    int result2 = __shfl_sync(0xb1111, data, 1);
}

// 通道 [0 - 15] 从通道 0 获取 'data'
// 通道 [16 - 31] 从通道 16 获取 'data'
int result3 = __shfl_sync(0xFFFFFFFF, value, warpSize / 2);

// 每个通道从其上方两个位置的通道获取 'data'
// 通道 30、31 获取其原始值
int result4 = __shfl_down_sync(0xFFFFFFFF, data, 2);

线程束洗牌错误用法的示例：

int laneId = threadIdx.x % warpSize;
int value = ...
// 未定义行为：通道 0 不参与该调用
int result = (laneId > 0) ? __shfl_sync(0xFFFFFFFF, value, 0) : 0;

if (laneId <= 4) {
    // 未定义行为：对于通道 3、4，目标通道 5、6 处于非活动状态
    result = __shfl_down_sync(0b11111, value, 2);
}

// 未定义行为：width 不是 2 的幂
__shfl_sync(0xFFFFFFFF, value, 0, /*width=*/31);

警告

这些内联函数不蕴含内存屏障（memory barrier），不保证任何内存序。

示例 1：在线程束内广播单个值

CUDA C++

#include <cassert>
#include <cuda/warp>

__global__ void warp_broadcast_kernel(int input) {
    int laneId = threadIdx.x % 32;
    int value;
    if (laneId == 0) { // 除通道 0 外的所有线程该变量均未使用
        value = input;
    }
    value = cuda::device::warp_shuffle_idx(value, 0); // 同步线程束中的所有线程，并从通道 0 获取 "value"

    assert(value == input);
}

int main() {
    warp_broadcast_kernel<<<1, 32>>>(1234);
    cudaDeviceSynchronize();
    return 0;
}

内联函数

#include <assert.h>

__global__ void warp_broadcast_kernel(int input) {
     int laneId = threadIdx.x % 32;
     int value;
     if (laneId == 0) { // 除通道 0 外的所有线程该变量均未使用
         value = input;
     }
     value = __shfl_sync(0xFFFFFFFF, value, 0); // 同步线程束中的所有线程，并从通道 0 获取 "value"

     assert(value == input);
}

int main() {
    warp_broadcast_kernel<<<1, 32>>>(1234);
    cudaDeviceSynchronize();
    return 0;
}

可在 Compiler Explorer 上查看该示例。

示例 2：在 8 线程子分区上执行包含式加法扫描（inclusive plus-scan）

提示

建议使用 cub::WarpScan 函数来实现高效且通用的线程束扫描函数。

CUDA C++

#include <cstdio>
#include <cub/cub.cuh>

__global__ void scan_sub_partition_with_8_threads_kernel() {
    using WarpScan    = cub::WarpScan<int, 8>;
    using TempStorage = typename WarpScan::TempStorage;
    __shared__ TempStorage temp_storage;

    int laneId = threadIdx.x % 32;
    int value = 31 - laneId; // 待累加的起始值
    int partial_sum;
    WarpScan(temp_storage).InclusiveSum(value, partial_sum);
    printf("Thread %d final value = %d\n", threadIdx.x, partial_sum);
}

int main() {
    scan_sub_partition_with_8_threads_kernel<<<1, 32>>>();
    cudaDeviceSynchronize();
    return 0;
}

内联函数

#include <stdio.h>

__global__ void scan_sub_partition_with_8_threads_kernel() {
     int laneId = threadIdx.x % 32;
     int value = 31 - laneId; // 待累加的起始值
     // 循环累加我所在分区内的扫描。
     // 8 个线程的扫描需要 log2(8) == 3 步
     for (int delta = 1; delta <= 4; delta *= 2) {
         int tmp         = __shfl_up_sync(0xFFFFFFFF, value, delta, /*width=*/8); // 从 laneId - delta 处读取

         int source_lane = laneId % 8 - delta;
         if (source_lane >= 0) // 'source_lane < 0' 的通道值保持不变
             value += tmp;
     }
     printf("Thread %d final value = %d\n", threadIdx.x, value);
}

int main() {
    scan_sub_partition_with_8_threads_kernel<<<1, 32>>>();
    cudaDeviceSynchronize();
    return 0;
}

可在 Compiler Explorer 上查看该示例。

示例 3：跨线程束归约

提示

建议使用 cub::WarpReduce 函数来实现高效且通用的线程束归约函数。

CUDA C++

#include <cstdio>
#include <cub/cub.cuh>
#include <cuda/warp>

__global__ void warp_reduce_kernel() {
    using WarpReduce = cub::WarpReduce<int>;
    using TempStorage = typename WarpReduce::TempStorage;
    __shared__ TempStorage temp_storage;

    int laneId     = threadIdx.x % 32;
    int value      = 31 - laneId; // 待累加的起始值
    auto aggregate = WarpReduce(temp_storage).Sum(value);
    aggregate      = cuda::device::warp_shuffle_idx(aggregate, 0);
    printf("Thread %d final value = %d\n", threadIdx.x, aggregate);
}

int main() {
    warp_reduce_kernel<<<1, 32>>>();
    cudaDeviceSynchronize();
    return 0;
}

内联函数

#include <stdio.h>

__global__ void warp_reduce_kernel() {
    int laneId = threadIdx.x % 32;
    int value = 31 - laneId; // 待累加的起始值
    // 使用 XOR 模式执行蝴蝶形归约
    // 全线程束归约需要 log2(32) == 5 步
    for (int i = 1; i <= 16; i *= 2)
        value += __shfl_xor_sync(0xFFFFFFFF, value, i);
    // "value" 此时包含所有线程之和
    printf("Thread %d final value = %d\n", threadIdx.x, value);
}

int main() {
    warp_reduce_kernel<<<1, 32>>>();
    cudaDeviceSynchronize();
    return 0;
}

可在 Compiler Explorer 上查看该示例。

5.4.6.6 线程束 __sync 内联函数约束

所有线程束 __sync 内联函数，例如：

__shfl_sync、__shfl_up_sync、__shfl_down_sync、__shfl_xor_sync
__match_any_sync、__match_all_sync
__reduce_add_sync、__reduce_min_sync、__reduce_max_sync、__reduce_and_sync、__reduce_or_sync、__reduce_xor_sync
__syncwarp

都使用 mask 参数指明哪些线程束线程参与该调用。该参数确保硬件执行内联函数前实现正确的汇聚。

mask 中的每一位对应一个线程的通道 ID（threadIdx.x % warpSize）。该内联函数会等待 mask 中指定的所有未退出线程束线程都到达该调用处。

正确执行必须满足以下约束：

每个调用线程在 mask 中都必须置位其对应的位。
每个非调用线程在 mask 中都必须将其对应的位清零。已退出的线程会被忽略。
mask 中指定的所有未退出线程都必须使用相同的 mask 值执行该内联函数。
线程束线程可以用不同的 mask 值并发调用该内联函数，前提是这些掩码互不相交（disjoint）。即使在分支分化（divergent）控制流中，这种情况也是合法的。

在以下情况下，线程束 __sync 函数的行为是无效的（例如核函数挂起）或未定义的：

调用线程在 mask 中未被指定。
mask 中指定的非退出线程未能最终退出，或未能在同一程序点以相同的 mask 值调用该内联函数。
在条件代码中，mask 中指定的所有未退出线程的所有条件都必须求值相同。

注

当所有线程束线程都参与该调用（即 mask 设置为 0xFFFFFFFF）时，内联函数的效率最高。

线程束内联函数正确用法的示例：

__global__ void valid_examples() {
     if (threadIdx.x < 4) {        // 线程 0、1、2、3 处于活动状态
         __all_sync(0b1111, pred); // 正确，线程 0、1、2、3 参与该调用
     }

     if (threadIdx.x == 0)
         return; // 退出
     // 正确，所有未退出的线程都参与该调用
     __all_sync(0xFFFFFFFF, pred);
}

互不相交掩码（disjoint mask）的示例：

__global__ void example_syncwarp_with_mask(int* input_data, int* output_data) {
     if (threadIdx.x < warpSize) {
          __shared__ int shared_data[warpSize];
          shared_data[threadIdx.x] = input_data[threadIdx.x];

          unsigned mask = threadIdx.x < 16 ? 0xFFFF : 0xFFFF0000; // 正确
          __syncwarp(mask);
          if (threadIdx.x == 0 || threadIdx.x == 16)
              output_data[threadIdx.x] = shared_data[threadIdx.x + 1];
     }
}

__global__ void example_syncwarp_with_mask_branches(int* input_data, int* output_data) {
     if (threadIdx.x < warpSize) {
         __shared__ int shared_data[warpSize];
         shared_data[threadIdx.x] = input_data[threadIdx.x];

          if (threadIdx.x < 16) {
              unsigned mask = 0xFFFF; // 正确
              __syncwarp(mask);
              output_data[threadIdx.x] = shared_data[15 - threadIdx.x];
          }
          else {
              unsigned mask = 0xFFFF0000; // 正确
              __syncwarp(mask);
              output_data[threadIdx.x] = shared_data[31 - threadIdx.x];
          }
     }
}

线程束内联函数错误用法的示例：

if (threadIdx.x < 4) {            // 线程 0、1、2、3 处于活动状态
     __all_sync(0b0000011, pred); // 错误，线程 2、3 处于活动状态但在 mask 中未置位

     __all_sync(0b1111111, pred); // 错误，线程 4、5、6 处于非活动状态但在 mask 中置位
}

// 错误，参与线程的掩码不同且重叠
__all_sync(threadIdx.x == 0 ? 1 : 0xFFFFFFFF, pred);

5.4.7 CUDA 特定的宏

5.4.7.1 CUDA_ARCH

宏 __CUDA_ARCH__ 表示正在编译的代码所针对的 NVIDIA GPU 的虚拟架构。其值可能与设备的实际计算能力不同。该宏使编写针对特定 GPU 架构特化的代码路径成为可能，这对于获得最优性能或使用架构特有的功能和指令可能是必要的。该宏还可用于区分主机端代码和设备端代码。

__CUDA_ARCH__ 仅在设备代码中定义，即在 __device__、__host__ __device__ 和 __global__ 函数中。该宏的值与 nvcc 选项 compute_ 相关，关系为 __CUDA_ARCH__ = <version> * 10。

示例：

nvcc --generate-code arch=compute_80,code=sm_90 prog.cu

将 __CUDA_ARCH__ 定义为 800。

CUDA_ARCH 约束

以下实体的类型签名不得依赖于 __CUDA_ARCH__ 是否定义，也不依赖于其值。
__global__ 函数和函数模板。
__device__ 和 __constant__ 变量。
纹理和表面（surface）。

示例：

```c #if !defined(CUDA_ARCH) typedef int my_type; #else typedef double my_type; #endif

device my_type my_var; // 错误：my_var 的类型依赖于 CUDA_ARCH

global void kernel(my_type in) { // 错误：kernel 的类型依赖于 CUDA_ARCH ... } ```

__global__ 函数模板的实例化不得依赖于 __CUDA_ARCH__ 是否定义或其值。也就是说，在所有设备程序和主机程序中，必须存在相同的实例化（具有等效的模板参数），而不论这些实例化是否在运行时被启动。

示例：

```c device int result;

template global void kernel(T in) { result = in; }

host device void host_device_function(void) { #if !defined(CUDA_ARCH) kernel<<<1, 1>>>(1); // 错误：仅在 CUDA_ARCH 未定义时才实例化 "kernel"！ #endif }

int main(void) { host_device_function(); cudaDeviceSynchronize(); return 0; } ```

要避免此问题，可以将实际计算从 __global__ 函数移入被其调用的 __device__ 函数模板中。在 __global__ 函数内，再使用 __CUDA_ARCH__ 条件性地用不同的参数实例化该 __device__ 函数模板。

在分离编译模式下，具有外部链接的函数或变量定义的存在与否，不得依赖于 __CUDA_ARCH__ 的定义或其值。

示例：

c #if !defined(__CUDA_ARCH__) void host_function(void) {} // 错误：host_function() 的定义仅在 __CUDA_ARCH__ // 未定义时存在 #endif

在分离编译中，为防止对象具有不同的行为，不得在头文件中使用预处理器宏 __CUDA_ARCH__。或者，所有对象都必须为相同的虚拟架构编译。如果一个弱（weak）函数或模板函数定义在头文件中，且其行为依赖于 __CUDA_ARCH__，那么当这些对象为不同的计算架构编译时，不同对象中该函数的实例可能发生冲突。

例如，如果头文件 a.h 包含：

c template<typename T> __device__ T* get_ptr() { #if __CUDA_ARCH__ == 900 return nullptr; /* 无地址 */ #else __shared__ T arr[256]; return arr; #endif }

那么，如果 a.cu 和 b.cu 都包含 a.h 并为相同类型实例化 get_ptr()，而 b.cu 期望一个非 NULL 地址，并使用以下命令编译：

bash nvcc -arch=compute_70 -dc a.cu nvcc -arch=compute_80 -dc b.cu nvcc -arch=sm_80 a.o b.o

在链接时只使用一个版本的 get_ptr() 函数，因此行为取决于选择了哪个版本。要避免此问题，要么将 a.cu 和 b.cu 为相同的计算架构编译，要么不在共享的头函数中使用 __CUDA_ARCH__。

编译器不保证会对上述 __CUDA_ARCH__ 的不支持用法生成诊断信息。

5.4.7.2 CUDA_ARCH_SPECIFIC 与 CUDA_ARCH_FAMILY_SPECIFIC

宏 __CUDA_ARCH_SPECIFIC__ 和 __CUDA_ARCH_FAMILY_SPECIFIC__ 分别用于标识具有架构特定和架构族特定功能的 GPU 设备。更多信息请参阅特性集编译器目标（Feature Set Compiler Targets）一节。

与 __CUDA_ARCH__ 类似，__CUDA_ARCH_SPECIFIC__ 和 __CUDA_ARCH_FAMILY_SPECIFIC__ 也仅在设备代码中定义，即在 __device__、__host__ __device__ 和 __global__ 函数中。这两个宏与 nvcc 选项 compute_a 和 compute_f 相关。

nvcc --generate-code arch=compute_100a,code=sm_100a prog.cu

__CUDA_ARCH__ == 1000。
__CUDA_ARCH_SPECIFIC__ == 1000。
__CUDA_ARCH_FAMILY_SPECIFIC__ == 1000。

nvcc --generate-code arch=compute_100f,code=sm_103f prog.cu

__CUDA_ARCH__ == 1000。
__CUDA_ARCH_FAMILY_SPECIFIC__ == 1000。
未定义 __CUDA_ARCH_SPECIFIC__。

nvcc -arch=sm_100 prog.cu

__CUDA_ARCH__ == 1000。
未定义 __CUDA_ARCH_FAMILY_SPECIFIC__。
未定义 __CUDA_ARCH_SPECIFIC__。

nvcc -arch=sm_100a prog.cu
# 等效于：
nvcc --generate-code arch=sm_100a,compute_100,compute_100a prog.cu

__CUDA_ARCH__ == 1000。
未定义 __CUDA_ARCH_FAMILY_SPECIFIC__。
__CUDA_ARCH_SPECIFIC__ == 1000 的版本和 __CUDA_ARCH_SPECIFIC__ 未定义的版本都会生成。

5.4.7.3 CUDA 特性测试宏

nvcc 提供以下预处理器宏用于特性测试。当 CUDA 前端编译器支持某一特性时，会定义相应的宏。

__CUDACC_DEVICE_ATOMIC_BUILTINS__：支持设备原子编译器内置函数（builtins）。
__NVCC_DIAG_PRAGMA_SUPPORT__：支持诊断控制 pragma。
__CUDACC_EXTENDED_LAMBDA__：支持扩展 lambda。由 --expt-extended-lambda 或 --extended-lambda 标志启用。
__CUDACC_RELAXED_CONSTEXPR__：支持放宽的 constexpr 函数。由 --expt-relaxed-constexpr 标志启用。

5.4.7.4 nv_pure 属性

在 C/C++ 中，纯函数（pure function）对其参数没有副作用，可以访问全局变量但不会修改它们。

CUDA 提供 __nv_pure__ 属性，主机和设备函数都支持。编译器将 __nv_pure__ 转换为 GNU 的 pure 属性或 Microsoft Visual Studio 的 noalias 属性。

__device__ __nv_pure__
int add(int a, int b) {
    return a + b;
}

5.4.8 CUDA 特定的函数

5.4.8.1 地址空间谓词函数

地址空间谓词函数用于确定指针的地址空间。

提示

建议使用 libcu++ 提供的 cuda::device::is_address_from() 和 cuda::device::is_object_from() 函数，作为地址空间谓词内联函数的可移植且更安全的替代方案。

__device__ unsigned __isGlobal      (const void* ptr);
__device__ unsigned __isShared      (const void* ptr);
__device__ unsigned __isConstant    (const void* ptr);
__device__ unsigned __isGridConstant(const void* ptr);
__device__ unsigned __isLocal       (const void* ptr);

如果 ptr 包含指定地址空间中对象的通用地址，则函数返回 1，否则返回 0。若参数是 NULL 指针，其行为未指定。

__isGlobal()：全局内存空间。
__isShared()：共享内存空间。
__isConstant()：常量内存空间。
__isGridConstant()：标注了 __grid_constant__ 的核函数参数。
__isLocal()：本地内存空间。

5.4.8.2 地址空间转换函数

CUDA 指针（T）可以访问对象而不论对象存储在哪里。例如，一个 int 既可以访问全局内存中的 int 对象，也可以访问共享内存中的 int 对象。

地址空间转换函数用于在通用地址与特定地址空间中的地址之间转换。当编译器无法确定指针的地址空间时（例如跨翻译单元或与 PTX 指令交互时），这些函数很有用。

__device__ size_t __cvta_generic_to_global (const void* ptr); // PTX: cvta.to.global
__device__ size_t __cvta_generic_to_shared (const void* ptr); // PTX: cvta.to.shared
__device__ size_t __cvta_generic_to_constant(const void* ptr); // PTX: cvta.to.const
__device__ size_t __cvta_generic_to_local   (const void* ptr); // PTX: cvta.to.local

__device__ void* __cvta_global_to_generic (size_t raw_ptr); // PTX: cvta.global
__device__ void* __cvta_shared_to_generic (size_t raw_ptr); // PTX: cvta.shared
__device__ void* __cvta_constant_to_generic(size_t raw_ptr); // PTX: cvta.const
__device__ void* __cvta_local_to_generic   (size_t raw_ptr); // PTX: cvta.local

作为与 PTX 指令互操作的示例，PTX 指令 ld.shared.s32 r0, [ptr]; 要求 ptr 引用共享内存地址空间。具有指向 __shared__ 内存中对象的 int* 指针的 CUDA 程序，在将指针传给 PTX 指令之前，需要调用 __cvta_generic_to_shared 将该指针转换为共享地址空间，如下所示：

__shared__ int smem_var;
smem_var        = 42;
size_t smem_ptr = __cvta_generic_to_shared(&smem_var);
int     output;
asm volatile("ld.shared.s32 %0, [%1];" : "=r"(output) : "l"(smem_ptr) : "memory");

assert(output == 42);

利用这些地址表示形式的一种常见优化是：利用共享、本地和常量空间的地址范围小于 32 位这一事实来减小数据结构的大小，从而可以存储 32 位地址而非 64 位指针，并节省寄存器。此外，32 位算术比 64 位算术更快。要获得这些地址的 32 位整数表示，将 64 位值从无符号 64 位整数强制转换为无符号 32 位整数进行截断即可：

__shared__ int smem_var;
uint32_t        smem_ptr_32bit = static_cast<uint32_t>(__cvta_generic_to_shared(&smem_var));

要从这样的 32 位表示恢复通用地址，将该地址零扩展回无符号 64 位整数，然后调用相应的地址空间转换函数：

size_t smem_ptr_64bit = static_cast<size_t>(smem_ptr_32bit); // 零扩展到 64 位

void* generic_ptr     = __cvta_shared_to_generic(smem_ptr_64bit);
assert(generic_ptr == &smem_var);

5.4.8.3 低级加载和存储函数

T __ldg(const T* address);

函数 __ldg() 执行只读的 L1/Tex 缓存加载。它支持所有 C++ 基本类型、CUDA 向量类型（x3 分量除外）以及扩展浮点类型，如 __half、__half2、__nv_bfloat16 和 __nv_bfloat162。

T __ldcg(const T* address);
T __ldca(const T* address);
T __ldcs(const T* address);
T __ldlu(const T* address);
T __ldcv(const T* address);

这些函数使用 PTX ISA 指南中指定的缓存操作符执行加载。它们支持所有 C++ 基本类型、CUDA 向量类型（x3 分量除外）以及扩展浮点类型，如 __half、__half2、__nv_bfloat16 和 __nv_bfloat162。

void __stwb(T* address, T value);
void __stcg(T* address, T value);
void __stcs(T* address, T value);
void __stwt(T* address, T value);

这些函数使用 PTX ISA 指南中指定的缓存操作符执行存储。它们支持所有 C++ 基本类型、CUDA 向量类型（x3 分量除外）以及扩展浮点类型，如 __half、__half2、__nv_bfloat16 和 __nv_bfloat162。

5.4.8.4 __trap()

提示

建议使用 libcu++ 提供的 cuda::std::terminate() 函数（C++ 参考）作为 __trap() 的可移植替代方案。

任何设备线程都可以调用 __trap() 函数发起陷阱（trap）操作。

void __trap();

核函数的执行被中止，主机程序中会产生中断。调用 __trap() 会导致 CUDA 上下文损坏，从而使后续的 CUDA 调用和核函数启动失败。

5.4.8.5 __nanosleep()

__device__ void __nanosleep(unsigned nanoseconds);

函数 __nanosleep(ns) 挂起线程约 ns 纳秒的睡眠时长。最大睡眠时长约为 1 毫秒。

示例：下面的代码实现了一个带指数退避（exponential back-off）的互斥锁。

__device__ void mutex_lock(unsigned* mutex) {
    unsigned ns = 8;
    while (atomicCAS(mutex, 0, 1) == 1) {
        __nanosleep(ns);
        if (ns < 256) {
            ns *= 2;
        }
    }
}

__device__ void mutex_unlock(unsigned *mutex) {
    atomicExch(mutex, 0);
}

5.4.8.6 动态规划扩展（DPX）指令

DPX 函数集支持求最多三个 16 位或 32 位有符号或无符号整数参数的最小值和最大值，以及融合的加法与最小/最大值。它具有可选的 ReLU（即钳制到零）功能。

比较函数：

三个参数。语义：max(a, b, c)，min(a, b, c)。

```c int __vimax3_s32 ( int, int, int); unsigned __vimax3_s16x2(unsigned, unsigned, unsigned); unsigned __vimax3_u32 (unsigned, unsigned, unsigned); unsigned __vimax3_u16x2(unsigned, unsigned, unsigned);

   int __vimin3_s32 (      int,      int,      int);

unsigned __vimin3_s16x2(unsigned, unsigned, unsigned); unsigned __vimin3_u32 (unsigned, unsigned, unsigned); unsigned __vimin3_u16x2(unsigned, unsigned, unsigned); ```

两个参数，带 ReLU。语义：max(a, b, 0)，max(min(a, b), 0)。

```c int __vimax_s32_relu ( int, int); unsigned __vimax_s16x2_relu(unsigned, unsigned);

   int __vimin_s32_relu (      int,      int);

unsigned __vimin_s16x2_relu(unsigned, unsigned); ```

三个参数，带 ReLU。语义：max(a, b, c, 0)，max(min(a, b, c), 0)。

```c int __vimax3_s32_relu ( int, int, int); unsigned __vimax3_s16x2_relu(unsigned, unsigned, unsigned);

   int __vimin3_s32_relu (      int,      int,      int);

unsigned __vimin3_s16x2_relu(unsigned, unsigned, unsigned); ```

两个参数，同时返回哪个参数更小/更大：

```c int __vibmax_s32 ( int, int, bool pred); unsigned __vibmax_u32 (unsigned, unsigned, bool pred); unsigned __vibmax_s16x2(unsigned, unsigned, bool pred); unsigned __vibmax_u16x2(unsigned, unsigned, bool pred);

   int __vibmin_s32 (      int,      int, bool* pred);

unsigned __vibmin_u32 (unsigned, unsigned, bool pred); unsigned __vibmin_s16x2(unsigned, unsigned, bool pred); unsigned __vibmin_u16x2(unsigned, unsigned, bool* pred); ```

融合的加法与最小/最大值：

三个参数，比较（第一 + 第二）与第三。语义：max(a + b, c)，min(a + b, c)。

```c int __viaddmax_s32 ( int, int, int); unsigned __viaddmax_s16x2(unsigned, unsigned, unsigned); unsigned __viaddmax_u32 (unsigned, unsigned, unsigned); unsigned __viaddmax_u16x2(unsigned, unsigned, unsigned);

   int __viaddmin_s32 (      int,     int,       int);

unsigned __viaddmin_s16x2(unsigned, unsigned, unsigned); unsigned __viaddmin_u32 (unsigned, unsigned, unsigned); unsigned __viaddmin_u16x2(unsigned, unsigned, unsigned); ```

三个参数，带 ReLU，比较（第一 + 第二）与第三和零。语义：max(a + b, c, 0)，max(min(a + b, c), 0)。

```c int __viaddmax_s32_relu ( int, int, int); unsigned __viaddmax_s16x2_relu(unsigned, unsigned, unsigned);

   int __viaddmin_s32_relu (      int,      int,      int);

unsigned __viaddmin_s16x2_relu(unsigned, unsigned, unsigned); ```

这些指令根据计算能力的不同由硬件加速或软件模拟。请参阅算术指令一节了解计算能力要求。完整的 API 可在 CUDA Math API 文档中找到。

DPX 是实现动态规划算法（如基因组学中的 Smith-Waterman 和 Needleman-Wunsch 算法，以及路线优化中的 Floyd-Warshall 算法）的极其有用的工具。

三个有符号 32 位整数的最大值，带 ReLU：

int a           = -15;
int b           = 8;
int c           = 5;
int max_value_0 = __vimax3_s32_relu(a, b, c); // max(-15, 8, 5, 0) = 8
int d           = -2;
int e           = -4;
int max_value_1 = __vimax3_s32_relu(a, d, e); // max(-15, -2, -4, 0) = 0

两个有符号 32 位整数之和与另一个有符号 32 位整数及零（ReLU）的最小值：

int a           = -5;
int b           = 6;
int c           = -2;
int max_value_0 = __viaddmax_s32_relu(a, b, c); // max(-5 + 6, -2, 0) = max(1, -2, 0) = 1

int d           = 4;
int max_value_1 = __viaddmax_s32_relu(a, d, c); // max(-5 + 4, -2, 0) = max(-1, -2, 0) = 0

两个无符号 32 位整数的最小值，并判断哪个值更小：

unsigned a = 9;
unsigned b = 6;
bool      smaller_value;
unsigned min_value = __vibmin_u32(a, b, &smaller_value); // min_value 为 6，smaller_value 为 true

三对无符号 16 位整数的最大值：

unsigned a         = 0x00050002;
unsigned b         = 0x00070004;
unsigned c         = 0x00020006;
unsigned max_value = __vimax3_u16x2(a, b, c); // max(5, 7, 2) 和 max(2, 4, 6)，因此 max_value 为 0x00070006

5.4.9 编译器优化提示

编译器优化提示用附加信息修饰代码，以帮助编译器优化生成的代码。

这些内置函数在 SIMT 设备代码中始终可用，但在 Tile 代码中不受支持。
主机代码支持取决于主机编译器。

5.4.9.1 #pragma unroll

默认情况下，编译器会展开行程数（trip count）已知的小循环。但是，可以使用 #pragma unroll 指令控制任意给定循环的展开。该指令必须紧挨着放在循环之前，且只适用于该循环。

指令后面可以可选地跟一个整数常量表达式。关于整数常量表达式有以下情况：

如果没有指定，当循环的行程数为常量时，循环将被完全展开。
如果其求值结果为 0 或 1，循环将不会展开。
如果是负整数或大于 INT_MAX，该 pragma 将被忽略，并发出警告。

示例：

struct MyStruct {
    static constexpr int value = 4;
};

inline constexpr int Count = 4;

__device__ void foo(int* p1, int* p2) {
    // 未指定参数，循环将被完全展开
    #pragma unroll
    for (int i = 0; i < 12; ++i)
        p1[i] += p2[i] * 2;

       // 展开值 = 5
       #pragma unroll (Count + 1)
       for (int i = 0; i < 12; ++i)
           p1[i] += p2[i] * 4;

       // 展开值 = 1，禁用循环展开
       #pragma unroll 1
       for (int i = 0; i < 12; ++i)
           p1[i] += p2[i] * 8;

       // 展开值 = 4
       #pragma unroll (MyStruct::value)
       for (int i = 0; i < 12; ++i)
           p1[i] += p2[i] * 16;

       // 负值，pragma unroll 被忽略
       #pragma unroll -1
       for (int i = 0; i < 12; ++i)
           p1[i] += p2[i] * 2;
}

可在 Compiler Explorer 上查看该示例。

5.4.9.2 __builtin_assume_aligned()

提示

建议使用 libcu++ 提供的 cuda::std::assume_aligned() 函数（C++ 参考）作为这些内置函数的可移植且更安全的替代方案。

void* __builtin_assume_aligned(const void* ptr, size_t align)
void* __builtin_assume_aligned(const void* ptr, size_t align, <integral type> offset)

这些内置函数使编译器能够假设返回的指针至少按 align 字节对齐。

三参数版本使编译器能够假设 (char*) ptr - offset 至少按 align 字节对齐。

align 必须是 2 的幂且为整数常量。

示例：

void* res1 = __builtin_assume_aligned(ptr, 32);      // 编译器可以假设 'res1' 至少按 32 字节对齐

void* res2 = __builtin_assume_aligned(ptr, 32, 8); // 编译器可以假设 'res2 = (char*) ptr - 8' 至少按 32 字节对齐

5.4.9.3 __builtin_assume() 与 __assume()

void __builtin_assume(bool predicate)
void __assume        (bool predicate) // 仅与 Microsoft 编译器搭配使用

该内置函数使编译器能够假设布尔参数为 true。如果参数在运行时为 false，则行为未定义。注意，如果参数有副作用，则行为未指定。

示例：

__device__ bool is_greater_than_zero(int value) {
    return value > 0;
}

__device__ bool f(int value) {
     __builtin_assume(value > 0);
     return is_greater_than_zero(value); // 返回 true，不求值条件
}

5.4.9.4 __builtin_constant_p()

int __builtin_constant_p(expression)

如果编译器能确定 expression 是编译时常量，该内置函数返回 1，否则返回 0。该表达式不会被求值，因此 expression 中的副作用不会发生。

示例：

__device__ int select(int value) {
    return __builtin_constant_p(value) ? value * 2 : value;
}

__global__ void kernel(int* out) {
    #pragma unroll
    for (int i = 0; i < 4; ++i) {
        out[i] = select(i); // out[i] = i * 2
    }
}

5.4.9.5 __builtin_expect()

long __builtin_expect(long input, long expected)

该内置函数告诉编译器 input 预期等于 expected，并返回 input 的值。它通常用于向编译器提供分支预测信息。其行为类似于 C++20 的 [[likely]] 和 [[unlikely]] 属性。

示例：

// 向编译器指明 "var == 0" 很可能成立
if (__builtin_expect(var, 0))
    doit();

5.4.9.6 __builtin_unreachable()

void __builtin_unreachable(void)

该内置函数告诉编译器，控制流永远不会到达调用该函数的位置。如果控制流在运行时到达该位置，程序的行为未定义。

该函数对于避免生成不可达分支的代码和禁用不可达代码的编译器警告很有用。

示例：

// 向编译器指明 default 标号永远不会到达。
switch (in) {
    case 1: return 4;
    case 2: return 10;
    default: __builtin_unreachable();
}

5.4.9.7 自定义 ABI pragma

#pragma nv_abi 指令使以分离编译模式编译的应用程序通过保留函数使用的寄存器数量，获得接近整体程序编译的性能。该 pragma 的使用语法如下，其中 EXPR 指任何整数常量表达式：

#pragma nv_abi preserve_n_data(EXPR) preserve_n_control(EXPR)

#pragma nv_abi 后面的参数是可选的，可以按任意顺序提供；但至少需要一个参数。
preserve_n 参数限制函数调用期间保留的寄存器数量：
preserve_n_data(EXPR) 限制数据寄存器的数量。
preserve_n_control(EXPR) 限制控制寄存器的数量。

#pragma nv_abi 指令可以紧挨着放在设备函数的声明或定义之前。

#pragma nv_abi preserve_n_data(16)
__device__ void dev_func();

#pragma nv_abi preserve_n_data(16) preserve_n_control(8)
__device__ int dev_func() {
      return 0;
}

或者，它也可以直接放在设备函数内 C++ 表达式语句中间接函数调用的前面。注意，虽然支持对自由函数（free function）的间接函数调用，但不支持对函数引用或类成员函数的间接调用。

__device__ int dev_func1();

struct MyStruct {
    __device__ int member_func2();
};

__device__ void test() {
    auto* dev_func_ptr = &dev_func1; // 类型：int (*)(void)
    #pragma nv_abi preserve_n_control(8)
    int v1 = dev_func_ptr();          // 正确，间接调用

    #pragma nv_abi preserve_n_control(8)
    int v2 = dev_func1();             // 错误，直接调用；pragma 不起作用

                                      // dev_func1 的类型：int(void)

      auto& dev_func_ref = &dev_func1; // 类型：int (&)(void)
      #pragma nv_abi preserve_n_control(8)
      int v3 = dev_func_ref();          // 错误，对引用的调用
                                        // pragma 不起作用

       auto member_function_ptr = &MyStruct::member_func2; // 类型：int (MyStruct::*)(void)
       #pragma nv_abi preserve_n_control(8)
       int v4 = member_function_ptr(); // 错误，对成员函数的间接调用
                                         // pragma 不起作用
}

当应用于设备函数的声明或定义时，该 pragma 会修改对该函数的所有调用的自定义 ABI 属性。当放在间接函数调用点时，它只影响该特定调用的 ABI 属性。注意，该 pragma 放在调用点时只影响间接函数调用，对直接函数调用没有影响。

#pragma nv_abi preserve_n_control(8)
__device__ int dev_func3();

__device__ int dev_func4();

__device__ void test() {
     int v1 = dev_func3();                       // 正确，pragma 影响直接调用

      auto* dev_func_ptr = &dev_func4; // 类型：int (*)(void)
      #pragma nv_abi preserve_n_control(8)
      int v2 = dev_func_ptr();          // 正确，pragma 影响间接调用

       int v3 = dev_func_ptr();                   // 错误，pragma 不起作用
}

注意，如果函数的声明与其对应定义中的 pragma 参数不匹配，则程序是病态（ill-formed）的。

5.4.9.8 MMA 吞吐 pragma

nv_mma_throughput pragma 是一个指令，用于启用专门针对矩阵乘累加操作调优的编译器优化。它仅在入口函数作用域有效，即只能用于 __global__ 函数。通过沿调用路径传播的优化，其效果可能延伸到从该入口函数调用的 __device__ 函数中。

示例：

#pragma nv_mma_throughput
__global__ void kernel(){
...
}

注

nv_mma_throughput pragma 仍处于实验阶段。它启用的设置是在有限的内部 NVIDIA 工作负载上调优的，不保证每个核函数都能获得性能提升。

5.4.10 调试与诊断

5.4.10.1 断言

#define assert(expression) /* 未指定 */

如果 expression 等于零，assert() 宏会停止核函数执行。如果程序在调试器中运行，会触发断点，允许使用调试器检查设备的当前状态。否则，在通过 cudaDeviceSynchronize()、cudaStreamSynchronize() 或 cudaEventSynchronize() 与主机端同步后，每个 expression 等于零的线程都会向 stderr 打印一条消息。该消息的格式如下：

<filename>:<line number>:<function>:
block: [blockIdx.x,blockIdx.y,blockIdx.z],
thread: [threadIdx.x,threadIdx.y,threadIdx.z]
Assertion `<expression>` failed.

核函数的执行被中止，主机程序中会产生中断。assert() 宏会导致 CUDA 上下文损坏，使任何后续的 CUDA 调用或核函数启动都以 cudaErrorAssert 失败。

如果 expression 不为零，核函数的执行不受影响。

例如，源文件 test.cu 中的以下程序：

#include <assert.h>

 __global__ void testAssert(void) {
     int is_one        = 1;
     int should_be_one = 0;

      // 这不会产生任何效果
      assert(is_one);

      // 这将中止核函数执行
      assert(should_be_one);
 }

 int main(void) {
     testAssert<<<1,1>>>();
     cudaDeviceSynchronize();
     return 0;
 }

将输出：

test.cu:11: void testAssert(): block: [0,0,0], thread: [0,0,0] Assertion `should_be_one` failed.

断言用于调试目的。由于它们可能影响性能，建议在生产代码中禁用它们。可以在包含 assert.h 或 <cassert> 之前定义 NDEBUG 预处理器宏，或使用编译器标志 -DNDEBUG，在编译时禁用断言。注意，表达式不应有副作用；否则，禁用断言会影响代码的功能。

assert() 宏在 __device__ 和 __tile__ 代码中都可用。

5.4.10.2 断点函数

任何设备线程都可以通过调用 __brkpt() 函数暂停核函数的执行。

void __brkpt();

5.4.10.3 诊断 pragma

以下 pragma 可用于管理触发特定诊断消息时错误的严重程度。

#pragma nv_diag_suppress
#pragma nv_diag_warning
#pragma nv_diag_error
#pragma nv_diag_default
#pragma nv_diag_once

这些 pragma 的用法如下：

#pragma nv_diag_xxx <error_number1>, <error_number2> ...

受影响的诊断使用警告消息中显示的错误号指定。任何诊断都可以改为错误，但只有警告可以在被改为错误后将其严重程度抑制或恢复。nv_diag_default pragma 将诊断的严重程度恢复到发布任何其他 pragma 之前生效的严重程度，即被任何命令行选项修改过的消息的正常严重程度。以下示例抑制了 foo() 的"已声明但从未引用"警告：

#pragma nv_diag_suppress 177 // "declared but never referenced"
void foo() {
    int i = 0;
}

#pragma nv_diag_default 177
void bar() {
    int i = 0;
}

以下 pragma 可用于保存和恢复当前的诊断 pragma 状态：

#pragma nv_diagnostic push
#pragma nv_diagnostic pop

示例：

#pragma nv_diagnostic push
#pragma nv_diag_suppress 177 // "declared but never referenced"
void foo() {
    int i = 0;
}

#pragma nv_diagnostic pop
void bar() {
    int i = 0; // 产生警告
}

注意，这些指令只影响 nvcc CUDA 前端编译器，对主机编译器没有影响。

当支持诊断 pragma 时，nvcc 定义宏 __NVCC_DIAG_PRAGMA_SUPPORT__。

5.4.11 线程束矩阵函数

C++ 线程束矩阵操作利用 Tensor Cores 加速形如 D=A*B+C 的矩阵问题。这些操作在计算能力 7.0 或更高的设备上支持混合精度浮点数据。这要求线程束中的所有线程协作。此外，这些操作仅在条件在整个线程束中求值相同时才允许出现在条件代码中，否则代码执行很可能挂起。

5.4.11.1 描述

以下所有函数和类型都定义在命名空间 nvcuda::wmma 中。子字节操作被视为预览版，即其数据结构和 API 可能会变化，且可能与未来的版本不兼容。此额外功能定义在 nvcuda::wmma::experimental 命名空间中。

template<typename Use, int m, int n, int k, typename T, typename Layout=void> class fragment;

void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned ldm);
void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned ldm, layout_t layout);

void store_matrix_sync(T* mptr, const fragment<...> &a, unsigned ldm, layout_t layout);

void fill_fragment(fragment<...> &a, const T& v);
void mma_sync(fragment<...> &d, const fragment<...> &a, const fragment<...> &b, const fragment<...> &c, bool satf=false);

fragment 一个重载类，包含分布在线程束中所有线程上的矩阵的一部分。矩阵元素到 fragment 内部存储的映射未指定，在未来的架构中可能会变化。

只允许模板参数的某些组合。第一个模板参数指定 fragment 将如何参与矩阵运算。Use 可接受的值为：

matrix_a，当 fragment 用作第一个乘数 A 时；
matrix_b，当 fragment 用作第二个乘数 B 时；或
accumulator，当 fragment 用作源或目标累加器（分别为 C 或 D）时。

m、n 和 k 尺寸描述了参与乘累加运算的线程束级矩阵分块（tile）的形状。每个分块的维度取决于其角色。matrix_a 的分块维度为 m x k；matrix_b 的分块维度为 k x n；累加器分块为 m x n。

对于乘数，数据类型 T 可以是 double、float、__half、__nv_bfloat16、char 或 unsigned char；对于累加器，可以是 double、float、int 或 __half。如"元素类型与矩阵尺寸"中所记录的，仅支持累加器与乘数类型的有限组合。matrix_a 和 matrix_b 的 fragment 必须指定 Layout 参数。row_major 或 col_major 分别表示矩阵行或列内的元素在内存中是连续的。累加器矩阵的 Layout 参数应保留默认值 void。只有在如下文所述加载或存储累加器时，才指定行或列布局。

load_matrix_sync 等待所有线程束通道到达 load_matrix_sync，然后从内存加载矩阵 fragment a。mptr 必须是一个 256 位对齐的指针，指向内存中矩阵的第一个元素。ldm 描述连续行（行主序布局）或连续列（列主序布局）之间以元素为单位的步长；对于 __half 元素类型必须是 8 的倍数，对于 float 元素类型必须是 4 的倍数（即两种情况下都是 16 字节的倍数）。如果 fragment 是累加器，则 layout 参数必须指定为 mem_row_major 或 mem_col_major。对于 matrix_a 和 matrix_b 的 fragment，布局从 fragment 的 layout 参数推断。线程束中所有线程的 mptr、ldm、layout 和 a 的所有模板参数的值必须相同。该函数必须由线程束中的所有线程调用，否则结果未定义。

store_matrix_sync 等待所有线程束通道到达 store_matrix_sync，然后将矩阵 fragment a 存储到内存。mptr 必须是一个 256 位对齐的指针，指向内存中矩阵的第一个元素。ldm 描述连续行（行主序布局）或连续列（列主序布局）之间以元素为单位的步长；对于 __half 元素类型必须是 8 的倍数，对于 float 元素类型必须是 4 的倍数（即两种情况下都是 16 字节的倍数）。输出矩阵的布局必须指定为 mem_row_major 或 mem_col_major。线程束中所有线程的 mptr、ldm、layout 和 a 的所有模板参数的值必须相同。

fill_fragment 用常量值 v 填充矩阵 fragment。由于矩阵元素到每个 fragment 的映射未指定，该函数通常由线程束中的所有线程以相同的 v 值调用。

mma_sync 等待所有线程束通道到达 mma_sync，然后执行线程束同步的矩阵乘累加操作 D=AB+C。也支持原地操作 C=AB+C。线程束中所有线程的 satf 值和每个矩阵 fragment 的模板参数必须相同。此外，fragment A、B、C 和 D 的模板参数 m、n 和 k 必须匹配。该函数必须由线程束中的所有线程调用，否则结果未定义。

如果 satf（饱和到有限值）模式为 true，则对目标累加器适用以下附加数值属性：

如果某个元素结果为 +Infinity，则对应的累加器将包含 +MAX_NORM。
如果某个元素结果为 -Infinity，则对应的累加器将包含 -MAX_NORM。
如果某个元素结果为 NaN，则对应的累加器将包含 +0。

由于矩阵元素到每个线程的 fragment 的映射未指定，调用 store_matrix_sync 后，必须从内存（共享或全局）访问单个矩阵元素。在线程束中的所有线程将对所有 fragment 元素均匀应用逐元素操作的特殊情况下，可以使用以下 fragment 类成员实现直接元素访问。

enum fragment<Use, m, n, k, T, Layout>::num_elements;
T fragment<Use, m, n, k, T, Layout>::x[num_elements];

例如，以下代码将累加器矩阵分块缩放一半。

wmma::fragment<wmma::accumulator, 16, 16, 16, float> frag;
float alpha = 0.5f; // 线程束中所有线程的值相同
/*...*/
for(int t=0; t<frag.num_elements; t++)
frag.x[t] *= alpha;

5.4.11.2 交替浮点（Alternate Floating Point）

Tensor Cores 在计算能力 8.0 及更高的设备上支持交替类型的浮点运算。

__nv_bfloat16 此数据格式是一种交替 fp16 格式，与 f32 具有相同的范围但精度降低（7 位）。您可以直接使用 cuda_bf16.h 中提供的 __nv_bfloat16 类型来使用此数据格式。具有 __nv_bfloat16 数据类型的矩阵 fragment 必须与 float 类型的累加器组合。支持的形状和操作与 __half 相同。

tf32 此数据格式是 Tensor Cores 支持的一种特殊浮点格式，与 f32 具有相同的范围但精度降低（≥10 位）。该格式的内部布局由实现定义。要在 WMMA 操作中使用此浮点格式，必须手动将输入矩阵转换为 tf32 精度。

为便于转换，提供了新的内联函数 __float_to_tf32。虽然该内联函数的输入和输出参数都是 float 类型，但输出在数值上为 tf32。此新精度仅用于 Tensor Cores，如果与其他浮点类型操作混合，结果的精度和范围将是未定义的。

一旦输入矩阵（matrix_a 或 matrix_b）转换为 tf32 精度，将具有 precision::tf32 精度的 fragment 与 float 数据类型的组合用于 load_matrix_sync，就能利用这一新能力。两个累加器 fragment 都必须是 float 数据类型。唯一支持的矩阵尺寸是 16x16x8（m-n-k）。

fragment 的元素表示为 float，因此 element_type<T> 到 storage_element_type<T> 的映射为：

precision::tf32 -> float

5.4.11.3 双精度

Tensor Cores 在计算能力 8.0 及更高的设备上支持双精度浮点运算。要使用此新功能，必须使用 double 类型的 fragment。mma_sync 操作将使用 .rn（舍入到最近偶数）舍入修饰符执行。

5.4.11.4 子字节操作

子字节 WMMA 操作提供了一种访问 Tensor Cores 低精度能力的方式。它们被视为预览功能，即其数据结构和 API 可能会变化，且可能与未来的版本不兼容。此功能通过 nvcuda::wmma::experimental 命名空间提供：

namespace experimental {
    namespace precision {
        struct u4; // 4 位无符号
        struct s4; // 4 位有符号
        struct b1; // 1 位
    }
    enum bmmaBitOp {
        bmmaBitOpXOR = 1, // 最低 compute_75
        bmmaBitOpAND = 2  // 最低 compute_80
    };
    enum bmmaAccumulateOp { bmmaAccumulateOpPOPC = 1 };
}

对于 4 位精度，可用的 API 保持不变，但必须将 experimental::precision::u4 或 experimental::precision::s4 指定为 fragment 数据类型。由于 fragment 的元素是打包在一起的，该 fragment 的 num_storage_elements 将小于 num_elements。因此，子字节 fragment 的 num_elements 变量返回子字节类型 element_type<T> 的元素数量。这对单比特精度同样成立，此时 element_type<T> 到 storage_element_type<T> 的映射如下：

experimental::precision::u4 -> unsigned（1 个存储元素中有 8 个元素）
experimental::precision::s4 -> int（1 个存储元素中有 8 个元素）
experimental::precision::b1 -> unsigned（1 个存储元素中有 32 个元素）
T -> T（所有其他类型）

子字节 fragment 允许的布局始终为：matrix_a 用 row_major，matrix_b 用 col_major。

对于子字节操作，load_matrix_sync 中 ldm 的值：对于元素类型 experimental::precision::u4 和 experimental::precision::s4 应为 32 的倍数，对于元素类型 experimental::precision::b1 应为 128 的倍数（即两种情况下都是 16 字节的倍数）。

注

以下 MMA 指令变体的支持已弃用，并将在 sm_90 中移除：

experimental::precision::u4

experimental::precision::s4

experimental::precision::b1 且 bmmaBitOp 设置为 bmmaBitOpXOR

bmma_sync 等待所有线程束通道执行 bmma_sync，然后执行线程束同步的位矩阵乘累加操作 D = (A op B) + C，其中 op 由逻辑操作 bmmaBitOp 和 bmmaAccumulateOp 定义的累加组成。可用操作为：

bmmaBitOpXOR：matrix_a 的一行与 matrix_b 的一列的 128 位异或；
bmmaBitOpAND：matrix_a 的一行与 matrix_b 的一列的 128 位与，在计算能力 8.0 及更高的设备上可用；
累加操作始终是 bmmaAccumulateOpPOPC，它统计置位（set）比特的数量。

5.4.11.5 限制

张量核（tensor cores）所需的特殊格式可能因每个设备架构的主版本号和次版本号而异。线程只持有整个矩阵的一个 fragment（不透明的、架构特定的 ABI 数据结构），开发者不得对各个参数如何映射到参与矩阵乘累加的寄存器做出假设，这进一步增加了复杂性。

由于 fragment 是架构特定的，如果函数 A 和函数 B 为不同的链接兼容架构编译并链接到同一个设备可执行文件中，那么在它们之间传递 fragment 是不安全的。这种情况下，fragment 的大小和布局将是某个特定架构的，在另一个架构中使用 WMMA API 会导致错误的结果，甚至可能造成损坏。两个链接兼容架构的一个示例是 sm_70 和 sm_75，它们的 fragment 布局不同。

fragA.cu: void foo() { wmma::fragment<...> mat_a; bar(&mat_a); }
fragB.cu: void bar(wmma::fragment<...> *mat_a) { // 操作 mat_a }

# sm_70 fragment 布局
$> nvcc -dc -arch=compute_70 -code=sm_70 fragA.cu -o fragA.o
# sm_75 fragment 布局
$> nvcc -dc -arch=compute_75 -code=sm_75 fragB.cu -o fragB.o
# 将两者链接在一起
$> nvcc -dlink -arch=sm_75 fragA.o fragB.o -o frag.o

这种未定义行为在编译时和运行时工具中也可能无法检测到，因此需要格外小心以确保 fragment 的布局一致。这种链接风险最可能出现在链接到一个遗留库时，该库既为不同的链接兼容架构构建，又期望被传入一个 WMMA fragment。

注意，在弱链接（weak linkage）的情况下（例如 CUDA C++ 内联函数），链接器可能选择任何可用的函数定义，这可能导致编译单元之间的隐式传递。

要避免这类问题，矩阵在经过外部接口传递时应始终先存储到内存（例如 wmma::store_matrix_sync(dst, …);），然后再以指针类型（例如 float *dst）安全地传给 bar()。

注意，由于 sm_70 可以在 sm_75 上运行，上面示例中的 sm_75 代码可以改为 sm_70 并在 sm_75 上正确工作。但是，当与为 sm_75 单独编译的其他二进制文件链接时，建议在应用程序中使用 sm_75 原生代码。

5.4.11.6 元素类型与矩阵尺寸

Tensor Cores 支持多种元素类型和矩阵尺寸。下表列出了支持的 matrix_a、matrix_b 和累加器矩阵的各种组合：

矩阵 A	矩阵 B	累加器	矩阵尺寸 (m-n-k)
`__half`	`__half`	float	16x16x16
`__half`	`__half`	float	32x8x16
`__half`	`__half`	float	8x32x16
`__half`	`__half`	`__half`	16x16x16
`__half`	`__half`	`__half`	32x8x16
`__half`	`__half`	`__half`	8x32x16
unsigned char	unsigned char	int	16x16x16
unsigned char	unsigned char	int	32x8x16
unsigned char	unsigned char	int	8x32x16
signed char	signed char	int	16x16x16
signed char	signed char	int	32x8x16
signed char	signed char	int	8x32x16

交替浮点支持：

矩阵 A	矩阵 B	累加器	矩阵尺寸 (m-n-k)
`__nv_bfloat16`	`__nv_bfloat16`	float	16x16x16
`__nv_bfloat16`	`__nv_bfloat16`	float	32x8x16
`__nv_bfloat16`	`__nv_bfloat16`	float	8x32x16
`precision::tf32`	`precision::tf32`	float	16x16x8

双精度支持：

矩阵 A	矩阵 B	累加器	矩阵尺寸 (m-n-k)
double	double	double	8x8x4

子字节操作的实验性支持：

矩阵 A	矩阵 B	累加器	矩阵尺寸 (m-n-k)
`precision::u4`	`precision::u4`	int	8x8x32
`precision::s4`	`precision::s4`	int	8x8x32
`precision::b1`	`precision::b1`	int	8x8x128

5.4.11.7 示例

以下代码在单个线程束中实现 16x16x16 矩阵乘法。

#include <mma.h>
using namespace nvcuda;

__global__ void wmma_ker(half *a, half *b, float *c) {
   // 声明 fragment
   wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::col_major> a_frag;
   wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::row_major> b_frag;
   wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;

   // 将输出初始化为零
   wmma::fill_fragment(c_frag, 0.0f);

   // 加载输入
   wmma::load_matrix_sync(a_frag, a, 16);
   wmma::load_matrix_sync(b_frag, b, 16);

   // 执行矩阵乘法
   wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

   // 存储输出
   wmma::store_matrix_sync(c, c_frag, 16, wmma::mem_row_major);
}

📘 CUDA 编程指南 · 中文 Wiki