CUDA C Programming Guide 在线教程学习笔记 Part 7

▶ 可缓存只读操作（Read-Only Data Cache Load Function），定义在 sm_32_intrinsics.hpp 中。从地址 adress 读取类型为 T 的函数返回，T 可以是 char，short，int，long longunsigned char，unsigned short，unsigned int，unsigned long long，int2，int4，uint2，uint4，float，float2，float4，doubledouble2 。

1 T __ldg(const T* address)

▶ 原子操作

● 原子操作只能在设备代码上使用。一台设备上的原子操作仅对该设备的内存体现原子性，跨设备原子操作（一台GPU对另一台GPU，或一台GPU对CPU）被视为普通读写操作。

● cc6.x引入限定范围的原子操作，如 atomicAdd_system() 限定原子操作对系统中主机和所有设备有效，atmoicAdd_block() 限定原子操作只对该线程块内所有线程有效等。代码举例：

 1 __global__ void mykernel(int *addr)
 2 {
 3     atomicAdd_system(addr, 10);// GPU端全局原子加法
 4 }
 5 
 6 
 7 void foo()
 8 {
 9     int *addr;
10     cudaMallocManaged(&addr, 4);
11     *addr = 0;
12     
13     mykernel << <... >> >(addr);
14 
15     __sync_fetch_and_add(addr, 10);// CPU端全局原子加法
16 }

● 所有原子操作均可以通过函数 atomicCAS() 来实现。代码举例（在 cc6.x 以下的系统中实现双精度原子加法）：

 1 __device__ double atomicAdd(double* address, double val)
 2 {
 3     unsigned long long int *address_as_ull = (unsigned long long int*)address;
 4     unsigned long long int old = *address_as_ull;
 5     unsigned long long int assumed;
 6     
 7     do
 8     {
 9         assumed = old;
10         old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));
11         // 将数据转化为 long long 来计算加法，防止 NaN 的比较和交换（NaN != NaN）
12     } while (assumed != old);
13 
14     return __longlong_as_double(old);
15 }

● 原子操作，定义于 device_atomic_functions.h 。

 1 // 原子加法， address 的值加上 val，返回 address 旧值
 2 int atomicAdd(int* address, int val); 
 3 unsigned int atomicAdd(unsigned int* address, unsigned int val);
 4 unsigned long long int atomicAdd(unsigned long long int* address, unsigned long long int val);
 5 float atomicAdd(float* address, float val); 
 6 double atomicAdd(double* address, double val); 
 7 
 8 // 原子减法， address 的值减去 val，返回 address 旧值
 9 int atomicSub(int* address, int val);
10 unsigned int atomicSub(unsigned int* address, unsigned int val);
11 
12 // 原子赋值， adress 赋值 val，返回 adress 旧值                                                                   
13 int atomicExch(int* address, int val);
14 unsigned int atomicExch(unsigned int* address, unsigned int val);
15 unsigned long long int atomicExch(unsigned long long int* address, unsigned long long int val);
16 float atomicExch(float* address, float val);
17 
18 // 原子 min， address 赋值 min(*address, val)，返回 adress 旧值
19 int atomicMin(int* address, int val);
20 unsigned int atomicMin(unsigned int* address, unsigned int val);
21 unsigned long long int atomicMin(unsigned long long int* address, unsigned long long int val);
22 
23 // 原子 max， address 赋值 max(*address, val)，返回 adress 旧值
24 int atomicMax(int* address, int val);
25 unsigned int atomicMax(unsigned int* address, unsigned int val);
26 unsigned long long int atomicMax(unsigned long long int* address, unsigned long long int val);
27 
28 // 原子自增， address 赋值 ((*address >= val) ? 0 : (*address + 1))，返回 adress 旧值
29 unsigned int atomicInc(unsigned int* address,unsigned int val);
30 
31 // 原子自减， address 赋值 (((*address == 0) | (*address > val)) ? val : (*address - 1))，返回 adress 旧值
32 unsigned int atomicDec(unsigned int* address, unsigned int val);
33 
34 // 原子交换，address 赋值 (*address == compare ? val : *address)，返回 adress 旧值
35 int atomicCAS(int* address, int compare, int val);
36 unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val);
37 unsigned long long int atomicCAS(unsigned long long int* address, unsigned long long int compare, unsigned long long int val);
38 
39 // 原子按位与，address 赋值 (*address & val )，返回 adress 旧值
40 int atomicAnd(int* address, int val); 
41 unsigned int atomicAnd(unsigned int* address, unsigned int val); 
42 unsigned long long int atomicAnd(unsigned long long int* address, unsigned long long int val);
43 
44 // 原子按位或，address 赋值 (*address | val )，返回 adress 旧值
45 int atomicOr(int* address, int val);
46 unsigned int atomicOr(unsigned int* address, unsigned int val); 
47 unsigned long long int atomicOr(unsigned long long int* address, unsigned long long int val);
48 
49 // 原子按位异或，address 赋值 (*address ^ val )，返回 adress 旧值
50 int atomicXor(int* address, int val);
51 unsigned int atomicXor(unsigned int* address, unsigned int val);
52 unsigned long long int atomicXor(unsigned long long int* address, unsigned long long int val);

▶ 线程束表决函数（Warp Vote Functions）见 part 8

▶ 线程束匹配函数（Warp Match Functions）见 part 8

▶ 线程束交织函数（Warp Shuffle Functions）见 part 8

▶ 线程束矩阵函数 Warp matrix functions [PREVIEW FEATURE]（略过）

▶ B.17. Profiler Counter Function（略过）

1 //device_functions.h
2 #define __prof_trigger(X) asm __volatile__ ("pmevent 	" #X ";")

▶ 警告函数 Assertion

● 代码举例在 Samples中，http://www.cnblogs.com/cuancuancuanhao/p/7775244.html 。

● 设备代码中触发 assert() 后，当主机中调用同步函数 cudaDeviceSynchronize()，cudaStreamSynchronize()，cudaEventSynchronize() 时将向 stderr 中写入错误信息，格式为：

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

● 设备代码中触发 assert() 后，主机调用该设备的任何调用都会返回 cudaErrorAssert（罢工），除非使用 cudaDeviceReset() 重新初始化该设备。

● 可以在预处理代码 #include assert.h 之前定义 #define NDUG 来使所有函数 assert() 无效化，减少该函数对性能造成的损失。

● 建议：在 assert() 的条件和内部不要使用会改变变量的值的操作，防止禁用 assert() 前后对结果的影响。

▶ 格式化输出函数 printf()

● 设备代码中的 printf() 返回输出的参数个数（不同于 C 中返回打印的字符数），上限32个。无参数时返回 0，输出表达式为 NULL 时返回 -1，内部错误返回 -2 。

● 设备代码中的 printf() 不会自己检查错误，而是交给主机完成最终的格式化和输出（注意格式兼容性问题）。

● 设备中 printf() 使用的缓冲区市固定大小的环形，若一次需要输出的内容太多有可能在缓冲区刷新之前就发生覆盖。以下过程可以刷新缓冲区：

■ 调用设备函数 <<< >>> 或 cuLaunchKernel()（调用前一定隐式刷新，如果环境变量 CUDA_LAUNCH_BLOCKING == 1，则调用后再次隐式刷新）

■ 使用同步函数 cudaDeviceSynchronize()，cuCtxSynchronize()，cudaStreamSynchronize()，cuStreamSynchronize()，cudaEventSynchronize()，cuEventSynchronize()

■ 使用内存拷贝函数 cudaMemcpy*()，cuMemcpy*()

■ 使用模块读取函数 cuModuleLoad()，cuModuleUnload()

■ 使用销毁上下文函数 cudaDeviceReset()，cuCtxDestroy()

■ 执行回调函数 cudaStreamAddCallback()，cuStreamAddCallback()

● 设备中 printf() 在程序退出时不会自动刷新，需要显式的调用 cudaDeviceReset()，cuCtxDestroy() 来强制刷新

● 设备中 printf() 使用共享数据结构，可能会改变线程时间的执行时间和顺序。

● 调整设备参数的函数。

 1 // driver_types.h
 2 enum __device_builtin__ cudaLimit
 3 {
 4     cudaLimitStackSize = 0x00,                      // 栈尺寸
 5     cudaLimitPrintfFifoSize = 0x01,                 // printf/fprintf 缓冲区尺寸
 6     cudaLimitMallocHeapSize = 0x02,                 // 堆内存尺寸
 7     cudaLimitDevRuntimeSyncDepth = 0x03,            // ？运行时同步深度
 8     cudaLimitDevRuntimePendingLaunchCount = 0x04    // ？运行时待办调用计数
 9 };
10 
11 //cuda_runtime.h
12 extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetLimit(size_t *pValue, enum cudaLimit limit);
13 extern __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t value);

● 类似的在 Driver API 中的函数。

 1 // cuda.h
 2 typedef enum CUlimit_enum {
 3     CU_LIMIT_STACK_SIZE = 0x00,                         // 栈尺寸
 4     CU_LIMIT_PRINTF_FIFO_SIZE = 0x01,                   // printf/fprintf 缓冲区尺寸
 5     CU_LIMIT_MALLOC_HEAP_SIZE = 0x02,                   // 堆内存尺寸
 6     CU_LIMIT_DEV_RUNTIME_SYNC_DEPTH = 0x03,             // ？运行时同步深度
 7     CU_LIMIT_DEV_RUNTIME_PENDING_LAUNCH_COUNT = 0x04,   // ？运行时待办调用计数
 8     CU_LIMIT_MAX
 9 } CUlimit;
10 
11 CUresult CUDAAPI cuCtxGetLimit(size_t *pvalue, CUlimit limit);
12 CUresult CUDAAPI cuCtxSetLimit(CUlimit limit, size_t value);

▶ 动态堆内存申请

● 堆内存申请失败时返回错误 CUDA_ERROR_SHARED_OBJECT_INIT_FAILED 。

● cc2.0 以上的设备，在设备代码中使用动态内存分配时，应该使用函数 malloc()，并胚胎使用 memset() 和 free() 。不能使用 cudaMalloc() 来申请（实验表明只能获得空指针）或 cudaMemset()（限定 __host__ 函数）和 cudaFree() （函数不配套）。

● 主机中使用 cudaMalloc() 仅受限于可使用的设备内存，而设备代码中中使用 malloc() 受限于设备堆内存申请上限参数 cudaLimitMallocHeapSize，可能需要在申请前临时修改（类似修改 printf() 的缓冲区）。

● 设备线程动态内存申请可以直接用 malloc()；设备线程块动态内存申请可以声明一个共享内存指针，使用其中一个线程来申请相应的内存。

● 代码举例

 1 #include <stdio.h>
 2 #include <malloc.h>
 3 #include <cuda_runtime.h>
 4 #include "device_launch_parameters.h" 
 5 
 6 __global__ void mallocTest()
 7 {
 8     size_t size = 1024;
 9     int *ptr = (int*)malloc(sizeof(int)*size);
10     memset(ptr, 0, size);
11     printf("Thread %d got pointer: %p
", threadIdx.x, ptr);
12     free(ptr);
13 }
14 
15 int main()
16 {
17     cudaDeviceSetLimit(cudaLimitMallocHeapSize, 4 * 1024 * 1024);// 设定申请的堆内存上限
18     mallocTest << <1, 4 >> >();
19     cudaDeviceSynchronize();
20 
21     getchar();
22     return 0;
23 }

● 输出结果：

Thread 0 got pointer: 0000000B017FF920
Thread 1 got pointer: 0000000B017F8020
Thread 2 got pointer: 0000000B017F7720
Thread 3 got pointer: 0000000B017F6F20

▶ 预编译命令 #pragma unroll 展开循环。

● #pragma unroll 命令加在循环之前。可以不另加参数，表示循环完全展开；也可以加整形常量表达式，如数字常量表达式或 const 变量。

● 代码举例

 1 // 完全展开
 2 #pragma unroll
 3 for (i = 0; i < m; i++)
 4     c[i] = a[i] + b[i];
 5 
 6 // 不展开
 7 #pragma unroll 1
 8 for (i = 0; i < m; i++)
 9     c[i] = a[i] + b[i];
10 
11 // 部分展开（这里展开了前 4 次迭代）
12 #pragma unroll 4
13 for (i = 0; i < m; i++)
14     c[i] = a[i] + b[i];
15 
16 // 部分展开，使用常量表达式
17 const int n = 4;
18 #pragma unroll n
19 for (i = 0; i < m; i++)
20     c[i] = a[i] + b[i];

▶ SIMD 视频指令

● cc3.0以上设备，汇编优化的PTX指令，同时操纵 4 个 8 bit 或 2 个 16 bit 数据。

● 指令举例：vadd2，vadd4，vsub2，vsub4，vavrg2，vavrg4，vabsdiff2，vabsdiff4，vmin2，vmin4，vmax2，vmax4，vset2，vset4

● asm() 基本语法

1 asm("template-string" : "constraint"(output) : "constraint"(input)"));

● 代码举例。使用指令 vabsdiff4 计算整形 4 字节 SIMD （理解成向量）A 和 B 绝对值差的和，放入 C 中。

1 asm("vabsdiff4.u32.u32.u32.add" " %0, %1, %2, %3;": "=r" (result):"r" (A), "r" (B), "r" (C));

● 其他参考资料："Using Inline PTX Assembly in CUDA"，"Parallel Thread Execution ISA Version 3.0" 。