Блог

Я пытаюсь сравнить общее время выполнения, затрачиваемое управляемыми и неуправляемыми версиями управления памятью CUDA.

В следующем примере кода у меня есть две функции, которые выполняют одно и то же. Единственное отличие заключается в управлении их памятью. Одна функция использует cudaMalloc() / cudaMemcpy() , а другой метод использует только cudaMallocManaged() .

Я использовал nvprof для вычисления разных времен и получил следующие результаты:

nvprof Вывод Управляемой Версии:

 == 29028 == Profiling result : Type Time(%) Time Calls Avg Min Max Name GPU activities : 100.00 % 59.425us 1 59.425us 59.425us 59.425us add(int, float*, float*) API calls : 78.08 % 296.49ms 2 148.24ms 1.7127ms 294.78ms cudaMallocManaged 19.61 % 74.451ms 1 74.451ms 74.451ms 74.451ms cuDevicePrimaryCtxRelease 1.55 % 5.8705ms 1 5.8705ms 5.8705ms 5.8705ms cudaLaunchKernel 0.67 % 2.5547ms 2 1.2774ms 974.40us 1.5803ms cudaFree 0.07 % 280.60us 1 280.60us 280.60us 280.60us cudaDeviceSynchronize 0.01 % 28.300us 3 9.4330us 3.0000us 13.300us cuModuleUnload 0.01 % 26.800us 1 26.800us 26.800us 26.800us cuDeviceTotalMem 0.00 % 17.700us 101 175ns 100ns 900ns cuDeviceGetAttribute 0.00 % 10.100us 3 3.3660us 300ns 8.8000us cuDeviceGetCount 0.00 % 3.2000us 1 3.2000us 3.2000us 3.2000us cuDeviceGetName 0.00 % 3.0000us 2 1.5000us 300ns 2.7000us cuDeviceGet 0.00 % 500ns 1 500ns 500ns 500ns cuDeviceGetLuid 0.00 % 200ns 1 200ns 200ns 200ns cuDeviceGetUuid  == 29028 == Unified Memory profiling result : Device "GeForce GTX 1070 (0)" Count Avg Size Min Size Max Size Total Size Total Time Name 64 128.00KB 128.00KB 128.00KB 8.000000MB 3.279000ms Host To Device 146 84.164KB 32.000KB 1.0000MB 12.00000MB 64.50870ms Device To Host  

Вывод неуправляемой версии nvprof :

 == 23864 == Profiling result : Type Time(%) Time Calls Avg Min Max Name GPU activities : 56.30 % 1.5032ms 2 751.60us 751.44us 751.76us[CUDA memcpy HtoD] 41.48 % 1.1075ms 1 1.1075ms 1.1075ms 1.1075ms[CUDA memcpy DtoH] 2.23 % 59.457us 1 59.457us 59.457us 59.457us add(int, float*, float*) API calls : 78.92 % 270.08ms 2 135.04ms 656.40us 269.43ms cudaMalloc 19.79 % 67.730ms 1 67.730ms 67.730ms 67.730ms cuDevicePrimaryCtxRelease 1.05 % 3.5796ms 3 1.1932ms 1.0106ms 1.4341ms cudaMemcpy 0.10 % 346.20us 2 173.10us 3.4000us 342.80us cudaFree 0.09 % 314.30us 1 314.30us 314.30us 314.30us cudaDeviceSynchronize 0.02 % 74.200us 1 74.200us 74.200us 74.200us cudaLaunchKernel 0.01 % 34.700us 3 11.566us 2.5000us 29.100us cuModuleUnload 0.01 % 24.100us 1 24.100us 24.100us 24.100us cuDeviceTotalMem 0.00 % 17.100us 101 169ns 100ns 900ns cuDeviceGetAttribute 0.00 % 9.0000us 3 3.0000us 300ns 8.0000us cuDeviceGetCount 0.00 % 3.2000us 1 3.2000us 3.2000us 3.2000us cuDeviceGetName 0.00 % 1.5000us 2 750ns 200ns 1.3000us cuDeviceGet 0.00 % 300ns 1 300ns 300ns 300ns cuDeviceGetUuid 0.00 % 300ns 1 300ns 300ns 300ns cuDeviceGetLuid  

Мой Код:

 int RunManagedVersion() {  int N = 1 lt;lt; 20;  float* x, * y;   // Allocate Unified Memory -- accessible from CPU or GPU  cudaMallocManaged(amp;x, N * sizeof(float));  cudaMallocManaged(amp;y, N * sizeof(float));   // initialize x and y arrays on the host  for (int i = 0; i lt; N; i  ) {  x[i] = 1.0f;  y[i] = 2.0f;  }   // Launch kernel on 1M elements on the GPU  int blockSize = 256;  int numBlocks = (N   blockSize - 1) / blockSize;  add lt;lt; lt;numBlocks, blockSize gt;gt; gt; (N, x, y);   // Wait for GPU to finish before accessing on host  cudaDeviceSynchronize();   // Check for errors (all values should be 3.0f)  float maxError = 0.0f;  for (int i = 0; i lt; N; i  )  maxError = fmax(maxError, fabs(y[i] - 3.0f));   std::cout lt;lt; "Max error: " lt;lt; maxError lt;lt; std::endl;   // Free memory  cudaFree(x);  cudaFree(y);   return 0; }  int RunUnmanagedVersion() {  int N = 1 lt;lt; 20;   //Declare pointers for input and output arrays  float* x = (float*)calloc(N, sizeof(float));  float* y = (float*)calloc(N, sizeof(float));   // initialize x and y arrays on the host  for (int i = 0; i lt; N; i  ) {  x[i] = 1.0f;  y[i] = 2.0f;  }   //Allocate device memory for input and output images  float* d_pX = 0;  float* d_pY = 0;  cudaMalloc(amp;d_pX, N * sizeof(float));  cudaMalloc(amp;d_pY, N * sizeof(float));   //Copy INPUT ARRAY data from host to device  cudaMemcpy(d_pX, x, N * sizeof(float), cudaMemcpyHostToDevice);  cudaMemcpy(d_pY, y, N * sizeof(float), cudaMemcpyHostToDevice);    // Launch kernel on 1M elements on the GPU  int blockSize = 256;  int numBlocks = (N   blockSize - 1) / blockSize;  add lt;lt; lt;numBlocks, blockSize gt;gt; gt; (N, d_pX, d_pY);   // Wait for GPU to finish before accessing on host  cudaDeviceSynchronize();   //Copy Results - Device to Host  cudaMemcpy(y, d_pY, N * sizeof(float), cudaMemcpyDeviceToHost);   // Check for errors (all values should be 3.0f)  float maxError = 0.0f;  for (int i = 0; i lt; N; i  )  maxError = fmax(maxError, fabs(y[i] - 3.0f));   std::cout lt;lt; "Max error: " lt;lt; maxError lt;lt; std::endl;   // device memory free  cudaFree(d_pX);  cudaFree(d_pX);   //host memory free  free(x);  free(y);   return 0; }  int main() {  RunUnmanagedVersion();   //RunManagedVersion();   return 0; }  

ВОПРОС: Я несколько раз повторял приведенный выше код и отметил, что время передачи данных из DeviceToHost значительно выше в случае управляемой версии (т. Е. Унифицированной памяти). Это нормально (и почему?) Или я делаю что-то не так в коде?

Я верю, что то, что я собираюсь здесь сказать, это:

1. без документов, следовательно, могут быть изменены, но
2. наблюдаемый при тщательном профилировании

Кроме того, все эти комментарии относятся к режиму единой системы обмена сообщениями, связанному с использованием единой системы обмена сообщениями Windows или Linux с графическим процессором до паскаля.

В этом режиме единой системы обмена сообщениями до паскаля и/или Windows передача данных с хоста на устройство инициируется в момент запуска ядра. Это проявится в виде задержки в процессе запуска ядра (с момента запроса запуска ядра до момента фактического начала выполнения кода ядра.

В этом случае система единой системы обмена сообщениями передает данные порциями фиксированного размера. Это видно из результатов вашего профилировщика:

 Count Avg Size Min Size Max Size Total Size Total Time Name 64 128.00KB 128.00KB 128.00KB 8.000000MB 3.279000ms Host To Device  

Мы пришли бы к выводу, что, поскольку на устройство действительно необходимо передать 8 МБ данных, а в случае без единой системы обмена сообщениями это, по-видимому, происходит за ~1,5 мс:

 GPU activities : 56.30 % 1.5032ms 2 751.60us 751.44us 751.76us[CUDA memcpy HtoD]  

что даже случай H-gt;D несколько менее эффективен в Windows по сравнению с случаем без единой системы обмена сообщениями. Я объясняю это необходимостью (по какой-либо причине) передавать 8 МБ данных относительно небольшими порциями по 128 КБ. Кроме того, WDDM имеет прямой контроль над этим графическим процессором в случае Windows, и CUDA фактически является «клиентом» WDDM для этих действий, особенно связанных с памятью. Вполне возможно, что WDDM решил, что хочет что-то сделать с/с графическим процессором во время передачи данных, и, возможно, добавил некоторые пробелы или неэффективность.

В случае D-gt;H, в Windows, ситуация, похоже, отличается и, возможно, хуже. Однако мы должны быть осторожны, чтобы оценить, что здесь происходит. Первый вопрос может быть:

Почему передается 12 МБ данных D-gt;H?

Кажется, есть несколько вещей, которые стоит отметить:

• Распределение единой системы обмена сообщениями, по-видимому, появляется первым в памяти устройства. (Это отличается от случая с вызовом по запросу!) Это означает, что если первое, что вы сделаете, это получите доступ к распределению единой системы обмена сообщениями в коде хоста, то распределение должно быть перенесено с устройства на хост. Это составляет 8 МБ из 12 МБ передачи, и если вы проведете тщательные эксперименты по профилированию, вы сможете убедиться в этом сами.
• Распределение единой системы обмена сообщениями на устройстве, по-видимому, передается коду хоста на основе действия кода хоста. Это самоочевидно, если внимательно рассмотреть первый пункт выше. Но даже если мы обратим внимание только на передачу данных после работы ядра, с помощью экспериментов с профилировщиками легко убедить себя, что если ни один код хоста после запуска ядра фактически не получит доступ к данным, там не произойдет никаких передач.

Вторая пуля выше означает, что мы могли бы предположить, что переводы D-gt;H могут быть:

1. «размазывается» в течение всего кода хоста, который на самом деле каким-то образом вызывает эти передачи
2. каким-то образом это происходит «одновременно» с кодом хоста.

Мы также могли бы сделать вывод, что только 1/3 сообщенной активности D-gt;H в единой системе обмена сообщениями фактически происходит после вызова ядра, и поэтому мы могли бы выбрать сравнение только этой части с отчетом D-gt;gt;H из случая, не относящегося к единой системе обмена сообщениями.

Суть всего этого в том, что я не думаю, что это тривиальный вопрос-сравнивать два случая, просто взглянув на тип данных, которые я привел выше. Да, случай с УМ, вероятно, работает хуже, чем случай без УМ. В документации CUDA нигде не указано, что ожидается, что они будут идентичны по производительности. Нет, вы не делаете ничего «неправильного».

FWIW, случай Максвелла/Кеплера в linux выглядит намного лучше, чем в Windows WDDM, поэтому я думаю, что WDDM, вероятно, также связан с менее эффективным поведением.