CUDA编程： GPU与CPU之间使用全局统一内存的完整代码及编译
最近碰到一个应用场景，需要从GPU访问host上创建的，一个很大的布隆过滤器（准确说是改进后的布谷鸟过滤器）。由于GPU卡上的显存有限，把整个过滤器复制到GPU卡显然不可能，于是想到用CUDA的全局统一内存来简化程序编写并提高性能。
由于以前没做过CUDA的编程，要从零开始学CUDA，还要进阶到用 统一虚拟内存寻址UVA，再到全局统一内存，甚至连CUDA的编译都是现学，碰到了不少问题。在参考这篇文章：
CPU与GPU的统一虚拟地址(CUDA UVA)原理
以及这篇文章：
CUDA全局内存
再到github上拉了一些源代码学习后，自己写了如一下完整的，使用统一内存，并可编译运行的简单例子，供正在入门CUDA编程的读者参考。
由于我只需要从GPU卡只读访问主机内存上的过滤器内容，因此这里除了演示cudaMallocManaged、cudaMemPrefetchAsync等与全局统一内存相关的函数用法外，还特别演示了cudaMemAdvise的用法。
//通知GPU只需读取a、b数组cudaMemAdvise(a, size, cudaMemAdviseSetReadMostly, deviceId);cudaMemAdvise(b, size, cudaMemAdviseSetReadMostly, deviceId);
这里也存在一个有待还没进一步验证的疑问：对于我的应用，过滤器应该在host内存上创建，所以使用cudaMallocHost在主机上创建固定不可分页内存，是否比用cudaMallocManaged创建全局内存+cudaMemAdvise指定为只读但可分页内存的效率要来的更高？
示例程序（文件名: uva_test.cu）完整的源代码如下:
/*********************************************************************************************************************************** 文件名 uva_test.cu* 编译命令： nvcc -o test_uva uva_test.cu* 一个测试CUDA的全局虚拟内存地址(UVA)的示范程序* author: Ryan***********************************************************************************************************************************/#include <iostream>#include <stdio.h>#include <cuda_runtime_api.h>using namespace std;// --------------------------------------------------------------------------------------------------------------------------------//计算GPU卡的SM数量int _ConvertSMVer2Cores(int major,int minor) {// Defines for GPU Architecture types (using the SM version to determine// the # of cores per SMtypedef struct {int SM; // 0xMm (hexidecimal notation), M = SM Major version,// and m = SM minor versionint Cores;} sSMtoCores;sSMtoCores nGpuArchCoresPerSM[] = {{0x20, 32 }, // Fermi Generation (SM 2.0) GF100 class{0x21, 48 }, // Fermi Generation (SM 2.1) GF10x class{0x30, 192 },{0x32, 192 },{0x35, 192 },{0x37, 192 },{0x50, 128 },{0x52, 128 },{0x53, 128 },{0x60, 64 },{0x61, 128 },{0x62, 128 },{0x70, 64 },{0x72, 64 },{0x75, 64 },{0x80, 64 },{0x86, 128 },{-1, -1 } };int index = 0;while(nGpuArchCoresPerSM[index].SM != -1) {if(nGpuArchCoresPerSM[index].SM == ((major << 4) + minor)) {return nGpuArchCoresPerSM[index].Cores;}index++;}return 0;}// --------------------------------------------------------------------------------------------------------------------------------//调用API显示GPU的硬件信息void PrintCudaInfo() {cudaError_t err;const char *sComputeMode[] ={"Multiple host threads","Only one host thread","No host thread","Multiple process threads","Unknown",NULL};int deviceCount = 0;cudaError_t error_id = cudaGetDeviceCount(&deviceCount);if(error_id != cudaSuccess) {printf("GPUEngine: CudaGetDeviceCount %s\n",cudaGetErrorString(error_id));return;}// 如果本机未安装GPU卡，diviceCount的值将为0.if(deviceCount == 0) {printf("GPUEngine: There are no available device(s) that support CUDA\n");return;}//当前系统安装的驱动版本int driver_Version=0;cudaDriverGetVersion(&driver_Version);printf("[+] 系统当前共检测到[%d]块GPU卡,安装的CUDA驱动版本为:%d.%d\n",deviceCount,driver_Version / 1000, (driver_Version % 100) / 10);for(int i = 0; i<deviceCount; i++) {err = cudaSetDevice(i);if(err != cudaSuccess) {printf("[E] 错误,调用cudaSetDevice(%d)时发生错误: %s\n",i,cudaGetErrorString(err));return;}cudaDeviceProp deviceProp;cudaGetDeviceProperties(&deviceProp,i);//如果要得到准确的流处理器数量，需要判断deviceProp.major，再乘以printf("[+] 第[%d]块GPU卡[%s] (共有:%dx%d=%d个流处理核心,主频:%d MHz) (算力: %d.%d) (设备内存:%.2f MB) (%s)\n",i+1,deviceProp.name,deviceProp.multiProcessorCount, _ConvertSMVer2Cores(deviceProp.major,deviceProp.minor),deviceProp.multiProcessorCount*_ConvertSMVer2Cores(deviceProp.major,deviceProp.minor),deviceProp.clockRate/1000,deviceProp.major,deviceProp.minor,(double)deviceProp.totalGlobalMem / 1048576.0,sComputeMode[puteMode]);printf("[+] 第[%d]块GPU卡[%s] (maxGridSize[(%d,%d,%d)] (本设备%s统一虚拟寻址UVA)\n",i+1,deviceProp.name,deviceProp.maxGridSize[0],deviceProp.maxGridSize[1],deviceProp.maxGridSize[2],deviceProp.unifiedAddressing ? "支持":"不支持");}}// --------------------------------------------------------------------------------------------------------------------------------void init_value(float num, float *a, int N){for(int i = 0; i < N; ++i){a[i] = num;}}// --------------------------------------------------------------------------------------------------------------------------------//检查计算结果void checkElementsAre(float target, float *vector, int N){for(int i = 0; i < N; i++){if(vector[i] != target){printf("FAIL: vector[%d] - %0.0f does not equal %0.0f\n", i, vector[i], target);exit(1);}}printf("[+] 经检验,[%d]个结果计算均正确.\n",N);}// --------------------------------------------------------------------------------------------------------------------------------//GPU的计算函数__global__ void addVectorsInto(float *result, float *a, float *b, int N){int index = threadIdx.x + blockIdx.x * blockDim.x;int stride = blockDim.x * gridDim.x;for(int i = index; i < N; i += stride){result[i] = a[i] + b[i];}}// --------------------------------------------------------------------------------------------------------------------------------//main入口函数int main(void){const int N = 2<<24;size_t size = N * sizeof(float);int count;int deviceId;int numberOfSMs; //GPU的内核数量//查询当前可用GPU卡数量cudaGetDeviceCount(&count);if (count == 0) {fprintf(stderr, "[E] 错误,当前系统未检测到GPU卡.\n");return -1;}//显示GPU卡的硬件参数PrintCudaInfo();//获得第一个GPU设备号cudaGetDevice(&deviceId); //获取GPU的内核数（注意仅是内核数，不是流处理器总数量）cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);float *a;float *b;float *ret;printf("[+] 开始演示CUDA 的UVA全局虚拟内存功能,将创建[%.2f] KB的UVA内存变量\n",(double)((size*3)/(1024*1024.0)));//调用UVA全局虚拟内存申请函数cudaMallocManaged(&a, size);cudaMallocManaged(&b, size);//通知GPU只需读取a数组cudaMemAdvise(a, size, cudaMemAdviseSetReadMostly, deviceId);cudaMemAdvise(b, size, cudaMemAdviseSetReadMostly, deviceId);cudaMallocManaged(&ret, size);init_value(3, a, N);init_value(4, b, N);init_value(0, ret, N);// 调用cudaMemPrefetchAsync 将数据预取到GPU，对于a及b,由于已经用cudaMemAdvise指定为只读,// 将只在GPU产生只读地址副本，中途不检查生缺页中断，// 也仅在GPU端需要时，GPU才从host内存读入数据cudaMemPrefetchAsync(a, size, deviceId);cudaMemPrefetchAsync(b, size, deviceId);cudaMemPrefetchAsync(ret, size, deviceId);size_t threadsPerBlock;size_t numberOfBlocks;threadsPerBlock = 256;//GPU并行的线程数量numberOfBlocks = 32 * numberOfSMs;cudaError_t addVectorsErr;cudaError_t asyncErr;//调用GPU进行计算addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(ret, a, b, N);//取得GPU计算的结果状态addVectorsErr = cudaGetLastError();if(addVectorsErr != cudaSuccess) printf("[E] GPU计算发生错误: %s\n", cudaGetErrorString(addVectorsErr));asyncErr = cudaDeviceSynchronize();if(asyncErr != cudaSuccess) printf("[E] GPU计算发生错误: %s\n", cudaGetErrorString(asyncErr));// 将GPU计算完成后的数据刷新回到CPUcudaMemPrefetchAsync(ret, size, cudaCpuDeviceId);checkElementsAre(7, ret, N);//销毁cudaMallocManaged创建的内存cudaFree(a);cudaFree(b);cudaFree(ret);return 0;}
上述代码，可以用以下命令编译为可执行文件：test_uva
$ nvcc -o test_uva uva_test.cu
在我的设备上运行的结果如下：
$ ./test_uva[+] 系统当前共检测到[1]块GPU卡,安装的CUDA驱动版本为:11.7[+] 第[1]块GPU卡[NVIDIA GeForce RTX 3070 Ti] (共有:48x128=6144个流处理核心,主频:1770 MHz) (算力: 8.6) (设备内存:7981.00 MB) (Multiple host threads)[+] 第[1]块GPU卡[NVIDIA GeForce RTX 3070 Ti] (maxGridSize[(2147483647,65535,65535)] (本设备支持统一虚拟寻址UAV)[+] 开始演示CUDA 的UVA全局虚拟内存功能,将创建[384.00] KB的UVA内存变量[+] 经检验,[33554432]个结果计算均正确.
（全文完，作于-07-30）
