两种方法使用零拷贝内存做简单的向量加和，并评估 GPU 计算结果与 CPU 计算结果的差。
▶ 源代码
1 #include <stdio.h> 2 #include <cuda.h> 3 #include <cuda_runtime.h> 4 #include "device_launch_parameters.h" 5 #include <helper_functions.h> 6 #include <helper_cuda.h> 7 8 #define MEMORY_ALIGNMENT 4096 9 #define ALIGN_UP(x,size) ( ((size_t)x+(size-1))&(~(size-1)) ) 10 11 __global__ void vectorAddGPU(float *a, float *b, float *c, int N) 12 { 13 int idx = blockIdx.x*blockDim.x + threadIdx.x; 14 if (idx < N) 15 c[idx] = a[idx] + b[idx]; 16 } 17 18 int main(int argc, char **argv) 19 { 20 printf(" Start. "); 21 22 // 设备检查 23 bool bMac; 24 cudaDeviceProp deviceProp; 25 cudaSetDevice(0); 26 cudaGetDeviceProperties(&deviceProp, 0); 27 if (CUDART_VERSION < 2020 || !deviceProp.canMapHostMemory)// CUDART_VERSION 为 CUDA Runtime API 版本，CUDA9.0 对应 9000 28 { 29 printf(" CUDA Runtime API not support MapHostMemory. "); 30 getchar(); 31 return 1; 32 } 33 cudaSetDeviceFlags(cudaDeviceMapHost);// MapHostFlag 功能正常，设置标志 34 #if defined(__APPLE__) || defined(MACOSX)// MacOS 系统不支持将普通堆内存设置为页锁定内存 35 bMac = true; 36 #else 37 bMac = false; 38 #endif 39 if (CUDART_VERSION < 4000 && !bMac)// 既不是 MacOS 系统，Runtime API 版本还不够高 40 { 41 printf(" CUDA Runtime API not support cudaHostRegister function. "); 42 getchar(); 43 return 1; 44 } 45 // 总体逻辑： 46 // CUDA Runtime version < 2200，不支持 MApHostMamory，退出 47 // CUDA Runtime version ∈[2200, 4000)，且为 MAcOS 系统，使用 cudaHostAlloc() + cudaHostAllocMapped 48 // CUDA Runtime version ∈[2200, 4000)，且不是 MAcOS 系统，退出 49 // CUDA Runtime version ≥ 4000，使用 malloc() + cudaHostRegister() 50 51 // 内存申请 52 int nelem = 1048576; 53 int bytes = nelem * sizeof(float); 54 float *a, *b, *c; 55 float *a_UA, *b_UA, *c_UA; 56 float *d_a, *d_b, *d_c; 57 if (CUDART_VERSION >= 4000 || bMac) 58 { 59 a_UA = (float *) malloc(bytes + MEMORY_ALIGNMENT); // 申请时多 4KB，用于滑动对齐，释放内存时以该指针为准 60 b_UA = (float *) malloc(bytes + MEMORY_ALIGNMENT); 61 c_UA = (float *) malloc(bytes + MEMORY_ALIGNMENT); 62 a = (float *) ALIGN_UP(a_UA, MEMORY_ALIGNMENT); // 指针指到 4K 对齐的位置上去，用于计算 63 b = (float *) ALIGN_UP(b_UA, MEMORY_ALIGNMENT); 64 c = (float *) ALIGN_UP(c_UA, MEMORY_ALIGNMENT); 65 cudaHostRegister(a, bytes, CU_MEMHOSTALLOC_DEVICEMAP); // 设置页锁定内存 66 cudaHostRegister(b, bytes, CU_MEMHOSTALLOC_DEVICEMAP); 67 cudaHostRegister(c, bytes, CU_MEMHOSTALLOC_DEVICEMAP); 68 } 69 else 70 { 71 cudaHostAlloc((void **)&a, bytes, cudaHostAllocMapped); // 使用函数 cudaHostAlloc() 一步到位 72 cudaHostAlloc((void **)&b, bytes, cudaHostAllocMapped); 73 cudaHostAlloc((void **)&c, bytes, cudaHostAllocMapped); 74 } 75 76 // 初始化和内存映射 77 for (int n = 0; n < nelem; n++) 78 { 79 a[n] = rand() / (float)RAND_MAX; 80 b[n] = rand() / (float)RAND_MAX; 81 } 82 cudaHostGetDevicePointer((void **)&d_a, (void *)a, 0); 83 cudaHostGetDevicePointer((void **)&d_b, (void *)b, 0); 84 cudaHostGetDevicePointer((void **)&d_c, (void *)c, 0); 85 86 // 调用内核 87 dim3 block(256, 1, 1); 88 dim3 grid((unsigned int)ceil(nelem / (float)block.x)); 89 vectorAddGPU << <grid, block >> > (d_a, d_b, d_c, nelem); 90 cudaDeviceSynchronize(); 91 92 // 检查结果 93 float errorNorm, refNorm, ref, diff; 94 errorNorm = 0.f; 95 refNorm = 0.f; 96 for (int n = 0; n < nelem; n++) 97 { 98 diff = c[n] - (ref = a[n] + b[n]);// ref 为 CPU 计算的和，diff 为 GPU 计算结果与 CPU 计算结果的差 99 errorNorm += diff*diff; // 向量 a + b 的两种计算结果的差的平方 100 refNorm += ref*ref; // 向量 a 与向量 b 的和的平方 101 } 102 errorNorm = (float)sqrt((double)errorNorm); 103 refNorm = (float)sqrt((double)refNorm); 104 printf(" Difference between GPU and CPU is %f, %f%% ", errorNorm, errorNorm / refNorm); 105 106 // 清理工作 107 if (CUDART_VERSION >= 4000 || bMac) 108 { 109 cudaHostUnregister(a); 110 cudaHostUnregister(b); 111 cudaHostUnregister(c); 112 free(a_UA); 113 free(b_UA); 114 free(c_UA); 115 } 116 else 117 { 118 cudaFreeHost(a); 119 cudaFreeHost(b); 120 cudaFreeHost(c); 121 } 122 printf(" Finish. "); 123 getchar(); 124 return 0; 125 }
▶ 输出结果：
1 Start. 2 Difference between GPU and CPU is 0.000000, 0.000000% 3 4 Finish.
▶ 涨姿势
● 两种使用零拷贝内存的方法，在代码的逻辑部分进行了说明
● 向上取整的宏函数，只对分母（size）为 2 的整数次幂的情况有效。
1 #define ALIGN_UP(x,size) ( ((size_t)x+(size-1))&(~(size-1)) )
e.g. size == 4096，则 ~ (size - 1) == 11111111 11111111 11110000 000000002，将其作为模板进行按位且操作，等价于取不低于 4096 的高位。