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

Question

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

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

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

提前感谢！

我的环境

• 驱动程序版本：418.87.00
• CUDA版本：10.1
• 乌班图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

Answer 1

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

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

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

你的另一个问题我没有答案。探查器通常无法完美地解决非常少量的活动。它不一定在分析运行期间对所有SM进行仪器化，并且它的一些结果可以根据GPC、TPC和/或整个GPU的大小进行缩放。这将是我的猜测，尽管这只是猜测。当分析几乎什么都不做的微小代码时，我通常不会期望分析器给出完全准确的结果。