CUDA - взаимозаменяемость памяти и взаимозаменяемость памяти - PullRequest
Новая
       37

CUDA - взаимозаменяемость памяти и взаимозаменяемость памяти

0 голосов
/ 05 мая 2018

Из того, что я понимаю, когда мы пытаемся перенести страничную память с хоста на устройство, cudamemcpy автоматически скопирует данные в закрепленную память (буфер), а затем перенесет на устройство.

Многие предложения по оптимизации кода предполагают использование закрепленной памяти вместо памяти с возможностью постраничного вывода. Я не понимаю, как это было бы быстрее. Хотя сама передача будет происходить быстрее, поскольку она напрямую из закрепленной памяти, а не для ее копирования перед передачей, вам все равно придется копировать содержимое из памяти с возможностью постраничной памяти в закрепленную память самостоятельно, что создает много накладных расходов. Я неправильно понимаю ситуацию? Может ли кто-нибудь объяснить мне, почему использование закрепленной памяти было бы быстрее, учитывая накладные расходы, которые она понесет при копировании, а также тот факт, что кажется, что мы просто вручную делаем то, что cudamemcpy может делать автоматически?

1 Ответ

0 голосов
/ 05 мая 2018

Закрепленная память требуется, если вы хотите дублировать копирование и вычисление.

В некоторых ситуациях закрепленная память также может повысить производительность. Это часто заметно, если мы можем повторно использовать буферы, которые используются для передачи данных между хостом и устройством.

вам все равно придется копировать содержимое из памяти с возможностью постраничного хранения в закрепленную память самостоятельно, что создает много накладных расходов.

Я не думаю, что вам нужно переносить данные из страничной памяти в закрепленную память во всех мыслимых случаях.

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

$ cat t113.cu
#include <stdio.h>
#include <stdlib.h>

typedef double my_T;
const int ds = 1024;
const int num_iter = 100;
const int block_dim = 16;

// C = A * B
// naive!!
template <typename T>
__global__ void mm(const T * __restrict__ A, const T * __restrict__ B, T * __restrict__ C, size_t d)
{
  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  int idy = threadIdx.y+blockDim.y*blockIdx.y;

  if ((idx < d) && (idy < d)){
    T temp = 0;
    for (int i = 0; i < d; i++)
      temp += A[idy*d + i]*B[i*d + idx];
    C[idy*d + idx] = temp;
    }
}

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

  int use_pinned = 0;
  if (argc > 1) use_pinned = atoi(argv[1]);
  if (use_pinned) printf("Using pinned memory\n");
  else printf("Using pageable memory\n");
  my_T *d_A, *d_B, *d_C, *h_A, *h_B, *h_C;
  int bs = ds*ds*sizeof(my_T);
  cudaMalloc(&d_A, bs);
  cudaMalloc(&d_B, bs);
  cudaMalloc(&d_C, bs);
  if (use_pinned){
    cudaHostAlloc(&h_A, bs, cudaHostAllocDefault);
    cudaHostAlloc(&h_B, bs, cudaHostAllocDefault);
    cudaHostAlloc(&h_C, bs, cudaHostAllocDefault);}
  else {
    h_A = (my_T *)malloc(bs);
    h_B = (my_T *)malloc(bs);
    h_C = (my_T *)malloc(bs);}
  cudaMemset(d_A, 0, bs);
  cudaMemset(d_B, 0, bs);
  memset(h_C, 0, bs);
  dim3 block(block_dim,block_dim);
  dim3 grid((ds+block.x-1)/block.x, (ds+block.y-1)/block.y);
  for (int iter = 0; iter<num_iter; iter++){
    mm<<<grid, block>>>(d_A, d_B, d_C, ds);
    if (iter > 1) if (h_C[0] != (my_T)((iter-2)*(iter-2)*ds)) printf("validation failure at iteration %d, was %f, should be %f\n", iter, h_C[0], (my_T) ((iter-2)*(iter-2)*ds));
    for (int i = 0; i < ds*ds; i++) {h_A[i] = iter; h_B[i] = iter;}
    cudaMemcpy(h_C, d_C, bs, cudaMemcpyDeviceToHost);
    cudaMemcpy(d_A, h_A, bs, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bs, cudaMemcpyHostToDevice);}
  printf("%s\n", cudaGetErrorString(cudaGetLastError()));
}
$ nvcc -arch=sm_60 -o t113 t113.cu
$ time ./t113
Using pageable memory
no error

real    0m1.987s
user    0m1.414s
sys     0m0.571s
$ time ./t113 1
Using pinned memory
no error

real    0m1.487s
user    0m0.903s
sys     0m0.579s
$

CUDA 9.1, CentOS 7.4, Tesla P100

Вкратце, этот код выполняет 100 «наивных» операций умножения матриц на GPU. На каждой итерации мы запускаем матричное умножение на GPU, и пока это делается, мы обновляем данные хоста (входные данные). Когда умножение матрицы завершено, мы передаем результаты на хост, затем передаем новые входные данные на устройство, затем выполняем еще одну итерацию.

Я не говорю, что этот код идеально оптимизирован. Например, ядро ​​является наивной реализацией (если вам нужно быстрое умножение матриц, вы должны использовать CUBLAS). И если вы серьезно относитесь к оптимизации, вы, вероятно, захотите перекрыть передачу данных в этом примере с выполнением кода устройства. В этом случае вы все равно будете вынуждены использовать закрепленные буферы. Но не всегда возможно добиться совпадения копирования и вычислений в каждом приложении, и в некоторых случаях (например, в приведенном примере) использование закрепленных буферов может помочь с точки зрения производительности.

Если вы настаиваете на сравнении со случаем, когда вы должны сначала скопировать данные из не закрепленного буфера в закрепленный буфер, то это может быть бесполезным. Но без конкретного примера того, что вы имеете в виду, для меня не очевидно, что вы не можете выполнять всю свою работу, используя только закрепленные буферы хоста (для данных, которые вы намереваетесь отправлять в / из GPU). Если вы читаете данные с диска или из сети, вы можете прочитать их в закрепленные буферы. Если вы сначала выполняете некоторые вычисления хоста, вы можете использовать закрепленные буферы. Затем отправьте эти закрепленные данные буфера в графический процессор.

...