CUDA 编程手册系列附录N – CUDA的统一内存（十五）

NVIDIA英伟达中国

已认证帐号

N.3. Performance Tuning

为了使用统一内存实现良好的性能，必须满足以下目标：

应避免错误：虽然可重放错误是启用更简单的编程模型的基础，但它们可能严重损害应用程序性能。故障处理可能需要几十微秒，因为它可能涉及 TLB 无效、数据迁移和页表更新。与此同时，应用程序某些部分的执行将停止，从而可能影响整体性能。
数据应该位于访问处理器的本地：如前所述，当数据位于访问它的处理器本地时，内存访问延迟和带宽明显更好。因此，应适当迁移数据以利用较低的延迟和较高的带宽。
应该防止内存抖动：如果数据被多个处理器频繁访问并且必须不断迁移以实现数据局部性，那么迁移的开销可能会超过局部性的好处。应尽可能防止内存抖动。如果无法预防，则必须进行适当的检测和解决。

为了达到与不使用统一内存相同的性能水平，应用程序必须引导统一内存驱动子系统避免上述陷阱。值得注意的是，统一内存驱动子系统可以检测常见的数据访问模式并自动实现其中一些目标，而无需应用程序参与。但是，当数据访问模式不明显时，来自应用程序的明确指导至关重要。 CUDA 8.0 引入了有用的 API，用于为运行时提供内存使用提示 ( cudaMemAdvise() ) 和显式预取 ( cudaMemPrefetchAsync() )。这些工具允许与显式内存复制和固定 API 相同的功能，而不会恢复到显式 GPU 内存分配的限制。

注意：Tegra 设备不支持 cudaMemPrefetchAsync() 。

N.3.1. Data Prefetching

数据预取意味着将数据迁移到处理器的内存中，并在处理器开始访问该数据之前将其映射到该处理器的页表中。数据预取的目的是在建立数据局部性的同时避免故障。这对于在任何给定时间主要从单个处理器访问数据的应用程序来说是最有价值的。由于访问处理器在应用程序的生命周期中发生变化，因此可以相应地预取数据以遵循应用程序的执行流程。由于工作是在 CUDA 中的流中启动的，因此预计数据预取也是一种流操作，如以下 API 所示：

cudaError_t cudaMemPrefetchAsync(const void *devPtr, 
                                     size_t count, 
                                     int dstDevice, 
                                     cudaStream_t stream);

其中由 devPtr 指针和 count 字节数指定的内存区域，ptr 向下舍入到最近的页面边界， count 向上舍入到最近的页面边界，通过在流中排队迁移操作迁移到 dstDevice 。为 dstDevice 传入 cudaCpuDeviceId 会导致数据迁移到 CPU 内存。考虑下面的一个简单代码示例：

void foo(cudaStream_t s) {
  char *data;
  cudaMallocManaged(&data, N);
  init_data(data, N);                                   // execute on CPU
  cudaMemPrefetchAsync(data, N, myGpuId, s);            // prefetch to GPU
  mykernel<<<..., s>>>(data, N, 1, compare);            // execute on GPU
  cudaMemPrefetchAsync(data, N, cudaCpuDeviceId, s);    // prefetch to CPU
  cudaStreamSynchronize(s);
  use_data(data, N);
  cudaFree(data);
}

如果没有性能提示，内核 mykernel 将在首次访问数据时出错，这会产生额外的故障处理开销，并且通常会减慢应用程序的速度。通过提前预取数据，可以避免页面错误并获得更好的性能。此 API 遵循流排序语义，即迁移在流中的所有先前操作完成之前不会开始，并且流中的任何后续操作在迁移完成之前不会开始。

N.3.2. Data Usage Hints

当多个处理器需要同时访问相同的数据时，单独的数据预取是不够的。在这种情况下，应用程序提供有关如何实际使用数据的提示很有用。以下咨询 API 可用于指定数据使用情况：

cudaError_t cudaMemAdvise(const void *devPtr, 
                              size_t count, 
                              enum cudaMemoryAdvise advice, 
                              int device);

其中，为从 devPtr 地址开始的区域中包含的数据指定的通知和计数字节的长度，四舍五入到最近的页面边界，可以采用以下值：

cudaMemAdviseSetReadMostly ：这意味着数据大部分将被读取并且只是偶尔写入。这允许驱动程序在处理器访问数据时在处理器内存中创建数据的只读拷贝。同样，如果在此区域上调用 cudaMemPrefetchAsync ，它将在目标处理器上创建数据的只读拷贝。当处理器写入此数据时，相应页面的所有副本都将失效，但发生写入的拷贝除外。此建议忽略设备参数。该建议允许多个处理器以最大带宽同时访问相同的数据，如以下代码片段所示：

char *dataPtr;
size_t dataSize = 4096;
// Allocate memory using malloc or cudaMallocManaged
dataPtr = (char *)malloc(dataSize);
// Set the advice on the memory region
cudaMemAdvise(dataPtr, dataSize, cudaMemAdviseSetReadMostly, 0);
int outerLoopIter = 0;
while (outerLoopIter < maxOuterLoopIter) {
    // The data is written to in the outer loop on the CPU
    initializeData(dataPtr, dataSize);
    // The data is made available to all GPUs by prefetching.
    // Prefetching here causes read duplication of data instead
    // of data migration
    for (int device = 0; device < maxDevices; device++) {
        cudaMemPrefetchAsync(dataPtr, dataSize, device, stream);
    // The kernel only reads this data in the inner loop
    int innerLoopIter = 0;