在TegraTK1上使用malloc()与cudaHostAlloc()分配的数据的CPU内存访问延迟

我正在执行一个简单的测试,它比较了使用malloc()分配的数据的访问延迟和使用cudaHostAlloc()从主机分配的数据(cpu正在执行访问).我注意到访问使用cudaHostAlloc()分配的数据比访问Jetson Tk1上的malloc()分配的数据要慢得多.

这不是离散GPU的情况,似乎只适用于TK1.经过一些调查,我发现用cudaHostAlloc()分配的数据是内存映射(mmap)到进程地址空间的/ dev/nvmap区域.对于映射在进程堆上的普通malloc数据,情况并非如此.我知道这个映射可能是允许GPU访问数据所必需的,因为必须从主机和设备都可以看到cudaHostAlloc的数据.

我的问题如下:从主机访问cudaHostAlloc数据的开销来自哪里？数据映射到CPU缓存上的/ dev/nvmap是否未缓存？

我相信我已经找到了这种行为的原因.在进一步调查之后(使用Linux跟踪事件并查看nvmap驱动程序代码),我发现开销的来源cudaHostAlloc()是使用该NVMAP_HANDLE_UNCACHEABLE标志将分配的数据标记为"不可缓存" 的事实.呼叫pgprot_noncached()是为了确保相关的PTE标记为不可缓存.

主机访问使用的数据的行为cudaMallocManaged()是不同的.数据将被缓存(使用标志NVMAP_HANDLE_CACHEABLE).因此,从主机访问此数据等同于malloc()'d数据.同样重要的是要注意,CUDA运行时不允许设备(GPU)访问cudaMallocManaged()与主机同时分配的任何数据,并且此类操作将生成段错误.但是,运行时允许并发访问cudaHostAlloc()'d设备和主机上的数据,我相信这是使cudaHostAlloc()'d数据不可缓存的原因之一.