1 资料

上一篇文章 CUDA C++ 高效入门第二章 – CUDA 线程组织之 Hello World，我们用了五个 Hello World 样例，由浅入深地为大家介绍了 CUDA 的线程组织。本文，我们将基于 CPU 和 GPU 分别实现两个数组相加，体会 CUDA 编程与 C++ 编程的区别。
本文参考资料如下：
（1）樊哲勇：CUDA编程 基础与实践 第二，三，四，五章
（2）cuda-programming-guide 1.2 编程模型
（3）cuda-programming-guide 2.1 CUDA C++ 入门
（4）cuda-programming-guide 2.2 编写 CUDA SIMT 内核
（5）cuda-programming-guide 2.5 NVCC 编译器
本系列博客汇总链接：CUDA C++ 高效入门系列。

2 正文

2.1 C++ 版本的两个数组相加

（1）下面是只使用 C++ 实现两个数组相加：02_array_add/add.cpp，基本逻辑是数组 a 与 数组 b 对应元素相加，复制给数组 c 的对应元素，并检查结果是否正确。

#include<cmath>#include<cstdlib>#include<cstdio>constdoubleEPSILON=1e-15;constdoublea=1.23;constdoubleb=2.34;constdoublec=3.57;voidadd(constdouble*px,constdouble*py,double*pz,constintN){for(inti=0;i<N;++i){pz[i]=px[i]+py[i];}}voidcheck(constdouble*pz,constintN){boolhas_error=false;for(inti=0;i<N;++i){if(fabs(pz[i]-c)>EPSILON){has_error=true;}}printf("%s\n",has_error?"Has errors":"No errors");}intmain(void){constintN=1e7;constintM=sizeof(double)*N;double*px=(double*)malloc(M);double*py=(double*)malloc(M);double*pz=(double*)malloc(M);for(inti=0;i<N;++i){px[i]=a;py[i]=b;}add(px,py,pz,N);check(pz,N);free(px);free(py);free(pz);return0;}

编译并运行：

ycao@Thinkpad-T14:~/cuda_junior$ ./run_demo.sh 02_array_add/add.cpp -------------------- g++-std=c++17-Wall-Wextra-O2-o/tmp/tmp.HHABvx7Hb5/run_demo_exec 02_array_add/add.cpp -------------------- No errors

2.2 CUDA 版本的两个数组相加

（1）使用 C++ 实现两个数组相加，对于 C++ 程序员来说，实在太简单了。但如果用 CUDA 实现，对于新手而言，就不一定那么容易了。因此，我们需要为大家引入简单 CUDA 程序的基本框架，请重点关注 main 函数中的逻辑。

头文件包含 常量定义（或者宏定义） C++自定义函数和 CUDA 核函数的定义（实现）intmain(void){分配主机与设备内存 初始化主机中的数据 将某些数据从主机复制到设备 调用核函数在设备中进行计算 将某些数据从设备复制到主机 释放主机与设备内存}

（2）下面是使用 CUDA 实现的两个数组相加：02_array_add/add1.cu，里面用到了若干 CUDA 运行时 API，细节直接读注释即可。

#include<cmath>#include<cstdlib>#include<cstdio>constdoubleEPSILON=1e-15;constdoublea=1.23;constdoubleb=2.34;constdoublec=3.57;// CUDA 设备函数由 __device__ 修饰符修饰，表示这是一个在 GPU 上运行的函数，// CUDA 设备函数由核函数或其他设备函数调用，不能由 CPU 端调用// CUDA 设备函数可以有返回值// CUDA 设备函数可以同时由 __host__ 和 __device__ 修饰符修饰，表示这个函数既可以在 CPU 上运行，也可以在 GPU 上运行，减少冗余代码// CUDA 设备函数不能同时由 __global__ 和 __device__ 修饰符修饰// CUDA 设备函数可以使用 __noinline__ 修饰符修饰，表示这个函数不进行内联展开（编译器不一定接受），减少编译器优化带来的问题// CUDA 设备函数可以使用 __forceinline__ 修饰符修饰，表示这个函数强制进行内联展开，减少函数调用开销__device__doubleadd1_device(constdoublex,constdoubley){returnx+y;}__global__voidadd1(constdouble*px,constdouble*py,double*pz,constintN){constinttid=blockDim.x*blockIdx.x+threadIdx.x;if(tid>=N){return;}pz[tid]=add1_device(px[tid],py[tid]);}__device__voidadd2_device(constdoublex,constdoubley,double&rz){rz=x+y;}__global__voidadd2(constdouble*px,constdouble*py,double*pz,constintN){constinttid=blockDim.x*blockIdx.x+threadIdx.x;if(tid>=N){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 中的一维全局索引。// 这个一维全局索引值与数组长度一致，因此可以直接用这个索引值访问数组元素，// 从而实现每个线程处理一个数组元素的计算，所有数组元素并发执行计算。constinttid=blockDim.x*blockIdx.x+threadIdx.x;// printf("tid = %d\n", tid);if(tid>=N){return;}pz[tid]=px[tid]+py[tid];}voidcheck(constdouble*pz,constintN){boolhas_error=false;for(inti=0;i<N;++i){if(fabs(pz[i]-c)>EPSILON){has_error=true;}}printf("%s\n",has_error?"Has errors":"No errors");}intmain(void){// 踩坑记录：对于 GeForce MX450，nvidia-smi 显示显存大小为 1878MiB，测试发现，如果设为 1e8，则超过显存大小，程序失败，1e7 可以。constintN=1e7;constintM=sizeof(double)*N;double*h_px=(double*)malloc(M);double*h_py=(double*)malloc(M);double*h_pz=(double*)malloc(M);for(inti=0;i<N;++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_size=128;constintgrid_size=(N%block_size==0)?(N/block_size):(N/block_size+1);printf("grid_size = %d, block_size = %d\n",grid_size,block_size);add<<<grid_size,block_size>>>(d_px,d_py,d_pz,N);// add1<<<grid_size, block_size>>>(d_px, d_py, d_pz, N);// add2<<<grid_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;}

编译运行：

ycao@Thinkpad-T14:~/cuda_junior$ ./run_demo.sh 02_array_add/add1.cu -------------------- nvcc-arch=sm_75-o/tmp/tmp.ePKavUQjc6/run_demo_exec 02_array_add/add1.cu -------------------- grid_size=78125, block_size=128No errors

（3）C++ 版本 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(inti=0;i<N;++i){pz[i]=px[i]+py[i];}}__global__voidadd(constdouble*px,constdouble*py,double*pz,constintN){constinttid=blockDim.x*blockIdx.x+threadIdx.x;// printf("tid = %d\n", tid);if(tid>=N){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__#include<cstdio>#defineCHECK_CUDA_CALL(call)\do{\constcudaError_t error_code=call;\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

#include<cmath>#include<cstdio>#include"cuda_error.cuh"constdoubleEPSILON=1e-15;constdoublea=1.23;constdoubleb=2.34;constdoublec=3.57;__global__voidadd(constdouble*px,constdouble*py,double*pz,constintN){constinttid=blockDim.x*blockIdx.x+threadIdx.x;// printf("tid = %d\n", tid);if(tid>=N){return;}pz[tid]=px[tid]+py[tid];}voidcheck(constdouble*pz,constintN){boolhas_error=false;for(inti=0;i<N;++i){if(fabs(pz[i]-c)>EPSILON){has_error=true;}}printf("%s\n",has_error?"Has errors":"No errors");}intmain(void){constintN=1e7;constintM=sizeof(double)*N;double*h_px=(double*)malloc(M);double*h_py=(double*)malloc(M);double*h_pz=(double*)malloc(M);for(inti=0;i<N;++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_size=128;constintgrid_size=(N%block_size==0)?(N/block_size):(N/block_size+1);printf("grid_size = %d, block_size = %d\n",grid_size,block_size);add<<<grid_size,block_size>>>(d_px,d_py,d_pz,N);// 这两句是检查核函数调用的错误，cudaGetLastError() 的作用是在主机与设备同步之前，捕捉最后一个错误// 通常情况下，主机异步调用核函数，调用后立即执行下一条语句，// 因此，捕捉核函数调用的错误需要在主机与设备同步之前，调用 cudaGetLastError() 函数。// 除了显示调用 cudaDeviceSynchronize() 实现主机与设备同步，还可以临时设置环境变量 CUDA_LAUNCH_BLOCKING=1，强制主机与设备同步。// 设置之后，所有核函数的调用都变为同步的，主机在调用核函数之后，必须等待设备执行完核函数之后，才能继续执行下一条语句。// 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;}

编译和运行：

ycao@Thinkpad-T14:~/cuda_junior$ ./run_demo.sh 03_check_api/check1.cu -------------------- nvcc-arch=sm_75-o/tmp/tmp.ZitlLOKBof/run_demo_exec 03_check_api/check1.cu -------------------- grid_size=78125, block_size=128No errors

3 总结

本文所有代码都托管在本人的 github 上：cuda_junior。下一篇，我们仍然基于两个数组相加这个例子，对 CUDA 的加速效果进行度量