第一个cuda程序详解

80 篇文章 3 订阅

CUDA是一个基于NVIDIA GPU的并行计算平台和编程模型,通过调用CUDA提供的API,可以开发高性能的并行程序。CUDA安装好之后,会自动配置好VS编译环境,按照UCDA模板新建一个工程“Hello CUDA”:



建好之后,发现该工程下已经存在一个项目 kernel.cu。这个是CUDA编程的入门示例,实现的功能是两个整型数组相加,代码如下:

[cpp]  view plain  copy
 print ?
  1. #include "cuda_runtime.h"  
  2. #include "device_launch_parameters.h"  
  3.   
  4. #include <stdio.h>  
  5.   
  6. cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);  
  7.   
  8. __global__ void addKernel(int *c, const int *a, const int *b)  
  9. {  
  10.     int i = threadIdx.x;  
  11.     c[i] = a[i] + b[i];  
  12. }  
  13.   
  14. int main()  
  15. {  
  16.     const int arraySize = 5;  
  17.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
  18.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
  19.     int c[arraySize] = { 0 };  
  20.   
  21.     // Add vectors in parallel.  
  22.     cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);  
  23.     if (cudaStatus != cudaSuccess) {  
  24.         fprintf(stderr, "addWithCuda failed!");  
  25.         return 1;  
  26.     }  
  27.   
  28.     printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",  
  29.         c[0], c[1], c[2], c[3], c[4]);  
  30.   
  31.     // cudaDeviceReset must be called before exiting in order for profiling and  
  32.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
  33.     cudaStatus = cudaDeviceReset();  
  34.     if (cudaStatus != cudaSuccess) {  
  35.         fprintf(stderr, "cudaDeviceReset failed!");  
  36.         return 1;  
  37.     }  
  38.   
  39.     return 0;  
  40. }  
  41.   
  42. // Helper function for using CUDA to add vectors in parallel.  
  43. cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)  
  44. {  
  45.     int *dev_a = 0;  
  46.     int *dev_b = 0;  
  47.     int *dev_c = 0;  
  48.     cudaError_t cudaStatus;  
  49.   
  50.     // Choose which GPU to run on, change this on a multi-GPU system.  
  51.     cudaStatus = cudaSetDevice(0);  
  52.     if (cudaStatus != cudaSuccess) {  
  53.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
  54.         goto Error;  
  55.     }  
  56.   
  57.     // Allocate GPU buffers for three vectors (two input, one output)    .  
  58.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
  59.     if (cudaStatus != cudaSuccess) {  
  60.         fprintf(stderr, "cudaMalloc failed!");  
  61.         goto Error;  
  62.     }  
  63.   
  64.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
  65.     if (cudaStatus != cudaSuccess) {  
  66.         fprintf(stderr, "cudaMalloc failed!");  
  67.         goto Error;  
  68.     }  
  69.   
  70.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
  71.     if (cudaStatus != cudaSuccess) {  
  72.         fprintf(stderr, "cudaMalloc failed!");  
  73.         goto Error;  
  74.     }  
  75.   
  76.     // Copy input vectors from host memory to GPU buffers.  
  77.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
  78.     if (cudaStatus != cudaSuccess) {  
  79.         fprintf(stderr, "cudaMemcpy failed!");  
  80.         goto Error;  
  81.     }  
  82.   
  83.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
  84.     if (cudaStatus != cudaSuccess) {  
  85.         fprintf(stderr, "cudaMemcpy failed!");  
  86.         goto Error;  
  87.     }  
  88.   
  89.     // Launch a kernel on the GPU with one thread for each element.  
  90.     addKernel << <1, size >> > (dev_c, dev_a, dev_b);  
  91.   
  92.     // Check for any errors launching the kernel  
  93.     cudaStatus = cudaGetLastError();  
  94.     if (cudaStatus != cudaSuccess) {  
  95.         fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));  
  96.         goto Error;  
  97.     }  
  98.   
  99.     // cudaDeviceSynchronize waits for the kernel to finish, and returns  
  100.     // any errors encountered during the launch.  
  101.     cudaStatus = cudaDeviceSynchronize();  
  102.     if (cudaStatus != cudaSuccess) {  
  103.         fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
  104.         goto Error;  
  105.     }  
  106.   
  107.     // Copy output vector from GPU buffer to host memory.  
  108.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
  109.     if (cudaStatus != cudaSuccess) {  
  110.         fprintf(stderr, "cudaMemcpy failed!");  
  111.         goto Error;  
  112.     }  
  113.   
  114. Error:  
  115.     cudaFree(dev_c);  
  116.     cudaFree(dev_a);  
  117.     cudaFree(dev_b);  
  118.   
  119.     return cudaStatus;  
  120. }  

程序首先定义了一个函数addWithCuda,它是调用GPU运算的入口函数,返回类型是cudaError_t。

cudaError_t是一个枚举类型,可以作为几乎所有CUDA函数的返回类型,用来检测函数执行期间发生的不同类型的错误,一共有80多个错误类型,可以在driver_types.h头文件中查看每一个整型对应的错误类型,如果返回0,代表执行成功。


第二个函数addKernel在最前有一个修饰符“__global__”,这个修饰符告诉编译器,被修饰的函数应该编译为在GPU而不是在CPU上运行,所以这个函数将被交给编译设备代码的编译器——NVCC编译器来处理,其他普通的函数或语句将交给主机编译器处理。


这里“设备”的概念可以理解为GPU和其显存组成的运算单元,“主机”可以理解为CPU和系统内存组成的运算单元。在GPU上执行的函数称为核函数。


addKernel函数定义:

[cpp]  view plain  copy
 print ?
  1. __global__ void addKernel(int *c, const int *a, const int *b)  
  2. {  
  3.     int i = threadIdx.x;  
  4.     c[i] = a[i] + b[i];  
  5. }  

这个核函数里有一个陌生的threadIdx.x,表示的是thread在x方向上的索引号,理解这个之前得先了解一下GPU线程的层次结构:



CUDA中的线程(thread)是设备中并行运算结构中的最小单位,类似于主机中的线程的概念,thread可以以一维、二维、三维的形式组织在一起,threadIdx.x表示的是thread在x方向的索引号,还可能存在thread在y和z方向的索引号threadIdx.y和threadIdx.z。


一维、二维或三维的thread组成一个线程块(Block),一维、二维或三维的线程块(Block)组合成一个线程块网格(Grid),线程块网格(Grid)可以是一维或二维的。通过网格块(Grid)->线程块(Block)->线程(thread)的 顺序可以定位到每一个并且唯一的线程


回到程序中的addKernel函数上来,这个函数会被GPU上的多个线程同时执行一次,线程间彼此没有通信,相互独立。到底会有多少个线程来分别执行核函数,是在“<<<>>>”符号里定义的。“<<<>>>”表示运行时配置符号,在本程序中的定义是<<<1,size>>>,表示分配了一个线程块(Block),每个线程块有分配了size个线程“<<<>>>”中的 参数并不是传递给设备代码的参数,而是定义主机代码运行时如何启动设备代码。以上定义的这些线程都是一个维度上的,可以通过thredaIdx.x来获取执行当前计算任务的线程的ID号。


cudaSetDevice函数用来设置要在哪个GPU上执行,如果只有一个GPU,设置为cudaSetDevice(0);


cudaMalloc函数用来为参与运算的数据分配显存空间,函数原型:cudaError_t cudaMalloc(void **p, size_t s);


cudaMemcpy函数用于主机内存和设备显存以及主机与主机之间,设备与设备之间相互拷贝数据,函数原型:

[cpp]  view plain  copy
 print ?
  1. cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind);  


第一个参数dst是目标数据地址,第二个参数src是源数据地址,第三个参数count是数据大小,第四个参数kind定义数据拷贝的类型,有如下几类枚举类型:

[cpp]  view plain  copy
 print ?
  1. /** 
  2.  * CUDA memory copy types 
  3.  */  
  4. enum __device_builtin__ cudaMemcpyKind  
  5. {  
  6.     cudaMemcpyHostToHost          =   0,      /**< Host   -> Host */  
  7.     cudaMemcpyHostToDevice        =   1,      /**< Host   -> Device */  
  8.     cudaMemcpyDeviceToHost        =   2,      /**< Device -> Host */  
  9.     cudaMemcpyDeviceToDevice      =   3,      /**< Device -> Device */  
  10.     cudaMemcpyDefault             =   4       /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */  
  11. };  

接下来在调用核函数时候添加了运行时配置符号“<<<>>>”,定义线程块和线程的数量,如<<<1,5>>>表示定义了一个线程块,每个线程块包含了5个线程。

cudaGetLastError函数用于返回最新的一个运行时调用错误,对于任何CUDA错误,都可以通过函数cudaGetErrorString函数来获取错误的详细信息。


cudaDeviceSynchronize函数提供了一个阻塞,用于等待所有的线程都执行完各自的计算任务,然后继续往下执行。


cudaFree函数用于释放申请的显存空间。


cudaDeviceReset函数用于释放所有申请的显存空间和重置设备状态;


第一个CUDA程序kernel.cu涉及的内容主要就是这些。CUDA的使用步骤如下:

  1. 主机代码执行
  2. 传输数据给GPU
  3. 确定Grid、Block大小
  4. 调用内核函数,GPU多线程运行程序
  5. 传输运算结果给CPU
  6. 继续主机代码执行

期间涉及到在设备上的一些显存空间申请、销毁等操作,从内存到显存上数据的相互拷贝是一个比较耗时的过程,应该尽量减少这种操作。

参与评论 您还未登录,请先 登录 后发表或查看评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
©️2022 CSDN 皮肤主题:大白 设计师:CSDN官方博客 返回首页
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、C币套餐、付费专栏及课程。

余额充值