Histogramme de programmation Cuda

Question

Je veux lancer un programme cuda, mais je suis un débutant. Je dois écrire un programme pour un histogramme. Mais avec des seaux. En fonction de la valeur maxValue (40 dans l'exemple), le nombre sera ajouté au compartiment approprié. Si nous avons 4 seaux:

histo: | 1 | 10 | 30 | 39 | 32 | 2 | 4 | 5 | 1 |

0-9 (1er seau)

10-19 (2e seau)

20-29 (3e seau)

30- 39 (4e seau)

Mon GPU a Capacité de calcul 1.1.

Je tentais de faire quelque chose comme ayant une température commune [] pour un bloc que chaque thread ajoute ses valeurs sur sa table temporaire:

__global__ void histo_kernel_optimized5(unsigned char *buffer, long size, 
           unsigned int *histo) 
{ 
    extern __shared__ unsigned int temp[]; 
    temp[threadIdx.x] = 0; 
    __syncthreads(); 

    int i = threadIdx.x + blockIdx.x * blockDim.x; 
    int offset = blockDim.x * gridDim.x; 
    int bucketID; 
    while (i < size) 
    { 
       bucketID = array[i]/Bwidth; 
       atomicAdd(&temp[bucketID], 1); 
       i += offset; 
    } 
    __syncthreads(); 


    atomicAdd(&(histo[threadIdx.x]), temp[threadIdx.x]); 
} 

histo_kernel_optimized <<<array_size/buckets, buckets,buckets*sizeof(unsigned int)>>>(buffer,SIZE, histogram) 

Mais le compilateur sais: Instruction « {atome , rouge} .shared 'nécessite. sm_12 cible ou plus

J'ai essayé aussi avoir une table temporaire pour chaque thread créé:

__global__ void histo_kernel_optimized5(unsigned char *buffer, long size, 
           unsigned int *histo) 
{ 
    unsigned int temp[buckets]; 
    int j; 
    for (j=0;j<buckets;j++){ 
     temp[j]=0; 
    } 

    int bucketID; 

    int i = threadIdx.x + blockIdx.x * blockDim.x; 
    int offset = blockDim.x * gridDim.x; 
    while (i < size) 
    { 
     bucketID = array[i]/Bwidth; 
     temp[bucketID]++; 
     i += offset; 
    } 


    for (j=0;j<buckets;j++){ 
     histo[j] += temp[j];  
    } 
} 

Mais compilateur ne me laisse pas le faire car il a besoin d'une constante pour créer la température table. Mais le problème dit que seaux sont dynamiquement donné pour la ligne de commande.

Existe-t-il une autre façon de le faire? Je ne sais pas comment le faire. Je suis confus.

Answer 1

Lorsque vous utilisez Atomics, le lancement moins de blocs va réduire les conflits (et donc d'améliorer la performance), car il ne faut coordonner entre moins de blocs. Lancez moins de blocs et faites en sorte que chaque bloc boucle sur plus d'éléments d'entrée.

for (unsigned tid = blockIdx.x*blockDim.x+threadIdx.x; 
       tid < size; tid += gridDim.x*blockDim.x) { 
    unsigned char value = array[tid]; // borrowing notation from another answer here 
    int bin = value % buckets; 
    atomicAdd(&histo[bin],1); 
} 

Answer 2

L'histogramme est vraiment facile à implémenter en utilisant des opérations atomiques. Je ne sais pas pourquoi vous écrivez un noyau aussi complexe. La motivation pour paralléliser l'opération est d'exploiter la nature parallèle de l'algorithme. Il n'est pas nécessaire d'itérer sur l'histogramme entier à l'intérieur du noyau. Voici un exemple de fonction noyau et wrapper CUDA pour calculer l'histogramme d'un tableau avec le nombre de cases spécifié. Je ne pense pas qu'il puisse être encore optimisé pour les appareils Compute 1.1. Mais pour Compute 1.2, la mémoire partagée peut être utilisée.

__global__ void kernel_getHist(unsigned char* array, long size, unsigned int* histo, int buckets) 
{ 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 

    if(tid>=size) return; 

    unsigned char value = array[tid]; 

    int bin = value % buckets; 

    atomicAdd(&histo[bin],1); 
} 

void getHist(unsigned char* array, long size, unsigned int* histo,int buckets) 
{ 
    unsigned char* dArray; 
    cudaMalloc(&dArray,size); 
    cudaMemcpy(dArray,array,size,cudaMemcpyHostToDevice); 

    unsigned int* dHist; 
    cudaMalloc(&dHist,buckets * sizeof(int)); 
    cudaMemset(dHist,0,buckets * sizeof(int)); 

    dim3 block(32); 
    dim3 grid((size + block.x - 1)/block.x); 

    kernel_getHist<<<grid,block>>>(dArray,size,dHist,buckets); 

    cudaMemcpy(histo,dHist,buckets * sizeof(int),cudaMemcpyDeviceToHost); 

    cudaFree(dArray); 
    cudaFree(dHist); 
} 

Answer 3

Il existe une solution pour les appareils sans opérations atomiques et montre une approche pour minimiser onchip collisions de mémoire, avec des subdivisions dans funes proproused par Podlozhnyuk à Histogram calculation in CUDA

Le code est à CUDASamples \ 3_Imaging \ histogramme (à partir Exemples CUDA)