第四章 CUDA编程模型

1、第四章第四章CUDACUDA编程模型编程模型本章内容本章内容2021-12-14并行计算4.1 CUDA4.1 CUDA并行线程组织结构并行线程组织结构4.2 4.2 存储器层次和显存分配存储器层次和显存分配4.3 CUDA4.3 CUDA变量和函数变量和函数上节内容回顾上节内容回顾2021-12-14并行计算扩展：如何实现两个向量的相加？上节内容回顾上节内容回顾2021-12-14并行计算 _global_ void add(int *a, int *b, int *c) cblockIdx.x = ablockIdx.x + bblockIdx.x;上节内容回顾上节内容回顾2021-12-

2、14并行计算#define N 512 int main() int *a, *b, *c; int *d_a, *d_b, *d_c; int size = N * sizeof(int); /分配GPU显存空间 cudaMalloc(void*)&d_a, size); cudaMalloc(void*)&d_a, size); cudaMalloc(void*)&d_a, size); /为输入数据分配内存空间 a = (int*)malloc(size); random_ints(a, N); b = (int*)malloc(size); random_int

3、s(b, N); c = (int*)malloc(size); /将内存中的输入数据拷贝到显存 cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice); /调用kernel add(d_a,d_b,d_c); /将显存中的计算结果回读到内存 cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost); /释放内存 free(a); free(b); free(c); cudaFree(d_a); cudaFree(

4、d_b); cudaFree(d_c); return 0; 上节内容回顾上节内容回顾2021-12-14并行计算如果调用kernel时指定启动1个线程块，此时，执行配置中参数应如何设置？上节内容回顾上节内容回顾2021-12-14并行计算 _global_ void add(int *a, int *b, int *c) cthreadIdx.x = athreadIdx.x + bthreadIdx.x;上节内容回顾上节内容回顾2021-12-14并行计算#define N 512 int main() int *a, *b, *c; int *d_a, *d_b, *d_c; int s

5、ize = N * sizeof(int); /分配GPU显存空间 cudaMalloc(void*)&d_a, size); cudaMalloc(void*)&d_a, size); cudaMalloc(void*)&d_a, size); /为输入数据分配内存空间 a = (int*)malloc(size); random_ints(a, N); b = (int*)malloc(size); random_ints(b, N); c = (int*)malloc(size); /将内存中的输入数据拷贝到显存 cudaMemcpy(d_a, a, size,

6、cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice); /调用并启动调用并启动kernel add(d_a,d_b,d_c); /将显存中的计算结果回读到内存 cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost); /释放内存 free(a); free(b); free(c); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); return 0; 4.1 并行线程组织结构并行线程组织结构4.1.1 并行线程组织结构并

7、行线程组织结构2021-12-14并行计算 线程(Thread)：kernel并行执行的最小单位 线程块(Block)：由多个线程组成 以1D、2D或3D组织 Block内的线程通过共享内存交换数据 最多包含512个线程 线程网格(Grid)：由多个相同大小和维度的线程块组成 以1D、2D或3D组织 各Block是并行执行的，没有执行顺序，共享全局内存 线程块数量不超过65535（硬件限制）4.1.1 并行线程组织结构并行线程组织结构2021-12-14并行计算线程束线程束2021-12-14并行计算 这32个线程的集合被“编织在一起”并且“步调一致”的形式执行。在程序中的每一行，线程束中的每

8、个线程都将在不同数据上执行相同的命令。 为了管理参与计算的数百个线程，SM采用了SIMT（Single Instruction，Multiple Thread，单指令多线程）执行模型。然后，创建、管理、调度和执行线程的工作都是由SIMT单元以warp块（32个连续的并行线程）为单位来完成的，且不同warp间的切换是零开销的。当一个SM指定了一个或者一个以上要执行计算的线程块时，这些线程块将会被分成多个warp块，然后等待SIMT单元调度。当SIMT单元收到指令单元发出的指令时，它就会选择一个已经处于等待状态的warp块，然后把下一条执行指令发送到已选择warp块的所有活动线程。线程束线程束20

9、21-12-14并行计算4.1.2 执行配置执行配置2021-12-14并行计算对于一个线程网格内的线程块数目和每个线程块内的线程数可通过执行配置来指定，其中gridDim和blockDim是CUDA定义的内置变量，分别表示了线程网格、线程块的维度（Dimension）与尺寸，这两个变量均为dim3数据类型。4.1.2 执行配置执行配置2021-12-14并行计算 执行配置指令，通过在内核函数即使用_global_限定符的函数名称和所有括号中的参数列表间插入表达式来实现，用于说明如何在CPU端启动内核以在GPU上执行并行计算。其中： Dg：指定线程网格维度、大小，类型为dim3。 Db：指定各

10、线程块维度、大小，类型为dim3。 Ns：可选参数，用于指定启动内核函数后需要为各线程块动态分配共享存储器的大小，默认的值为0，类型为size_t。 s ： 可 选 参 数 ， 用 来 指 定 相 关 流 ， 默 认 的 值 为 0 ， 类 型 为cudaStream_t。4.1.2 执行配置执行配置2021-12-14并行计算 dim3结构类型 dim3是基亍uint3定义的矢量类型，相当于由3个unsigned int型组成的结构体。uint3类型有三个数据成员unsigned int x; unsigned int y; unsigned int z。 可使用于1D、2D或3D的索引。

11、dim3结构类型变量用在核函数调用的中。4.1.2 执行配置执行配置2021-12-14并行计算Block和Grid的维度 1D、2D或3D（变量类型为dim3） blockDim，gridDim blockDim.x，blockDim.y，blockDim.z gridDim.x，gridDim.y，gridDim.z4.1.2 执行配置执行配置2021-12-14并行计算Block和Thread具有各自的ID 1D、2D或3D（变量类型为uint3） 线程块和线程ID索引blockIdx，threadIdx blockIdx.x，blockIdx.y，blockIdx.z threadId

12、x.x，threadIdx.y，threadIdx.z4.1.2 执行配置执行配置2021-12-14并行计算如何定义grid和block？4.1.2 执行配置执行配置2021-12-14并行计算如果结合Blocks和Threads，则执行配置中参数应如何设置？4.1.2 执行配置执行配置2021-12-14并行计算 _global_ void add(int *a, int *b, int *c) int index = blockIdx.x * bloxkDim.x + threadIdx.x; cindex = aindex + bindex;4.1.2 执行配置执行配置2021-12-

13、14并行计算假定每个block有8个thread：4.1.2 执行配置执行配置2021-12-14并行计算 #define N (2048*2048) #define THREADS_PER_BLOCK 512 int main() int *a, *b, *c; int *d_a, *d_b, *d_c; int size = N * sizeof(int); /分配GPU显存空间 cudaMalloc(void*)&d_a, size); cudaMalloc(void*)&d_a, size); cudaMalloc(void*)&d_a, size); /为输入

14、数据分配内存空间 a = (int*)malloc(size); random_ints(a, N); b = (int*)malloc(size); random_ints(b, N); c = (int*)malloc(size); /将内存中的输入数据拷贝到显存 cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice); /调用并启动调用并启动kernel add(d_a,d_b,d_c); /将显存中的计算结果回读到内存 cudaMemcpy(c,

15、 d_c, size, cudaMemcpyDeviceToHost); /释放内存 free(a); free(b); free(c); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); return 0; 举例：递增数组元素举例：递增数组元素2021-12-14并行计算4.2 存储器层次和显存分配存储器层次和显存分配4.2.1 存储器层次存储器层次2021-12-14并行计算4.2.1 存储器层次存储器层次2021-12-14并行计算全局内存全局内存 通俗意义上的设备内存，位于GPU片外。共享内存共享内存 1. 位置：位置：位于GPU片内的高速缓存。

16、 2. 形式：形式：关键字_shared_添加到变量声明中。如_shared_ float cache10。 3. 目的：目的：对于GPU上启动的每个线程块，CUDA C编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存，但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程能够在计算上通信和协作。4.2.1 存储器层次存储器层次2021-12-14并行计算常量内存常量内存 位置：位置：设备内存，GPU片外 形式：形式：关键字_constant_添加到变量声明中。如_constant_ float s10;。 目的：目的：为了提升性能。常量内存采取了不

17、同于标准全局内存的处理方式。在某些情况下，用常量内存替换全局内存能有效地减少内存带宽。 特点：特点：常量内存用于保存在核函数执行期间不会发生变化的数据。变量的访问限制为只读。NVIDIA硬件提供了64KB的常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时，静态地分配空间。 要 求 ：要 求 ： 当 我 们 需 要 拷 贝 数 据 到 常 量 内 存 中 应 该 使 用cudaMemcpyToSymbol()，而cudaMemcpy()会复制到全局内存。 性能提升的原因：性能提升的原因：对常量内存的单次读操作可以广播到其他的“邻近”线程。这将节约15次读取操作。（

18、为什么是15，因为“邻近”指半个线程束，一个线程束包含32个线程的集合。）常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生额外的内存通信量。4.2.1 存储器层次存储器层次2021-12-14并行计算纹理内存纹理内存 1. 位置：位置：设备内存 2. 目的：目的：能够减少对内存的请求并提供高效的内存带宽。是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序设计，意味着一个线程读取的位置可能与邻近线程读取的位置“非常接近”。如下图： 3. 纹理变量（引用）必须声明为文件作用域内的全局变量。 4. 形式：形式：分为一维纹理内存 和 二维纹理内存。4.2.2 GPU 显存分配显

19、存分配2021-12-14并行计算 cudaMalloc() 分配GPU显存空间 函数原型： cudaMalloc (void *devPtr, size_t size)对象数组指针数组大小 函数用处：与C语言中的malloc函数一样，分配GPU显存空间注意： 可以将cudaMalloc()分配的指针传递给在设备上执行的函数； 可以在设备代码中使用cudaMalloc()分配的指针进行设备内存读写操作； 可以将cudaMalloc()分配的指针传递给在主机上执行的函数； 不可以在主机代码中使用cudaMalloc()分配的指针进行主机内存读写操作。4.2.2 GPU 显存分配显存分配2021-

20、12-14并行计算 cudaFree() 释放GPU显存空间 函数原型： cudaFree ( void* devPtr )对象数组指针 函数用处：与c语言中的free()函数一样，只是此函数释放的是cudaMalloc()分配的内存。如：int BLOCK_SIZE = 64;float* Md;int size = BLOCK_SIZE * BLOCK_SIZE * sizeof(float); cudaMalloc(void*)&Md, size);cudaFree(Md);4.2.3 Host-Device数据交换数据交换2021-12-14并行计算 cudaMemcpy()

21、主机与设备端进行数据传输 函数原型： cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind) 函数用处：与c语言中的memcpy函数一样，只是此函数可以在主机内存和GPU内存之间互相拷贝数据。 函数参数：cudaMemcpyKind kind表示数据拷贝方向。FHostToHostFHostToDeviceFDeviceToHostFDeviceToDevice 与C中的memcpy()一样，以同步方式执行，即当函数返回时，复制操作就已经完成了，并且在输出缓冲区中包含了复制进去的内容。4.3 CUDA变量和函数变量和函数4.3.1 变量类型限定符变量类型限定符变量类型限定符变量类型限定符说明变量在说明变量在GPU上的存储器位置上的存储器位置 _device_：使用该限定符声明的变量位于全局存储器中，参与并行计算的所有线程都可以访问，与整个应用程序具有相同的生存周期。 _shared_：使用该限定符声明的变量位于共享存储器中，一个线程块中的所有线程都可以访问，与线程块具有相同的生存周期。 _constant_：使用该限定符声明的变量位于常数存储器中，参与并行计