diff --git a/01-introduction.pdf b/cours/01-introduction.pdf
similarity index 100%
rename from 01-introduction.pdf
rename to cours/01-introduction.pdf
diff --git a/02-clusters-slurm.pdf b/cours/02-clusters-slurm.pdf
similarity index 100%
rename from 02-clusters-slurm.pdf
rename to cours/02-clusters-slurm.pdf
diff --git a/cours/03-gpu-cuda.pdf b/cours/03-gpu-cuda.pdf
new file mode 100644
index 0000000000000000000000000000000000000000..989ec9fa82a13d99d72e7b44dba5441a55c400a0
Binary files /dev/null and b/cours/03-gpu-cuda.pdf differ
diff --git a/exercices/01-somme-vecteurs-gpu/fig/.gitkeep b/exercices/01-somme-vecteurs-gpu/fig/.gitkeep
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/exercices/01-somme-vecteurs-gpu/fig/1-D-array-in-CUDA-kernels.png b/exercices/01-somme-vecteurs-gpu/fig/1-D-array-in-CUDA-kernels.png
new file mode 100644
index 0000000000000000000000000000000000000000..79e4bb1a0b2ee8e9d2124007c57d8ac0861fa2c6
Binary files /dev/null and b/exercices/01-somme-vecteurs-gpu/fig/1-D-array-in-CUDA-kernels.png differ
diff --git a/exercices/01-somme-vecteurs-gpu/fig/CUDA-GridBlockThread-Structure.png b/exercices/01-somme-vecteurs-gpu/fig/CUDA-GridBlockThread-Structure.png
new file mode 100644
index 0000000000000000000000000000000000000000..ff10e1130c5e113e382033bd6cfa3d43f57c42f6
Binary files /dev/null and b/exercices/01-somme-vecteurs-gpu/fig/CUDA-GridBlockThread-Structure.png differ
diff --git a/exercices/01-somme-vecteurs-gpu/fig/CUDABThreadNumCalc.png b/exercices/01-somme-vecteurs-gpu/fig/CUDABThreadNumCalc.png
new file mode 100644
index 0000000000000000000000000000000000000000..94997f5cc6aec58f31e3f39253b289f149308a44
Binary files /dev/null and b/exercices/01-somme-vecteurs-gpu/fig/CUDABThreadNumCalc.png differ
diff --git a/exercices/01-somme-vecteurs-gpu/fig/Software-Perspective_for_thread_block.jpg b/exercices/01-somme-vecteurs-gpu/fig/Software-Perspective_for_thread_block.jpg
new file mode 100644
index 0000000000000000000000000000000000000000..f4aabed8c637b6adb11cdeed87916d5a0f4d8b79
Binary files /dev/null and b/exercices/01-somme-vecteurs-gpu/fig/Software-Perspective_for_thread_block.jpg differ
diff --git a/exercices/01-somme-vecteurs-gpu/fig/Warp-Scheduler-Gpu.jpg b/exercices/01-somme-vecteurs-gpu/fig/Warp-Scheduler-Gpu.jpg
new file mode 100644
index 0000000000000000000000000000000000000000..c25d4eff9495fb315a57a61fb5a59b6b44445b1d
Binary files /dev/null and b/exercices/01-somme-vecteurs-gpu/fig/Warp-Scheduler-Gpu.jpg differ
diff --git a/exercices/01-somme-vecteurs-gpu/fig/clip_image004.jpg b/exercices/01-somme-vecteurs-gpu/fig/clip_image004.jpg
new file mode 100644
index 0000000000000000000000000000000000000000..3e576e743df8d51fd5f1b63c86fc03febf1fb58f
Binary files /dev/null and b/exercices/01-somme-vecteurs-gpu/fig/clip_image004.jpg differ
diff --git a/exercices/01-somme-vecteurs-gpu/readme.md b/exercices/01-somme-vecteurs-gpu/readme.md
new file mode 100644
index 0000000000000000000000000000000000000000..48bcd3fbc95e0dd3605e459579b0fadfd78cfcb5
--- /dev/null
+++ b/exercices/01-somme-vecteurs-gpu/readme.md
@@ -0,0 +1,162 @@
+# Série 1: addition de vecteurs sur GPU
+
+[[_TOC_]]
+
+Vous avez vu en cours que sur GPU Il y a beaucoup plus de coeurs disponible que sur un CPU.
+Ceci même si les CPU sont groupés sur du multi-socket ou en noeuds.
+Chacun de ces coeurs peut donc exécuter une suite d'instructions spécifée par votre code.
+
+Pour mettre en oeuvre du code sur GPU vous utiliserez le language CUDA C qui est un C avec "extensions" pour pouvoir écrire, lancer un kernel et effectuer mouvements mémoire entre le CPU, que l'on appelle l'hôte (host), et le GPU, que l'on appelle le dispositif (device).
+Si vous connaissez le C, il n'y a que peu d'ajout à faire pour réaliser une addition de vecteurs par rapport à un code C "classique".
+
+Le but s'agit de réaliser une addition de deux vecteurs sur GPU et l'exécuter sur les GPU d'un cluster académique.
+Cet exercice est considéré comme le "hello world" du CUDA.
+
+## Réservation d'un GPU sur un cluster
+
+Réserver un noeud GPU n'est pas plus compliqué que réservé un noeud CPU sur un cluster.
+Il faut néanmoins penser à trois choses:
+1. utiliser une partition (file d'attente) GPU
+2. penser à demander le nombre de GPU requis pour son calcul avec `--gpus=n`
+3. charger un module permettant d'utiliser CUDA (par exemple `fosscuda/2020b`)
+
+Le point deux est important, sinon vous ne "verrez" pas le GPU du noeud de calcul.
+Le `--gpus=n` s'utilise comme les autres options SLURM, au début du script de soumission:
+```bash
+#SBATCH partition=debug-gpu
+#SBATCH gpus=1
+...
+```
+
+## Parallélisme et addition de vecteurs
+
+On rappelle que l'addition de deux vecteurs $`\textbf{v}`$ et $`\textbf{u}`$ appartenant à $`R^{n}`$ est définie comme:
+
+$`\textbf{v} + \textbf{u} = \begin{pmatrix} v_1 \\ v_2 \\ \vdots \\ v_n \end{pmatrix} + \begin{pmatrix} u_1 \\ u_2 \\ \vdots \\ u_n \end{pmatrix} = \begin{pmatrix} v_1 + u_1 \\ v_2 + u_2 \\ \vdots \\ v_n + u_n \end{pmatrix}`$
+
+ce qu'on pourrait implémenter simplement de la manière suivante (en pseudo C) :
+
+```c
+for (int i = 0; i < n; i++) {
+ c[i] = v[i] + u[i];
+}
+
+```
+où le vecteur (tableau) `c` stockerait la valeur du résultat de l'addition de `u` et de `v`.
+
+Cette opération peut déjà être accélérée par le compilateur en vectorisant l'opération avec `VADDSD` provenant de l'ensemble d'instruction (instruction set) [AVX](https://en.wikipedia.org/wiki/Advanced_Vector_Extensions) (advanced vector extensions).
+
+<!--
+Vous pouvez aussi utliser MPI (~MIMD) pour additionner de très grand vecteurs, avec le pseudo-code suivant:
+```c
+MPI_INIT // donne n_proc (rangs) à disposition
+
+double* V, U, C // de longueur N
+double sub_a, sub_b, sub_c
+
+scatter(A, N/n_proc, MPI_DOUBLE, sub_a, N/n_proc, MPI_DOUBLE, 0, COMM)
+scatter(B, N/n_proc, MPI_DOUBLE, sub_b, N/n_proc, MPI_DOUBLE, 0, COMM)
+
+for (int k = 0; k < N/n_proc; k++) sub_a[k] = sub_v[k] + sub_u[k]
+
+gather(C, N/n_proc, MPI_DOUBLE, sub_c, N/n_proc, MPI_DOUBLE, 0, COMM)
+
+MPI_FINALIZE
+```
+-->
+
+Ou vous pouvez utiliser le GPU pour résoudre le problème avec une apporche quasi-SIMD.
+
+### L'approche (quasi) SIMD sur GPU (block and threads)
+
+Contrairement à du parallélisme basé sur des processus, CUDA ajoute une organisation en bloc des threads, et ces blocs sont eux-même organisés en grille.
+Comme ce qui vous a été présenté lors du cours:
+
+
+
+**Source de la figure:** wikipedia.org
+
+Cette organisation par grille de threads fait que CUDA se prête assez bien au problème sur des domaines carrés, ce qui est le cas de nombreuses applications scientifiques.
+
+Imaginons que nous souhaitons additionner des vecteurs contenant huit éléments.
+On peut donc définir quatre blocs de deux threads selon la figure suivante:
+
+<img src="fig/CUDABThreadNumCalc.png" alt="add-vec-cuda" width="500"/>
+
+**Source de la figure:** Libby Shoop du Macalester College
+
+Dans ce cas il est très intuitif de penser en grille (à une dimension).
+CUDA offre des accesseurs à la dimension `x` de la grille de blocs et de threads, il s'agit de:
+ - `threadIdx.x`: l'indice du thread sur la dimension `x` au sein de son bloc,
+ - `blockIdx.x`: l'indice du bloc sur la dimension `x` au sein de la grille,
+ - `blockDim.x`: le nombre de thread par bloc dans la direction `x`.
+
+Dans notre exemple de vecteur, avec un total de huit threads, indéxé de 0 à 7, chaque thread accède "naturellement" à la donnée suivante au sein du kernel:
+```
+int t_id = blockDim.x*blockIdx.x + threadIdx.x;
+```
+On peut illustrer le type de calcul ci-dessus avec:
+
+
+
+**Source de la figure:** Documentation CUDA de NVIDIA.
+
+Ces explications et illustrations donnent quasiment la solution de l'exercice.
+Mais pour rendre l'exercice plus stimulant, nous vous laissons trouver la solution complète par vous-même.
+
+Notez que pour des problème plus complexes, notez qu'il est possible d'utiliser des indices en deux ou trois dimensions en accédant au valeur `y` ou `z`.
+Par exemple: `threadIdx.x` ou `blockIdx.z`.
+
+
+
+**Source de la figure:** 3dgep.com.
+
+### Du SIMD, oui mais par groupe
+
+On a parlé de quasi SIMD dans cette série (et dans le cours) parce que tous les threads du GPU n'exécutent pas la même instruction.
+Concrètement sur chaque SM (straming multiprocessor) du GPU un thread block est composé de warps.
+Un warp est un ensemble de 32 threads au sein d'un thread block qui exécutent la **même instruction**.
+D'où le pseudo-SIMD, car c'est le seul moment où l'on a un garantie qu'un groupe de threads exécutent la même instruction.
+
+
+
+**Source de la figure:** wikipedia.org
+
+On note aussi que les SM contiennent 8 coeurs.
+Il faut donc 4 cycles pour que les 32 threads d'un warp exécturent la même instruction.
+
+### Le kernel CUDA
+
+Un kernel CUDA est simplement une fonction qui peut être appelée depuis le host et exécutée sur le device.
+Ces fonctions ont une particularité syntaxique parce qu'elles sont déclarée avec le mot clé `__global__`:
+```c
+__global__ void ker(type arg_1, type arg_2, ...)
+```
+
+où `type` est un type quelconque C (comme `int`, `double*`, `void`, etc ...).
+
+et sont appelées avec le nombre de blocs et le nombre de threads par bloc:
+```
+ker<<<N_BLOCS, N_THRDS>>>(arg_1, arg_2, ...);
+```
+c'est à dire, votre grille.
+Voir schéma précédent.
+
+Notez que pour modulariser votre code, vous pouvez créer des fonctions `__device__`:
+```c
+__device__ void f(type arg_1, type arg_2, ...);
+```
+
+Mais ces fonctions ne peuvent être appelées que depuis un kernel/fonction qui s'exécute sur le device.
+
+La structure du code sera la suivante:
+- écrire le kernel (le code) qui va s'exécuter sur le device (le GPU) dans une ou plusieures fonctions
+- dans le programme principal du host (le CPU), vous devez:
+ - allouer la mémoire des tableaux (càd les vecteur) sur le host,
+ - allouer la mémoire des tableaux sur le device,
+ - copier les données du host sur le device
+- éxecuter le kernel (la/les fonctions) pour calculer la somme, càd la nouvelle donnée à partir des données copiées sur le device. Ceci en spécifiant combien de blocs et threads par block vous voulez utiliser.
+- une fois le kernel terminé, copier les résultats de la mémoire du device vers le host.
+- vérifier la validité des résultats !
+- libérer la mémoire sur le device **ET** le host.
+