CUDA C++ 高效入门第二章 -- CUDA 线程组织之两个数组相加

CUDA C 高效入门第二章 -- CUDA 线程组织之两个数组相加1 资料2 正文2.1 C 版本的两个数组相加2.2 CUDA 版本的两个数组相加2.3 为 CUDA 版本的两个数组相加添加运行时错误检查3 总结1 资料上一篇文章 CUDA C 高效入门第二章 – CUDA 线程组织之 Hello World我们用了五个 Hello World 样例由浅入深地为大家介绍了 CUDA 的线程组织。本文我们将基于 CPU 和 GPU 分别实现两个数组相加体会 CUDA 编程与 C 编程的区别。本文参考资料如下1樊哲勇CUDA编程 基础与实践 第二三四五章2cuda-programming-guide 1.2 编程模型3cuda-programming-guide 2.1 CUDA C 入门4cuda-programming-guide 2.2 编写 CUDA SIMT 内核5cuda-programming-guide 2.5 NVCC 编译器本系列博客汇总链接CUDA C 高效入门系列。2 正文2.1 C 版本的两个数组相加1下面是只使用 C 实现两个数组相加02_array_add/add.cpp基本逻辑是数组 a 与 数组 b 对应元素相加复制给数组 c 的对应元素并检查结果是否正确。#includecmath#includecstdlib#includecstdioconstdoubleEPSILON1e-15;constdoublea1.23;constdoubleb2.34;constdoublec3.57;voidadd(constdouble*px,constdouble*py,double*pz,constintN){for(inti0;iN;i){pz[i]px[i]py[i];}}voidcheck(constdouble*pz,constintN){boolhas_errorfalse;for(inti0;iN;i){if(fabs(pz[i]-c)EPSILON){has_errortrue;}}printf(%s\n,has_error?Has errors:No errors);}intmain(void){constintN1e7;constintMsizeof(double)*N;double*px(double*)malloc(M);double*py(double*)malloc(M);double*pz(double*)malloc(M);for(inti0;iN;i){px[i]a;py[i]b;}add(px,py,pz,N);check(pz,N);free(px);free(py);free(pz);return0;}编译并运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 02_array_add/add.cpp -------------------- g-stdc17-Wall-Wextra-O2-o/tmp/tmp.HHABvx7Hb5/run_demo_exec 02_array_add/add.cpp -------------------- No errors2.2 CUDA 版本的两个数组相加1使用 C 实现两个数组相加对于 C 程序员来说实在太简单了。但如果用 CUDA 实现对于新手而言就不一定那么容易了。因此我们需要为大家引入简单 CUDA 程序的基本框架请重点关注 main 函数中的逻辑。头文件包含 常量定义或者宏定义 C自定义函数和 CUDA 核函数的定义实现intmain(void){分配主机与设备内存 初始化主机中的数据 将某些数据从主机复制到设备 调用核函数在设备中进行计算 将某些数据从设备复制到主机 释放主机与设备内存}2下面是使用 CUDA 实现的两个数组相加02_array_add/add1.cu里面用到了若干 CUDA 运行时 API细节直接读注释即可。#includecmath#includecstdlib#includecstdioconstdoubleEPSILON1e-15;constdoublea1.23;constdoubleb2.34;constdoublec3.57;// CUDA 设备函数由 __device__ 修饰符修饰表示这是一个在 GPU 上运行的函数// CUDA 设备函数由核函数或其他设备函数调用不能由 CPU 端调用// CUDA 设备函数可以有返回值// CUDA 设备函数可以同时由 __host__ 和 __device__ 修饰符修饰表示这个函数既可以在 CPU 上运行也可以在 GPU 上运行减少冗余代码// CUDA 设备函数不能同时由 __global__ 和 __device__ 修饰符修饰// CUDA 设备函数可以使用 __noinline__ 修饰符修饰表示这个函数不进行内联展开编译器不一定接受减少编译器优化带来的问题// CUDA 设备函数可以使用 __forceinline__ 修饰符修饰表示这个函数强制进行内联展开减少函数调用开销__device__doubleadd1_device(constdoublex,constdoubley){returnxy;}__global__voidadd1(constdouble*px,constdouble*py,double*pz,constintN){constinttidblockDim.x*blockIdx.xthreadIdx.x;if(tidN){return;}pz[tid]add1_device(px[tid],py[tid]);}__device__voidadd2_device(constdoublex,constdoubley,doublerz){rzxy;}__global__voidadd2(constdouble*px,constdouble*py,double*pz,constintN){constinttidblockDim.x*blockIdx.xthreadIdx.x;if(tidN){return;}add2_device(px[tid],py[tid],pz[tid]);}// CUDA 核函数的返回值必须是 void, 不能有返回值但核函数里面可以使用 return 语句提前结束函数的执行return 后面不能跟任何值。// CUDA 核函数必须使用 __global__ 修饰符修饰表示这是一个在 GPU 上运行的函数// CUDA 核函数不能同时由 __host__ 和 __global__ 修饰符修饰// CUDA 核函数支持 C 的重载overload// CUDA 核函数的参数个数必须确定不支持可变参数// CUDA 核函数支持传递非指针变量内容对每个线程可见// CUDA 核函数不能是类的成员函数。// CUDA 核函数由 CPU 端调用由 GPU 端执行CC 3.5 之前核函数不能调用核函数// CUDA 核函数被调用时必须使用 ... 语法指定 grid 和 block 的尺寸__global__voidadd(constdouble*px,constdouble*py,double*pz,constintN){// 这个公式得到是线程在整个 Grid 中的 x 维度的坐标// 但由于本样例中只使用了一维网格和一维线程块因此 x 维度的坐标就是线程在整个 Grid 中的一维全局索引。// 这个一维全局索引值与数组长度一致因此可以直接用这个索引值访问数组元素// 从而实现每个线程处理一个数组元素的计算所有数组元素并发执行计算。constinttidblockDim.x*blockIdx.xthreadIdx.x;// printf(tid %d\n, tid);if(tidN){return;}pz[tid]px[tid]py[tid];}voidcheck(constdouble*pz,constintN){boolhas_errorfalse;for(inti0;iN;i){if(fabs(pz[i]-c)EPSILON){has_errortrue;}}printf(%s\n,has_error?Has errors:No errors);}intmain(void){// 踩坑记录对于 GeForce MX450nvidia-smi 显示显存大小为 1878MiB测试发现如果设为 1e8则超过显存大小程序失败1e7 可以。constintN1e7;constintMsizeof(double)*N;double*h_px(double*)malloc(M);double*h_py(double*)malloc(M);double*h_pz(double*)malloc(M);for(inti0;iN;i){h_px[i]a;h_py[i]b;}// CUDA 编程的传统是将 CPU 端的变量命名为 h_ 开头表示 host 端的变量// 将 GPU 端的变量命名为 d_ 开头表示 device 端的变量。double*d_px,*d_py,*d_pz;// 所有的 CUDA 运行时 API 都是以 cuda 开头// cudaMalloc() 函数在 GPU 显存上分配内存API 接口// cudaError_t cudaMalloc(void **address, size_t size);// address 是待分配内存的指针是一个二重指针size 是待分配内存的大小单位是字节。// cudaError_t 是一个枚举类型表示 CUDA 运行时 API 的返回值cudaSuccess 表示成功其他值表示失败。cudaMalloc((void**)d_px,M);cudaMalloc((void**)d_py,M);cudaMalloc((void**)d_pz,M);// cudaMemcpy() 函数的作用是拷贝数据API 接口// cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind);// dst 是目标地址src 是源地址count 是拷贝的字节数// kind 是拷贝的方向cudaMemcpyHostToDevice 表示从 CPU 端拷贝到 GPU 端cudaMemcpyDeviceToHost 表示从 GPU 端拷贝到 CPU 端cudaMemcpyDeviceToDevice 表示在 GPU 端拷贝。// cudaMemcpyDefault 表示由 CUDA 运行时 API 自动选择拷贝的方向通常情况下CUDA 运行时 API 会根据 dst 和 src 的地址判断拷贝的方向。// cudaMemcpyHostToHost 表示从 CPU 端拷贝到 CPU 端。cudaMemcpy(d_px,h_px,M,cudaMemcpyHostToDevice);cudaMemcpy(d_py,h_py,M,cudaMemcpyHostToDevice);constintblock_size128;constintgrid_size(N%block_size0)?(N/block_size):(N/block_size1);printf(grid_size %d, block_size %d\n,grid_size,block_size);addgrid_size,block_size(d_px,d_py,d_pz,N);// add1grid_size, block_size(d_px, d_py, d_pz, N);// add2grid_size, block_size(d_px, d_py, d_pz, N);cudaMemcpy(h_pz,d_pz,M,cudaMemcpyDeviceToHost);check(h_pz,N);free(h_px);free(h_py);free(h_pz);cudaFree(d_px);cudaFree(d_py);cudaFree(d_pz);return0;}编译运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 02_array_add/add1.cu -------------------- nvcc-archsm_75-o/tmp/tmp.ePKavUQjc6/run_demo_exec 02_array_add/add1.cu -------------------- grid_size78125, block_size128No errors3C 版本 add 函数与 CUDA 版本的核心区别重中之重C 版本的 add 函数就是一个循环通过逐个遍历数组元素完成求和以及赋值操作。这个循环就是前面文章所说的 CPU 顺序执行的思维方式见 CUDA C 高效入门第一章 – CUDA 导学之为什么要学习 GPU CUDA 。到了 CUDA 编程时代默认的思维方式就是并行就是想办法分解任务然后使用足够多的线程去执行这是所有 CUDA 程序设计的出发点。非如此不足以发挥 CUDA 成千上万个处理核的威力。针对两个数组相加这个例子使用 CUDA 编程下意识的就要想到数组有多长就用多少个 CUDA 线程一个线程只处理一个元素即可。因此CUDA 版本的 add 核函数就需要先获得线程的一维索引然后使用这个索引去访问数组这样就能做到一个线程处理一个元素。至于求和以及赋值则与 C 版本是一样的。总结一下看似长得差不多的两个 add 函数其思维方式是有本质区别的而这个区别又来自于 CPU 和 GPU 设计原理的的本质区别。voidadd(constdouble*px,constdouble*py,double*pz,constintN){for(inti0;iN;i){pz[i]px[i]py[i];}}__global__voidadd(constdouble*px,constdouble*py,double*pz,constintN){constinttidblockDim.x*blockIdx.xthreadIdx.x;// printf(tid %d\n, tid);if(tidN){return;}pz[tid]px[tid]py[tid];}2.3 为 CUDA 版本的两个数组相加添加运行时错误检查1上一节使用到的几个 CUDA 运行时 API 都有 cudaError_t 类似的返回值因此我们可以写一个宏在调用这些 API 时对返回值进行检查。我们将这个宏写在一个头文件中03_check_api/cuda_error.cuh#ifndef__ERROR_CUH__#define__ERROR_CUH__#includecstdio#defineCHECK_CUDA_CALL(call)\do{\constcudaError_t error_codecall;\if(error_code!cudaSuccess){\printf(CUDA Error: \n);\printf( File: %s\n,__FILE__);\printf( Line: %d\n,__LINE__);\printf( Error code: %d\n,error_code);\printf( Error string: %s\n,cudaGetErrorString(error_code));\exit(1);\}\}while(0)#endif// __ERROR_CUH__2下面是对 CUDA 版本的两个数组相加添加运行时错误检查细节看注释就行03_check_api/check1.cu#includecmath#includecstdio#includecuda_error.cuhconstdoubleEPSILON1e-15;constdoublea1.23;constdoubleb2.34;constdoublec3.57;__global__voidadd(constdouble*px,constdouble*py,double*pz,constintN){constinttidblockDim.x*blockIdx.xthreadIdx.x;// printf(tid %d\n, tid);if(tidN){return;}pz[tid]px[tid]py[tid];}voidcheck(constdouble*pz,constintN){boolhas_errorfalse;for(inti0;iN;i){if(fabs(pz[i]-c)EPSILON){has_errortrue;}}printf(%s\n,has_error?Has errors:No errors);}intmain(void){constintN1e7;constintMsizeof(double)*N;double*h_px(double*)malloc(M);double*h_py(double*)malloc(M);double*h_pz(double*)malloc(M);for(inti0;iN;i){h_px[i]a;h_py[i]b;}double*d_px,*d_py,*d_pz;CHECK_CUDA_CALL(cudaMalloc((void**)d_px,M));CHECK_CUDA_CALL(cudaMalloc((void**)d_py,M));CHECK_CUDA_CALL(cudaMalloc((void**)d_pz,M));// 这里故意将 cudaMemcpyHostToDevice 错写成 cudaMemcpyDeviceToHost测试 CHECK_CUDA_CALL 宏的功能// CHECK_CUDA_CALL(cudaMemcpy(d_px, h_px, M, cudaMemcpyDeviceToHost));CHECK_CUDA_CALL(cudaMemcpy(d_px,h_px,M,cudaMemcpyHostToDevice));CHECK_CUDA_CALL(cudaMemcpy(d_py,h_py,M,cudaMemcpyHostToDevice));// 由于 block 最大为 1024, 这里故意写成 1280, 测试核函数调用失败的情况// const int block_size 1280;constintblock_size128;constintgrid_size(N%block_size0)?(N/block_size):(N/block_size1);printf(grid_size %d, block_size %d\n,grid_size,block_size);addgrid_size,block_size(d_px,d_py,d_pz,N);// 这两句是检查核函数调用的错误cudaGetLastError() 的作用是在主机与设备同步之前捕捉最后一个错误// 通常情况下主机异步调用核函数调用后立即执行下一条语句// 因此捕捉核函数调用的错误需要在主机与设备同步之前调用 cudaGetLastError() 函数。// 除了显示调用 cudaDeviceSynchronize() 实现主机与设备同步还可以临时设置环境变量 CUDA_LAUNCH_BLOCKING1强制主机与设备同步。// 设置之后所有核函数的调用都变为同步的主机在调用核函数之后必须等待设备执行完核函数之后才能继续执行下一条语句。// CHECK_CUDA_CALL(cudaGetLastError());// CHECK_CUDA_CALL(cudaDeviceSynchronize());CHECK_CUDA_CALL(cudaMemcpy(h_pz,d_pz,M,cudaMemcpyDeviceToHost));check(h_pz,N);free(h_px);free(h_py);free(h_pz);CHECK_CUDA_CALL(cudaFree(d_px));CHECK_CUDA_CALL(cudaFree(d_py));CHECK_CUDA_CALL(cudaFree(d_pz));return0;}编译和运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 03_check_api/check1.cu -------------------- nvcc-archsm_75-o/tmp/tmp.ZitlLOKBof/run_demo_exec 03_check_api/check1.cu -------------------- grid_size78125, block_size128No errors3 总结本文所有代码都托管在本人的 github 上cuda_junior。下一篇我们仍然基于两个数组相加这个例子对 CUDA 的加速效果进行度量