Параллельность потоков cuFFT

Question

Итак, я использую cuFFT в сочетании с функцией потока CUDA. У меня проблема в том, что я не могу заставить ядра cuFFT работать с полным параллелизмом. Ниже приведены результаты, полученные из nvvp. В каждом потоке работает ядро ​​2D пакетного БПФ на 128 изображений размером 128х128. Я настраиваю 3 потока для запуска 3 независимых плана партии FFT.

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

for (unsigned int j = 0; j < NUM_IMAGES; j++) { 
     gpuErrchk(cudaMemcpyAsync(dev_pointers_in[j], 
           image_vector[j], 
           NX*NY*NZ*sizeof(SimPixelType), 
           cudaMemcpyHostToDevice, 
           streams_fft[j])); 
     gpuErrchk(cudaMemcpyAsync(dev_pointers_out[j], 
           out, 
           NX*NY*NZ*sizeof(cufftDoubleComplex), 
           cudaMemcpyHostToDevice, 
           streams_fft[j])); 
     cufftExecD2Z(planr2c[j], 
        (SimPixelType*)dev_pointers_in[j], 
        (cufftDoubleComplex*)dev_pointers_out[j]); 

    } 

Затем я изменил код так, что я закончил все копии памяти (синхронизировать) и отправить все ядра на потоки сразу, и я получил следующий результат профилирования:

Тогда я подтвердил, что ядра не работают одновременно.

Я посмотрел на один link, который объясняет в деталях, как настроить, чтобы использовать полный параллелизм либо проходящий «-default-потоком каждого потока» аргумент командной строки или #define CUDA_API_PER_THREAD_DEFAULT_STREAM, прежде чем #include или в вашем коде. Это функция, введенная в CUDA 7. Я проверил образец кода в приведенной выше ссылке на моем MacBook Pro Retina 15 'с GeForce GT750M (тот же самый компьютер, что и в приведенной выше ссылке), и я смог получить параллельные прогоны ядра. Но я не смог запустить ядра cuFFT параллельно.

Тогда я нашел это link с кем-то, говорящим, что ядро ​​cuFFT будет занимать весь графический процессор, поэтому нет двух ядер cuFFT, работающих параллельно. Тогда я застрял. Поскольку я не нашел никакой официальной документации, касающейся того, позволяет ли CUFFT использовать параллельные ядра. Это правда? Есть ли способ обойти это?

Answer 1

Предполагаете, что вы назвали cufftSetStream() перед кодом, который вы указали, подходящий для каждого planr2c[j], так что каждый план связан с отдельным потоком. Я не вижу этого в коде, который вы опубликовали. Если вы действительно хотите, чтобы ядра cufft перекрывались с другими ядрами cufft, для необходимо запустить эти ядра для разделения потоков. Таким образом, вызов cufft exec для изображения 0 должен быть запущен в другой поток, чем вызов exec cufft для изображения 1, например.

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

Сказав, что параллельные копии памяти с выполнением ядра, но не с параллельными ядрами, - это то, что я ожидал бы для БПФ разумного размера.

A 128x128 FFT в первом приближении будет вращаться ~ 15 000 потоков, поэтому, если мои потоковые блоки составляют ~ 500 нитей, это будет 30 потоков, что будет держать GPU довольно занятым, оставляя не много «комнаты», для дополнительных ядер. (Фактически вы можете обнаружить общие блоки и потоки для ядра в самом профилировщике.) Ваш GT750m probably has 2 Kepler SMs с a maximum of 16 blocks per SM, поэтому максимальная мгновенная емкость составляет 32 блока. И этот номер емкости может быть уменьшен для конкретного ядра из-за использования общей памяти, использования регистра или других факторов.

Мгновенная пропускная способность любого GPU, на котором вы работаете (максимальные блоки на количество SM SM SM), определит потенциал перекрытия (параллелизма) ядер. Если вы превысите эту емкость при запуске одного ядра, то это будет «заполнять» графический процессор, предотвращая параллелизм ядра в течение некоторого периода времени.

Теоретически возможно, чтобы ядра CUFFT запускались одновременно. Но, как и любой сценарий параллелизма ядра, CUFFT или иначе, использование ресурсов этими ядрами должно быть довольно низким, чтобы фактически свидетельствовать о параллелизме. Обычно, когда у вас мало ресурсов, это подразумевает ядра с относительно небольшим количеством потоков/потоков. Эти ядра обычно не требуют много времени для выполнения, что еще более затрудняет фактическое подтверждение параллелизма (поскольку задержка запуска и другие факторы задержки могут мешать). Самый простой способ наблюдать параллельные ядра - это иметь ядра с необычно низкими требованиями к ресурсам в сочетании с необычно долгим временем выполнения. Обычно это не типичный сценарий для ядер CUFFT или любых других ядер.

Перекрытие копии и вычисления по-прежнему является полезной функцией потоков с CUFFT. И идея параллелизма, без основы понимания емкости машины и ограничений ресурсов, несколько неразумна сама по себе. Например, если параллелизм ядра был произвольным достижимым («я должен иметь возможность запускать любые 2 ядра одновременно»), без учета емкости или специфики ресурсов, то после того, как вы одновременно запускаете два ядра, следующим логическим шагом будет перейти к 4, 8, 16 ядрам одновременно. Но реальность такова, что машина не может справиться с такой работой одновременно. Как только вы обнаружите достаточный параллелизм (свободно переводимый как «достаточный поток») в одном запуске ядра, выставляя дополнительный параллелизм работы через дополнительные запуска ядра, обычно не может заставить машину работать быстрее или быстрее обрабатывать работу.