small benchmark for testing worst-case UVM overhead in CUDA

uvm_perf_test.cu

#include <vector>
#include <iomanip>
#include <iostream>
#include <cinttypes>
#include <thread>

////////////////////////////////////////////////////////////////////////////////
////////////////////////////////////////////////////////////////////////////////

#include <chrono>

class timer {
  typedef std::chrono::high_resolution_clock::time_point time_point;
  typedef std::chrono::duration<double>                  duration_type;

public:
  void start() { then = std::chrono::high_resolution_clock::now(); }
  void stop() { now = std::chrono::high_resolution_clock::now(); }
  double elapsed() { return std::chrono::duration_cast<duration_type>(now - then).count(); }

private:
  time_point then, now;
};

template < typename callable >
double time(callable f) {
  timer stopwatch;
  stopwatch.start();
  f();
  stopwatch.stop();
  return stopwatch.elapsed();
}

////////////////////////////////////////////////////////////////////////////////
////////////////////////////////////////////////////////////////////////////////

void daxpy_cpu(double * x, double * y, double * z, int n) {
  for (int i = 0; i < n; i++) {
    z[i] = x[i] + y[i];
  }
}

__global__ void daxpy_gpu(double * x, double * y, double * z, int n) {
  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  if (tid < n) {
    z[tid] = x[tid] + y[tid];
  }
}

////////////////////////////////////////////////////////////////////////////////
////////////////////////////////////////////////////////////////////////////////

int num_runs = 20;

double perf_test_host(int n) {

  std::vector< double > x(n);
  std::vector< double > y(n);
  std::vector< double > z(n);

  for (int i = 0; i < n; i++) {
    x[i] = 1;
    y[i] = 2;
    z[i] = 0;
  }

  timer stopwatch;
  stopwatch.start();
  for (int i = 0; i < num_runs; i++) {
    saxpy_cpu(&x[0], &y[0], &z[0], n);
    saxpy_cpu(&z[0], &x[0], &y[0], n);
  }
  stopwatch.stop();

  for (int i = 0; i < n; i++) {
    if (z[i] != 41) {
      std::cout << "wrong answer!" << std::endl;
      break;
    }
  }

  return stopwatch.elapsed();
}

double perf_test_uvm(int n) {

  double * x; 
  double * y; 
  double * z; 

  cudaMallocManaged(&x, sizeof(double) * n);
  cudaMallocManaged(&y, sizeof(double) * n);
  cudaMallocManaged(&z, sizeof(double) * n);

  for (int i = 0; i < n; i++) {
    x[i] = 1;
    y[i] = 2;
    z[i] = 0;
  }

  int blocks = 256;
  int blocksize = (n + blocks - 1) / blocks;

  timer stopwatch;
  stopwatch.start();
  for (int i = 0; i < num_runs; i++) {
    daxpy_cpu(x, y, z, n);
    daxpy_gpu<<<blocksize, blocks>>>(z, x, y, n);
    cudaDeviceSynchronize();
  }
  stopwatch.stop();

  for (int i = 0; i < n; i++) {
    if (z[i] != 41) {
      std::cout << "wrong answer!" << std::endl;
      break;
    }
  }

  cudaFree(x);
  cudaFree(y);
  cudaFree(z);

  return stopwatch.elapsed();
}

double perf_test_manual(int n) {

  std::vector< double > x_h(n);
  std::vector< double > y_h(n);
  std::vector< double > z_h(n);

  double * x_d; 
  double * y_d; 
  double * z_d; 

  cudaMalloc(&x_d, sizeof(double) * n);
  cudaMalloc(&y_d, sizeof(double) * n);
  cudaMalloc(&z_d, sizeof(double) * n);

  for (int i = 0; i < n; i++) {
    x_h[i] = 1;
    y_h[i] = 2;
    z_h[i] = 0;
  }

  int blocks = 256;
  int blocksize = (n + blocks - 1) / blocks;

  timer stopwatch;
  stopwatch.start();
  for (int i = 0; i < num_runs; i++) {

    if (i > 0) {
      cudaMemcpy(&x_h[0], x_d, sizeof(double) * n, cudaMemcpyDeviceToHost);
      cudaMemcpy(&y_h[0], y_d, sizeof(double) * n, cudaMemcpyDeviceToHost);
      cudaMemcpy(&z_h[0], z_d, sizeof(double) * n, cudaMemcpyDeviceToHost);
    }

    daxpy_cpu(&x_h[0], &y_h[0], &z_h[0], n);

    cudaMemcpy(x_d, &x_h[0], sizeof(double) * n, cudaMemcpyHostToDevice);
    cudaMemcpy(y_d, &y_h[0], sizeof(double) * n, cudaMemcpyHostToDevice);
    cudaMemcpy(z_d, &z_h[0], sizeof(double) * n, cudaMemcpyHostToDevice);

    daxpy_gpu<<<blocksize, blocks>>>(z_d, x_d, y_d, n);

  }
  cudaDeviceSynchronize();
  stopwatch.stop();

  cudaMemcpy(&z_h[0], z_d, sizeof(double) * n, cudaMemcpyDeviceToHost);
  cudaDeviceSynchronize();

  for (int i = 0; i < n; i++) {
    if (z_h[i] != 41) {
      std::cout << "wrong answer!" << std::endl;
      break;
    }
  }

  cudaFree(x_d);
  cudaFree(y_d);
  cudaFree(z_d);

  return stopwatch.elapsed();
}

double perf_test_device(int n) {

  std::vector< double > x_h(n);
  std::vector< double > y_h(n);
  std::vector< double > z_h(n);

  double * x_d; 
  double * y_d; 
  double * z_d; 

  cudaMalloc(&x_d, sizeof(double) * n);
  cudaMalloc(&y_d, sizeof(double) * n);
  cudaMalloc(&z_d, sizeof(double) * n);

  for (int i = 0; i < n; i++) {
    x_h[i] = 1;
    y_h[i] = 2;
    z_h[i] = 0;
  }

  int blocks = 256;
  int blocksize = (n + blocks - 1) / blocks;

  cudaMemcpy(x_d, &x_h[0], sizeof(double) * n, cudaMemcpyHostToDevice);
  cudaMemcpy(y_d, &y_h[0], sizeof(double) * n, cudaMemcpyHostToDevice);
  cudaMemcpy(z_d, &z_h[0], sizeof(double) * n, cudaMemcpyHostToDevice);

  timer stopwatch;
  stopwatch.start();
  for (int i = 0; i < num_runs; i++) {
    daxpy_gpu<<<blocksize, blocks>>>(x_d, y_d, z_d, n);
    daxpy_gpu<<<blocksize, blocks>>>(z_d, x_d, y_d, n);
  }
  cudaDeviceSynchronize();
  stopwatch.stop();

  cudaMemcpy(&z_h[0], z_d, sizeof(double) * n, cudaMemcpyDeviceToHost);
  cudaDeviceSynchronize();

  for (int i = 0; i < n; i++) {
    if (z_h[i] != 41) {
      std::cout << "wrong answer!" << std::endl;
      break;
    }
  }

  cudaFree(x_d);
  cudaFree(y_d);
  cudaFree(z_d);

  return stopwatch.elapsed();
}

double perf_test_device_uvm(int n) {

  double * x; 
  double * y; 
  double * z; 

  cudaMallocManaged(&x, sizeof(double) * n);
  cudaMallocManaged(&y, sizeof(double) * n);
  cudaMallocManaged(&z, sizeof(double) * n);

  for (int i = 0; i < n; i++) {
    x[i] = 1;
    y[i] = 2;
    z[i] = 0;
  }

  int blocks = 256;
  int blocksize = (n + blocks - 1) / blocks;

  cudaMemPrefetchAsync(x, sizeof(double) * n, 0);
  cudaMemPrefetchAsync(y, sizeof(double) * n, 0);
  cudaMemPrefetchAsync(z, sizeof(double) * n, 0);
  cudaDeviceSynchronize();

  timer stopwatch;
  stopwatch.start();
  for (int i = 0; i < num_runs; i++) {
    daxpy_gpu<<<blocksize, blocks>>>(x, y, z, n);
    daxpy_gpu<<<blocksize, blocks>>>(z, x, y, n);
  }
  cudaDeviceSynchronize();
  stopwatch.stop();

  for (int i = 0; i < n; i++) {
    if (z[i] != 41) {
      std::cout << "wrong answer!" << std::endl;
      break;
    }
  }

  cudaFree(x);
  cudaFree(y);
  cudaFree(z);

  return stopwatch.elapsed();
}

int main() {

  // ignore the timings from the first run, 
  // to avoid timing first-kernel-launch overhead
  perf_test_uvm(1000);

  std::cout << std::setw(70) << "performance in GBps (higher is better)" << std::endl;

  std::cout << std::setw(15) << "n";
  std::cout << std::setw(15) << "host only";
  std::cout << std::setw(24) << "no cudaMemcpy (uvm)";
  std::cout << std::setw(20) << "manual cudaMemcpy";
  std::cout << std::setw(15) << "device only";
  std::cout << std::setw(20) << "device only (uvm)";
  std::cout << std::endl;

  for (int i = 15; i < 21; i++) {

    int n = 1<<i;

    uint64_t bytes = 3 * 2 * num_runs * n * sizeof(double);

    double GB = bytes / 1000000000.0;

    std::cout << std::setw(15) << n;
    std::cout << std::setw(15) << GB / perf_test_host(n);
    std::cout << std::setw(24) << GB / perf_test_uvm(n);
    std::cout << std::setw(20) << GB / perf_test_manual(n);
    std::cout << std::setw(15) << GB / perf_test_device(n);
    std::cout << std::setw(20) << GB / perf_test_device_uvm(n);
    std::cout << std::endl;
 
  }

}

output from running on a GV100 gpu machine with intel 13900k (compiled with nvcc uvm_perf_test.cu -O3 -arch=sm_70):

                                performance in GBps (higher is better)
              n      host only     no cudaMemcpy (uvm)   manual cudaMemcpy    device only   device only (uvm)
          32768        177.679                 1.14639             6.19351        233.669             275.465
          65536         151.93                  1.6168             6.74325        354.508             497.911
         131072         64.261                 3.20241               6.991        516.309             826.193
         262144        60.6726                 6.34253             8.32662        469.368             665.075
         524288        61.4274                 6.60688               9.178        468.544             544.338
        1048576        57.8731                 6.67996             9.92983        535.682             574.515

output from running on a hip version (rocm 6.0) of this test on a Radeon VII machine with intel 12700k:

                                performance in GBps (higher is better)
              n      host only      no hipMemcpy (uvm)    manual hipMemcpy    device only   device only (uvm)
          32768        113.508                  18.729             4.27342        234.866             6.98121
          65536        94.2211                 4.67749             5.79793        401.405             7.75288
         131072        69.9229                 8.83035             6.24912        61.5186             12.1142
         262144        69.1119                 17.9187             10.2904        268.848             32.6475
         524288        70.1276                  23.666             10.2775        412.232             144.353
        1048576        47.1457                  25.077              10.749        540.449             107.227

I don't understand these timings:

• the naive uvm approach in column 2 appears to be significantly faster than manually copying (??)
• allocating with hipMallocManaged(), and never copying back to the host is 5-30x slower than allocating hipMalloc() and never copying to host (??)

Note: AMD's higher throughput on no ***Memcpy (uvm) and manual ***Memcpy data is related to a faster PCIe bus. The Titan V released in 2017, so it has PCIe 3.0x16 (~16 GBps max), while the Radeon VII came out in 2019 with PCIe 4.0x16 (~32 GBps max).

from a different AMD card with XNACK

env HSA_XNACK=1 srun -N 1 -p pdebug a.out 
flux-job: f3LrLQbtazKH started                                                                                              00:00:02
                                performance in GBps (higher is better)
              n      host only      no hipMemcpy (uvm)    manual hipMemcpy    device only   device only (uvm)
          32768        51.2754                0.742181             5.53053        234.478             281.885
          65536        48.0435                  1.3787             9.00739        473.872             508.158
         131072        53.2398                 1.77426             10.1386        929.575              689.69
         262144        60.0587                 1.77001             11.6738        1223.98             1079.84
         524288         60.918                  1.9941             13.2752        1005.71             1172.49
        1048576        58.4539                  2.0073             12.5879        1106.58             1172.99

from a different AMD card without XNACK

$ srun -N 1 -p pdebug a.out 
                                performance in GBps (higher is better)
              n      host only      no hipMemcpy (uvm)    manual hipMemcpy    device only   device only (uvm)
          32768        52.6523                 28.0241             5.29573        238.559             22.1148
          65536        47.9916                 30.9428             9.10463        494.242             75.6288
         131072        51.8506                 34.8372             10.7828        917.958             175.658
         262144        59.4912                 38.2315             16.5798        1231.05             195.093
         524288        60.7979                 40.3616             17.2538        1032.72             338.708
        1048576        63.7128                 41.3376             17.6276        1099.86             531.526