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 SM
typedef struct {
int SM; // 0xMm (hexidecimal notation), M = SM Major version,
// and m = SM minor version
int 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计算完成后的数据刷新回到CPU
cudaMemPrefetchAsync(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）
4788
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
