Сокращение блоков в CUDA

Я пытаюсь сделать сокращение в CUDA, и я действительно новичок. В настоящее время я изучаю пример кода от NVIDIA.

Я думаю, я действительно не уверен, как настроить размер блока и размер сетки, особенно когда мой входной массив больше (512 X 512), чем размер одного блока.

Вот код.

template <unsigned int blockSize>
__global__ void reduce6(int *g_idata, int *g_odata, unsigned int n)
{
    extern __shared__ int sdata[];
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*(blockSize*2) + tid;
    unsigned int gridSize = blockSize*2*gridDim.x;
    sdata[tid] = 0;

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

    __syncthreads();

    if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); }
    if (blockSize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); }
    if (blockSize >= 128) { if (tid < 64) { sdata[tid] += sdata[tid + 64]; } __syncthreads(); }

    if (tid < 32) 
    {
        if (blockSize >= 64) sdata[tid] += sdata[tid + 32];
        if (blockSize >= 32) sdata[tid] += sdata[tid + 16];
        if (blockSize >= 16) sdata[tid] += sdata[tid + 8];
        if (blockSize >= 8) sdata[tid] += sdata[tid + 4];
        if (blockSize >= 4) sdata[tid] += sdata[tid + 2];
        if (blockSize >= 2) sdata[tid] += sdata[tid + 1];
    }

    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

Однако мне кажется, что g_odata[blockIdx.x] сохраняет частичные суммы из всех блоков, и, если я хочу получить окончательный результат, мне нужно просуммировать все члены в массиве g_odata[blockIdx.x].

Мне интересно: есть ли ядро ​​для всего суммирования? или я неправильно понимаю вещи здесь? Я был бы очень признателен, если бы кто-нибудь мог научить меня этому. Большое спасибо.


algorithm cuda reduction cub
person Ono    schedule 08.04.2014    source источник
comment
Также обратите внимание, что __shared__ данные должны быть volatile в приведенном выше коде, иначе правильный окончательный результат не может быть гарантирован. Это можно увидеть по ссылке, предоставленной @Robert.   -  person Farzad    schedule 08.04.2014

Ответы (3)


arrow_upward
7
arrow_downward

Чтобы лучше понять эту тему, вы можете посмотреть это PDF от NVIDIA, который графически объясняет все стратегии, которые вы использовали в своем коде.

person Leos313    schedule 18.11.2015

arrow_upward
14
arrow_downward

Ваше понимание правильное. Сокращения, продемонстрированные здесь, заканчиваются последовательностью блочные суммы хранятся в глобальной памяти.

Чтобы суммировать все эти суммы блоков вместе, требуется некоторая форма глобальной синхронизации. Вы должны дождаться завершения всех блоков, прежде чем складывать их суммы вместе. На этом этапе у вас есть несколько вариантов, некоторые из которых:

  1. запустить новое ядро ​​после основного ядра, чтобы суммировать суммы блоков вместе
  2. добавить суммы блоков на хост
  3. используйте atomics, чтобы сложить суммы блоков вместе, в конце основного ядра
  4. используйте такой метод, как уменьшение числа потоков, чтобы сложить суммы блоков вместе в основном ядре.
  5. Используйте совместные группы CUDA, чтобы синхронизировать всю сетку в коде ядра. Суммируйте суммы блоков после синхронизации всей сетки (возможно, в одном блоке).

Если поискать по тегу CUDA, можно найти примеры всего этого и обсуждение их плюсов и минусов. Чтобы увидеть, как основное ядро, которое вы опубликовали, используется для полной редукции, посмотрите ссылку пример кода параллельного сокращения.

person Robert Crovella    schedule 08.04.2014

arrow_upward
8
arrow_downward

Роберт Кровелла уже ответил на этот вопрос, который в основном касается понимания, а не производительности.

Тем не менее, для всех, кто сталкивается с этим вопросом, я просто хочу подчеркнуть, что CUB делает доступными функции сокращения блоков. . Ниже я привожу простой рабочий пример использования BlockReduce CUB.

#include <cub/cub.cuh>
#include <cuda.h>

#include "Utilities.cuh"

#include <iostream>

#define BLOCKSIZE   32

const int N = 1024;

/**************************/
/* BLOCK REDUCTION KERNEL */
/**************************/
__global__ void sum(const float * __restrict__ indata, float * __restrict__ outdata) {

    unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;

    // --- Specialize BlockReduce for type float. 
    typedef cub::BlockReduce<float, BLOCKSIZE> BlockReduceT; 

    // --- Allocate temporary storage in shared memory 
    __shared__ typename BlockReduceT::TempStorage temp_storage; 

    float result;
    if(tid < N) result = BlockReduceT(temp_storage).Sum(indata[tid]);

    // --- Update block reduction value
    if(threadIdx.x == 0) outdata[blockIdx.x] = result;

    return;  
}

/********/
/* MAIN */
/********/
int main() {

    // --- Allocate host side space for 
    float *h_data       = (float *)malloc(N * sizeof(float));
    float *h_result     = (float *)malloc((N / BLOCKSIZE) * sizeof(float));

    float *d_data;      gpuErrchk(cudaMalloc(&d_data, N * sizeof(float)));
    float *d_result;    gpuErrchk(cudaMalloc(&d_result, (N / BLOCKSIZE) * sizeof(float)));

    for (int i = 0; i < N; i++) h_data[i] = (float)i;

    gpuErrchk(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice));

    sum<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_data, d_result);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    gpuErrchk(cudaMemcpy(h_result, d_result, (N / BLOCKSIZE) * sizeof(float), cudaMemcpyDeviceToHost));

    std::cout << "output: ";
    for(int i = 0; i < (N / BLOCKSIZE); i++) std::cout << h_result[i] << " ";
    std::cout << std::endl;

    gpuErrchk(cudaFree(d_data));
    gpuErrchk(cudaFree(d_result));

    return 0;
}

В этом примере создается массив длиной N, результатом которого является сумма 32 последовательных элементов. Так

result[0] = data[0] + ... + data[31];
result[1] = data[32] + ... + data[63];
....
person Vitality    schedule 30.07.2015