CUDA - несколько ядер для вычисления одного значения

Question

Эй, я пытаюсь написать ядро, чтобы по существу сделать следующее в C

 float sum = 0.0;
 for(int i = 0; i < N; i++){
   sum += valueArray[i]*valueArray[i];
 }
 sum += sum / N;

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

int i0 = blockIdx.x * blockDim.x + threadIdx.x;

   for(int i=i0; i<N; i += blockDim.x*gridDim.x){
        *d_sum += d_valueArray[i]*d_valueArray[i];
    }

  *d_sum= __fdividef(*d_sum, N);

Код, используемый для вызова ядра:

  kernelName<<<64,128>>>(N, d_valueArray, d_sum);
  cudaMemcpy(&sum, d_sum, sizeof(float) , cudaMemcpyDeviceToHost);

Я думаю, что каждое ядро ​​вычисляет частичную сумму, но окончательный оператор деления не учитывает накопленное значение из каждого потока. Каждое ядро ​​производит свое собственное окончательное значение для d_sum?

Кто-нибудь знает, как я могу сделать это эффективным способом? Может быть, использование общей памяти между потоками? Я очень новичок в программировании на GPU. Приветствия

Answer 1

Вы обновляете d_sum из нескольких потоков.

См. Следующий пример SDK:

http://developer.download.nvidia.com/compute/cuda/sdk/website/samples.html

Вот код из этого примера. Обратите внимание, как это двухэтапный процесс. Суммируйте каждый блок потока и затем __syncthreads, прежде чем пытаться накопить окончательный результат.

#define ACCUM_N 1024
__global__ void scalarProdGPU(
    float *d_C,
    float *d_A,
    float *d_B,
    int vectorN,
    int elementN
){
    //Accumulators cache
    __shared__ float accumResult[ACCUM_N];

    ////////////////////////////////////////////////////////////////////////////
    // Cycle through every pair of vectors,
    // taking into account that vector counts can be different
    // from total number of thread blocks
    ////////////////////////////////////////////////////////////////////////////
    for(int vec = blockIdx.x; vec < vectorN; vec += gridDim.x){
        int vectorBase = IMUL(elementN, vec);
        int vectorEnd  = vectorBase + elementN;

        ////////////////////////////////////////////////////////////////////////
        // Each accumulator cycles through vectors with
        // stride equal to number of total number of accumulators ACCUM_N
        // At this stage ACCUM_N is only preferred be a multiple of warp size
        // to meet memory coalescing alignment constraints.
        ////////////////////////////////////////////////////////////////////////
        for(int iAccum = threadIdx.x; iAccum < ACCUM_N; iAccum += blockDim.x){
            float sum = 0;

            for(int pos = vectorBase + iAccum; pos < vectorEnd; pos += ACCUM_N)
                sum += d_A[pos] * d_B[pos];

            accumResult[iAccum] = sum;
        }

        ////////////////////////////////////////////////////////////////////////
        // Perform tree-like reduction of accumulators' results.
        // ACCUM_N has to be power of two at this stage
        ////////////////////////////////////////////////////////////////////////
        for(int stride = ACCUM_N / 2; stride > 0; stride >>= 1){
            __syncthreads();
            for(int iAccum = threadIdx.x; iAccum < stride; iAccum += blockDim.x)
                accumResult[iAccum] += accumResult[stride + iAccum];
        }

        if(threadIdx.x == 0) d_C[vec] = accumResult[0];
    }
}