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

恭喜!你遇到了更新的CUDA版本的2.x计算能力+ TCC + 64位操作系统功能:)

阅读其余内容以了解更多信息!

首先是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(统一虚拟寻址).这意味着:使用类似零拷贝的访问自动增强代码.

这也是"主机分配的主机内存的自动映射"段落中当前版本的CUDA文档中的内容

Answer 2

映射内存是一种固定内存.它是在您固定内存并传递cudaHostAllocMapped标志时创建的. 但是,即使您已指定cudaHostAllocDefault,在某些情况下内存也会"映射".我相信TCC模式与64位操作系统相结合足以满足"自动映射"功能所需的环境.

关键问题是UVA是否有效.在你的情况下,它是.

关于为何具有显式能力的问题,它适用于UVA无效的情况(例如在32位主机OS中).

从文档中(当UVA生效时):

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

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