Как правильно использовать шаг в cuda для уменьшения количества блоков? - PullRequest
Новая
       65

Как правильно использовать шаг в cuda для уменьшения количества блоков?

0 голосов
/ 31 марта 2020

Привет всем, я пытаюсь использовать метод grid-stride и функции atomi c для выполнения многоблочного сокращения.
Я знаю, что обычный способ сделать это - запустить два ядра или использовать метод lastblock как направлено в этой заметке. (или в этом учебном пособии)

Однако я подумал, что это также можно сделать, используя grid-stepde с кодом atomi c.
Как я проверял, это работало очень хорошо ...
, пока для некоторого числа не получится неправильный ответ. (что очень странно)

Я проверил некоторые "n" и обнаружил, что я получаю неправильный ответ для n = 1234565, 1234566, 1234567.
Это весь мой код выполнения n суммы 1. Таким образом, ответ должен быть n.
Любая помощь или комментарий приветствуется. 

#include<iostream>

__global__ void stride_sum(const double* input,
                           const int size,
                           double* sumOut){
    extern __shared__ double sm[];

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

    //doing grid loop using stride method.
    for(unsigned int s=i;
            s<size;
            s+=blockDim.x*gridDim.x){
        sm[tid] = input[i];
        __syncthreads();

        //doing parallel reduction.
        for(unsigned int ss = blockDim.x/2;ss>0;ss>>=1){
            if(tid<ss && tid+ss<size) sm[tid] += sm[tid+ss];
            __syncthreads();
        }

        //atomically add results to sumOut.
        if(tid==0) atomicAdd(sumOut, sm[0]);
    }
}

int main(){

    unsigned int n = 1234567;
    int blockSize = 4;
    int nBlocks = (n + blockSize - 1) / blockSize;
    int sharedMemory = sizeof(double)*blockSize;

    double *data, *sum;

    cudaMallocManaged(&data, sizeof(double)*n);
    cudaMallocManaged(&sum, sizeof(double));

    std::fill_n(data,n,1.);
    std::fill_n(sum,1,0.);

    stride_sum<<<nBlocks, blockSize, sharedMemory>>>(data,n,sum);

    cudaDeviceSynchronize();

    printf("res: 10.f \n",sum[0]);

    cudaFree(data);
    cudaFree(sum);

    return 0;
}

1 Ответ

1 голос
/ 31 марта 2020

Вы ошиблись в своей реализации. Это будет работать: 

__global__ void stride_sum(const double* input,
                           const int size,
                           double* sumOut)
{
    extern __shared__ volatile double sm[];

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

    //doing grid loop using stride method.
    double val = 0.;
    for(unsigned int s=i; s<size; s+=blockDim.x*gridDim.x){
        val += input[i]; 
    }

    // Load partial sum to memory
    sm[tid] = val; 
    __syncthreads();

    //doing parallel reduction.
    for(unsigned int ss = blockDim.x/2;ss>0;ss>>=1){
        if(tid<ss && tid+ss<size) sm[tid] += sm[tid+ss];
        __syncthreads();
    }

   //atomically add results to sumOut.
   if(tid==0) atomicAdd(sumOut, sm[0]);
}

[Никогда не компилируйте и не запускайте, используйте собственный риск]

Короче - делайте суммирование с шагом по сетке, затем один общий ресурс уменьшение памяти, , затем обновление одного атома c. Ваша реализация имеет неопределенное поведение в некоторых местах, особенно условно выполняемые вызовы __syncthreads и использование неинициализированной разделяемой памяти, когда некоторые потоки выпадают из суммы l oop.

...