CUDA 统一内存情况下的 DeviceToHost 和 HostToDevice 时间

Posted 2023-03-23

【中文标题】CUDA 统一内存情况下的 DeviceToHost 和 HostToDevice 时间

【英文标题】：DeviceToHost and HostToDevice times in case of CUDA Unified Memory

【发布时间】：2022-01-06 14:18:40

【问题描述】：

我正在尝试比较 CUDA 内存管理的托管和非托管版本所花费的总执行时间。

在下面的示例代码中，我有两个功能完全相同。唯一不同的是它们的内存管理。一个函数使用cudaMalloc()/cudaMemcpy()，另一种方法只使用cudaMallocManaged()。

我使用nvprof计算不同的时间，得到以下输出：

托管版本nvprof 输出：

== 29028 == Profiling result :
Type  Time(%)      Time     Calls       Avg       Min       Max  Name
GPU activities : 100.00 % 59.425us         1  59.425us  59.425us  59.425us  add(int, float*, float*)
API calls : 78.08 % 296.49ms         2  148.24ms  1.7127ms  294.78ms  cudaMallocManaged
19.61 % 74.451ms         1  74.451ms  74.451ms  74.451ms  cuDevicePrimaryCtxRelease
1.55 % 5.8705ms         1  5.8705ms  5.8705ms  5.8705ms  cudaLaunchKernel
0.67 % 2.5547ms         2  1.2774ms  974.40us  1.5803ms  cudaFree
0.07 % 280.60us         1  280.60us  280.60us  280.60us  cudaDeviceSynchronize
0.01 % 28.300us         3  9.4330us  3.0000us  13.300us  cuModuleUnload
0.01 % 26.800us         1  26.800us  26.800us  26.800us  cuDeviceTotalMem
0.00 % 17.700us       101     175ns     100ns     900ns  cuDeviceGetAttribute
0.00 % 10.100us         3  3.3660us     300ns  8.8000us  cuDeviceGetCount
0.00 % 3.2000us         1  3.2000us  3.2000us  3.2000us  cuDeviceGetName
0.00 % 3.0000us         2  1.5000us     300ns  2.7000us  cuDeviceGet
0.00 % 500ns         1     500ns     500ns     500ns  cuDeviceGetLuid
0.00 % 200ns         1     200ns     200ns     200ns  cuDeviceGetUuid

== 29028 == Unified Memory profiling result :
Device "GeForce GTX 1070 (0)"
Count  Avg Size  Min Size  Max Size  Total Size  Total Time  Name
64  128.00KB  128.00KB  128.00KB  8.000000MB  3.279000ms  Host To Device
146  84.164KB  32.000KB  1.0000MB  12.00000MB  64.50870ms  Device To Host

非托管版本nvprof 输出：

== 23864 == Profiling result :
Type  Time(%)      Time     Calls       Avg       Min       Max  Name
GPU activities : 56.30 % 1.5032ms         2  751.60us  751.44us  751.76us[CUDA memcpy HtoD]
41.48 % 1.1075ms         1  1.1075ms  1.1075ms  1.1075ms[CUDA memcpy DtoH]
2.23 % 59.457us         1  59.457us  59.457us  59.457us  add(int, float*, float*)
API calls : 78.92 % 270.08ms         2  135.04ms  656.40us  269.43ms  cudaMalloc
19.79 % 67.730ms         1  67.730ms  67.730ms  67.730ms  cuDevicePrimaryCtxRelease
1.05 % 3.5796ms         3  1.1932ms  1.0106ms  1.4341ms  cudaMemcpy
0.10 % 346.20us         2  173.10us  3.4000us  342.80us  cudaFree
0.09 % 314.30us         1  314.30us  314.30us  314.30us  cudaDeviceSynchronize
0.02 % 74.200us         1  74.200us  74.200us  74.200us  cudaLaunchKernel
0.01 % 34.700us         3  11.566us  2.5000us  29.100us  cuModuleUnload
0.01 % 24.100us         1  24.100us  24.100us  24.100us  cuDeviceTotalMem
0.00 % 17.100us       101     169ns     100ns     900ns  cuDeviceGetAttribute
0.00 % 9.0000us         3  3.0000us     300ns  8.0000us  cuDeviceGetCount
0.00 % 3.2000us         1  3.2000us  3.2000us  3.2000us  cuDeviceGetName
0.00 % 1.5000us         2     750ns     200ns  1.3000us  cuDeviceGet
0.00 % 300ns         1     300ns     300ns     300ns  cuDeviceGetUuid
0.00 % 300ns         1     300ns     300ns     300ns  cuDeviceGetLuid

我的代码：

int RunManagedVersion()

    int N = 1 << 20;
    float* x, * y;

    // Allocate Unified Memory -- accessible from CPU or GPU
    cudaMallocManaged(&x, N * sizeof(float));
    cudaMallocManaged(&y, N * sizeof(float));

    // initialize x and y arrays on the host
    for (int i = 0; i < N; i++) 
        x[i] = 1.0f;
        y[i] = 2.0f;
    

    // Launch kernel on 1M elements on the GPU
    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    add << <numBlocks, blockSize >> > (N, x, y);

    // Wait for GPU to finish before accessing on host
    cudaDeviceSynchronize();

    // Check for errors (all values should be 3.0f)
    float maxError = 0.0f;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(y[i] - 3.0f));

    std::cout << "Max error: " << maxError << std::endl;

    // Free memory
    cudaFree(x);
    cudaFree(y);

    return 0;


int RunUnmanagedVersion()

    int N = 1 << 20;

    //Declare pointers for input and output arrays
    float* x = (float*)calloc(N, sizeof(float));
    float* y = (float*)calloc(N, sizeof(float));

    // initialize x and y arrays on the host
    for (int i = 0; i < N; i++) 
        x[i] = 1.0f;
        y[i] = 2.0f;
    

    //Allocate device memory for input and output images
    float* d_pX = 0;
    float* d_pY = 0;
    cudaMalloc(&d_pX, N * sizeof(float));
    cudaMalloc(&d_pY, N * sizeof(float));

    //Copy INPUT ARRAY data from host to device
    cudaMemcpy(d_pX, x, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_pY, y, N * sizeof(float), cudaMemcpyHostToDevice);


    // Launch kernel on 1M elements on the GPU
    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    add << <numBlocks, blockSize >> > (N, d_pX, d_pY);

    // Wait for GPU to finish before accessing on host
    cudaDeviceSynchronize();

    //Copy Results - Device to Host
    cudaMemcpy(y, d_pY, N * sizeof(float), cudaMemcpyDeviceToHost);

    // Check for errors (all values should be 3.0f)
    float maxError = 0.0f;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(y[i] - 3.0f));

    std::cout << "Max error: " << maxError << std::endl;

    // device memory free
    cudaFree(d_pX);
    cudaFree(d_pX);

    //host memory free
    free(x);
    free(y);

    return 0;


int main()

    RunUnmanagedVersion();

    //RunManagedVersion();

    return 0;

问题：我多次使用上述代码，并注意到在托管版本（即统一内存）的情况下，来自DeviceToHost 的数据传输时间要长得多。这是正常的（为什么？）还是我在代码中做错了什么？

【问题讨论】：

我不相信任何这些都被记录或指定，所以关于“为什么”的权威答案是不可能的，我不相信。你没有做错什么。您对 D->H 时间的解释与非托管情况不能直接比较，因为这些传输是与主机代码同时执行的。这些细节可以从探查器中观察到，但确切的特征没有记录，AFAIK。在 H->D 的情况下，内核启动会触发传输。在 H->D 情况下，传输由cudaDeviceSynchronize() 启用，但由主机代码活动驱动

【参考方案1】：

我相信我在这里要说的是：

未记录，因此可能会更改，但 通过仔细分析可以观察到

此外，所有这些 cmets 都特定于与 Windows UM 使用相关的 UM 机制，或者与 pre-pascal GPU 的 linux 使用相关联。

在这种 pre-pascal 和/或 windows UM 机制中，从主机到设备的数据传输是在内核启动时启动的。这将表现为内核启动过程中的延迟（从请求内核启动到内核代码实际开始执行）。

在这种情况下，UM 系统以固定大小的块传输数据。从您的分析器输出中可以看出这一点：

Count  Avg Size  Min Size  Max Size  Total Size  Total Time  Name
64  128.00KB  128.00KB  128.00KB  8.000000MB  3.279000ms  Host To Device

我们会得出结论，因为确实有 8MB 的数据需要传输到设备，而在非 UM 的情况下，它似乎在大约 1.5 毫秒内发生：

GPU activities : 56.30 % 1.5032ms         2  751.60us  751.44us  751.76us[CUDA memcpy HtoD]

与非 UM 情况相比，即使是 H->D 情况在 Windows 上的性能也有所降低。我将此归因于需要（无论出于何种原因）以相对较小的 128KB 块传输 8MB 数据。另外，在windows的情况下，WDDM对这块GPU有直接的控制权，而CUDA实际上是WDDM对这些活动的“客户端”，尤其是与内存相关的活动。 WDDM 完全有可能决定在数据传输期间对 GPU 做一些事情，并且可能插入了一些空白或效率低下。

在 D->H 的情况下，在 windows 上，情况似乎有所不同，而且可以说更糟。但是，我们必须小心评估这里发生的情况。第一个问题可能是：

为什么要传输 12MB 的数据 D->H？

似乎有几点需要注意：

UM 分配似乎首先出现在设备内存中。 （这与请求分页的情况不同！）这意味着如果您要做的第一件事是访问主机代码中的 UM 分配，那么分配必须从设备转移到主机。这占 12MB 传输量中的 8MB，如果您进行仔细的分析实验，您可以相信这一点。 设备上的 UM 分配似乎根据主机代码活动转移到主机代码。如果仔细考虑上面的第一个项目符号，这是不言而喻的。但即使我们只关注内核活动后的数据传输，也很容易通过分析器实验来说服自己，如果内核启动后没有主机代码真正访问数据，那么那里就不会发生传输。

上面的第二条表示我们可以假设 D->H 转移可能是：

在实际上以某种方式导致这些传输的主机代码的持续时间内“被抹去” 不知何故与主机代码“同时”发生。

我们还可以得出结论，报告的 UM D->H 活动中只有 1/3 实际发生在内核调用之后，因此我们可能选择仅将该部分与来自非UM 案例。

所有这一切的结果是，我认为仅通过查看我在上面摘录的数据类型来比较这两种情况并不是一件小事。是的，UM 案例的性能可能确实比非 UM 案例差。 CUDA 文档中没有任何地方说明这些预期性能相同。不，你没有做任何“错误”的事情。

FWIW，Linux 上的 maxwell/kepler UM 案例看起来很多比它在 Windows WDDM 上的表现要好，所以我认为 WDDM 可能也涉及效率较低的行为。