CUDA: cufftExecR2C - ненужные копии памяти - PullRequest
Новая
       59

CUDA: cufftExecR2C - ненужные копии памяти

0 голосов
/ 01 июля 2018

Я экспериментирую с cuda и наблюдаю, как данные копируются с хоста на устройство при вызове 

cufftExecR2C(plan, src, dst);

, который я не предпринимаю, поскольку мой указатель src является допустимым дескриптором памяти устройства, которую я хотел бы преобразовать. Перед cufftExecR2C (...) я инициализировал аргументы следующим образом: 

  float* src;
  cudaMalloc((&src),  image_rows * image_cols  * sizeof(float) );
  cudaMemcpy(src, image.data()  ,  image_rows * image_cols  * sizeof(float)  , cudaMemcpyHostToDevice);


cufftComplex* dst;
cudaMalloc((void**)&dst    , image_rows * (image_cols/2+1) * sizeof(cufftComplex) );


 cufftHandle plan;
 cufftPlan2d(&plan, image_rows, image_cols, CUFFT_R2C))

Запуск профилировщика nvidia (nvprof) - только с учетом fft - я получаю следующий результат 

...
cudaProfilerStart();
cufftExecR2C(plan, src, dst);
cudaProfilerStop();
...

enter image description here

Я бы хотел избежать 3 ненужных звонков с устройства на устройство. Я не понимаю, почему cuda выполняет эти дополнительные копии (особенно, почему хост на устройство - данные уже находятся в памяти устройства)?

Программа выполняется на GeForce GT 540M с использованием Cuda 8.0.

Спасибо!

1 Ответ

0 голосов
/ 06 июля 2018

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

Рассмотрим следующий пример, собранный из фрагментов кода, представленных в вашем вопросе: 

#include "cufft.h"
#include "cuda_profiler_api.h"
#include <random>
#include <algorithm>
#include <iterator>
#include <iostream>
#include <functional>

int main()
{
  const int image_rows = 1600, image_cols = 2048;

  std::random_device rnd_device;
  std::mt19937 mersenne_engine {rnd_device()};
  std::uniform_real_distribution<float> dist {0.0, 255.0};

  auto gen = [&dist, &mersenne_engine](){
                 return dist(mersenne_engine);
             };

  std::vector<float> image(image_rows * image_cols);
  std::generate(std::begin(image), std::end(image), gen);

  float* src;
  cudaMalloc((&src),  image_rows * image_cols  * sizeof(float) );
  cudaMemcpy(src, &image[0],  image_rows * image_cols  * sizeof(float)  , cudaMemcpyHostToDevice);
  cufftComplex* dst;
  cudaMalloc((void**)&dst    , image_rows * (image_cols/2+1) * sizeof(cufftComplex) );

  cufftHandle plan;
  cufftPlan2d(&plan, image_rows, image_cols, CUFFT_R2C);

  cudaProfilerStart();
  cufftExecR2C(plan, src, dst);
  cudaProfilerStop();

  return 0;
}

Я подставил массив случайных значений для вашего изображения. Теперь давайте скомпилируем и профилируем это: 

$ nvcc -std=c++11 -o unecessary unecessary.cu -lcufft
$ nvprof ./unecessary
==10314== NVPROF is profiling process 10314, command: ./unecessary
==10314== Profiling application: ./unecessary
==10314== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   74.39%  2.2136ms         1  2.2136ms  2.2136ms  2.2136ms  [CUDA memcpy HtoD]
                    6.66%  198.30us         1  198.30us  198.30us  198.30us  void spRadix0064B::kernel1Mem<unsigned int, float, fftDirection_t=-1, unsigned int=32, unsigned int=4, CONSTANT, ALL, WRITEBACK>(kernel_parameters_t<fft_mem_radix1_t, unsigned int, float>)
                    6.50%  193.47us         1  193.47us  193.47us  193.47us  void spRadix0025B::kernel1Mem<unsigned int, float, fftDirection_t=-1, unsigned int=64, unsigned int=4, CONSTANT, ALL, WRITEBACK>(kernel_parameters_t<fft_mem_radix1_t, unsigned int, float>)
                    6.25%  185.98us         1  185.98us  185.98us  185.98us  void spVector1024C::kernelMem<unsigned int, float, fftDirection_t=-1, unsigned int=2, unsigned int=5, LUT, ALL, WRITEBACK>(kernel_parameters_t<fft_mem_t, unsigned int, float>)
                    6.20%  184.38us         1  184.38us  184.38us  184.38us  __nv_static_45__32_spRealComplex_compute_70_cpp1_ii_1f28721c__ZN13spRealComplex24postprocessC2C_kernelMemIjfL9fftAxii_t3EEEvP7ComplexIT0_EPKS4_T_15coordDivisors_tIS8_E7coord_tIS8_ESC_S8_S3_10callback_t

[Вызовы API удалены для краткости]

Похоже, ты прав! Огромный memcpy прямо в сводной статистике GPU!

Итак, давайте профилируем это снова правильно : 

$ nvprof --profile-from-start off ./unecessary
==11674== NVPROF is profiling process 11674, command: ./unecessary
==11674== Profiling application: ./unecessary
==11674== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   25.96%  196.28us         1  196.28us  196.28us  196.28us  void spRadix0064B::kernel1Mem<unsigned int, float, fftDirection_t=-1, unsigned int=32, unsigned int=4, CONSTANT, ALL, WRITEBACK>(kernel_parameters_t<fft_mem_radix1_t, unsigned int, float>)
                   25.25%  190.91us         1  190.91us  190.91us  190.91us  void spRadix0025B::kernel1Mem<unsigned int, float, fftDirection_t=-1, unsigned int=64, unsigned int=4, CONSTANT, ALL, WRITEBACK>(kernel_parameters_t<fft_mem_radix1_t, unsigned int, float>)
                   24.65%  186.39us         1  186.39us  186.39us  186.39us  void spVector1024C::kernelMem<unsigned int, float, fftDirection_t=-1, unsigned int=2, unsigned int=5, LUT, ALL, WRITEBACK>(kernel_parameters_t<fft_mem_t, unsigned int, float>)
                   24.15%  182.59us         1  182.59us  182.59us  182.59us  __nv_static_45__32_spRealComplex_compute_70_cpp1_ii_1f28721c__ZN13spRealComplex24postprocessC2C_kernelMemIjfL9fftAxii_t3EEEvP7ComplexIT0_EPKS4_T_15coordDivisors_tIS8_E7coord_tIS8_ESC_S8_S3_10callback_t

[Опять же, вызовы API удалены для краткости]

memcpy ушел. Профилировщик сообщает только о четырех запусках ядра, связанных с выполнением преобразования. Нет передачи памяти. Передача памяти, о которой сообщается в исходном выводе профилировщика, является передачей от хоста к устройству в начале программы и не связана с вызовом cuFFT. Причина, по которой он включен, состоит в том, что по умолчанию nvprof включает профилирование с начала выполнения программы, и первоначальный вызов cudaProfilerStart не имеет никакого эффекта, поскольку профилирование уже включено. О правильном способе кодирования профиля вы можете прочитать в документации по здесь .

Я предложу свою собственную гипотезу в отсутствие обещанного MCVE - вы не правильно использовали профилировщик, и сообщенные переносы, по сути, являются переносами, которые происходят в другом месте вашего кода и которые включены в профилировщик вывод, но совершенно не связаны с работой cuFFT.

...