Как скопировать память между разными gpus в cuda

Перенос данных с одного графического процессора на другой часто требует промежуточной обработки через память хоста. Исключение составляют случаи, когда графические процессоры и топология системы поддерживают одноранговый (P2P) доступ, а P2P был явно включен. В этом случае передача данных может осуществляться напрямую по шине PCIE от одного графического процессора к другому.

В любом случае (с доступностью/включенностью P2P или без нее) типичный вызов API среды выполнения cuda будет cudaMemcpyPeer/cudaMemcpyPeerAsync, как показано в cuda p2pBandwidthLatencyTest пример кода.

В Windows одним из требований P2P является поддержка обоих устройств драйвером в режиме TCC. Режим TCC, по большей части, недоступен для графических процессоров GeForce (недавно было сделано исключение для графических процессоров семейства GeForce Titan с использованием драйверов и среды выполнения, доступных в наборе инструментов CUDA 7.5RC).

Поэтому в Windows эти графические процессоры не смогут использовать преимущества прямой передачи P2P. Тем не менее, почти идентичная последовательность может использоваться для передачи данных. Среда выполнения CUDA обнаружит характер передачи и выполнит скрытое выделение для создания промежуточного буфера. Затем передача будет завершена в 2 этапа: передача с исходного устройства на промежуточный буфер и передача из промежуточного буфера на целевое устройство.

Ниже приведен полностью проработанный пример, показывающий, как передавать данные с одного графического процессора на другой, используя преимущества доступа P2P, если он доступен:

$ cat t850.cu
#include <stdio.h>
#include <math.h>

#define SRC_DEV 0
#define DST_DEV 1

#define DSIZE (8*1048576)

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)


int main(int argc, char *argv[]){

  int disablePeer = 0;
  if (argc > 1) disablePeer = 1;
  int devcount;
  cudaGetDeviceCount(&devcount);
  cudaCheckErrors("cuda failure");
  int srcdev = SRC_DEV;
  int dstdev = DST_DEV;
  if (devcount <= max(srcdev,dstdev)) {printf("not enough cuda devices for the requested operation\n"); return 1;}
  int *d_s, *d_d, *h;
  int dsize = DSIZE*sizeof(int);
  h = (int *)malloc(dsize);
  if (h == NULL) {printf("malloc fail\n"); return 1;}
  for (int i = 0; i < DSIZE; i++) h[i] = i;
  int canAccessPeer = 0;
  if (!disablePeer) cudaDeviceCanAccessPeer(&canAccessPeer, srcdev, dstdev);
  cudaSetDevice(srcdev);
  cudaMalloc(&d_s, dsize);
  cudaMemcpy(d_s, h, dsize, cudaMemcpyHostToDevice);
  if (canAccessPeer) cudaDeviceEnablePeerAccess(dstdev,0);
  cudaSetDevice(dstdev);
  cudaMalloc(&d_d, dsize);
  cudaMemset(d_d, 0, dsize);
  if (canAccessPeer) cudaDeviceEnablePeerAccess(srcdev,0);
  cudaCheckErrors("cudaMalloc/cudaMemset fail");
  if (canAccessPeer) printf("Timing P2P transfer");
  else printf("Timing ordinary transfer");
  printf(" of %d bytes\n", dsize);
  cudaEvent_t start, stop;
  cudaEventCreate(&start); cudaEventCreate(&stop);
  cudaEventRecord(start);
  cudaMemcpyPeer(d_d, dstdev, d_s, srcdev, dsize);
  cudaCheckErrors("cudaMemcpyPeer fail");
  cudaEventRecord(stop);
  cudaEventSynchronize(stop);
  float et;
  cudaEventElapsedTime(&et, start, stop);
  cudaSetDevice(dstdev);
  cudaMemcpy(h, d_d, dsize, cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy fail");
  for (int i = 0; i < DSIZE; i++) if (h[i] != i) {printf("transfer failure\n"); return 1;}
  printf("transfer took %fms\n", et);
  return 0;
}

$ nvcc -arch=sm_20 -o t850 t850.cu
$ ./t850
Timing P2P transfer of 33554432 bytes
transfer took 5.135680ms
$ ./t850 disable
Timing ordinary transfer of 33554432 bytes
transfer took 7.274336ms
$

Примечания:

1. Передача любого параметра командной строки отключит использование P2P, даже если он доступен.
2. Приведенные выше результаты относятся к системе, в которой возможен одноранговый доступ, и оба графических процессора подключены через канал PCIE Gen2, обеспечивающий пропускную способность около 6 ГБ/с в одном направлении. Время передачи P2P соответствует этому (32 МБ/5 мс ~= 6 ГБ/с). Время передачи без P2P больше, но не удваивается. Это связано с тем, что для передач в/из промежуточного буфера после передачи некоторых данных в промежуточный буфер может начаться исходящая передача. Драйвер/среда выполнения использует это для частичного перекрытия передачи данных.

Обратите внимание, что в целом поддержка P2P может различаться в зависимости от графического процессора или семейства графических процессоров. Возможность запуска P2P на одном типе или семействе графических процессоров не обязательно означает, что он будет работать на другом типе или семействе графических процессоров, даже в той же системе/настройке. Последним фактором, определяющим поддержку GPU P2P, являются предоставляемые инструменты, которые запрашивают среду выполнения через cudaDeviceCanAccessPeer. Поддержка P2P может зависеть от системы и других факторов. Никакие заявления, сделанные здесь, не являются гарантией поддержки P2P для любого конкретного графического процессора в любой конкретной конфигурации.

Примечание. Требования к драйверу TCC в Windows были смягчены в последних версиях драйверов. С последними драйверами должна быть возможность обмена данными P2P между устройствами в режиме WDDM, если выполняются остальные требования.

Заявление о поддержке TCC является общим. Поддерживаются не все графические процессоры. Последним фактором, определяющим поддержку TCC (или нет) на конкретном графическом процессоре, является инструмент nvidia-smi. Ничто здесь не должно рассматриваться как гарантия поддержки TCC на вашем конкретном графическом процессоре.

Этот ответ был написан до появления NVLINK и других изменений в технологиях ЦП и ГП. Для любой данной системы окончательным арбитром/детерминантом того, доступен ли P2P между любыми двумя графическими процессорами в системе, является результат, возвращаемый cudaDeviceCanAccessPeer(). Другие общие утверждения об узлах NUMA и других характеристиках топологии системы являются второстепенными и не должны использоваться для принятия окончательного решения.