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

Question

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

独立 GPU 并非如此，似乎仅适用于 TK1。经过一些调查，我发现用 cudaHostAlloc() 分配的数据被内存映射 (mmap) 到进程地址空间的/dev/nvmap 区域。对于映射到进程堆上的普通 malloc 数据，情况并非如此。我知道此映射可能是允许 GPU 访问数据所必需的，因为 cudaHostAlloc 的数据必须从主机和设备可见。

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

Answer 1

我相信我已经找到了这种行为的原因。经过进一步调查(使用 Linux trace events 并查看 nvmap driver code )我发现开销的来源来自于使用 cudaHostAlloc() 分配的数据被标记为“不可缓存”的事实 NVMAP_HANDLE_UNCACHEABLE 标志。调用 pgprot_noncached() 以确保相关的 PTE 被标记为不可缓存。

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