10.2 页锁定主机内存

在前9章的各个示例中，都是通过CUDAAlloc()在GPU上分配内存，以及通过标准的C库函数malloc()在主机上分配内存。除此之外，CUDA运行时还提供了自己独有的机制来分配主机内存：CUDAHostAlloc()。如果malloc()已经能很好地满足C程序员的需求，那么为什么还要使用这个函数？

事实上，malloc()分配的内存与CUDAHostAlloc()分配的内存之间存在着一个重要差异。C库函数malloc()将分配标准的，可分页的（Pagable）主机内存，而CUDAHostAlloc()将分配页锁定的主机内存。页锁定内存也称为固定内存（PinnedMemory）或者不可分页内存，它有一个重要的属性：操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中。因此，操作系统能够安全地使某个应用程序访问该内存的物理地址，因为这块内存将不会被破坏或者重新定位。

由于GPU知道内存的物理地址，因此可以通过“直接内存访问（Direct Memory Access, DMA）”技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU的介入，这也就同样意味着，CPU很可能在DMA的执行过程中将目标内存交换到磁盘上，或者通过更新操作系统的可分页表来重新定位目标内存的物理地址。CPU可能会移动可分页的数据，这就可能对DMA操作造成延迟。因此，在DMA复制过程中使用固定内存是非常重要的。事实上，当使用可分页内存进行复制时，CUDA驱动程序仍然会通过DAM把数据传输给GPU。因此，复制操作将执行两遍，第一遍从可分页内存复制到一块“临时的”页锁定内存，然后再从这个页锁定内存复制到GPU上。

因此，每当从可分页内存中执行复制操作时，复制速度将受限于PCIE传输速度和系统前端总线速度相对较低的一方。在某些系统中，这些总线在带宽上有着巨大的差异。因此当在GPU和主机间复制数据时，这种差异会使页锁定主机内存的性能比标准可分页内存的性能要高大约2倍。即使PCIE的速度与前端总线的速度相等，由于可分页内存需要更多一次由CPU参与的复制操作，因此会带来额外的开销。

然而，你也不能进入另一个极端：查找每个malloc调用并将其替换为cudaHostAlloc()调用。

固定内存是一把双刃剑。当使用固定内存时，你将失去虚拟内存的所有功能。特别是，在应用程序中使用每个页锁定内存时都需要分配物理内存，因为这些内存不能交换到磁盘上。这意味着，与使用标准的malloc()调用相比，系统将更快地耗尽内存。因此，应用程序在物理内存较少的机器上会运行失败，而且意味着应用程序将影响在系统上运行的其他应用程序的性能。

这些情况并不是说不使用cudaHostAlloc()，而是提醒你应该清楚页锁定内存的隐含作用。我们建议，仅对cudaMemcpy()调用中的源内存或者目标内存，才使用页锁定内存，并且在不再需要使用它们时立即释放，而不是等到应用程序关闭时才释放。cudaHostAlloc()与到目前为止学习的其他内容一样简单，我们来看一个示例，这个示例很好地说明了如何分配固定内存，以及它相对于标准可分页内存的性能优势。

这个应用程序非常简单，主要是测试cudaMemcpy()在可分配内存和页锁定内存上的性能。我们要做的就是分配一个GPU缓冲区，以及一个大小相等的主机缓冲区，然后在这两个缓冲区之间执行一些复制操作。我们允许用户指定复制的方向，例如为“上”（从主机到设备）或者为“下”（从设备到主机）。你还将注意到，为了获得精确的时间统计，我们为复制操作的起始时刻和结束时刻分别设置了CUDA事件。你可能会记得之前性能测试示例中是如何实现这些操作的，但如果忘了也不要紧，下面这些代码会帮你回忆它们的用法：

float CUDA_malloc_test(int size, bool up) {
    CUDAEvent_t start, stop;
    int *a, *dev_a;
    float elapsedTime;
    HANDLE_ERROR( CUDAEventCreate( &start ) );
    HANDLE_ERROR( CUDAEventCreate( &stop ) );
    a = (int*)malloc( size * sizeof( *a ) );
    HANDLE_NULL( a );
    HANDLE_ERROR( CUDA malloc( (void**)&dev_a,
                          size * sizeof( *dev_a ) ) );
};

首先为size个整数分别分配主机缓冲区和GPU缓冲区，然后执行100次复制操作，并由参数up来指定复制方向，在完成复制操作后停止计时器。

HANDLE_ERROR(udaEventRecord( start, 0 ) );   
for (int i=0; i<100; i++) { if (up)   
    HANDLE_ERROR(udaMemcpy(dev_a, a, size * sizeof(*dev_a),udaMemcpyHostToDevice ) ); else   
    HANDLE_ERROR(udaMemcpy(a, dev_a,

size \* sizeof( \*dev_a ),udaMemcpyDeviceToHost ) );}   
HANDLE_ERROR(udaEventRecord(stop,0)）;   
HANDLE_ERROR(udaEventSynchronize(stop))；   
HANDLE_ERROR(udaEventElapsedTime(& elapsedTime, start，stop)）;

在执行了100次复制操作后，释放主机缓冲区和GPU缓冲区，并且销毁计时事件。

free(a);   
HANDLE_ERROR(udaFree(dev_a)）;   
HANDLE_ERROR(udaEventDestroy（start））;   
HANDLE_ERROR(udaEventDestroy（stop））;   
return elapsedTime;

函数cuda_malloc_test()通过标准的C函数malloc()来分配可分页主机内存，在分配固定内存时则使用了cudaHostAlloc()。

float CUDA_host_alloc_test(int size, bool up) {
    CUDAEvent_t start, stop;
    int *a, *dev_a;
    float elapsedTime;
    HANDLE_ERROR( CUDAEventCreate( &start ) );
    HANDLE_ERROR( CUDAEventCreate( &stop ) );
    HANDLE_ERROR( CUDAHostAlloc( (void**)&a, size * sizeof(*a), CUDAHostAllocDefault ) );
    HANDLE_ERROR( CUDAMalloc( (void**)&dev_a, size * sizeof(*dev_a) ) );
    HANDLE_ERROR( CUDAEventRecord( start, 0 ) );
    for (int i=0; i<100; i++) {
        if (up)
            HANDLE_ERROR( CUDAMemcpy( dev_a, a, size * sizeof(*a), CUDAMemcpyHostToDevice ) );
        else
            HANDLE_ERROR( CUDAMemcpy( a, dev_a, size * sizeof(*a), CUDAMemcpyDeviceToHost ) );
    }
};

}   
HANDLE_ERROR(udaEventRecord(stop,0)）;   
HANDLE_ERROR(udaEventSynchronize(stop))；   
HANDLE_ERROR(udaEventElapsedTime(& elapsedTime, start，stop）);   
HANDLE_ERROR(udaFreeHost(a)）;   
HANDLE_ERROR(udaFree(dev_a)）;   
HANDLE_ERROR(udaEventDestroy（start））;   
HANDLE_ERROR(udaEventDestroy(stop)）;   
return elapsedTime;   
}

你可以看到，cudaHostAlloc()分配的内存与malloc()分配的内存在使用方式上是相同的。与malloc()的不同之处在于最后一个参数cudaHostAllocDefault。最后一个参数的取值范围是一组标志，我们可以通过这些标志来修改cudaHostAlloc()的行为，并分配不同形式的固定主机内存。在第11章中，我们将看到其他的标志值，但就目前而言，只需使用默认值，因此将参数指定为cudaHostAllocDefault以获得默认的行为。要释放cudaHostAlloc()分配的内存，必须使用cudaFreeHost()。也就是说，正如每次调用malloc()后都需要调用一次free()一样，在每次调用cudaHostAlloc()后，也需要调用一次cudaFreeHost()。

main()函数的代码如下所示。

include"../common/book.h"   
#defineSIZE (10\*1024\*1024)   
int main(void){ float elapsedTime; float MB  $=$  (float)100\*SIZE\*sizeof(int)/1024/1024; elapsedTime  $\equiv$  CUDA_malloc_test(SIZE, true); printf("Time using CUDAMalloc:  $\text{忍} 3 . 1 f$  ms\n", elapsedTime）; printf("\tMB/s during copy up:  $\text{忍} 3 . 1 f \backslash \mathsf { n } "$  ， MB/( elapsedTime/1000))；

由于cuda_malloc_test()的参数up为true，因此前一次调用将测试从主机到设备（或者说“上升”到设备）的复制性能。要测试相反方向的性能，可以执行相同的调用，只需将第二个参数指定为false。

elapsedTime = CUDA_malloc_test( SIZE, false );
printf( "Time using CUDAAlloc: %3.1f ms\n",

elapsedTime ); printf( "\tMB/s during copy down: %3.1f\n", MB/( elapsedTime/1000 ) );

我们执行了相同的步骤来测试cudaHostAlloc()的性能，将cuda_host_alloc_test()调用两次，一次将up指定为true，另一次指定为false。

elapsedTime =uda_host_alloc_test(SIZE, true); printf("Time using daemonHostAlloc: %3.1f ms\n", elapsedTime); printf("\tMB/s during copy up: %3.1f\n", MB/( elapsedTime/1000)); elapsedTime =uda_host_alloc_test(SIZE, false); printf("Time using daemonHostAlloc: %3.1f ms\n", elapsedTime); printf("\tMB/s during copy down: %3.1f\n", MB/( elapsedTime/1000));

在GeForce GTX 285上，当使用固定内存而不是可分页内存时，我们观察到从主机复制到设备的性能从2.77GB/s提升到5.11GB/s。当从设备复制到主机时，获得的性能提升是类似的，即从2.43GB/s提升到5.46GB/s。因此，对于大多数PCIE带宽有限的应用程序，当使用固定内存而不是标准可分页内存时，可以观察到显著的性能提升。但页锁定内存的作用并不仅限于性能提升。在后面的章节中会看到，在一些特殊情况中也需要使用页锁定内存。

10.2_页锁定主机内存

10.2 页锁定主机内存