▶ 可缓存只读操作（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 \t" #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\n", 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" 。
