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[deviceProp.computeMode]);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]个结果计算均正确.

（全文完，作于2022-07-30）

CUDA编程： GPU与CPU之间使用全局统一内存的完整代码及编译相关推荐

CUDA编程——GPU架构，由sp，sm，thread，block，grid，warp说起
目录 1.从硬件看 2.从软件看 3.对应关系 4.SIMT和SIMD 掌握部分硬件知识,有助于程序员编写更好的CUDA程序,提升CUDA程序性能,本文目的是理清sp,sm,thread,block, ...
Java网络编程 Socket、ServerSocket 详解，方法介绍及完整代码示例
Java网络编程 Socket.ServerSocket 详解,方法介绍及完整代码示例概念什么是网络编程? 网络编程是指编写运行在多个设备(计算机)的程序,这些设备通过网络连接起来.当这些通过网络 ...
cuda编程（7）：实现LK稀疏光流算法--完整的cuda程序
一.原理二.算法的假设(其实这样的条件比较苛刻) 三.具体实现 main.cc #include <opencv2/opencv.hpp> #include <opencv2/co ...
CUDA编程（一）：GPU计算与CUDA编程简介
CUDA编程(一):GPU计算与CUDA编程简介 GPU计算 GPU硬件资源 GPU软件资源 GPU存储资源 CUDA编程 GPU计算 NVIDIA公司发布的CUDA是建立在GPU上的一个通用并行计算 ...
CUDA刷新器：CUDA编程模型
CUDA刷新器:CUDA编程模型 CUDA Refresher: The CUDA Programming Model CUDA,CUDA刷新器,并行编程这是CUDA更新系列的第四篇文章,它的目标是 ...
cuda编程_CUDA刷新器：CUDA编程模型
CUDA刷新器:CUDA编程模型 CUDA Refresher: The CUDA Programming Model CUDA,CUDA刷新器,并行编程这是CUDA更新系列的第四篇文章,它的目标是 ...
CUDA 编程学习
0 简单的CUDA简介 1.简单教程 CUDA C ++只是使用CUDA创建大规模并行应用程序的方法之一.它允许您使用功能强大的C ++编程语言来开发由GPU上运行的数千个并行线程加速的高性能算法.许 ...
【CS231n 课程笔记】第八讲-----常见的深度学习框架，以及GPU和CPU介绍。
目录 0.写在前面 1.CPU VS GPU 1.1GPU 与CPU的比较 1.2性能测试 1.3实践中的问题 2.深度学习框架 2.1 计算图思想和深度学习框架 2.2 tensorflow 2.2 ...
c语言数天数类型,GPU和CPU到底有什么区别
描述显卡的发展可以说是非常的快,人们对于视觉化上的要求也越来越高,随着用户对于图像处理上面的要求不断超出处理器的计算能力.另一方面CPU处理能力也不断强大,但在进入3D时代后,人们发现庞大的3D图像 ...

CUDA编程： GPU与CPU之间使用全局统一内存的完整代码及编译

CUDA编程： GPU与CPU之间使用全局统一内存的完整代码及编译

CUDA编程： GPU与CPU之间使用全局统一内存的完整代码及编译相关推荐

最新文章

热门文章