当 CPU 尝试读取由 GPU 初始化的托管内存时，为啥数据会从主机迁移到设备？

Posted 2023-03-23

【中文标题】当 CPU 尝试读取由 GPU 初始化的托管内存时，为啥数据会从主机迁移到设备？

【英文标题】：Why data migrate from Host to Device when CPU try to read a managed memory initialized by GPU?当 CPU 尝试读取由 GPU 初始化的托管内存时，为什么数据会从主机迁移到设备？

【发布时间】：2020-12-07 08:05:39

【问题描述】：

在下面的测试代码中，我们通过 GPU 初始化数据，然后通过 CPU 访问数据。我对 nvprof 的分析结果有 2 个问题。

为什么会有一个数据从主机迁移到设备？据我了解，它应该是设备到主机。

为什么 H->D 计数是 2？我认为应该是 1，因为数据在一页中。

提前致谢！

我的环境

驱动版本：418.87.00 CUDA 版本：10.1 ubuntu 18.04

#include <cuda.h>
#include <iostream>

using namespace std;

__global__ void setVal(char* data, int idx)

    data[idx] = 'd';



int main()

    const int count =  8;

    char* data;
    cudaMallocManaged((void **)&data, count);

    setVal<<<1,1>>>(data, 0); //GPU page fault

    cout<<" cpu read " << data[0] <<endl; 
  
    cudaFree(data);

    return 0;

==28762== Unified Memory profiling result:
Device "GeForce GTX 1070 (0)"
   Count  Avg Size  Min Size  Max Size  Total Size  Total Time  Name
       2  32.000KB  4.0000KB  60.000KB  64.00000KB  11.74400us  Host To Device
       1         -         -         -           -  362.9440us  Gpu page fault groups
Total CPU Page faults: 1

【参考方案1】：

为什么会有一个数据从主机迁移到设备？据我了解，它应该是设备到主机。

您正在主机和设备之间颠簸数据。因为 GPU 内核启动是异步的，所以您的主机代码在内核启动之后实际上是在访问 GPU 代码之前之前的数据。

在你的内核调用之后放置一个cudaDeviceSynchronize()，这样CPU代码在内核完成之前不会尝试读取数据。

对于您的其他问题，我没有答案。探查器通常无法完美解析非常少量的活动。它不一定会在分析运行期间检测所有 SM，并且它的一些结果可能会针对 GPC、TPC 和/或整个 GPU 的大小进行缩放。这将是我的猜测，尽管这只是猜测。在分析几乎什么都不做的微小代码时，我通常不期望分析器提供完全准确的结果。