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 。

T __ldg(const T* address)

▶ 原子操作

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

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

__global__ void mykernel(int *addr)
{
    atomicAdd_system(addr, 10);// GPU端全局原子加法
}


void foo()
{
    int *addr;
    cudaMallocManaged(&addr, 4);
    *addr = 0;
    
    mykernel << <... >> >(addr);

    __sync_fetch_and_add(addr, 10);// CPU端全局原子加法
}

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

__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int *address_as_ull = (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull;
    unsigned long long int assumed;
    
    do
    {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));
        // 将数据转化为 long long 来计算加法，防止 NaN 的比较和交换（NaN != NaN）
    } while (assumed != old);

    return __longlong_as_double(old);
}

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

// 原子加法， address 的值加上 val，返回 address 旧值
int atomicAdd(int* address, int val); 
unsigned int atomicAdd(unsigned int* address, unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address, unsigned long long int val);
float atomicAdd(float* address, float val); 
double atomicAdd(double* address, double val); 

// 原子减法， address 的值减去 val，返回 address 旧值
int atomicSub(int* address, int val);
unsigned int atomicSub(unsigned int* address, unsigned int val);

// 原子赋值， adress 赋值 val，返回 adress 旧值                                                                   
int atomicExch(int* address, int val);
unsigned int atomicExch(unsigned int* address, unsigned int val);
unsigned long long int atomicExch(unsigned long long int* address, unsigned long long int val);
float atomicExch(float* address, float val);

// 原子 min， address 赋值 min(*address, val)，返回 adress 旧值
int atomicMin(int* address, int val);
unsigned int atomicMin(unsigned int* address, unsigned int val);
unsigned long long int atomicMin(unsigned long long int* address, unsigned long long int val);

// 原子 max， address 赋值 max(*address, val)，返回 adress 旧值
int atomicMax(int* address, int val);
unsigned int atomicMax(unsigned int* address, unsigned int val);
unsigned long long int atomicMax(unsigned long long int* address, unsigned long long int val);

// 原子自增， address 赋值 ((*address >= val) ? 0 : (*address + 1))，返回 adress 旧值
unsigned int atomicInc(unsigned int* address,unsigned int val);

// 原子自减， address 赋值 (((*address == 0) | (*address > val)) ? val : (*address - 1))，返回 adress 旧值
unsigned int atomicDec(unsigned int* address, unsigned int val);

// 原子交换，address 赋值 (*address == compare ? val : *address)，返回 adress 旧值
int atomicCAS(int* address, int compare, int val);
unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val);
unsigned long long int atomicCAS(unsigned long long int* address, unsigned long long int compare, unsigned long long int val);

// 原子按位与，address 赋值 (*address & val )，返回 adress 旧值
int atomicAnd(int* address, int val); 
unsigned int atomicAnd(unsigned int* address, unsigned int val); 
unsigned long long int atomicAnd(unsigned long long int* address, unsigned long long int val);

// 原子按位或，address 赋值 (*address | val )，返回 adress 旧值
int atomicOr(int* address, int val);
unsigned int atomicOr(unsigned int* address, unsigned int val); 
unsigned long long int atomicOr(unsigned long long int* address, unsigned long long int val);

// 原子按位异或，address 赋值 (*address ^ val )，返回 adress 旧值
int atomicXor(int* address, int val);
unsigned int atomicXor(unsigned int* address, unsigned int val);
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（略过）

//device_functions.h
#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() 使用共享数据结构，可能会改变线程时间的执行时间和顺序。

● 调整设备参数的函数。

// driver_types.h
enum __device_builtin__ cudaLimit
{
    cudaLimitStackSize = 0x00,                      // 栈尺寸
    cudaLimitPrintfFifoSize = 0x01,                 // printf/fprintf 缓冲区尺寸
    cudaLimitMallocHeapSize = 0x02,                 // 堆内存尺寸
    cudaLimitDevRuntimeSyncDepth = 0x03,            // ？运行时同步深度
    cudaLimitDevRuntimePendingLaunchCount = 0x04    // ？运行时待办调用计数
};

//cuda_runtime.h
extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetLimit(size_t *pValue, enum cudaLimit limit);
extern __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t value);

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

// cuda.h
typedef enum CUlimit_enum {
    CU_LIMIT_STACK_SIZE = 0x00,                         // 栈尺寸
    CU_LIMIT_PRINTF_FIFO_SIZE = 0x01,                   // printf/fprintf 缓冲区尺寸
    CU_LIMIT_MALLOC_HEAP_SIZE = 0x02,                   // 堆内存尺寸
    CU_LIMIT_DEV_RUNTIME_SYNC_DEPTH = 0x03,             // ？运行时同步深度
    CU_LIMIT_DEV_RUNTIME_PENDING_LAUNCH_COUNT = 0x04,   // ？运行时待办调用计数
    CU_LIMIT_MAX
} CUlimit;

CUresult CUDAAPI cuCtxGetLimit(size_t *pvalue, CUlimit limit);
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()；设备线程块动态内存申请可以声明一个共享内存指针，使用其中一个线程来申请相应的内存。

● 代码举例

#include <stdio.h>
#include <malloc.h>
#include <cuda_runtime.h>
#include "device_launch_parameters.h" 

__global__ void mallocTest()
{
    size_t size = 1024;
    int *ptr = (int*)malloc(sizeof(int)*size);
    memset(ptr, 0, size);
    printf("Thread %d got pointer: %p
", threadIdx.x, ptr);
    free(ptr);
}

int main()
{
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, 4 * 1024 * 1024);// 设定申请的堆内存上限
    mallocTest << <1, 4 >> >();
    cudaDeviceSynchronize();

    getchar();
    return 0;
}

● 输出结果：

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 变量。

● 代码举例

// 完全展开
#pragma unroll
for (i = 0; i < m; i++)
    c[i] = a[i] + b[i];

// 不展开
#pragma unroll 1
for (i = 0; i < m; i++)
    c[i] = a[i] + b[i];

// 部分展开（这里展开了前 4 次迭代）
#pragma unroll 4
for (i = 0; i < m; i++)
    c[i] = a[i] + b[i];

// 部分展开，使用常量表达式
const int n = 4;
#pragma unroll n
for (i = 0; i < m; i++)
    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() 基本语法

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

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

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