如何从设备端预取CUDA统一内存至主机端？性能对比分析

背景与问题 #

我正在对比cudaMalloc与cudaMallocManaged的性能，应用场景是一个对用户隐藏GPU使用的矩阵库（用户可像使用普通库一样操作，部分运算自动调用GPU）。

当算法仅使用GPU时，可通过cudaMemPrefetch将内存预取至GPU端。测试结果显示：

• cost[0][2]与cost[1][2]性能相当
• cost[0][3]的速度慢很多
  但反向预取（从设备到主机）似乎无法正常工作，cost[1][4]比cost[0][3]+cost[0][4]慢约10-15%。

请问是否存在将CUDA统一内存从设备端预取至主机端的方法？

测试代码 #

#include <cuda_runtime.h>
#include <thrust/execution_policy.h>
#include <thrust/sort.h>
#include <thrust/device_ptr.h>
#include <string>
#include <chrono>
#include <random>
using namespace std;

class MyTimer {
    std::chrono::time_point<std::chrono::system_clock> start;

public:
    void startCounter() {
        start = std::chrono::system_clock::now();
    }

    int64_t getCounterNs() {
        return std::chrono::duration_cast<std::chrono::nanoseconds>(std::chrono::system_clock::now() - start).count();
    }

    int64_t getCounterMs() {
        return std::chrono::duration_cast<std::chrono::milliseconds>(std::chrono::system_clock::now() - start).count();
    }

    double getCounterMsPrecise() {
        return std::chrono::duration_cast<std::chrono::nanoseconds>(std::chrono::system_clock::now() - start).count()
                / 1000000.0;
    }
};


int N = 10000000;

void GenData(int N, float* a)
{
  for (int i = 0; i < N; i ++) a[i] = float(rand() % 1000000) / (rand() % 100 + 1);
}

__global__
void HelloWorld()
{
  printf("Hello world\n");
}

constexpr int npoints = 6;
const string costnames[] = {"allocate", "H2D", "sort", "D2H", "hostsum", "free"};
double cost[3][npoints];
volatile double dummy = 0;

void Test1()
{
  MyTimer timer;

  timer.startCounter();
  float *h_a = new float[N];
  float *d_a;
  cudaMalloc(&d_a, N * sizeof(float));
  cudaDeviceSynchronize();
  cost[0][0] += timer.getCounterMsPrecise();

  GenData(N, h_a);
  dummy = h_a[rand() % N];

  timer.startCounter();
  cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);
  cudaDeviceSynchronize();
  cost[0][1] += timer.getCounterMsPrecise();

  timer.startCounter();
  thrust::device_ptr<float> dev_ptr = thrust::device_pointer_cast(d_a);
  thrust::sort(dev_ptr, dev_ptr + N);
  cudaDeviceSynchronize();
  cost[0][2] += timer.getCounterMsPrecise();

  timer.startCounter();
  cudaMemcpy(h_a, d_a, N * sizeof(float), cudaMemcpyDeviceToHost);  
  cudaDeviceSynchronize();
  dummy = h_a[rand() % N];
  cost[0][3] += timer.getCounterMsPrecise();
  
  timer.startCounter();
  float sum = 0;
  for (int i = 0; i < N; i++) sum += h_a[i];
  dummy = sum;
  cost[0][4] += timer.getCounterMsPrecise();

  timer.startCounter();
  delete[] h_a;
  cudaFree(d_a);
  cudaDeviceSynchronize();
  cost[0][5] += timer.getCounterMsPrecise();

  for (int i = 0; i < npoints; i++) dummy += cost[0][i];
}

void Test2()
{
  MyTimer timer;

  timer.startCounter();
  float *a;
  cudaMallocManaged(&a, N * sizeof(float));  
  cost[1][0] += timer.getCounterMsPrecise();

  GenData(N, a);
  dummy = a[rand() % N];

  timer.startCounter();
  cudaMemPrefetchAsync(a, N * sizeof(float), 0, 0);
  cudaDeviceSynchronize();
  cost[1][1] += timer.getCounterMsPrecise();

  timer.startCounter();
  thrust::device_ptr<float> dev_ptr = thrust::device_pointer_cast(a);
  thrust::sort(dev_ptr, dev_ptr + N);
  cudaDeviceSynchronize();
  cost[1][2] += timer.getCounterMsPrecise();

  timer.startCounter();
  cudaMemPrefetchAsync(a, N * sizeof(float), 0, 0);
  cudaDeviceSynchronize();
  dummy = a[rand() % N];
  cost[1][3] += timer.getCounterMsPrecise();

  timer.startCounter();
  float sum = 0;
  for (int i = 0; i < N; i++) sum += a[i];
  dummy = sum;
  cost[1][4] += timer.getCounterMsPrecise();

  timer.startCounter();  
  cudaFree(a);
  cudaDeviceSynchronize();
  cost[1][5] += timer.getCounterMsPrecise();

  for (int i = 0; i < npoints; i++) dummy += cost[1][i];
}

void Test3()
{
  MyTimer timer;

  timer.startCounter();
  float *a;
  cudaMallocManaged(&a, N * sizeof(float));  
  cost[2][0] += timer.getCounterMsPrecise();

  GenData(N, a);
  dummy = a[rand() % N];

  timer.startCounter();
  //cudaMemPrefetchAsync(a, N * sizeof(float), 0, 0);
  //cudaDeviceSynchronize();
  cost[2][1] += timer.getCounterMsPrecise();

  timer.startCounter();
  thrust::device_ptr<float> dev_ptr = thrust::device_pointer_cast(a);
  thrust::sort(dev_ptr, dev_ptr + N);
  cudaDeviceSynchronize();
  cost[2][2] += timer.getCounterMsPrecise();

  timer.startCounter();
  // cudaMemPrefetchAsync(a, N * sizeof(float), 0, 0);
  // cudaDeviceSynchronize();
  dummy = a[rand() % N];
  cost[2][3] += timer.getCounterMsPrecise();

  timer.startCounter();
  float sum = 0;
  for (int i = 0; i < N; i++) sum += a[i];
  dummy = sum;
  cost[2][4] += timer.getCounterMsPrecise();

  timer.startCounter();  
  cudaFree(a);
  cudaDeviceSynchronize();
  cost[2][5] += timer.getCounterMsPrecise();

  for (int i = 0; i < npoints; i++) dummy += cost[2][i];
}

int main()
{
  srand(time(NULL));
  HelloWorld<<<1,1>>>();

  // warmup
  Test1();
  Test2();
  for (int i = 0; i < 3; i++)
  for (int j = 0; j < npoints; j++) cost[i][j] = 0;  

  int ntest = 10;
  for (int t = 1; t <= ntest; t++) {
    Test1();
    Test2();
    Test3();
  }

  for (int i = 0; i < npoints; i++) {
    cout << "cost " << costnames[i] << " = " << (cost[0][i] / ntest) << " , " << (cost[1][i] / ntest) << " , " << (cost[2][i] / ntest) << "\n";
  }

  return 0;
}

测试结果（2080ti） #

Hello world
cost allocate = 0.245438 , 0.0470603 , 0.029834
cost H2D = 6.25315 , 6.36215 , 3.71e-05
cost sort = 2.61625 , 2.6077 , 14.5418
cost D2H = 8.74573 , 0.0520719 , 0.0759482
cost hostsum = 6.98815 , 17.9619 , 18.3188
cost free = 2.82205 , 3.8711 , 4.12887

解决方案 #

要将统一内存从设备预取至主机，关键是在cudaMemPrefetchAsync中指定主机设备ID，而非GPU设备ID（0）。主机的设备ID可通过CUDA定义的cudaCpuDeviceId常量直接获取。

修改代码 #

将Test2中的反向预取代码替换为：

cudaMemPrefetchAsync(a, N * sizeof(float), cudaCpuDeviceId, 0);

原理说明 #

• cudaMemPrefetchAsync的第三个参数是目标设备ID，cudaCpuDeviceId是CUDA官方定义的主机设备标识
• 显式预取到主机后，后续主机端的内存访问（如hostsum循环）无需再触发隐式数据迁移，消除了额外的同步延迟和拷贝开销
• 需确保预取操作完成后再执行主机访问，可通过cudaDeviceSynchronize()实现同步

修改后预期效果 #

修改后cost[1][4]的耗时会接近cost[0][3]+cost[0][4]的总和，显式预取替代了隐式迁移，性能与cudaMalloc方案的主机访问耗时基本持平。

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