2016-02-18 7 views
0

Я использую CUDA 7.0 с nVidia 980 GTX для некоторой обработки изображений. В конкретной итерации несколько фрагментов обрабатываются независимо через 15-20 вызовов ядра и многочисленные вызовы API cuFFT FFT/IFFT.Поведение API CUDA cuFFT в параллельных потоках

Из-за этого я поместил каждую плиту в свой собственный поток CUDA, чтобы каждая плитка выполняла свою последовательность операций асинхронно по отношению к хосту. Каждая плитка имеет одинаковый размер в рамках итерации, поэтому они разделяют план cuFFT. Хост-хост быстро перемещается по командам, пытаясь сохранить работу графического процессора. Я испытываю периодическое состояние гонки, в то время как эти операции обрабатываются параллельно, и, в частности, вопрос о cuFFT. Если я положу план cuFFT в потоке 0, используя cuFFTSetStream() для tile 0, и FFT для фрагмента 0 фактически не был выполнен на графическом процессоре еще до того, как хост установит поток общего потока cuFFT в поток 1 для фрагмента 1 перед он обрабатывает работу плитки 1 на графическом процессоре, каково поведение cuFFTExec() для этого плана?

С другой стороны, вызов cufftExec() выполняется в потоке, который был установлен в момент вызова cufftExec() независимо от того, используется ли cuFFTSetStream() для изменения потока для последующих фрагментов перед предыдущим FFT звонки действительно начаты/завершены?

Приносим извинения, что не отправлял код, но я не могу опубликовать свой фактический источник.

ответ

2

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(), если вы внимательно об этом подумаете, Придется прийти к выводу, что так должно быть. Поток запуска ядра в указан во время запуска ядра. (И, конечно, я считаю, что это считается «предпочтительным» поведением.)

+0

В этом случае (тот же план повторно использовался в 2 потоках), если бы ядра были запущены одновременно, был бы конфликт рабочей зоны? Что останавливает ядра в потоке 14 от перезаписи промежуточных результатов, используемых потоком 13? – KQS

+0

Да, это правильно. Обновлен мой ответ. В этом конкретном случае, похоже, не существует какого-либо перекрытия фактического исполнения на основе вывода профилировщика, и это, как правило, мой опыт с БПФ любого разумного размера, но ваше наблюдение является правильным, если есть перекрытие (фактическое) или (для правильности). –

 Смежные вопросы

  • Нет связанных вопросов^_^