CUDA实战笔记(14)——页锁定(Page-Locked)主机内存

本系列代码托管于：https://github.com/chintsan-code/cuda_by_example

本篇使用的项目为：copy_timed

之前，使用cudaMalloc()在GPU上分配内存，使用malloc()在主机分配内存。除此之外，还可以使用cudaHostAlloc()来分配主机内存。

malloc()和cudaHostAlloc()的区别：

• malloc()分配的是标准的、可分页的(Pagable)内存
• cudaHostAlloc()分配的是页锁定(Page-Locked)内存，也成为固定内存(Pinned Memory)或者不可分页内存。

页锁定(Page-Locked)内存：

操作系统不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中。因此操作系统能够安全地使某个应用程序访问该内存的物理地址，因为这块内存将不会被破坏或者重新定位。

由于GPU知道内存的物理地址，因此可以通过”直接内存访问(Direct Memory Access，DMA)“技术在GPU和主机之间复制数据。

CPU很可能会在DMA执行过程中将目标内存交换到磁盘上，或者通过更新系统的可分页来重新定位目标内存的物理地址。CPU可能会移动可分页的数据，这就可能对DMA操作造成延迟。因此在DMA复制过程中使用固定内存是非常重要的。（事实上，当使用可分页内存进行复制时，CUDA驱动程序仍然会通过DMA把数据传输给CPU。因此复制会执行两遍：第一遍从可分页内存复制一份到”临时的”页锁定内存，然后再从这个页锁定内存复制到GPU）

因此当在GPU和主机之间复制数据时，使用页锁定主机内存的性能比标准可分页内存的性能要高大约2倍。

但是，使用页锁定内存时，将会失去虚拟内存的所有功能。特别是，在应用程序中使用每个页锁定内存时都需要分配物理内存，因为这些内存不能交换到磁盘上。这意味着，与使用标准的malloc()调用相比，系统将更快地耗尽内存。

因此，应该只对cudaMemcpy()调用中的源内存或目标内存使用用页锁定内存，并且在不使用时立即释放。

使用malloc()分配内存：

int* a;
a = (int*)malloc(size * sizeof(int));

使用cudaHostAlloc()分配内存：

int* a;
cudaHostAlloc((void**)&a, size * sizeof(int), cudaHostAllocDefault);

释放使用malloc()分配的内存：

free(a);

释放使用cudaHostAlloc()分配的内存：

cudaFreeHost(a);

完整代码如下：

// copy_timed

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include "../../common/book.h"

#include <stdio.h>

#define SIZE    (64*1024*1024)

float cuda_malloc_test(int size, bool up) {
    cudaEvent_t start, end;
    int *a, *dev_a;
    float elapsedTime;

    HANDLE_ERROR(cudaEventCreate(&start));
    HANDLE_ERROR(cudaEventCreate(&end));
    
    // 分配主机缓冲区和GPU缓冲区
    a = (int*)malloc(size * sizeof(int));  // 使用标准C函数malloc()来分配可分页主机内存
    HANDLE_NULL(a);
    HANDLE_ERROR(cudaMalloc((void**)&dev_a, size * sizeof(int)));

    HANDLE_ERROR(cudaEventRecord(start, 0));
    // 执行100次复制操作，并由参数up指定复制方向
    for (int i = 0; i < 100; i++) {
        if (up) {
            // cudaMemcpyHostToDevice
            HANDLE_ERROR(cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice));
        }else{
            // cudaMemcpyDeviceToHost
            HANDLE_ERROR(cudaMemcpy(a, dev_a, size * sizeof(int), cudaMemcpyDeviceToHost));
        }
    }
    HANDLE_ERROR(cudaEventRecord(end, 0));
    HANDLE_ERROR(cudaEventSynchronize(end));
    HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, end));

    free(a);
    HANDLE_ERROR(cudaFree(dev_a));
    HANDLE_ERROR(cudaEventDestroy(start));
    HANDLE_ERROR(cudaEventDestroy(end));

    return elapsedTime;
}

float cuda_host_alloc_test(int size, bool up) {
    cudaEvent_t start, end;
    int* a, * dev_a;
    float elapsedTime;

    HANDLE_ERROR(cudaEventCreate(&start));
    HANDLE_ERROR(cudaEventCreate(&end));

    // 分配主机缓冲区和GPU缓冲区
    HANDLE_ERROR(cudaHostAlloc((void**)&a, size * sizeof(int), cudaHostAllocDefault));  // 使用cudaHostAlloc()来分配固定内存
    HANDLE_NULL(a);
    HANDLE_ERROR(cudaMalloc((void**)&dev_a, size * sizeof(int)));

    HANDLE_ERROR(cudaEventRecord(start, 0));
    // 执行100次复制操作，并由参数up指定复制方向
    for (int i = 0; i < 100; i++) {
        if (up) {
            // cudaMemcpyHostToDevice
            HANDLE_ERROR(cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice));
        }
        else {
            // cudaMemcpyDeviceToHost
            HANDLE_ERROR(cudaMemcpy(a, dev_a, size * sizeof(int), cudaMemcpyDeviceToHost));
        }
    }
    HANDLE_ERROR(cudaEventRecord(end, 0));
    HANDLE_ERROR(cudaEventSynchronize(end));
    HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, end));

    HANDLE_ERROR(cudaFreeHost(a));  // 使用cudaFreeHost()释放由cudaHostAlloc()分配的内存
    HANDLE_ERROR(cudaFree(dev_a));
    HANDLE_ERROR(cudaEventDestroy(start));
    HANDLE_ERROR(cudaEventDestroy(end));

    return elapsedTime;
}

int main() {
    float elapsedTime;
    float MB = (float)100 * SIZE * sizeof(int) / 1024 / 1024;

    // 测试从Host到Device的复制性能(使用malloc分配的内存)
    elapsedTime = cuda_malloc_test(SIZE, true);
    printf("Time using malloc: %3.1f ms\n", elapsedTime);
    printf("\tMB/s during copy up: %3.1f\n", MB / (elapsedTime / 1000));

    // 测试从Device到Host的复制性能(使用malloc分配的内存)
    elapsedTime = cuda_malloc_test(SIZE, false);
    printf("Time using malloc: %3.1f ms\n", elapsedTime);
    printf("\tMB/s during copy down: %3.1f\n", MB / (elapsedTime / 1000));

    // 测试从Host到Device的复制性能(使用cudaHostAlloc分配的内存)
    elapsedTime = cuda_host_alloc_test(SIZE, true);
    printf("Time using cudaHostAlloc: %3.1f ms\n", elapsedTime);
    printf("\tMB/s during copy up: %3.1f\n", MB / (elapsedTime / 1000));

    // 测试从Device到Host的复制性能(使用cudaHostAlloc分配的内存)
    elapsedTime = cuda_host_alloc_test(SIZE, false);
    printf("Time using cudaHostAlloc: %3.1f ms\n", elapsedTime);
    printf("\tMB/s during copy down: %3.1f\n", MB / (elapsedTime / 1000));


    return 0;
}