为啥在内核启动之前开始数据迁移（D 到 H）？

Posted 2023-03-23

【中文标题】为啥在内核启动之前开始数据迁移（D 到 H）？

【英文标题】：Why data migration (D to H) started before kernel launch?为什么在内核启动之前开始数据迁移（D 到 H）？

【发布时间】：2022-01-07 03:37:30

【问题描述】：

以下是基于 CUDA 托管内存的代码的 nvvp 分析结果：

问题： 如上面的屏幕截图所示，数据迁移（D 到 H）甚至在内核启动之前就开始了。尽管与需要传输的总数据（大约 4 MB）相比，第一个数据迁移块（在内核启动之前开始）的大小非常小（大约 450kb）。那么，这些数据究竟是什么？如何在发布前进行数据迁移（D 到 H）？

我的代码：

#include <iostream>
#include <math.h>

#include "device_launch_parameters.h"
#include "cuda_runtime.h"

// CUDA kernel to add elements of two arrays
__global__
void add(int n, float* x, float* y)

    int index = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride)
        y[i] = x[i] + y[i];



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 main()

    //RunUnmanagedVersion();

    RunManagedVersion();

    return 0;

【参考方案1】：

分析器不会将这种类型的 D->H 转移定位到特定的时间点。

相反，它将时间线分成大致固定的时间段，并确定在该段期间是否发生了任何 D->H 传输。如果是这样，则该段将被报告为“其中”有 D->H 传输，并且该段将用颜色编码以指示该段内传输的相对“密度”。

因此，您所指的绿色条的开头不一定是发生 D->H 转移的时刻，而是发生某些 D->H 转移活动的测量段的开始.

由于绿色条在内核完成后很好地延伸，显然是内核完成后 D->H 传输活动的原因，这是预期的。

（可以在内核启动之前、期间和之后就 H->D 线和蓝条提出类似的问题。显然不需要 H->D 活动之后 内核启动。同样，根据固定的时间段测量和报告活动。）

顺便说一下（对于未来的读者），在内核启动之前的数据迁移 (D->H) 已在here 进行了解释。

【讨论】：

谢谢，我也认为时间线不能准确表示数据传输的实际开始/停止，但想确认我没有遗漏任何东西。