cuda - 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

Answer 1

Congratulations! You're encountering a 2.x compute capability + TCC + 64-bit OS feature with newer CUDA versions :)

阅读其余部分以了解更多信息!

首先是CUDA教给我们的一个小理论总结:

固定内存 不是零拷贝，因为 GPU 无法访问它(它没有映射到它的地址空间中)并且它被用来有效地从主机传输到 GPU。它是页面锁定(宝贵的内核资源)内存，与可分页的普通内存相比具有一些性能优势。

固定零拷贝内存 是页锁定内存(通常用 cudaHostAllocMapped 标志分配)，它也被 GPU 使用，因为它映射到它的地址空间。

为什么要从设备访问主机分配的内存而没有明确指定？

查看 CUDA 4.0(及更高版本)的发行说明:

• (Windows and Linux) Added support for unified virtual address space.

Devices supporting 64-bit and compute 2.0 and higher capability now share a single unified address space between the host and all devices. This means that the pointer used to access memory on the host is the same as the pointer to used to access memory on the device. Therefore, the location of memory may be queried directly from its pointer value; the direction of a memory copy need not be specified.

总结一下:如果您的卡是 2.0+(并且是:https://developer.nvidia.com/cuda-gpus)，您运行的是 64 位操作系统并且在 Windows 上您打开了 TCC 模式，那么您将自动使用 UVA ( Unified Virtual Addressing ) 主机和设备之间 .这意味着:通过类似零复制的访问自动增强您的代码。

这也在CUDA documentation for the current version在“主机分配的主机内存的自动映射”段落中