11.2_零拷贝主机内存

11.2 零拷贝主机内存

在第10章中，我们介绍了固定内存（或者说页锁定内存），这种新型的主机内存能够确保不会交换出物理内存。我们通过调用cudahostAlloc()来分配这种内存，并且传递参数cudahostAllocDefault来获得默认的固定内存。在前面曾提到，本章会介绍在分配固定内存时可以使用的其他参数值。除了cudahostAllocDefault外，还可以传递的标志之一是cudahostAllocMapped。通过cudahostAllocMapped分配的主机内存也是固定的，它与通过cudahostAllocDefault分配的固定内存有着相同的属性，特别是当它不能从物理内存中交换出去或者重新定位时。但这种内存除了可以用于主机与GPU之间的内存复制外，还可以打破第3章中的主机内存规则之一：可以在CUDAC核函数中直接访问这种类型的主机内存。由于这种内存不需要复制到GPU，因此也称为零拷贝内存。

11.2.1 通过零拷贝内存实现点积运算

通常，GPU只能访问GPU内存，而CPU也只能访问主机内存。但在某些环境中，打破这种规则或许能带来更好的效果。为了说明由GPU访问主机内存将带来哪些好处，我们来重新回顾前面的归约运算：矢量点积运算。如果你已经完整地读了本书，那么肯定还记得第一个点积运算版本：将两个输入矢量复制到GPU，对相应的元素执行乘积计算，然后将中间结果复制回到主机，并在CPU上完成求和计算。

在这个版本中，我们不需要将输入矢量显式复制到GPU，而是使用零拷贝内存从GPU中直接访问数据。这个版本的点积运算非常类似于对固定内存的性能测试程序。我们将编写两个函数，其中一个函数是对标准主机内存的测试，另一个函数将在GPU上执行归约运算，并使用零拷贝内存作为输入缓冲区和输出缓冲区。首先，我们来看看点积运算的主机内存版本。按照惯例，首先创建计时事件，然后分配输入缓冲区和输出缓冲区，并用数据填充输入缓冲区。

float malloc_test( int size ) {udaEvent_t start, stop;float \*a,\*b,c,\*partial_c;float \*dev_a,\*dev_b,\*dev_partial_c;
float elapsedTime;   
HANDLE_ERROR(udaEventCreate(&start）);   
HANDLE_ERROR(udaEventCreate(&stop）);   
//在CPU上分配内存   
a=(float\*）malloc（size\*sizeof(float）);   
b=(float\*）malloc（size\*sizeof(float）);   
partial_c=(float\*）mallocBlocksPerGrid\*sizeof(float）);   
//在GPU上分配内存   
HANDLE_ERROR(udaMalloc（void\*\&dev_a, size\*sizeof(float））;   
HANDLE_ERROR(udaMalloc（void\*\&dev_b, size\*sizeof(float））;   
HANDLE_ERROR(udaMalloc（void\*\&dev_partial_c, blocksPerGrid\*sizeof(float））;   
//用数据填充主机内存   
for（int  $\mathrm{i} = 0$  ：i<size；i++）{ a[i]  $=$  i; b[i]  $=$  i\*2;

在分配好内存并且创建完数据后，就可以开始计算。启动计时器，将输入数据复制到GPU，执行点积核函数，并将中间计算结果复制回主机。

HANDLE_ERROR(udaEventRecord( start, 0 ) ); //将数组‘a’和‘b’复制到GPU   
HANDLE_ERROR(udaMemcpy( dev_a, a, size*sizeof(float),udaMemcpyHostToDevice)）;   
HANDLE_ERROR(udaMemcpy( dev_b, b, size*sizeof(float),udaMemcpyHostToDevice)）;   
dot<<<blocksPerGrid, threadsPerBlock>>>(size,dev_a,dev_b, dev_partial_c);   
//将数组‘c’从GPU复制到CPU   
HANDLE_ERROR(udaMemcpy( partial_c, dev_partial_c, blocksPerGrid*size(of float),udaMemcpyDeviceToHost)）;

现在，我们需要像第5章中那样结束CPU上的计算。在执行这个操作前，首先要停止事件计时器，因为它只需测量在GPU上完成的工作：

HANDLE_ERROR(udaEventRecord(stop,0)）;   
HANDLE_ERROR(udaEventSynchronize(stop))；   
HANDLE_ERROR(udaEventElapsedTime(& elapsedTime, start，stop））;

最后，将中间计算结果相加起来，并释放输入缓冲区和输出缓冲区。

//结束CPU上的计算  
c = 0;  
for (int i = 0; i < blocksPerGrid; i++) {  
    c += partial_c[i];  
}  
HANDLE_ERROR(udaFree(dev_a));  
HANDLE_ERROR(udaFree(dev_b));  
HANDLE_ERROR(udaFree(dev_partial_c));  
//释放CPU上的内存  
free(a);  
free(b);  
free(partial_c);  
//释放事件  
HANDLE_ERROR(udaEventDestroy(start));  
HANDLE_ERROR(udaEventDestroy(stop));  
printf("Value calculated: %f\n", c);  
return elapsedTime;

使用零拷贝内存的版本是非常类似的，只是在内存分配上有所不同。我们首先分配输入缓冲区和输出缓冲区，并同样用数据来填充输入内存：

float CUDA_host_alloc_test(int size) {
   udaEvent_t start, stop;
    float *a, *b, c, *partial_c;
    float *dev_a, *dev_b, *dev_partial_c;
    float elapsedTime;
    HANDLE_ERROR(cudaEventCreate(&start));
    HANDLE_ERROR(cudaEventCreate(&stop));
    //在CPU上分配内存
    HANDLE_ERROR(cudaHostAlloc((void**)&a, size*sizeof(float),
varaHostAllocWriteCombined |  
varaHostAllocMapped）);  
HANDLE_ERROR(纵深Alloc( (void**)&b,  
size*sizeof(float),  
纵深AllocWriteCombined |  
纵深AllocMapped)）;  
HANDLE_ERROR(纵深Alloc( (void**)&partial_c,  
blocksPerGrid*sizeof(float),  
纵深AllocMapped)）;  
//用数据填充主机内存  
for（int  $\mathrm{i} = 0$  ：i  $<$  size；i++）{a[i]  $\equiv$  i;  
b[i]  $\equiv$  i\*2;

与第10章一样，我们再次使用了cudaHostAlloc()，只是通过参数flags来指定内存的其他行为。cudaHostAllocMapped这个标志告诉运行时将从GPU中访问这块内存。换句话说，这个标志意味着分配零拷贝内存。对于两个输入缓冲区，我们还指定了标志cudaHostAllocWriteCombined。这个标志表示，运行时应该将内存分配为“合并式写入（Write-Combined）”内存。这个标志并不会改变应用程序的功能，但却可以显著地提升GPU读取内存时的性能。然而，当CPU也要读取这块内存时，“合并式写入”会显得很低效，因此在决定是否使用这个标志之前，必须首先考虑应用程序的可能访问模式。

在使用标志cudaHostAllocMapped来分配主机内存后，就可以从GPU中访问这块内存。然而，GPU的虚拟内存空间与CPU是不同的，因此在GPU上访问它们与在CPU上访问它们有着不同的地址。调用cudaHostAlloc()将返回这块内存在CPU上的指针，因此需要调用cudaHostGetDevicePointer()来获得这块内存在GPU上的有效指针。这些指针将被传递给核函数，并在随后由GPU对这块内存执行读取和写入等操作：

HANDLE_ERROR(udaHostGetDevicePointer(&dev_a, a, 0));  
HANDLE_ERROR(udaHostGetDevicePointer(&dev_b, b, 0));  
HANDLE_ERROR(udaHostGetDevicePointer(&dev_partial_c, partial_c, 0));

在获得了有效的设备指针后，就可以启动计时器以及核函数。

HANDLE_ERROR(udaEventRecord( start, 0 ) );  
dot<<<blocksPerGrid, threadsPerBlock>>>( size, dev_a, dev_b, dev_partial_c );  
HANDLE_ERROR(udaThreadSynchronize() );

即使指针dev_a、dev_b和dev_partial_c都位于主机上，但对于核函数来说，它们看起来与GPU内存一样，这正是由于调用了cudaHostGetDevicePointer()。由于部分计算结果已经位于主机上，因此就不再需要通过cudaMemcpy()将它们从设备上复制回来。然而，你可能注意到了在程序中调用了cudaThreadSynchronize()将CPU与GPU同步。如果在核函数中会修改零拷贝内存的内容，那么在核函数的执行期间，零拷贝内存的内容是未定义的。在同步完成后，就可以确信核函数已经完成，并且在零拷贝内存中包含了计算好的结果，因此就可以停止计时器并结束在CPU上的计算。

HANDLE_ERROR(udaEventRecord(stop,0)）;   
HANDLE_ERROR(udaEventSynchronize(stop)）;   
HANDLE_ERROR(udaEventElapsedTime(& elapsedTime, start，stop）);   
//结束CPU上的操作   
 $\mathbf{c} = \mathbf{0}$  ·   
for（int  $i = 0$  ；i<blocksPerGrid;i++) {  $\mathrm{c + =}$  partial_c[i];   
1

在使用cudahostAlloc()的点积运算代码中，唯一剩下的事情就是执行释放操作。

HANDLE_ERROR(udaFreeHost(a)）;   
HANDLE_ERROR(udaFreeHost(b)）;   
HANDLE_ERROR(udaFreeHost(partial_c）);   
//释放事件   
HANDLE_ERROR(udaEventDestroy（start））;   
HANDLE_ERROR(udaEventDestroy(stop）);   
printf("Value calculated：%f\n"，c）;   
return elapsedTime;

需要注意的是，无论在cudaHostAlloc()中使用什么标志，总是按照相同的方式来释放内存，即只需调用cudaFreeHost()。

基本内容就是这样！剩下的工作就是观察main()如何将这些代码片段组合在一起。在main()中，首先要判断设备是否支持映射主机内存。我们使用第10章中判断设备是否支持重叠的方法，即调用cudaGetDeviceProperties()。

int main(void) {
   udaDeviceProp prop;
int whichDevice;   
HANDLE_ERROR( CUDAGetDevice( &whichDevice ) );   
HANDLE_ERROR( CUDAGetDeviceProperties( &prop, whichDevice ) ); if(prop.canMapHostMemory != 1) { printf("Device cannot map memory.\n"); return 0; }

如果设备支持零拷贝内存，那么接下来就是将运行时置入能分配零拷贝内存的状态。通过调用CUDASetDeviceFlags()来实现这个操作，并且传递标志值CUDADeviceMapHost来表示我们希望设备映射主机内存：

HANDLE_ERROR(udaSetDeviceFlags(udaDeviceMapHost));

这就是main()中的主要操作。我们运行两个测试，分别显示二者的执行时间，并退出应用程序：

float elapsedTime  $=$  malloc_test(N); printf("Time using CUDAAlloc:  $\text{日} 3 . 1 \mathrm { f ~ m s } \backslash \mathrm { n } "$  elapsedTime）; elapsedTime  $=$  CUDA_host_alloc_test(N); printf("Time using CUDAHostAlloc:  $\text{日} 3 . 1 \mathrm { f ~ m s } \backslash \mathrm { n } "$  elapsedTime）;

核函数本身与第5章中的核函数没有区别，但我们还是给出了完整的核函数：

define imin(a,b) (a<b?a:b)   
const int N = 33 * 1024 * 1024;   
const int threadsPerBlock = 256;   
const int blocksFerGrid = imin(32, (N+threadsPerBlock-1) / threadsPerBlock);   
__global__ void dot(int size, float *a, float *b, float *c) { _shared__ float cache[threadsPerBlock]; int tid = threadIdx.x + blockIdx.x * blockDim.x; int cacheIndex = threadIdx.x; float temp = 0; while (tid < size) { temp += a[tid] * b[tid]; tid += blockDim.x * gridDim.x; }
//设置cache中的值  
cache［cacheIndex]  $\equiv$  temp;  
//同步这个线程块中的线程  
__syncthreads();  
//对于归约运算，threadsPerBlock必须为2的幂  
int i  $=$  blockDim.x/2;  
while（i！=0）{if（cacheIndex< i）cache[cacheIndex]  $+ =$  cache[cacheIndex+i];_syncthreads();i/=2;  
}  
if（cacheIndex  $\equiv = 0$  ）c[blockIdx.x]  $\equiv$  cache[0];

11.2.2 零拷贝内存的性能

零拷贝内存能够带来哪些好处？对于独立GPU和集成GPU，答案是不同的。独立GPU自己拥有专门的DRAM，通常位于CPU之外的电路板上。例如，如果在计算机中已经安装了一块图形卡，那么这个GPU就是一个独立GPU。集成GPU是系统芯片组中内置的图形处理器，通常与CPU共享系统内存。在许多基于NVIDIA nForce媒体与通信处理器（Media and Communications Processor, MCP）构建的现代系统中，都包含了支持CUDA的集成GPU。除了nForce MCP外，基于NVIDIA新推出的ION平台的上网本、笔记本以及桌面计算机都包含了集成的和支持CUDA的GPU。对于集成GPU，使用零拷贝内存通常都会带来性能提升，因为内存在物理上与主机是共享的。将缓冲区声明为零拷贝内存的唯一作用就是避免不必要的数据复制。但要记住，天下没有免费的午餐，所有类型的固定内存都存在一定的局限性，零拷贝内存同样也不例外：每个固定内存都会占用系统的可用物理内存，这最终将降低系统的性能。

当输入内存和输出内存都只能使用一次时，那么在独立GPU上使用零拷贝内存将带来性能提升。由于GPU在设计时考虑了隐藏内存访问带来的延迟，因此这种机制在某种程度上将减轻PCIE总线上读取和写入等操作的延迟，从而会带来可观的性能提升。但由于GPU不会缓存零拷贝内存的内容，如果多次读取内存，那么最终将得不偿失，还不如一开始就将数据复制到GPU。

如何判断某个GPU是集成的还是独立的？当然，你可以打开计算机的机箱来观察，但这种方法对于CUDA C应用程序来说是不可行的。在代码中可以通过CUDAGetDeviceProperties()返回

的结构来判断GPU的这个属性。在该结构中有一个域integrated，如果设备是集成GPU，那么这个域的值为true，否则为false。

由于点积运算应用程序满足“仅读取/写入一次”这个约束条件，因此在使用零拷贝内存时能获得性能提升。事实上，程序在性能上确实有一定程度的提升。在GeForce GTX 285上，当使用零拷贝内存时，程序的执行时间减少了 45%45\% ，从98.1毫秒降到52.1毫秒。在GeForce GTX 280上同样能获得性能提升，执行时间减少 34%34\% ，从143.9毫秒降为94.7毫秒。当然，由于计算量与带宽的比值不同，以及芯片组之间PCIE总线带宽的不同，不同的GPU将表现出不同的性能特性。

前の章11.1_本章目标
次の章11.3_使用多个GPU
11.2_零拷贝主机内存 - CUDA by Example | OpenTech