CPU 在 Tegra TK1 上使用 malloc() 与 cudaHostAlloc() 分配的数据的内存访问延迟
CPU memory access latency of data allocated with malloc() vs. cudaHostAlloc() on Tegra TK1
我正在执行一个简单的测试，比较使用 malloc() 分配的数据和来自主机的 cudaHostAlloc() 分配的数据的访问延迟（cpu 正在执行访问）。
我注意到访问使用 cudaHostAlloc() 分配的数据比访问使用 Jetson Tk1 上的 malloc() 分配的数据慢得多。
离散 GPU 并非如此，似乎仅适用于 TK1。经过一些调查，我发现用 cudaHostAlloc() 分配的数据被内存映射 (mmap) 到进程地址 space 的 /dev/nvmap 区域。对于映射到进程堆上的普通 malloc 数据，情况并非如此。我知道此映射可能是允许 GPU 访问数据所必需的，因为 cudaHostAlloc 的数据必须从主机和设备都可见。
我的问题如下：
从主机访问 cudaHostAlloc 数据的开销从何而来？
映射到 /dev/nvmap 的数据是否在 CPU 缓存中未缓存？
我相信我已经找到了这种行为的原因。经过进一步调查（使用 Linux trace events and looking at the nvmap driver code）后，我发现开销的来源是使用 cudaHostAlloc() 分配的数据使用 NVMAP_HANDLE_UNCACHEABLE 标志标记为 "uncacheable"。调用 pgprot_noncached() 以确保相关 PTE 被标记为不可缓存。
主机访问使用cudaMallocManaged()分配的数据的行为是不同的。数据将被缓存（使用标志 NVMAP_HANDLE_CACHEABLE）。因此从主机访问此数据相当于 malloc()'d 数据。同样重要的是要注意，CUDA 运行时不允许设备 (GPU) 访问与主机同时使用 cudaMallocManaged() 分配的任何数据，这样的操作会产生段错误。但是，运行时允许并发访问设备和主机上的 cudaHostAlloc()'d 数据，我相信这是使 cudaHostAlloc()'d 数据不可缓存的原因之一。
我正在执行一个简单的测试，比较使用 malloc() 分配的数据和来自主机的 cudaHostAlloc() 分配的数据的访问延迟（cpu 正在执行访问）。 我注意到访问使用 cudaHostAlloc() 分配的数据比访问使用 Jetson Tk1 上的 malloc() 分配的数据慢得多。
离散 GPU 并非如此，似乎仅适用于 TK1。经过一些调查，我发现用 cudaHostAlloc() 分配的数据被内存映射 (mmap) 到进程地址 space 的 /dev/nvmap 区域。对于映射到进程堆上的普通 malloc 数据，情况并非如此。我知道此映射可能是允许 GPU 访问数据所必需的，因为 cudaHostAlloc 的数据必须从主机和设备都可见。
我的问题如下： 从主机访问 cudaHostAlloc 数据的开销从何而来？ 映射到 /dev/nvmap 的数据是否在 CPU 缓存中未缓存？
我相信我已经找到了这种行为的原因。经过进一步调查（使用 Linux trace events and looking at the nvmap driver code）后，我发现开销的来源是使用 cudaHostAlloc() 分配的数据使用 NVMAP_HANDLE_UNCACHEABLE 标志标记为 "uncacheable"。调用 pgprot_noncached() 以确保相关 PTE 被标记为不可缓存。
主机访问使用cudaMallocManaged()分配的数据的行为是不同的。数据将被缓存（使用标志 NVMAP_HANDLE_CACHEABLE）。因此从主机访问此数据相当于 malloc()'d 数据。同样重要的是要注意，CUDA 运行时不允许设备 (GPU) 访问与主机同时使用 cudaMallocManaged() 分配的任何数据，这样的操作会产生段错误。但是，运行时允许并发访问设备和主机上的 cudaHostAlloc()'d 数据，我相信这是使 cudaHostAlloc()'d 数据不可缓存的原因之一。