Параллельность потоков cuFFT
Поэтому я использую 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]);
}
Затем я изменил свой код так, чтобы я закончил все копии памяти (синхронизировал) и отправил все ядра сразу в потоки, и я получил следующий результат профилирования:
Затем мне подтвердили, что ядра не работают одновременно.
Я посмотрел на одного ссылка на сайт в котором подробно объясняется, как настроить использование полного параллелизма путем передачи аргумента командной строки «–default-stream per-thread» или #define CUDA_API_PER_THREAD_DEFAULT_STREAM перед #include или в вашем коде. Это функция, представленная в CUDA 7. Я запустил пример кода по приведенной выше ссылке на моем MacBook Pro Retina 15 ‘с GeForce GT750M (тот же компьютер, что и в приведенной выше ссылке), и я смог получить параллельные запуски ядра. Но я не смог запустить мои ядра cuFFT параллельно.
Тогда я нашел это ссылка на сайт когда кто-то говорит, что ядро cuFFT будет занимать весь графический процессор, поэтому два ядра cuFFT не будут работать параллельно. Тогда я застрял. Поскольку я не нашел никакой формальной документации, касающейся того, разрешает ли CUFFT параллельные ядра. Это правда? Есть ли способ обойти это?
Решение
Я полагаю, вы звонили cufftSetStream() до кода, который вы показали, подходит для каждого planr2c[j], так что каждый план связан с отдельным потоком. Я не вижу этого в коде, который вы разместили. Если вы действительно хотите, чтобы ядра cufft перекрывались с другими ядрами cufft, это необходимо для те ядра, которые будут запущены в отдельных потоках. Таким образом, вызов cufft exec для изображения 0 должен быть запущен, например, в потоке, отличном от вызова cufft exec для изображения 1.
Для того чтобы любые две операции CUDA чтобы иметь возможность перекрываться, они должен запускаться в разные потоки.
Сказав это, одновременные копии памяти с выполнением ядра, но не одновременные ядра, — это то, что я ожидал бы для разумных размеров FFT.
БПФ 128×128 в приближении первого порядка будет раскручивать ~ 15 000 потоков, так что, если у моих блоков потоков будет ~ 500 потоков каждый, это будет 30 потоковых блоков, что будет держать GPU достаточно занятым, оставляя немного места для дополнительных ядер. (Вы можете узнать общее количество блоков и потоков для ядра в самом профилировщике.) Ваш GT750m вероятно, имеет 2 Kepler SM с максимум 16 блоков на СМ таким образом, максимальная мгновенная емкость 32 блоков. И это число емкости может быть уменьшено для конкретного ядра из-за использования общей памяти, использования регистра или других факторов.
Мгновенная емкость любого графического процессора, на котором вы работаете (максимальное количество блоков на SM * количество SM), будет определять вероятность перекрытия (параллелизма) ядер. Если вы превысите эту емкость за один запуск ядра, это будет «заполнять» графический процессор, предотвращая параллелизм ядра в течение некоторого периода времени.
Теоретически возможно, чтобы ядра CUFFT работали одновременно. Но, как и в любом сценарии параллелизма ядра, CUFFT или иного, использование ресурсов этими ядрами должно быть довольно низким, чтобы фактически обеспечить параллелизм. Обычно, когда вы используете мало ресурсов, это подразумевает ядра с относительно небольшим количеством потоков / потоковых блоков. Эти ядра обычно не требуют много времени для выполнения, что еще более затрудняет фактическое наблюдение параллелизма (потому что задержка запуска и другие факторы задержки могут мешать). Самый простой способ наблюдать параллельные ядра — это иметь ядра с необычно низкими требованиями к ресурсам в сочетании с необычно долгим временем работы. Обычно это не типичный сценарий для ядер CUFFT или любых других ядер.
Наложение копирования и вычисления — все еще полезная функция потоков с CUFFT. И идея параллелизма, без основы понимания производительности машины и ограничений ресурсов, сама по себе несколько неразумна. Например, если параллелизм ядра был произвольно достижимым («я мог бы заставить любые 2 ядра работать одновременно»), без учета емкости или специфики ресурса, то после того, как вы запустили два ядра одновременно, следующим логическим шагом было бы перейти к 4, 8, 16 ядрам одновременно. Но реальность такова, что машина не может справиться с такой большой работой одновременно. После того, как вы продемонстрировали достаточный параллелизм (в широком смысле переводится как «достаточное количество потоков») в одном запуске ядра, раскрытие дополнительной параллелизации работы посредством дополнительных запусков ядра обычно не может заставить машину работать быстрее или быстрее обрабатывать работу.
Другие решения
Других решений пока нет …