11.4_可移动的固定内存

11.4 可移动的固定内存

在使用多个GPU的程序中，最后一个重要的部分就是固定内存的使用。在第10章中已经介绍了，固定内存实际上是主机内存，只是该内存页锁定在物理内存中，以便防止被换出或者重定位。然而，这些内存页仅对于单个CPU线程来说是“固定的”。也就是说，如果某个线程分配了固定内存，那么这些内存只是对于分配它们的线程来说是页锁定的。如果在线程之间共享指向这块内存的指针，那么其他的线程将把这块内存视为标准的、可分页的内存。

这种行为的副作用之一就是，当其他线程（不是分配固定内存的线程）试图在这块内存上执行cudaMemcpy()时，将按照标准的可分页内存速率来执行复制操作。在第10章中曾介绍过，这种速率大约为最高传输速度的 50%50\% 。更糟糕的是，如果线程试图将一个cudaMemcpyAsync()调用放入CUDA流的队列中，那么将失败，因为cudaMemcpyAsync()需要使用固定内存。由于这块内存对于除了分配它的线程以外的其他线程来说似乎是可分页的，因此这个调用会失败，甚至导致任何后续操作都无法进行。

然而，对于这个问题有一种补救方案。我们可以将固定内存分配为可移动的，这意味着可以在主机线程之间移动这块内存，并且每个线程都将其视为固定内存。要达到这个目的，需要使用cudaHostAlloc()来分配内存，并且在调用时使用一个新的标志：cudaHostAllocPortable。这个标志可以与其他标志一起使用，例如cudaHostAllocWriteCombined和cudaHostAllocMapped。这意味着在分配主机内存时，可将其作为可移动、零拷贝以及合并式写入等的任意组合。

为了说明可移动固定内存的作用，我们将进一步修改使用多GPU的点积运算应用程序。我们将修改最初使用零拷贝内存的点积运算程序，因此在这个版本中将融合零拷贝版本的代码和

多GPU版本的代码。首先，我们需要验证至少有两个支持CUDA的GPU，并且二者都能处理零拷贝缓冲区。

int main(void) { int deviceCount; HANDLE_ERROR(udaGetDeviceCount( &deviceCount）); if (deviceCount<2）{ printf(“We need at least two compute 1.0 or greater " devices,but only found  $\text{d}\backslash \text{n}"$  deviceCount）; return 0; } JudaDeviceProp prop; for(int  $i = 0$  ；  $i <   2$  ；  $i + +$  ）{ HANDLE_ERROR(udaGetDeviceProperties( &prop，i））; if(prop.canMapHostMemory！=1）{ printf(“Device  $\text{d}$  cannot map memory.\n”，i）; return 0; }

在前面的示例中，已经准备好了在主机上分配内存以便保存输入矢量。然而，为了分配可移动的固定内存，首先需要设置将在哪个CUDA设备上运行。由于我们还希望在这个设备上分配零拷贝内存，因此在调用CUDASetDevice()之后，接着调用了CUDASetDeviceFlags()，和11.1.1节中一样。

float \*a,\*b;   
HANDLE_ERROR( CUDASetDevice(0)）;   
HANDLE_ERROR( CUDASetDeviceFlags( CUDADeviceMapHost))；   
HANDLE_ERROR( CUDAHostAlloc( (void\*\*)&a,N\*sizeof(float), CUDAHostAllocWriteCombined | CUDAHostAllocPortable | CUDAHostAllocMapped))；   
HANDLE_ERROR( CUDAHostAlloc( (void\*\*)&b,N\*sizeof(float), CUDAHostAllocWriteCombined | CUDAHostAllocPortable | CUDAHostAllocMapped))；

在本章的前面，我们等到已经分配了内存并且创建了线程后才调用CUDASetDevice()。然而，在使用CUDAHostAlloc()分配页锁定内存时，首先要通过调用CUDASetDevice()来初始化设备。你还将注意到，我们将新介绍的标志CUDAHostAllocPortable传递给这两个内存分配操作。由于这些内存是在调用了CUDASetDevice(0)之后才分配的，因此，如果没有将这些内存指定为可移动的内存，那么只有第0个CUDA设备会把这些内存视为固定内存。

继续之前的应用程序，为输入矢量生成数据，并采用11.2节中多GPU示例中的方式来准备DataStruct结构。

//用数据填充主机内存  
for（int  $\mathbf{i} = 0$  ：i<N；i++）{a[i]  $=$  i;b[i]  $=$  i\*2;1

//为使用多线程做好准备

DataStruct data[2];  
data[0].deviceID = 0;  
data[0].offset = 0;  
data[0].size = N/2;  
data[0].a = a;  
data[0].b = b;  
data[1].deviceID = 1;  
data[1].offset = N/2;  
data[1].size = N/2;  
data[1].a = a;  
data[1].b = b;

然后，我们创建第二个线程，并调用routine()开始在每个设备上执行计算。

CUTThread thread = start_thread( routine, &(data[1]) );  
routine( &(data[0]) );  
end_thread( thread );

由于主机内存是由CUDA运行时分配的，因此需要用CUDAFreeHost()而不是free()来释放它。

// 释放CPU上的内存  
HANDLE_ERROR(udaFreeHost(a));  
HANDLE_ERROR(udaFreeHost(b));  
printf("Value calculated: %f\n", data[0].returnValue + data[1].returnValue);  
return 0;

为了在多GPU应用程序中支持可移动的固定内存和零拷贝内存，我们需要对routine()的代码进行两处修改。第一处修改有些微妙。

void\*routine(void\*pvoidData）{
DataStruct *data = (DataStruct*)pvoidData;  
if (data->deviceID != 0) {  
    HANDLE_ERROR(udaSetDevice(data->deviceID));  
    HANDLE_ERROR(udaSetDeviceFlags(udaDeviceMapHost));  
}

你可能还记得，在多GPU版本的代码中，我们需要在routine()中调用CUDASetDevice()，从而确保每个线程控制一个不同的GPU。另一方面，在这个示例中，我们已经在主线程中调用了一次CUDASetDevice()。这么做的原因是为了在main()中分配固定内存。因此，我们只希望在还没有调用CUDASetDevice()的设备上调用CUDASetDevice()和CUDASetDeviceFlags()。也就是，如果deviceID不是0，那么将调用这两个函数。虽然在第0个设备上再次这些函数调用将产生更整洁的代码，但事实上这种做法是错误的。一旦在某个线程上设置了这个设备，那么将不能再次调用CUDASetDevice()，即便传递的是相同的设备标识符。粗体显示if()语句将帮助我们避免CUDA运行时中的这个问题，因此我们继续讨论routine()的下一处重要修改。

除了使用可移动的固定内存外，我们还使用了零拷贝内存，以便从GPU中直接访问这些内存。因此，我们不再像之前的应用程序那样使用cudaMemcpy()，而是使用cudaHostGetDevicePointer()来获得主机内存的有效设备指针，这与前面零拷贝示例中采用的方法一样。然而，你可能会注意到使用了标准的GPU内存来保存临时计算结果。这块内存同样是通过CUDAAlloc()来分配的。

int size = data->size;  
float *a, *b, c, *partial_c;  
float *dev_a, *dev_b, *dev_partial_c;  
// 在CPU上分配内存  
a = data->a;  
b = data->b;  
partial_c = (float*)malloc( blocksPerGrid* sizeof(float));  
HANDLE_ERROR(udaHostGetDevicePointer( &dev_a, a, 0));  
HANDLE_ERROR(udaHostGetDevicePointer( &dev_b, b, 0));  
HANDLE_ERROR(udaMalloc( (void**)&dev_partial_c, blocksPerGrid* sizeof(float) ));  
// 计算GPU读取数据的偏移量‘a’和‘b’  
dev_a += data->offset;  
dev_b += data->offset;

此时，我们已经完全做好了准备，因此可以启动核函数并且将结果从GPU中复制回来。

$\mathrm{dot} <   <   <   \mathrm{blocksPerGrid,threadsPerBlock} > > >$  (size，dev_a，dev_b， dev_partial_c）;
// 将数组 'c' 从GPU复制回CPU  
HANDLE_ERROR(udaMemcpy( partial_c, dev_partial_c, blocksPerGrid*size(float),udaMemcpyDeviceToHost));

最后，像之前的点积示例一样，在CPU上将临时和值相加起来，释放内存，并返回到main()。

//结束CPU上的操作  
c = 0;  
for (int i = 0; i < blocksPerGrid; i++) {  
    c += partial_c[i];  
}  
HANDLE_ERROR(udaFree(dev_partial_c));  
//释放CPU上的内存  
free(partial_c);  
data->returnValue = c;  
return 0;
Previous11.3_使用多个GPU
Next11.5_本章小结