NVIDIA CUDA 高性能计算笔记（一）cuda编程简介及矩阵赋值案例

NVIDIA CUDA 高性能计算笔记（一）

​       CUDA （Compute  Unified Device Architecture）是NIVIDIA 推出的通用并行计算平台，支持C，C++，Python等语言，实现CPU和GPU协同计算。其架构采用Grid-Blocks-Threads线程层次结构和SIMT并行模式，在给出CUDA的编程实例之前，需要给出模型的基础知识做个简单的介绍。
1.1CUDA编程模型简介

​      CUDA编程模型是一个异构模型，需要GPU和CPU协同工作。在CUDA架构中，我们用host端指代CPU及其内存的，用device指代GPU及其内存。CUDA程序中即包含Host程序，又包含device程序，它们分别在CPU与GPU上运行。同时，host与device之间进行通信，这样它们之间可以进行数据拷贝。典型的CUDA程序的执行的程序的流程为：

• 分配host内存，并进行数据初始化；
  
• 分配device内存（显存、共享内存），并从host端将数据拷贝到device端；
  
• 调用CUDA的核函数在device函数上完成指定的运算；
  
• 将device上的运算结果拷贝到host上；
  
• 释放device和host上分配的内存。
  

​        由于CUDA编程模型实际上是异构编程模型，所以需要区分host和device上的代码，在CUDA中是通过函数类型限定词区别开host和device上的函数，主要的三个函数类型限定词如下：

• __global__: 在device端上执行，从host中调用（一些特定的GPU也可以从device上调用），返回类型必须为 void , 不支持可变参数，不能成为类成员函数。注意__global__ 定义的kernel是异步的，这意味着host端不会等待kernel执行完就执行下一步；
  
• __device__: 在device端上执行，但仅可以从device中调用，不可以和 __global__ 同时用；
  
• __host__: 在host上执行，仅可以从host中调用，一般省略不写，不可以和 __global__同时用，但可以和 __device__，此时函数会在device和host都编译。
  

​         上面的流程中最重要的一个过程是调用CUDA的核函数来执行并行计算，kernel是CUDA中的一个重要的概念，kernel是在device上线程中并行执行的函数，在调用时需要用   来指定kernel要执行的线程数量，在CUDA中，每个线程都要执行核函数，并且每个线程会分配一个唯一的\(thread\space ID\) ,这个\(ID\) 值可以通过核函数的内置变量 thread Idx 来获得。
​       要深刻理解\(kernel\)，必须要对\(kernel\) 的线程层次结构有一个清晰的认识。首先，\(GPU\)上很多并形化的轻量级线程。\(kernel\) 在device上执行时实际上是启动很多线程，一个\(kernel\) 所启动的所有线程称为网格\(grid\) ，同一个网格的线程共享相同的全局内存空间，grid是线程结构的第一个层次，而网格又可以分为很多线程块(block)，一个线程块里面包含很多线程，这是第二个层次。  为了编程方便，\(grid\) 和\(block\) 都是定义为 dim3 类型的变量，dim3 可以看成是包含三个无符号整数\((x,y,z)\) 成员的结构体变量，在定义时，缺失值初始化为1。因此，grid和block可以灵活地定义为1-dim，2-dim以及3-dim的结构，对于，\(knernel\)在定义调用时也必须通过执行配置 来指定kernel所使用的线程数及结构。

​    所以，为了方便编程，CUDA中使用了 dim3 类型（dim3 是基于unit3定义的矢量类型，相当于由3个 unsigned int类型组成的结构体）的内建变量 threadIdx 和 blockIdx。这样，就可以使用一维、二维或三维的索引来标识线程，构成 一维、二维或三维线程块。使得线程组织形式对各种域（向量、矩阵，或者高维张量）中数据的划分变得直观、自然。

• 对于一维的block，线程的\(threadID\)就是\(threadId.x\);
  
• 对于大小为\((Dx,Dy)\)的二维线程块block，线程的\(threadID\) 是 \((threadIdx.x+threadIdx.x\times{Dx})\);
  
• 对于大小为\((Dx,Dy,Dz)\)的三维线程块block, 线程的\(threadID\)是（\(threadIdx.x+threadIdx.y\times{Dx}+threadIdx.z\times{Dx}\times{Dy}\)）;
  

另外，线程还有内置变量gridDim，用于获取网格块各个维度的大小。
​       此外，这里简单介绍一下CUDA的内存模块，如图所示。可以看到，每个线程有自己的私有本地内存（\(Local Memory\)）, 而每个线程块有包含共享内存（\(Shared \space Memory\)）。还可以访问一些只读内存块：常用内存（\(Constant \space Memory\)）和纹理内存 （\(Texture \space Memory\)）。内存结构涉及到程序优化，这里就过多讨论。

​       还有重要一点，你需要对\(GPU\)的硬件实现有一个基本的认识。上面说到了\(kernel\)的线程组织层次，那么一个\(kernel\) 实际上会启动很多线程，这些线程是逻辑上是并行的，但是在物理层也是无法却并不一定。这其实和CPU的多线程有类似之处，多线程如果没有多核支持，在物理层也无法实现并行的。但是好在\(GPU\) 存在很多CUDA核心，充分利用CUDA核心可以充分发挥GPU的并行计算能力。GPU硬件的一个核心组件是SM，前面已经说过，SM是Streaming Multiprocessor，SM的核心组件包括的CUDA核心、共享内存、寄存器等，SM可以并发的执行上，一个线程块只能在一个SM上被调度。SM一般可以调度多个线程块，这要看SM本身的能力。那么有可能一个kernel的各个线程块被分配多个SM，所以grid只是逻辑层，而SM才是执行的物理层。SM采用的是SIMT(Single-Instruction, Multiple-Thread，单指令多线程)架构，基本的执行单元是线程束（warps)，线程束包含32个线程，这些线程同时执行相同的指令，但是每个线程都包含自己的指令地址计数器和寄存器状态，也有自己独立的执行路径。所以尽管线程束中的线程同时从同一程序地址执行，但是可能具有不同的行为，比如遇到了分支结构，一些线程可能进入这个分支，但是另外一些有可能不执行，它们只能死等，因为GPU规定线程束中所有线程在同一周期执行相同的指令，线程束分化会导致性能下降。当线程块被划分到某个SM上时，它将进一步划分为多个线程束，因为这才是SM的基本执行单元，但是一个SM同时并发的线程束数是有限的。这是因为资源限制，SM要为每个线程块分配共享内存，而也要为每个线程束中的线程分配独立的寄存器。所以SM的配置会影响其所支持的线程块和线程束并发数量。总之，就是网格和线程块只是逻辑划分，一个kernel的所有线程其实在物理层是不一定同时并发的。所以kernel的grid和block的配置不同，性能会出现差异，这点是要特别注意的。还有，由于SM的基本执行单元是包含32个线程的线程束，所以block大小一般要设置为32的倍数。
内存类型：内存作用：全局内存（Global Memory）容量最大（通常数GB），所有线程可访问，但延迟高（400-800周期）共享内存（shared Memory）片上内存，速度比全局内存快100倍，但容量有限（每SM通常16-64KB）寄存器（Registers）最快的存储，每个线程私有常量内存（Constant Memory）只读缓存，适合广播数据纹理内存（Texture Memory）专为图形处理优化的特殊缓存​     内存访问特性比较：
内存类型物理位置作用域带宽、速度使用场景显式控制关键字寄存器GPU芯片寄存器线程私有最高（1周期）高频访问的私有变量（如循环计数器）自动分配（局部变量）共享内存GPU芯片上的SM处理器线程块共享高（1-32周期）线程协作（如规约运算、矩阵分块）__share__本地内存实际在全局内存中内存线程私有中低（\(\approx\)全局内存）大数组或寄存器不足时的溢出变量编译器自动分配全局内存GPU设备显存所有线程+主机中（400~800周期）大规模数据存储，需要频繁访问时需合并访问优化cudaMalloc分配常量内存GPU芯片上的缓存所有线程只读中（缓存加速）需要广播给所有线程的至多__constant__纹理内存GPU专用缓存所有线程中 （优化访存）图形处理、具有空间局部性的非对齐访问纹理API绑定主机内存CPU内存主机+设备（需要拷贝）最低（PCLe瓶颈）CPU-GPU数据传输的中间存储malloc、cudaHostAlloc下面我将详细地介绍CUDA中各种内存管理函数的功能、参数和使用方法。
CUDA是一种用于异构并行计算的编程模型，经常需要在主机端（host）和设备端（Device）之间进行数据传输。这是因为CUDA核函数传入的必须是指向其中处理GPU显存的三个关键的API：cudaMalloc,cudaMemcpy和 cudaFree。

• cudaMalloc：
  

其接口API形式：cudaError_t  cudaMalloc(void ** devPtr,size_t size )函数功能：在设备上分配线性内存size字节，并通过指针返回分配的内存devPtr。分配的内存对应任何类型的变量。记忆没有被清除。失败时返回 cudaErrorMemoryAllocation。参数：devPtr 设备内存分配指针；size ：分配的字节数返回值：cudaSuccess , cudaErrorMemoryAllocation注意事项：
分配的内存

• cudaMemcpy ：
  其接口形式：cudaError_t cudaMemcpy(void * dist, const void * src,size_t count,CudaMemcpyKind kind)函数功能：将指向的内存区域的字节复制到指向的存储区域参数：dist-目的存储地址；src -源内存地址；count-复制内存的字节数； kind-传输类型返回值：cudaSuccess，cudaErrorInvalidValue，cudaErrorInvalidDevicePointer，cudaErrorInvalidMemcpyDirection
  
• cudaFree：
  其接口形式：cudaError_t cudaFree(void * devPtr)函数功能：释放由 指向的内存空间，该空间必须是之前调用cudaMalloc()或cudaMallocPitch()时返回过的。否则，或者如果cudaFree()之前已被调用过，则返回错误。如果 为 0，则不执行作。cudaFree() 在失败时返回cudaErrorInvalidDevicePointer。参数:devPtr -设备指针指向内存释放返回值：cudaSuccess，cudaErrorInvalidDevicePointer, cudaErrorInitialization​
  

1.2 CUDA的第一个程序—矩阵赋值(Matrix Assign)

​            在本节通过一个矩阵赋值（matrix Assign）例子开始真正的CUDA程序实现，本例是在SDK中template程序的基础上修改得到的。\(template\) 是 \(NVIDIA\) 公司提供的CUDA程序模板，也就是CUDA程序最基本的框架。要创建一个CUDA程序，可以把整个template文件复制一份。在一个CUDA程序中，基本的主机端代码主要完成以下的功能：

• 启动CUDA，使用多卡时应该时应该加上设备号，或使用\(cudaSetDevice()\)设备GPU设备；
  
• 为输入数据分配内存空间；
  
• 初始化输入数据；
  
• 为GPU分配内存，用于存放输入数据；
  
• 将内存中的输入数据拷贝到显存；
  
• 为GPU分配显存，用于存放输出数据；
  
• 调用device端的kernel进行计算，将结果写到显存中的对应区域；
  
• 为CPU分配内存，用于存放GPU传回来的输出数据；
  
• 将显存中的结果读取到内存；
  
• 释放内存和显存空间；
  
• 退出CUDA；
  

最简单的设备端代码主要完成以下功能：

• 从显存读取数据到GPU片内；
  
• 对数据进行处理；
  
• 将处理后的数据写回显存；
  其整个工程包含了三：
  （1）主程序文件CPU-Host端程序（example1main.cu）；
  

​       （2）GPU设备端函数的处理函数头文件（example_matrixassign_kernel.cuh）；
​       （3）GPU设备端函数的处理函数文件（example_matrixassign_kernel.cu）；
File1：主程序文件CPU-Host端程序（example1main.cu）；

1. #include<stdio.h> //系统头文件
2. #include<stdlib.h>
3. #include<string.h>
4. #include<math.h>
6. #include"cuda_runtime.h" //cuda项目头文件
7. #include"device_launch_parameters.h"
8. #include"example_matrixassign_kernel.cuh"  //核函数的数据的头文件
11. void runTest(int argc, char** argv);
13. int main(int argc,char** argv){
15.         runTest(argc,argv);
17. }
19. void runTest(int argc, char** argv){
21.         unsigned int num_blocks = 4;  //定义网格中的线程块数量
22.         unsigned int num_threads= 4;  //定义每个线程块中的线程数量
24.         unsigned int mem_size = sizeof(float) * num_blocks * num_threads; //为了数据分配的存储器大小，这里每一个人线程计算一个flaot
26.         //在host端分配内存，h_表示host端，i表示input，o表示output
27.         float* h_idata = nullptr;
28.         float* h_odata = nullptr;
30.         h_idata =(float *)malloc(mem_size);
31.         h_odata = (float*)malloc(mem_size);
33.         if(h_idata != nullptr) {
34.            memset(h_idata, 0, mem_size);
35.         }else{
36.                 return;
37.         }
38.         if(h_odata!=nullptr){
39.                 memset(h_odata, 0, mem_size);
40.         }else{
41.                 return;
42.         }
43.        
44.         //在device端分配显存，d_表示device端，i表示input，o表示output
45.         float* d_idata = nullptr;
46.         float* d_odata = nullptr;
48.         cudaError_t cudaStatus;  //cuda状态判断
50.         cudaStatus=cudaMalloc((void**)&d_idata, mem_size);
51.         if(cudaStatus != cudaSuccess){
52.                 printf("d_idata is cudaMalloc failed!\n");
53.                 return;
54.         }
55.         cudaStatus=cudaMalloc((void**)&d_odata, mem_size);
56.         if(cudaStatus!=cudaSuccess){
57.                 printf("d_odata is cudaMalloc failed!\n");
58.                 return;
59.         }
60.        
61.         //初始化内存中的值
62.         for(unsigned int i = 0; i < num_threads * num_blocks;i++){
63.                 h_idata[i] =1.0f;
64.         }//end for(unsigned int i = 0; i < num_threads * num_blocks;i++)
66.         //将内存中的输入数据读入设备端显存，这样就完成了主机对设备的数据写入
67.         cudaStatus=cudaMemcpy(d_idata,h_idata,mem_size,cudaMemcpyHostToDevice);
69.         //设置运行参数，即网格的形状和线程块的形状
70.         dim3 grid(num_blocks,1,1);
71.         dim3 block(num_threads,1,1);
73.         // 运行核函数，调用GPU进行运算
74.         testMatrixAssignKernel <<<grid, block>>> (d_idata,d_odata);
76.         //将结果从显存写入内存
77.         cudaStatus = cudaMemcpy(h_odata,d_odata,mem_size,cudaMemcpyDeviceToHost);
79.         //打印结果
80.         printf("赋值前的矩阵：\n");
81.         for (unsigned int iblock = 0; iblock < num_blocks; iblock++) {
82.                 for (unsigned int ithread = 0; ithread < num_threads; ithread++) {
83.                         printf("%5.0f", h_idata[iblock * num_threads + ithread]);
84.                 }//end for(unsigned int ithread = 0; ithread < num_threads; ithread++)
85.                 printf("\n");
86.         }//end for(unsigned int iblock = 0; iblock < num_blocks; iblock++)
88.         printf("赋值后的矩阵：\n");
89.         for(unsigned int iblock = 0; iblock < num_blocks; iblock++){
90.                 for(unsigned int ithread = 0; ithread < num_threads; ithread++){
91.                         printf("%5.0f",h_odata[iblock*num_threads+ithread]);
92.                 }//end for(unsigned int ithread = 0; ithread < num_threads; ithread++)
93.                 printf("\n");
94.         }//end for(unsigned int iblock = 0; iblock < num_blocks; iblock++)
96.         //输出存储器指针
97.         free(h_idata);
98.         free(h_odata);
99.         cudaFree(d_idata);
100.         cudaFree(d_odata);
101. }

复制代码

从代码中看出，CUDA的主机端代码与C语言非常相似。但也有一部分C语言中没有的语句，下面逐一进行分析。
​      （1）cudaMalloc(size)在显存global memory上分配大小为size字节的线性空间。需要注意的是，与malloc和free一样，cudaMalloc() 也必须与cudaFree()成对使用，否则无法释放显存空间，运行几次程序以后显卡上就没有显存可供分配，程序也就无法正常运行了。另外，为了杜绝指针指费的情况现象，最好在程序结束前将指针赋空并摧毁。
​      （2） cudaMemcpy()用于拷贝存储器中的数据，其中第二参数是指向目标的指针，第二个参数是指向源的指针，第三个参数是需要拷贝的字节数，第四个参数是拷贝操作的类型。拷贝操作类型共有三种：

• cudaMemcpyDeviceToHost 将显存中的数据拷贝内存中；
  
• cudaMemcpyHostToDevice 将内存中的数据拷贝到显存中；
  
• cudaMemcpyDeviceToDevice将global memory中的数据拷贝到同一个CUDA上下文的global的另一个区域中；
  

​      （3）运算符对kernel函数完整的执行参数配置形式是，其中各个参数的含义是：

• 参数Dg用于定义整个grid的维度和尺寸，为dim3类型，但实际上只有前两维可以不为1。Dim3 Dg(Dg.x,Dg.y,1)中每行有Dg.x个block，每列有Dg.y个block的维度，第三维恒为1。
  
• 参数Db为dim3类型，用于定义每个block的维度与尺寸。Dim3 Db(Db.x,Db.y,Db.z) 中每行有Db.x个thread，每列Db.y个thread，高为Db.z，可以定义三维尺寸。整个block中共有Db.x*Db.y*Db.z 个线程；
  
• 参数Ns是一个可选参数，用于设置每个block的共享内存shared memory以外，最多能够动态分配的shared memory大小，单位为Byte。
  
• 参数\(s\)是一个cudaStream_t类型的可选参数，初始值为0。在本案例中没有用到Stream的相关内容因此这个参数不填，默认为0号流。
  

File2：主程序文件CPU-Host端程序（example1main.cu）；
​

1. #pragma once
2. #ifndef EXAMPLE_MATRIXASSIGN_KERNEL_H
3. #define EXAMPLE_MATRIXASSIGN_KERNEL_H
5. #include<stdio.h>
6. #include"cuda_runtime.h"
8. __global__ void testMatrixAssignKernel(float* data_input, float* data_output);
11. #endif // !_EXAMPLE_MATRIXASSIGN_KERNEL_H_

复制代码

File2：主程序文件CPU-Host端程序（example1main.cu）；

1. __global__ void testMatrixAssignKernel(float *data_input,float *data_output){
3.         //shared memory,extern表示大小由host端的Ns参数确定
4.         extern __shared__ float sdata[];
6.         const unsigned int bid = blockIdx.x; //线程所在的block的索引号
7.         const unsigned int tid_in_block = threadIdx.x; //线程在block中的位置
8.         const unsigned int tid_in_grid = blockDim.x * blockIdx.x + threadIdx.x;
10.         //按行划分任务时，线程在整个grid中的位置
12.   // 将数据从global memory读入shared memory
13.         sdata[tid_in_block] = data_input[tid_in_grid];
14.         //读入数据后进行一次同步，保证计算时所有数据均已到位
15.         __syncthreads();
17.         // 计算
18.         sdata[tid_in_block] = (float)tid_in_grid;
19.         //  sdata[tid_in_block] *= (float)tid_in_block;
20.         //  sdata[tid_in_block] *= (float)tid_in_grid;
22.           //进行同步，确保要写入的数据已经被更新
23.         __syncthreads();
25.         // 将shared memory中的数据写到global memory
26.         data_output[tid_in_grid] = sdata[tid_in_block];
29. }

复制代码

由上可知，最简单的__gloabal__程序由以下的过程组成：

• 分配\(shared \space memory\)；
  
• 将\(global\space memory\) 中的数据读入\(shared \space memory\);
  
• 将进行计算，将结果写到\(shared \space memory\);
  
• 将\(shared\)中的结果写到\(global \space memory\) ;
  

​         进行一次GPU计算，要在多种存储器进行几次数据传输，要消耗相当多的时间。这导致了较大的延迟，这导致使\(GPU\) 不适合处理一些实时性要求很高的应用。不同存储器间的数据传输速率和使用方法有很大差异，开发人员需要根据硬件的特点来设计算法，以优化存储器访问。在理想情况下，在所有的存储器传输进行的同时，GPU的各个核心也始终在进行计算，这样就能够很好的隐藏各种访问延迟。CUDA 并不是一种完全硬件透明的语言，程序员需要根据硬件特征将任务进行合理的分解，在编程时对数据传输和寄存器访问进行优化。
​       __global__前缀表示这一段代码是cuda GPU端内核函数。内核函数运行在设备上，其返回类型必须为void。__global__函数中是每一个线程要执行的语句，但由于\(shared\space memory\)和同步的存在，在最好将__global__函数理解为对每一个block的行为的描述。
​        在这一端内核函数中，首先定义了\(shared \space memory\) 中的变量；然后根据内建变量定义每一个block和thread的索引，对任务进行划分；最后，每一个线程执行了相同的求和运算，但处理数据不同，由线程的索引决定的。程序员在编写__global__函数之前，要先对任务进行划分，设计各个block的工作流程后，做到成竹在胸。
​        由于CUDA采用了两层并行，因此本例在划分任务时，每个thread在grid中的索引\(tid\_in\_grid\) 是由thread所在block内编号tid计算得来的。计算出每个线程的索引后，就可以根据索引处理线程中不同的数据，请读者好好体会这一点。
​      extern __shared__ float sdata[] 在shared  memory中为数组data动态分配了空间。extern 在设备端和主机端有不同的含义：__device__和__global__ 函数中表示动态分配，而在主机端函数中表示外部变量。如果要静态分配一块 shared memory，那么在__shared__之前就不加extern，还必须在[]中写上要分配的字节数。动态分配的shared memory大小，是的执行参数中第三个参数规定的大小。关于shared memory大小。
​        CUDA定义了一些内建变量如下：

• gridDim： 网格的维度的变量，dim3类型
  
• blockIdx:    块的索引变量，unit3类型
  
• blockDim：块的维度变量，dim3类型
  
• threadIdx：块内的线程索引变量，unit3类型
  
• warpSize：线程中的warp大小，int类型
  

其输出结果：

​

来源：程序园用户自行投稿发布，如果侵权，请联系站长删除
免责声明：如果侵犯了您的权益，请联系站长，我们会及时删除侵权内容，谢谢合作！