ajout TP correction gamma

exercices/02-correction-gamma/readme.md

+# Kernel grid en deux dimensions
+
+[[_TOC_]]
+
+Le but de cet exercice est que vous arriver à travailler avec une grille 2D de block et threads sur GPU. 
+En effet, la limite de blocks que vous pouvez avoir sur la direction x est de 2'147'483'647.
+Si vous souhaitez aller au-dela de cette limite, ceci même pour traiter de données en 1D, vous deverez utiliser les autres dimension à disposition.
+
+Ce qu'il faut savoir c'est que:
+- le nombre maximum de blocks en dim x est 2'147'483'647
+- la taille maximum d'un block en dim x est 1'024
+- le nombre maximum de blocks en dim y est 65'535
+- la taille maximum d'un block en dim y est 1'024
+- le nombre maximum de blocks en dim z est 65'535
+- la taille maximum d'un block en dim z est 64
+
+Le nombre maximum de threads par block est de 1'024.
+
+Ceci en date de l'exercice et peut varier en fonction de compute capability de CUDA.
+La compute capability de CUDA définit l'ensemble d'attribut que vous avez avec votre device.
+Que ce soit hardware ou software.
+Ceci dépend de la génération du GPU utilisé.
+
+Pour cet exercice, il ne sera pas nécessaire de pousser ces limites en terme de dimension sur le GPU.
+Mais il faut être capable d'exploiter les dimensions offerte par CUDA pour:
+- aller au dela de la limite 2'147'483'647 si nécessaire
+- être capable de travailler aisément sur une structure régulière en 2 ou 3D
+
+Rappelez vous également que les GPU on une mémoire limitée.
+Jetez un oeil aux [spécifications des machines à l'UniGE](https://doc.eresearch.unige.ch/hpc/hpc_clusters#for_advanced_users) ou de votre [propore GPU](https://developer.nvidia.com/cuda-gpus).
+
+## Exécuter un kernel sur une grille 2D
+
+Pour exécuter un kernel 2D, cela se fait fondamentalement comme une grille 1D: avec les paramètre d'exécution du kernel (`my_kernel<<<GRID_SIZE,BLOCK_SIZE>>>`).
+Vous avez probablement passé de simple entier lors de la série précédente.
+Mais pour pouvoir utiliser des dimensions supérieur il faut instancier un `dim3`.
+Un `dim3` vous permet de spécifier les tailles en `x`, `y` et `z` (dans cet ordre), comme par exemple:
+```c
+const dim3 BLOCK_SIZE = dim3(3, 2);
+```
+où je définis un grille de taille 3 x 2 (= 6 blocks).
+Notez que le `dim3` n'est pas une erreur et est équivalent à:
+```c
+const dim3 BLOCK_SIZE = dim3(3, 2, 1);
+```
+Vous pouvez faire de même avec la taille des blocks:
+```c
+const dim3 GRID_SIZE = dim3(4, 3);
+```
+et lancer le kernel de la sorte:
+```c
+my_kernel<<<GRID_SIZE, BLOCK_SIZE>>>(...);
+```
+
+Vous pouvez ensuite utiliser ces dimensions dans votre kernel, comme en 1D:
+```c
+int rows = blockIdx.y*blockDim.y + threadIdx.y;
+int cols = blockIdx.x*blockDim.x + threadIdx.x; 
+```
+
+Si la notion grille 2D n'est pas claire, voici un schéma pour vous aider :
+
+![2D grid](fig/CUDA-GridBlockThread-Structure.png)
+
+où, probablement, on a:
+```c
+const dim3 BLOCK_SIZE = dim3(3, 3);
+const dim3 GRID_SIZE = dim3(3, 3);
+```
+
+## Exercice 1: addition de vecteurs
+
+Adaptez l'exercice de la série de l'addition de vecteurs de manière à ce que vos vecteurs soient projetés sur une kernel grid en 2D en place d'une grille en 1D.
+Imaginez que les vecteurs sont projeté rangée par rangée ou colone par colone.
+
+C'est-à-dire que vous passerez d'une répartition sur une grille 1D comme:
+
+![vect 1D grid](fig/g31567.png)
+
+à une répartition sur une grille 2D comme:
+
+![vect 2D grid](fig/g31611.png)
+
+## Exercice 2: correction de luminosité d'une image
+
+Nous allons à présent utliser la kernel grid 2D pour travailler "naturellement" avec des données en 2D.
+Nous appliquerons ceci avec le réhaussement d'une image.
+Notez que pour des raisons de simplicité, l'image ne sera qu'une simple matrice d'entiers. 
+Nous ne lirons pas d'image au format PNG, JPG, ou autre format existant.
+
+(Vous avez vu, lors du cours de VISNUM ???), que l'on obtient une bonne correction de la luminosité
+en appliquant une transformation non-linéaire, appelée [correction gamma](https://en.wikipedia.org/wiki/Gamma_correction).
+
+Cette fonction dépend d'un paramètre nommé gamma, et voici à quoi ressemble cette transformation pour quelques valeurs de gamma choisie sur des niveau de gris:
+![gamma](fig/Gamma-transformation-of-grey-level.png)
+
+Attention sur cette figure les `L` sont les différent niveau de gris or les niveaux de gris sont souvent normalisé entre 0 (noir) et 1 (blanc) dans les implémentation.
+
+Voici un pseudo-code illustrant l'application de la transformation gamma sur une image:
+```python
+# habituellement un pixel est codé sur un unsigned int 8,
+# un entier de 8 bits (noir: 0 -> blanc: 255)
+def func_gamma(img, gamma):
+  X,Y = dim(img)
+  G = zeros(X,Y)
+  for i in X:
+    for j in Y:
+      pixel_norm = img[i,j].as_double()/255.0
+      new_pixel = 255 * pow(pixel_norm, gamma)
+      G[i,j] <- new_pixel.as_uint8()
+  return G
+```
+
+### Travail à réaliser
+
+Implémenter la transformation gamma sur GPU.
+Il n'y a pas besoin de viser les performances pour cet exercice.
+Assurez-vous surtout que vous ayez bien compris le modèle de programmation CUDA (mouvement mémoire, exécution de kernels et synchronization).
+
+### Lecture de l'image
+
+Pour l'image, pas besoin de lire une vrai image.
+Générez simplement une matrice d'entiers en 2D contenant uniquement de pixels en niveau de gris.
+Ce sont génàralement des entiers allant de 0 à 255.
+Cette image doit être générée sur le host et copiée sur le device.
+
+Attention, à la connaissance de l'auteur, CUDA ne support pas les entiers non signé sur 8 bits.
+Mais comme ce n'est pas très important pour cet exercice, vous pouvez utiliser de simple entiers non 
+signés (`unsigned int`). Pensez à vos conversions de type pour le calcul dans votre fonction de transformation.
+
+Voici un exemple d'image qui génère un dégradé de gris. 
+Mais vous êtes libre de tester avec la matrice qui vous plait.
+Encore, pour cet exercice, 
+
+```c
+// non tested (bad teacher, bad...), but should be fine 
+int* gray_Scale(int n, int m) {
+  int n_pixels = n/m;
+  char *img = malloc(n_pixels);
+  if (!img) return NULL;
+
+  int extra_pix = n_pixels % 255;
+  int per_gray = n_pixels / 255;
+  int g = 0;
+  while (i < count) {
+    for (int k = 0; k < per_gray; ++k) {
+      img[i] = g;
+      i++;
+    }
+
+    if (extra_pix > 0) {
+      img[i] = g;
+      i++;
+      extra_pix--;
+    }
+
+    g++;
+  }
+  return img;
+}
+```
+
+Comparez également le résultat obtenu par votre code GPU avec une implémentation sur CPU.
+
+Si vous souhaitez **vraiment** lire et écrire une image sur CPU en **C++**, sachez que [openimageio (OIO)](https://sites.google.com/site/openimageio/home) est disponible sous forme de module sur les clusters:
+
+```
+OpenImageIO: OpenImageIO/1.7.17
+
+Description:
+  OpenImageIO is a library for reading and writing images, and a bunch of related 
+  classes, utilities, and applications.
+
+You will need to load all module(s) on any one of the lines below before the 
+"OpenImageIO/1.7.17" module is available to load.
+
+  icc/2017.1.132-GCC-6.3.0-2.27  impi/2017.1.132
+  ifort/2017.1.132-GCC-6.3.0-2.27  impi/2017.1.132
+```