CUDA固定内存是否为零拷贝？

Question

固定内存旨在增加从主机到设备的传输速率(api 参考文献)。然而，我发现我不需要调用cuMemcpyHtoD使内核可以访问值，或者调用cuMemcpyDtoA使主机可以读取值。我认为这种方法不会起作用，但实际上它确实有效：

__global__ void testPinnedMemory(double * mem)
{
    double currentValue = mem[threadIdx.x];
    printf("Thread id: %d, memory content: %f\n", threadIdx.x, currentValue);
    mem[threadIdx.x] = currentValue+10;
}

void test() 
{
    const size_t THREADS = 8;
    double * pinnedHostPtr;
    cudaHostAlloc((void **)&pinnedHostPtr, THREADS, cudaHostAllocDefault);

    //set memory values
    for (size_t i = 0; i < THREADS; ++i)
        pinnedHostPtr[i] = i;

    //call kernel
    dim3 threadsPerBlock(THREADS);
    dim3 numBlocks(1);
    testPinnedMemory<<< numBlocks, threadsPerBlock>>>(pinnedHostPtr);

    //read output
    printf("Data after kernel execution: ");
    for (int i = 0; i < THREADS; ++i)
        printf("%f ", pinnedHostPtr[i]);    
    printf("\n");
}

输出：

Data after kernel execution: 10.000000 11.000000 12.000000 13.000000 14.000000 15.000000 16.000000 17.000000
Thread id: 0, memory content: 0.000000
Thread id: 1, memory content: 1.000000
Thread id: 2, memory content: 2.000000
Thread id: 3, memory content: 3.000000
Thread id: 4, memory content: 4.000000
Thread id: 5, memory content: 5.000000
Thread id: 6, memory content: 6.000000
Thread id: 7, memory content: 7.000000

我的问题是：

• 固定内存是否是零拷贝？我认为只有映射的固定内存才是零拷贝。
• 如果它是零拷贝，为什么还需要一种显式的方式将其映射到设备（cudaHostAlloc使用cudaHostAllocMapped选项）

我正在使用CUDA Toolkit 5.5，Quadro 4000驱动程序设置为TCC模式，并编译选项sm_20，compute_20

- Budric

Answer 1

恭喜! 您正在遇到一个2.x计算能力+TCC+64位操作系统功能，使用新的CUDA版本 :) 阅读其余部分以了解更多信息！首先，作为CUDA教给我们的一个小理论总结：固定内存不是零拷贝，因为GPU无法访问它（它未映射到其地址空间），并且用于高效地从主机传输到GPU。它是页面锁定（有价值的内核资源）内存，并且比可分页的常规内存具有一些性能优势。固定的零拷贝内存是页面锁定内存（通常使用cudaHostAllocMapped标志分配），也被GPU使用，因为映射到其地址空间。为什么您在未明确指定的情况下从设备访问主机分配的内存？请查看CUDA 4.0（及更高版本）的发行说明：

(Windows和Linux) 添加了对统一虚拟地址空间的支持。 支持64位和计算能力2.0及更高版本的设备现在在主机和所有设备之间共享单个统一的地址空间。这意味着用于访问主机上的内存的指针与用于访问设备上的内存的指针相同。因此，可以直接从其指针值查询内存的位置；不需要指定内存复制的方向。 总之：如果您的卡是2.0+（并且它是：https://developer.nvidia.com/cuda-gpus），您正在运行64位操作系统并且在Windows上打开了TCC模式，则自动使用主机和设备之间的UVA（Unified Virtual Addressing）。这意味着：自动提高您的代码零拷贝式访问的能力。

这也在当前版本的CUDA文档中，在段落"自动映射主机分配的主机内存"中提到。

- Marco A.

Answer 2

映射内存是一种固定内存的类型。当您固定内存并传递cudaHostAllocMapped标志时，它会被创建。然而，即使您指定了cudaHostAllocDefault，也会在某些情况下将内存“映射”。我相信TCC模式结合64位操作系统足以满足“自动映射”功能所需的条件。

核心问题是UVA是否生效。在您的情况下，它生效了。

关于为什么需要明确的能力的问题，这是为了在UVA无效的情况下使用（例如在32位主机操作系统中）。

来自文档（当UVA生效时）：

自动映射主机分配的主机内存

通过cudaMallocHost()和cudaHostAlloc()分配的所有主机内存始终可以直接从支持统一寻址的所有设备访问。无论是否指定了标志cudaHostAllocPortable和cudaHostAllocMapped，都是如此。在支持统一寻址的所有设备上，在内核中访问已分配的主机内存的指针值与在主机上访问该内存的指针值相同。不需要调用cudaHostGetDevicePointer()来获取这些分配的设备指针。

- Robert Crovella