CUDA: Не параллельный редукционный образец Марка Харриса, просто суммирующий каждый блок нити?

Question

Ссылка на его слайдах:

Вот его код первой версии параллельного сокращения:

__global__ void reduce0(int *g_idata, int *g_odata) { 
extern __shared__ int sdata[]; 

// each thread loads one element from global to shared mem 
unsigned int tid = threadIdx.x; 
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; 
sdata[tid] = g_idata[i]; 
__syncthreads(); 

// do reduction in shared mem 
for(unsigned int s=1; s < blockDim.x; s *= 2) { 
if (tid % (2*s) == 0) { 
sdata[tid] += sdata[tid + s]; 
} 
__syncthreads(); 
} 

// write result for this block to global mem 
if (tid == 0) g_odata[blockIdx.x] = sdata[0]; 
} 

, который позже он оптимизирует. Как это не просто суммировать все int для каждого потока и помещать ответ в другой вектор? Это то, что он должен делать? Разве не * g_odata сам вектор, так как он помещает сумму в каждую «blockIdx.x» точку в векторе? Как вы получаете вектор g_idata для суммирования на один номер?

Answer 1

Как это не просто суммирование всех значений для каждого блока потоков и размещение ответа в другом векторе?

Выполняется именно это.

Это то, что он должен делать?

Да.

Это не g_odata сам вектор, так как он поместил сумму на каждую точку «blockIdx.x» в векторе?

Да, это вектор, содержащий суммы уровня блока.

Как вы можете получить вектор g_idata для одного единственного числа?

Вызвать ядро ​​дважды. Однажды на исходном наборе данных и один раз на векторном выходе из предыдущего вызова (суммы уровня блока). Обратите внимание, что этот второй шаг использует только один блок и требует, чтобы вы могли запускать достаточно потоков на каждый блок, чтобы покрыть весь вектор, по одному потоку на сумму с предыдущего шага. Если вы просмотрите cuda sample code, который предназначен для сопровождения этой презентации, которую вы связали, вы найдете такую ​​вызывающую последовательность, например, в строках 304 и 333 из reduction.cpp. Второй вызов reduce<T> выполняет сокращение, который суммирует частичный блок сумму, как указан в комментарии на линии 324:

304:reduce<T>(n, numThreads, numBlocks, whichKernel, d_idata, d_odata); 

    // check if kernel execution generated an error 
    getLastCudaError("Kernel execution failed"); 

    if (cpuFinalReduction) 
    { 
     // sum partial sums from each block on CPU 
     // copy result from device to host 
     checkCudaErrors(cudaMemcpy(h_odata, d_odata, numBlocks*sizeof(T), cudaMemcpyDeviceToHost)); 

     for (int i=0; i<numBlocks; i++) 
     { 
      gpu_result += h_odata[i]; 
     } 

     needReadBack = false; 
    } 
    else 
    { 
324: // sum partial block sums on GPU 
     int s=numBlocks; 
     int kernel = whichKernel; 

     while (s > cpuFinalThreshold) 
     { 
      int threads = 0, blocks = 0; 
      getNumBlocksAndThreads(kernel, s, maxBlocks, maxThreads, blocks, threads); 

333:  reduce<T>(s, threads, blocks, kernel, d_odata, d_odata); 

отметить, что выходной сигнал d_odata от первого сокращения на линии 304 передаются в качестве ввода ко второму сокращению на линии 333.

отметить также, что необходимость, и этот метод ядра-разложения покрыта в presentation you linked на слайдах 3 - 5.