CI1 : De la couleur au noir et blanc
Maitriser l'environnement de développement et exécuter un premier kernel.
Utilisation de SLURM (∼20mn, – facile)
Nous allons accéder à des GPUs en utilisant un logiciel appelé SLURM. Slurm permet des ressources partagées sur plusieurs serveurs comme des GPUs ou des CPUs.
Il va falloir créer un compte sur le SLURM de l'école. Générer une clef SSH (si vous n'en avez pas déjà une), et envoyer à votre enseignant la clef PUBLIQUE par mail, ainsi que votre pseudo de connexion aux machines de l'école.
Pour accéder aux machines contenant les GPUs, nous allons utiliser SSH. Comme la machine se trouve dans le réseau de l'école, on ne peut y accéder que depuis l'école, ou à travers le VPN, ou en passant par le portail SSH publique. Demander à votre enseignant l'adresse IP. Si vous utiliser Linux, vous pouvez configurer le fichier ~/.ssh/config que la manière suivante pour un accès facilité :
Vous pourrez alors vous connecter depuis n'importe où en faisant ssh tsp-client. Si vous êtes sous Windows, une solution consiste à utiliser WSL2 ou Putty.
Host LAB_GATEWAY
User VOTRE_IDENTIFIANT_TSP
Hostname ssh1.imtbs-tsp.eu
IdentitiesOnly yes
IdentityFile CHEMIN_VERS_CLEF_SSH_PRIVEE
ServerAliveInterval 60
ServerAliveCountMax 2
Host tsp-client
User VOTRE_IDENTIFIANT_TSP
Hostname ADRESSE_IP
IdentitiesOnly yes
IdentityFile CHEMIN_VERS_CLEF_SSH_PRIVEE
ServerAliveInterval 60
ServerAliveCountMax 2
ProxyJump LAB_GATEWAY
La machine sur laquelle vous vous connectez est peu puissante : n'y exécutez pas de programmes couteux en ressources !
Dans Slurm, vous avez la possibilité de demander des ressources. Si celles-ci sont disponibles, elles vous seront attribuées et vous serez connectés aux ressources en question dans un nouveau terminal. Commencez par exécuter la commande nvidia-smi. Cette commande donne un aperçu des GPUs et de leur utilisation. Normalement, sur la machine client configurée plus haut cette commande ne fonctionne pas (il n'y a pas de GPU). Maintenant, demandez des ressources en mode interactif à l'aide de la commande srun --time=10:00:00 --gres=gpu:1 --pty bash. Les commandes bash commencent par un s. Ici, nous demandons un GPU (--gres=gpu:1) pour une durée de 10 heures with bash as a pseudo-terminal (--pty bash). Si des ressources sont disponibles, vous devriez arriver dans un nouveau terminal avec un GPU. Exécutez nvidia-smi pour vérifier que c'est bien le cas. Pour quitter le terminal, vous pouvez faire Ctrl + D, ou taper exit.
Une commande utile de Slurm est squeue (squeue -u $USER permet de visualiser uniquement vos jobs). Elle permet de regarder l'utilisation des ressources et d'avoir des éléments clefs à propos des jobs (allocation de ressource dans Slurm). En particulier, on peut voir le JobID qui nous permet d'annuler un job (i.e., de le terminer) à l'aide de la commande scancel MON_JOB_ID.
De la couleur à l'échelle des gris (∼60mn, – facile)
Dans cette exercice, nous allons écrire notre premier kernel, puis compiler et exécuter notre code sur GPU. Le but de notre programme sera de transformer une image en couleur en son équivalent en noir et blanc.
Les exercices se trouvent sur un dépôt Git. Vous pouvez y accéder ici. Cloner ce dépôt sur le client slurm (c'est-à-dire la première machine sur laquelle vous vous connecter avant de faire srun). Votre dossier HOME se trouve dans le répertoire /array/shared/home/$USER. Ce dossier est partagé avec la machine contenant le GPU : tout se que vous y mettez sera accessible après avoir fait un srun.
Pour éditer votre code, plusieurs solutions s'offrent à vous.
- Éditer le code sur machine via SSH directement dans le terminal à l'aide d'un programme tel que nano ou vim.
- Éditer le code sur votre machine puis le transférer sur la machine distante. Cela peut passer par SCP, rsync, ou un autre dépôt Git.
- Développer depuis votre IDE (par exemple, CLion ou VSCode) en mode Remote SSH. Il faudra configurer la connexion.
Reprenez le code vu dans le cours calculant le carré d'un nombre lesson_code_snippets/lesson1_code_snippets/square.cu. Nous allons le compiler et l'exécuter. Pour cela, demander un GPU. Ensuite, déplacez vous dans le dossier de square.cu. Pour compiler le code, nous n'utilisons pas gcc mais nvcc avec la commande suivante : nvcc square.cu -o square. Cette commande génère un exécutable square que vous pouvez exécuter (uniquement sur la machine avec le GPU) avec ./square. Vérifiez que vous avez bien le bon résultat.
Modifiez le code pour maintenant calculer le cube de chaque nombre de la même manière que vue dans le cours.
Passons maintenant à l'implémentation d'un filtre transformant une image couleur en sa version en noir et blanc. Implémentez le kernel dans problem_sets/problem_set1/student_func.cu appelé rgba_to_greyscale. Les instructions sont dans le code. Notez que la représentation des images données en paramètre est faite par un tableau en une seule dimension. Dans ce cas, chaque ligne de l'image se suit. Cela permet d'assurer que l'image est représentée sur une région continue de la mémoire.
Commencez par déterminer la position x et y dans l'image en fonction de la position du block (blockIdx et blockDim) et de la position dans ce block (threadIdx). Ensuite, transformer cette représentation 2D en une représentation 1D (en considérant que toutes les lignes se suivent).
__global__ void rgba_to_greyscale(const uchar4* const rgbaImage,
unsigned char* const greyImage,
int numRows, int numCols)
{
int pos_x = blockDim.x * blockIdx.x + threadIdx.x;
int pos_y = blockDim.y * blockIdx.y + threadIdx.y;
if (pos_x < numRows && pos_y < numCols){
int position = numRows * pos_y + pos_x;
uchar4 color = rgbaImage[position];
unsigned char grey = (unsigned char)(0.299f * color.x + 0.587f * color.y + 0.114f * color.z);
greyImage[position] = grey;
}
}
Le kernel étant implémenté, il faut configurer les tailles des blocks et la taille de la grille. Modifiez la fonction your_rgba_to_greyscale pour donner une taille aux blocs et à la grille. À noter que même si nous représentons l'image en 2D, cela ne nous empêche pas d'utiliser un bloc en 2D.
Commencez par un block de taille carré plutôt petit. La taille de la grille correspond au nombre de blocks pouvant entrer horizontalement et verticalement.
void your_rgba_to_greyscale(const uchar4 * const h_rgbaImage, uchar4 * const d_rgbaImage,
unsigned char* const d_greyImage, size_t numRows, size_t numCols)
{
//You must fill in the correct sizes for the blockSize and gridSize
//currently only one block with one thread is being launched
int block_size = 32;
const dim3 blockSize(block_size, block_size, 1); //TODO
int n_blocks_x = numRows / block_size + 1;
int n_blocks_y = numCols / block_size + 1;
const dim3 gridSize(n_blocks_x, n_blocks_y, 1); //TODO
rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols);
cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
}
Pour compiler le programme, il nous faut d'abord générer un Makefile. Pour cela, tapez la commande cmake . dans le dossier dans lequel se trouve le programme. Si tout se passe bien, vous devrier voir apparaitre un fichier Makefile. Ensuite, nous allons simplement exécuter la commande make. Make et le Makefile sont configurés pour appelés les bons compilateurs et combiner les résultats. Ici, si tout se passe bien, vous devriez générer un fichier nommé HW1 permettant d'exécuter le programme. Par exemple, vous pouvez taper ./HW1 cinque_terre_small.jpg. Cela génèrera l'image en noir et blanc et une image de comparaison avec l'image de référence. S'il y a très peu de points blancs, c'est que vous avez bien implémenté la fonction (des erreurs de précisions créent ces points blancs).
Jouez avec la taille des blocs. Observez des la vitesse d'exécution augmente avec la taille des blocs (n'oubliez pas de recompiler avec make !). Y-a-t'il des limites à la taille des blocs ? Pourquoi ?
Le nombre de threads étant limité à 1024, on ne peut pas augmenter la taille des blocs indéfiniement (32x32 au maximum pour un bloc carré).
