Сокращение CUDA, подход для больших массивов

Question

У меня есть следующий код сокращения суммы «Франкенштейна», взятый частично из common CUDA reduction slices, частично из образцов CUDA.

__global__ void reduce6(float *g_idata, float *g_odata, unsigned int n) 
{ 
    extern __shared__ float sdata[]; 

    // perform first level of reduction, 
    // reading from global memory, writing to shared memory 
    unsigned int tid = threadIdx.x; 
    unsigned int i = blockIdx.x*blockSize*2 + threadIdx.x; 
    unsigned int gridSize = blockSize*2*gridDim.x; 
    sdata[tid] = 0; 
    float mySum = 0; 

    while (i < n) { 
     sdata[tid] += g_idata[i] + g_idata[i+MAXTREADS]; 
     i += gridSize; 
    } 
    __syncthreads(); 


    // do reduction in shared mem 
    if (tid < 256) 
     sdata[tid] += sdata[tid + 256]; 
    __syncthreads(); 

    if (tid < 128) 
     sdata[tid] += sdata[tid + 128]; 
    __syncthreads(); 

    if (tid < 64) 
     sdata[tid] += sdata[tid + 64]; 
    __syncthreads(); 


#if (__CUDA_ARCH__ >= 300) 
    if (tid < 32) 
    { 
     // Fetch final intermediate sum from 2nd warp 
     mySum = sdata[tid]+ sdata[tid + 32]; 
     // Reduce final warp using shuffle 
     for (int offset = warpSize/2; offset > 0; offset /= 2) 
      mySum += __shfl_down(mySum, offset); 
    } 
    sdata[0]=mySum; 
#else 

    // fully unroll reduction within a single warp 
    if (tid < 32) { 
     sdata[tid] += sdata[tid + 32]; 
     sdata[tid] += sdata[tid + 16]; 
     sdata[tid] += sdata[tid + 8]; 
     sdata[tid] += sdata[tid + 4]; 
     sdata[tid] += sdata[tid + 2]; 
     sdata[tid] += sdata[tid + 1]; 
    } 
#endif 
    // write result for this block to global mem 
    if (tid == 0) g_odata[blockIdx.x] = sdata[0]; 
    } 

Я буду использовать это, чтобы уменьшить развернутый массив большого размера (например, 512^3 = 134217728 = n) на Теслах K40 GPU.

У меня есть некоторые вопросы относительно переменной blockSize и ее значения.

С этого момента, я попытаюсь объяснить свое понимание (или правильно или неправильно) о том, как это работает:

Чем больше я выбираю blockSize, тем быстрее этот код будет выполняться, так как он будет тратить меньше времени весь цикл, но он не завершит сокращение всего массива, но он вернет меньший массив размером dimBlock.x, правильно? Если я использую blockSize=1, этот код вернется в 1 вызов значения уменьшения, но он будет очень медленным, потому что он не использует силу CUDA практически ничем. Поэтому мне нужно вызвать ядро ​​сокращения несколько раз, каждый раз с меньшим blokSize и уменьшая результат предыдущего вызова, чтобы уменьшить, пока не дойду до наименьшей точки.

что-то вроде (pesudocode)

blocks=number; //where do we start? why? 
while(not the min){ 

    dim3 dimBlock(blocks); 
    dim3 dimGrid(n/dimBlock.x); 
    int smemSize = dimBlock.x * sizeof(float); 
    reduce6<<<dimGrid, dimBlock, smemSize>>>(in, out, n); 

    in=out; 

    n=dimGrid.x; 
    dimGrid.x=n/dimBlock.x; // is this right? Should I also change dimBlock? 
} 

В какое значение я должен начать? Я думаю, это зависит от GPU. Какие значения shoudl это для Tesla k40 (просто для меня, чтобы понять, как эти значения выбраны)?

Является ли моя логика некорректной? как?

Answer 1

Существует инструмент CUDA для получения хороших размеров сетки и блоков для вас: Cuda Occupancy API.

В ответ на «Чем больше я выбираю BLOCKSIZE, тем быстрее этот код будет выполняться» - Не обязательно, как вы хотите, чтобы размеры, которые дают максимальную occupancy (соотношение активных перекосов к общему числу возможных активным перекосов).

См. Этот ответ для получения дополнительной информации How do I choose grid and block dimensions for CUDA kernels?.

Наконец, для графических процессоров Nvidia, поддерживающих Kelper или более поздних версий, есть shuffle intrinsics, чтобы упростить и ускорить сокращения. Вот статья о том, как использовать перетасовку: Faster Parallel Reductions on Kepler.

Обновление для выбора количества потоков:

Вы не можете использовать максимальное количество потоков, если это приводит к менее эффективное использование регистров. От ссылки на заполнение:

Для расчета количества посещений число регистров, используемых каждым потоком, является одним из ключевых факторов. Например, устройства с вычислительной способностью 1.1 имеют 8 192 32-разрядных регистра на один процессор и могут иметь не более 768 одновременных потоков (24 перелома x 32 потока на дефис). Это означает, что в одном из этих устройств для многопроцессорного устройства, имеющего 100% -ное заполнение, каждый поток может использовать не более 10 регистров. Однако этот подход определения того, как количество регистров влияет на занятость, не учитывает гранулярность распределения регистров. Например, на устройстве вычислительной возможности 1.1, ядро ​​с 128-потоковыми блоками с использованием 12 регистров на поток приводит к заполнению 83% с 5 активными 128-потоковыми блоками на каждый процессор, тогда как ядро ​​с 256-потоковыми блоками, использующее те же 12 регистров на поток, приводит к занятость 66%, потому что только два блока с 256 потоками могут находиться на мультипроцессоре.

Таким образом, я понимаю, что увеличение числа потоков имеет потенциал для ограничения производительности из-за способа распределения регистров. Однако это не всегда так, и вам нужно сделать расчет (как в приведенном выше описании) самостоятельно, чтобы определить оптимальное количество потоков на блок.