为何CUDA托管内存代码内核启动前会出现D to H数据迁移

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

提问内容 #

如上述截图所示，设备到主机（D to H）的数据迁移在内核启动前就已开始，且首个启动的数据迁移块大小约450KB，远小于总需传输的约4MB数据。请问该部分迁移的数据具体是什么？为什么D to H数据迁移会在内核启动前发生？

测试代码 #

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

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

// 实现两个数组元素相加的CUDA内核
__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;

    // 分配统一内存，CPU和GPU均可访问
    cudaMallocManaged(&x, N * sizeof(float));
    cudaMallocManaged(&y, N * sizeof(float));

    // 在主机端初始化x和y数组
    for (int i = 0; i < N; i++) {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }

    // 在GPU上启动内核处理1M个元素
    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    add << <numBlocks, blockSize >> > (N, x, y);

    // 等待GPU执行完成，再在主机端访问数据
    cudaDeviceSynchronize();

    // 校验错误，所有值应为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;

    // 释放内存
    cudaFree(x);
    cudaFree(y);

    return 0;
}

int main()
{
    //RunUnmanagedVersion();

    RunManagedVersion();

    return 0;
}

问题解答 #

1. 450KB迁移数据的内容
   这部分数据和你代码中定义的x、y业务数组完全无关，属于CUDA驱动和运行时的内部控制数据，具体包括：

• 当前GPU设备上下文的状态同步元数据
• 统一内存（UM）的页表映射管理结构
• add内核的二进制镜像元信息、参数校验配置
• nvvp剖析工具注入的性能采集控制块

2. D2H迁移在内核启动前触发的原因
   CUDA程序执行到首次内核启动前，会完成驱动侧的初始化收尾工作：

• 设备端初始化完成后，需要把上下文控制信息回传到主机运行时，用来匹配主机侧的进程地址空间
• nvvp做性能剖析时，需要先从设备读取采集模块的初始状态，保证后续时间线统计的准确性
• 统一内存管理模块需要先同步设备端的页错误处理回调配置，才能在后续内核执行过程中正确处理按需页迁移

补充说明：你代码中的4MB业务数据只会产生两类和业务相关的迁移：内核启动时主机写过的x、y页从主机到设备（H2D）的预迁移，以及内核执行完成后主机端校验时y数组从设备到主机（D2H）的迁移，这两部分在nvvp时间线上会出现在内核执行的前后，和你提前看到的450KB内部数据迁移没有关联。

内容的提问来源于stack exchange，提问作者skm