一、概述

1. CUDA编程结构

CUDA6.0开始提出统一寻址（Unified Memory），它连接了主机内存和设备内存空间，可使用单个指针访问CPU和GPU内存。

主机代码按照ANSI C标准进行编写，而设备代码使用CUDA C进行编写。

一个典型的CUDA程序实现流程遵循以下模式。

1. 把数据从CPU内存拷贝到GPU内存。

2.调用核函数对存储在GPU内存中的数据进行操作。

3.将数据从GPU内存传送回到CPU内存。

2. 内存管理

CUDA编程模型假设系统是由一个主机和一个设备组成，而且各自拥有独立的内存。

①cudaMalloc函数负责执行GPU内存分配：其函数原型为：

        cudaError_t cudaMalloc (void** devPtr, size_t size)

该函数负责向设备分配一定字节的线性内存，并以devPtr的形式返回指向所分配内存的指针。

②cudaMemcpy函数负责主机和设备之间的数据传输，其函数原型为：

        函数 cudaError_t cudaMemcpy (void* dst, const void* src, size_t count, cudaMemcpyKind kind)

此从src指向的源存储区复制一定数量的字节到dst指向的目标存储区.复制方向由kind指定，其中的kind有以下几种：

内存层级结构

在GPU的内存层次结构中，最主要的两种内存是全局内存和共享内存。全局内存类似于CPU的系统内存，而共享内存类似于CPU的缓存。然而GPU的共享内存可以由CUDA C的内核直接控制。

一旦内核被调用，控制权立刻被传回主机。当核函数在GPU上运行时，主机可以执行其他函数。因此，内核与主机是异步的。

cudaMemcpy的调用会导致主机运行阻塞。

3. 线程管理

当核函数在主机端启动时，它的执行会移动到设备上，此时设备中会产生大量的线程，并且每个线程都执行由核函数指定的语句。

由一个内核启动所产生的所有线程统称为一个网格。同一网格中的所有线程共享相同的全局内存空间。一个网格由多个线程块构成，一个线程块包含一组线程，同一个线程块内线程协作可以通过同步和共享内存的方式实现。

不同块内的线程不能协作。

线程依靠以下两个坐标变量区分彼此。

· blockIdx（线程块在线程格内的索引）

· threadIdx（块内的线程索引）

CUDA可以组织三维的网格和块。网格和块的维度由下列两个内置变量指定。

· blockDim（线程块的维度，用每个线程块中的线程数来表示）

· gridDim（线程格的维度，用每个线程格中的线程数来表示）

它们是dim3类型的变量，是基于uint3定义的整数型向量，用来表示维度。当定义一个dim3类型的变量时，所有未指定的元素都被初始化为1且忽略不计。dim3类型变量中的每个组件可以通过它的x、y、z字段获得。

网格和线程块的维度

通常，一个线程格会被组织成线程块的二维数组形式，一个线程块会被组织成线程的三维数组形式。

对于块大小的一个主要限制因素就是可利用的计算资源，如寄存器，共享内存等。

4. 启动一个CUDA核函数

CUDA内核调用是对C语言函数调用语句的延伸，<<<>>>运算符内时核函数的执行配置。

kernel_name <<<grid, block>>>(argument list);

执行配置中的第一个值是网格维度，也就是启动块的数目。第二个值是块维度，也就是每个块中线程的数目。

核函数的调用与主机线程是异步的。核函数调用结束后，控制权立刻返回给主机端。可以调用以下函数来强制主机端程序等待所有的核函数执行结束：

        cudaError_t cudaDeviceSynchronize(void);

一些CUDA运行时API在主机和设备之间是隐式同步的当使用cudaMemcpy函数在主机和设备之间拷贝数据时，主机端隐式同步，即主机端程序必须等待数据拷贝完成后才能继续执行程序。cudaMemcpy之前所有的核函数调用完成后开始拷贝数据。当拷贝完成后，控制权立刻返回给主机端。

异步行为

不同于C语言的函数调用，所有的CUDA核函数的启动都是异步的。CUDA内核调用完成后，控制权立刻返回给CPU。

5. 编写核函数

核函数是在设备端执行的代码。在核函数中，需要为一个线程规定要进行的计算以及要进行的数据访问。当核函数被调用时，许多不同的CUDA线程并行执行同一个计算任务。以下是用__global__声明定义核函数：

        __global__ void kernel_name(argument list);

核函数必须有一个void返回类型。

CUDA函数的限制

· 只能访问设备内存

· 必须具有void返回类型

· 不支持可变数量的参数

· 不支持静态变量

· 显示异步行为

6. 验证核函数

验证核函数代码

两个简单实用的方法：① 在核函数中使用printf函数。② 可以将执行参数设置为<<<1,1>>>，强制用一个块核一个线程执行核函数，模拟串行执行程序。

7. 处理错误

由于许多CUDA调用是异步的，所以有时可能很难确定某个错误是由哪一步程序引起的。定义一个错误处理宏封装所有的CUDA API调用，这简化了错误检查过程：

二、给核函数计时

衡量核函数性能的方法有很多，最简单的方法是在主机端使用一个CPU或GPU计时器来计算内核的执行时间。

1. 用CPU计时器计时

2. 用nvprof工具计时

将应用程序的测量值与理论峰值进行比较，可以判定你的应用程序时受限于算法还是受限于内存带宽。大多数HPC工作负载受内存带宽的限制。

3. 组织并行线程

· 改变执行配置对内核性能有影响

· 传统的核函数实现一般不能获得最佳性能

· 对于一个给定的核函数，尝试使用不同的网格和线程块大小可以获得更好的性能

四、 设备管理

在本节，你将通过以下两种方法学习查询和管理GPU设备：

· CUDA运行时API函数

· NVIDIA系统管理界面（nvidia-smi）命令行实用程序

对于一个有N个GPU的系统，nvidia-smi从0到N-1标记设备ID。使用环境变量CUDA_VISIBLE_DEVICES，就可以在运行时指定所选的GPU且无须更改应用程序。例如，设置CUDA_VISIBLE_DEVICES=2,3。然后在运行时，nvidia驱动程序将只使用ID为2和3的设备，并会将设备分别映射为0和1。