EDIT: Как было отмечено в комментариях, если тот же самый план (тот же созданный дескриптор) используется для одновременного выполнения БПФ на одном устройстве с помощью потоков, то the user is responsible for managing separate work areas for each usage of such plan. Казалось, что вопрос фокусируется на самом поведении потока, и мой оставшийся ответ также фокусируется на этом, но это важный момент.
Если я помещаю план CUFFT в потоке 0, используя cuFFTSetStream() для плитки 0, и FFT для плитки 0 не был фактически выполнен на GPU еще до хоста устанавливает поток в общем плане CUFFT к поток 1 для плитки 1, прежде чем он выдает работу плитки 1 на графическом процессоре, каково поведение cuFFTExec() для этого плана?
Позвольте мне притвориться, что вы сказали поток 1 и поток 2, просто чтобы избежать возможной путаницы вокруг потока NULL.
CUFFT должен уважать поток, который был определен для плана во время передачи плана CUFFT через cufftExecXXX()
. Последующие изменения плана с помощью cufftSetStream()
не должны влиять на поток, используемый для ранее выпущенных вызовов cufftExecXXX()
.
Мы можем проверить это с помощью довольно простого теста, используя профилировщик. Рассмотрим следующий код теста:
$ cat t1089.cu
// NOTE: this code omits independent work-area handling for each plan
// which is necessary for a plan that will be shared between streams
// and executed concurrently
#include <cufft.h>
#include <assert.h>
#include <nvToolsExt.h>
#define DSIZE 1048576
#define BATCH 100
int main(){
const int nx = DSIZE;
const int nb = BATCH;
size_t ws = 0;
cufftHandle plan;
cufftResult res = cufftCreate(&plan);
assert(res == CUFFT_SUCCESS);
res = cufftMakePlan1d(plan, nx, CUFFT_C2C, nb, &ws);
assert(res == CUFFT_SUCCESS);
cufftComplex *d;
cudaMalloc(&d, nx*nb*sizeof(cufftComplex));
cudaMemset(d, 0, nx*nb*sizeof(cufftComplex));
cudaStream_t s1, s2;
cudaStreamCreate(&s1);
cudaStreamCreate(&s2);
res = cufftSetStream(plan, s1);
assert(res == CUFFT_SUCCESS);
res = cufftExecC2C(plan, d, d, CUFFT_FORWARD);
assert(res == CUFFT_SUCCESS);
res = cufftSetStream(plan, s2);
assert(res == CUFFT_SUCCESS);
nvtxMarkA("plan stream change");
res = cufftExecC2C(plan, d, d, CUFFT_FORWARD);
assert(res == CUFFT_SUCCESS);
cudaDeviceSynchronize();
return 0;
}
$ nvcc -o t1089 t1089.cu -lcufft -lnvToolsExt
$ cuda-memcheck ./t1089
========= CUDA-MEMCHECK
========= ERROR SUMMARY: 0 errors
$
Мы просто делаем два вперед БПФ подряд, переключение потоков в промежутке между ними. Мы будем использовать nvtx marker, чтобы четко определить точку, с которой происходит запрос изменения связи потока плана. Теперь давайте посмотрим на nvprof --print-api-trace
выходе (удаление пространное запуска до преамбулы):
983.84ms 617.00us cudaMalloc
984.46ms 21.628us cudaMemset
984.48ms 37.546us cudaStreamCreate
984.52ms 121.34us cudaStreamCreate
984.65ms 995ns cudaPeekAtLastError
984.67ms 996ns cudaConfigureCall
984.67ms 517ns cudaSetupArgument
984.67ms 21.908us cudaLaunch (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>) [416])
984.69ms 349ns cudaGetLastError
984.69ms 203ns cudaPeekAtLastError
984.70ms 296ns cudaConfigureCall
984.70ms 216ns cudaSetupArgument
984.70ms 8.8920us cudaLaunch (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>) [421])
984.71ms 272ns cudaGetLastError
984.71ms 177ns cudaPeekAtLastError
984.72ms 314ns cudaConfigureCall
984.72ms 229ns cudaSetupArgument
984.72ms 9.9230us cudaLaunch (void spRadix0256B::kernel3Mem<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=2, L1, ALL, WRITEBACK>(kernel_parameters_t<fft_mem_radix3_t, unsigned int, float>) [426])
984.73ms 295ns cudaGetLastError
984.77ms - [Marker] plan stream change
984.77ms 434ns cudaPeekAtLastError
984.78ms 357ns cudaConfigureCall
984.78ms 228ns cudaSetupArgument
984.78ms 10.642us cudaLaunch (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>) [431])
984.79ms 287ns cudaGetLastError
984.79ms 193ns cudaPeekAtLastError
984.80ms 293ns cudaConfigureCall
984.80ms 208ns cudaSetupArgument
984.80ms 7.7620us cudaLaunch (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>) [436])
984.81ms 297ns cudaGetLastError
984.81ms 178ns cudaPeekAtLastError
984.81ms 269ns cudaConfigureCall
984.81ms 214ns cudaSetupArgument
984.81ms 7.4130us cudaLaunch (void spRadix0256B::kernel3Mem<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=2, L1, ALL, WRITEBACK>(kernel_parameters_t<fft_mem_radix3_t, unsigned int, float>) [441])
984.82ms 312ns cudaGetLastError
984.82ms 152.63ms cudaDeviceSynchronize
$
Мы видим, что каждая операция FFT требует 3 вызовы ядра. В промежутке мы видим маркер nvtx, указывающий, когда был сделан запрос на изменение потока плана, и неудивительно, что это происходит после запуска первых 3 ядер, но до последнего 3. Наконец, отметим, что по существу все время исполнения поглощается в окончательном вызове cudaDeviceSynchronize()
.Все предыдущие вызовы являются асинхронными и поэтому выполняются более или менее «немедленно» в течение первой миллисекунды выполнения. Окончательная синхронизация поглощает все время обработки 6 ядер, составляя около 150 миллисекунд.
Так что, если cufftSetStream
должны были иметь эффект на первой итерации cufftExecC2C()
вызова, можно было бы ожидать, чтобы увидеть некоторые или все из первых 3-х ядер, запускаемыми в том же потоке, который использовался в течение последних 3-х ядер. Но когда мы смотрим на nvprof --print-gpu-trace
выход:
$ nvprof --print-gpu-trace ./t1089
==3757== NVPROF is profiling process 3757, command: ./t1089
==3757== Profiling application: ./t1089
==3757== Profiling result:
Start Duration Grid Size Block Size Regs* SSMem* DSMem* Size Throughput Device Context Stream Name
974.74ms 7.3440ms - - - - - 800.00MB 106.38GB/s Quadro 5000 (0) 1 7 [CUDA memset]
982.09ms 23.424ms (25600 2 1) (32 8 1) 32 8.0000KB 0B - - Quadro 5000 (0) 1 13 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>) [416]
1.00551s 21.172ms (25600 2 1) (32 8 1) 32 8.0000KB 0B - - Quadro 5000 (0) 1 13 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>) [421]
1.02669s 27.551ms (25600 1 1) (16 16 1) 61 17.000KB 0B - - Quadro 5000 (0) 1 13 void spRadix0256B::kernel3Mem<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=2, L1, ALL, WRITEBACK>(kernel_parameters_t<fft_mem_radix3_t, unsigned int, float>) [426]
1.05422s 23.592ms (25600 2 1) (32 8 1) 32 8.0000KB 0B - - Quadro 5000 (0) 1 14 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>) [431]
1.07781s 21.157ms (25600 2 1) (32 8 1) 32 8.0000KB 0B - - Quadro 5000 (0) 1 14 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>) [436]
1.09897s 27.913ms (25600 1 1) (16 16 1) 61 17.000KB 0B - - Quadro 5000 (0) 1 14 void spRadix0256B::kernel3Mem<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=2, L1, ALL, WRITEBACK>(kernel_parameters_t<fft_mem_radix3_t, unsigned int, float>) [441]
Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.
$
мы видим, что на самом деле первые 3 ядра выпускаются в первый поток, и последние 3 ядра выпускаются во второй поток, так же, как просили. (И общее время выполнения всех ядер составляет приблизительно 150 мс, как и предполагалось выходом трассировки api.) Поскольку запуски ядра ядра асинхронны и выдаются до возврата вызова cufftExecC2C()
, если вы внимательно об этом подумаете, Придется прийти к выводу, что так должно быть. Поток запуска ядра в указан во время запуска ядра. (И, конечно, я считаю, что это считается «предпочтительным» поведением.)
В этом случае (тот же план повторно использовался в 2 потоках), если бы ядра были запущены одновременно, был бы конфликт рабочей зоны? Что останавливает ядра в потоке 14 от перезаписи промежуточных результатов, используемых потоком 13? – KQS
Да, это правильно. Обновлен мой ответ. В этом конкретном случае, похоже, не существует какого-либо перекрытия фактического исполнения на основе вывода профилировщика, и это, как правило, мой опыт с БПФ любого разумного размера, но ваше наблюдение является правильным, если есть перекрытие (фактическое) или (для правильности). –