CSC 8610 – Introduction au traitement d'image

CI6 : Histogrammes rapides

Écrire un algorithme rapide d'histogramme.

Algorithme rapide histogramme (∼90mn, – moyen)

Dans cet exercice, nous allons implémenter un algorithme rapide d'histogramme en utilisant la mémoire partagée et des accès intelligents à la mémoire.

Nous allons calculer un histogramme sur des données dont la valeur correspond au numéro de la bin. Par exemple, la valeur 452 ira dans la bin 452. La distribution des données suit une loi normale. Ce qui signifie que pour des valeurs comprises entre 0 et numBins (le nombre de bins), il y aura plus de valeurs vers le centre que sur les bords. Cette propriété va nous être utile par la suite car si nous arrivons à réduire le nombre d'écritures atomiques sur les valeurs centrales les plus fréquentes, nous serons en mesure d'améliorer la vitesse de notre algorithme.

Commencez par faire un git pull pour être sur d'être sur la bonne version du code. Rendez-vous ensuite dans problem_sets/problem_set5 où vous trouverez le fichier student.cu. Exécutez cmake . pour générer le Makefile puis make pour compiler le code.

Écrivez une première version de l'histogramme dans la fonction computeNaiveHistogram. Pour cette version, écrivez un kernel utilisant une opération atomique pour mettre à jour la mémoire globale. Exécutez le code et notez que le temps de calcul est relativement long.

__global__
void naiveHisto(const unsigned int* const vals,
               unsigned int* const histo,
               int numVals)
{
  int global_id = threadIdx.x + blockDim.x*blockIdx.x;
  if (global_id >= numVals) return;
  atomicAdd(&(histo[vals[global_id]]), 1);
}

void computeNaiveHistogram(const unsigned int* const d_vals,
                      unsigned int* const d_histo,
                      const unsigned int numBins,
                      const unsigned int numElems)
{
  int blocks = ceil(numElems / N_THREADS);

  naiveHisto <<< blocks, N_THREADS >>> (d_vals, d_histo, numElems);

  cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
}

Nous proposons maintenant d'améliorer les performances de l'histogramme en utilisant la mémoire partagée. Pour ce faire, vous allez devoir implémenter la fonction computeHistogram. Nous proposons ici d'utiliser une mémoire partagée par block qui contiendra un histogramme local au block. Une fois cet histogramme local calculé, nous mettrons à jour la mémoire globale. Le but est d'utiliser le moins d'opérations atomiques possible, en particulier pour la mémoire globale et sur les données les plus fréquentes. Rappelez-vous que les données suivent une loi normale et que donc certains valeurs sont plus fréquentes que d'autres.

__global__
void perBlockHisto(const unsigned int* const vals, //INPUT
               unsigned int* const histo,      //OUPUT
               int numVals,
               int numBins)
{
  extern __shared__ unsigned int sharedHisto[]; //size as original histo

  // Initialisation
  // As numBin can be more than blockDim, each thread initialize several indexes
  for (int i = threadIdx.x; i < numBins; i += blockDim.x) {
    sharedHisto[i] = 0;
  }

  __syncthreads();

  int global_id = threadIdx.x + blockIdx.x * blockDim.x;
  atomicAdd(&sharedHisto[vals[global_id]], 1);

  __syncthreads();

  for (int i = threadIdx.x; i < numBins; i += blockDim.x) {
    atomicAdd(&histo[i], sharedHisto[i]);
  }
}

void computeHistogram(const unsigned int* const d_vals, //INPUT
                      unsigned int* const d_histo,      //OUTPUT
                      const unsigned int numBins,
                      const unsigned int numElems)
{
  int blocks = ceil(numElems / N_THREADS);
  perBlockHisto <<<blocks, N_THREADS, sizeof(unsigned int)*numBins>>>(d_vals, d_histo, numElems, numBins);
  cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
}

        