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

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）