Bloc de réduction de CUDA

Je suis en train de faire de la réduction de CUDA et je suis vraiment un débutant. Je suis actuellement en étude d'un exemple de code à partir de NVIDIA.

Je suppose que je ne suis pas vraiment sûr de savoir comment configurer le bloc de la taille et de la taille de la grille, en particulier quand mon tableau d'entrée est plus grande (512 X 512) qu'un seul bloc de taille.

Voici le code.

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];
}

Cependant, il me semble que la g_odata[blockIdx.x] enregistre les sommes partielles de tous les blocs, et, si je veux obtenir le résultat final, j'ai besoin de la somme de tous les termes à l'intérieur de la g_odata[blockIdx.x] tableau.

Je me demande: est-il un noyau à faire la sommation? ou suis-je incompréhension des choses ici? J'apprécierais vraiment si quelqu'un peut m'instruire avec ce. Merci beaucoup.

Veuillez également noter __shared__ les données doivent être volatile dans le code ci-dessus, sinon corriger résultat final ne peut pas être garantie. Il peut être vu dans le lien @Robert fourni.
OK. Merci man:)

OriginalL'auteur Ono | 2014-04-08

  1. 4

    Afin d'avoir une meilleure idée de cette rubrique, vous pouvez avoir un regard sur ce pdf de NVIDIA qui explique, graphiquement, toutes les stratégies que vous avez utilisées dans votre code.

    OriginalL'auteur Leos313

  2. 11

    Votre compréhension est correcte. Les réductions démontré ici jusqu'à la fin avec une séquence de bloc-sommes déposées dans la mémoire globale.

    À la somme de toutes ces blocs sommes ensemble, exige une certaine forme de synchronisation globale. Vous devez attendre jusqu'à ce que tous les blocs sont complets avant de les ajouter leur sommes ensemble. Vous avez un certain nombre d'options à ce stade, certains sont:

    1. le lancement d'un nouveau noyau après le principal noyau de la somme de la bloquer-les sommes ensemble
    2. ajouter le bloc sommes sur l'hôte
    3. utilisation atomics pour ajouter le bloc sommes ensemble, à la fin de la principale du noyau
    4. utiliser une méthode comme threadfence de réduction pour ajouter le bloc sommes ensemble dans le principal noyau.

    Si vous effectuez une recherche autour du CUDA de la marque que vous pouvez trouver des exemples de toutes ces choses, et les discussions de leurs avantages et inconvénients. Pour voir comment le principal noyau que vous avez posté est utilisé pour une réduction complète, regardez la réduction parallèle des exemples de code.

    Je vois... merci!
    Un supplément pour la boucle dans le code de CPU a l'ajouter de l'emploi.

    OriginalL'auteur Robert Crovella

  3. 6

    Robert Crovella a déjà répondu à cette question, qui est principalement sur la compréhension plutôt que sur la performance.

    Cependant, pour tous ceux qui se cogner dans cette question, je veux juste souligner que CUB fait de bloquer les fonctions de réduction disponibles. Ci-dessous, je donne un simple exemple sur la manière dont l'utilisation de la CUB BlockReduce.

    #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;
}

    Dans cet exemple, un tableau de longueur N est créé et le résultat est la somme de 32 consécutives éléments. Donc

    result[0] = data[0] + ... + data[31];
result[1] = data[32] + ... + data[63];
....

    OriginalL'auteur JackOLantern