TP 2 : Programmation CUBLAS
5ème année ingénieur de Polytech Paris-Sud
Stéphane
Vialle, CentraleSupélec & LISN
Stéphane.Vialle@centralesupelec.fr
Objectifs du TP :
Ce TP a pour objectif de pratiquer la programmation d'un GPU au sein d'un noeud de calcul CPU+GPU : il consiste à
implanter un produit de matrices denses sur un GPU en
utilisant la bibliothèque CUBLAS. On expérimentera différents types de
données et de précision des calculs, et on comparera les précisions
des résultats obtenus et les vitesses de calculs mesurées.
Plate-forme
de développement :
Les
machines utilisées seront celles des clusters Tx ou Cameron du DCE de
CentraleSupélec :
- Tx : chaque machine contient un CPU Intel XEON quad-core
hyperthreadés, et un GPU NVIDIA GTX-2080Ti (architecture Turing)
- Cameron
: chaque machine contient un CPU Intel XEON hexa-core
hyperthreadés, et un GPU NVIDIA GTX-1080 (architecture Pascal)
L'environnement
CUDA C et C++ est disponible sur chaque machine (et donc le
compilateur "nvcc"
et les drivers pour utiliser le GPU).
Vous utiliserez les
comptes de TP "23ppsgpu_i", où i est une valeur entière
entre 1 et 10.
Depuis votre poste de travail en mode graphique avec dcejs :
- allocation de ressource : demandez 1 machine en mode exclusif (configuration par défaut)
- action : démarrez VNC
OU BIEN depuis votre poste de travail en mode alphanumérique :
- vous vous connecterez
par ssh sur la machine chome.metz.supelec.fr en indiquant le login 23ppsgpu_i
- par un simple terminal (xterm depuis Linux, ou powershell depuis Windows), puis en entrant la commande suivante :
ssh -l 23ppsgpu_i chome.metz.supelec.fr
- une fois un terminal ouvert sur chome.metz.supelec.fr vous réserverez UN nœud sur le cluster à l'aide d'une commande slurm:
- Pendant le TP : srun --reservation=XXX -N 1 --exclusive --pty bash
--reservation=XXX : pré-réservation de machines, de nom XXX (demander ce nom à l'enseignant)
-N 1 : UN noeud
--exclusive : être seul sur le noeud et pouvoir utiliser tous les coeurs CPU
--pty bash : lancer un shell (bash) pour une session interactive
- Après le TP (si besoin) :
- pour obtenir un noeud Tx (GTX 2080 Ti): srun -p gpu_tp -C tx -N 1 --exclusive --pty bash
- pour obtenir un noeud Cameron (GTX 1080) : srun -p gpu_tp -C cam -N 1 --exclusive --pty bash
Travail à effectuer :
Remarques préliminaires :
- Le squelette de programme que vous utiliserez contient un code de produit de matrices denses en OpenMP et CUDA.
- La
partie OpenMP est complète, et est destinée à permettre de vérifier les résultats obtenus en CUDA.
- La partie CUDA est en partie développée, mais il vous reste à compléter le fichier gpu.cu :
- Le squelette est compilable et contient une aide intégrée : exécutez 'make' puis './MatrixProduct -h'.
- Pour valider votre premier code vous compilerez en Double Précision (le type "T_real" devient le type "double") avec "-DDP" dans le Makefile, les résultats seront identiques sur CPU et sur GPU, mais les performances des GPU
s'effondreront (car il s'agit de cartes GPU grand public non adaptées à la Double Précision).
- Pour
faire vos mesures de performances vous compilerez en Simple Précision
(le type "T_real" devient le type "float") avec "#-DDP" dans le
Makefile. La simple précision est adaptée
aux capacités des GeForce GTX1080 et RTX2080, mais il se peut que vous observiez des
différences entre les calculs sur CPU et sur GPU!!
1 - Implantation d'un appel CUBLAS "gemm" suivi du lancement d'un kernel de transposition
- Récupérez et compilez le squelette de programme OpenMP+CUDA.
- Implantez le pseudo-kernel K4 dans le fichier 'gpu.cu' pour calculer C = AxB :
- utilisez la routine CUBLAS_GEMM(...) (voir main.h),
il s'agit du renommage de cublasSgemm ou de cublasDgemm, selon que l'on
compile en simple ou double précision (-DDP dans le Makefile),
- utilisez la macro CHECK_CUBLAS_SUCCESS (voir gpu.h) pour lancer vos appels CUBLAS et tester leurs codes de retour,
- et lancez une grille de blocs du kernel TransposeKernel_v1,
associé à des blocs carrés de BLOCK_SIZE_XY_KT1 x BLOCK_SIZE_XY_KT1
threads, pous transposer le résultat de l'appel CUBLAS. Lancez une
grille de blocs adaptée à toutes les tailles de matrices (SIZE x SIZE).
- Testez
votre implantation sur une matrice de 4096x4096 DOUBLE (option -DDP active dans le 'Makefile' et changement de SIZE dans le 'main.h'). Obtenez-vous les mêmes valeurs qu'avec les BLAS sur
CPU : MatrixProduct -t CPU -cpu-k 1 -cpu-nt 4, sur Tx (et -cpu-nt 6 sur Cameron) ?
- Mesurez les performances du pseudo-kernel K4 obtenues sur une matrice de 4096x4096 FLOAT éléments (#-DDP dans le Makefile).
- Récupérez le fichier Excel de saisi des résultats, et complétez-le au fur et à mesure du TP.
- Comparez aux meilleures performances obtenues sur CPU multi-coeurs en OpenMP avec un kernel de même niveau (kernel 1 - BLAS sur CPU),
- MatrixProduct -t CPU -cpu-k 1 -cpu-nt 8 sur Tx
- MatrixProduct -t CPU -cpu-k 1 -cpu-nt 12 sur Cameron (sauf certaines machines Cam... vérifiez le nombre de coeurs de votre CPU Cam)
- Calculez le speedup GPU vs best-config-CPU.
- Mesurez les performances du pseudo-kernel K4 obtenues SANS verrouiller la mémoire sur le CPU
- Au début du fichier gpu.cu commentez les instructions de
verrouillage mémoire de la routine gpuInit(), et les instructions de
déverrouillage de la routine gpuFinalize().
- Refaites rune mesure de performances du pseudo-kernel K4, et notez la Bw mesurée. Comparez là à celle obtenue précédemment.
- Réactivez les opérations de verrouillage et déverrouillage.
2 - Implantation d'un appel CUBLAS "gemm" suivi d'un appel CUBLAS "geam"
- Créez
le pseudo-kernel K5 dans le fichier 'gpu.cu' :
- utilisez la routine CUBLAS_GEMM(...) : voir main.h
- et utilisez la routine CUBLAS_GEAM(...) : voir main.h
- Testez
et validez votre implantation sur une matrice de 4096x4096 FLOAT.
- Vérifiez que vous obtenez les mêmes valeurs qu'avec le pseudo-kernel K4.
- Mesurez les performances obtenues sur une matrice de 4096x4096 FLOAT.
- Comparez-les aux performances obtenues avec le pseudo-kernel K4
-
Comparez-les aux performances de la meilleure configuration sur le CPU
3 - Implantation d'un unique appel CUBLAS "gemm" permettant d'obtenir C = AxB
- Créez
le pseudo-kernel K6 dans le fichier 'gpu.cu' :
- utilisez la routine CUBLAS_GEMM(...) : voir main.h
- pour qu'elle calcule C = AxB à elle seule
- Testez
et validez votre implantation sur une matrice de 4096x4096 FLOAT.
- Vérifiez que vous obtenez les mêmes valeurs qu'avec le pseudo-kernel K5.
- Mesurez les performances obtenues sur une matrice de 4096x4096 FLOAT.
- Comparez-les aux performances obtenues avec le pseudo-kernel K5
-
Comparez-les aux performances de la meilleure configuration sur le CPU
4 - Utilisation des TensorCores avec des types de données et de calculs "standard"
- Implantez le pseudo-kernel K7 dans le fichier 'gpu.cu' pour exploiter les TensorCores :
- Utilisez la routine cublasGemmEx(...)
- Utilisez le type T_CUBLAS_real défini dans 'main.h'
- Utilisez le type T_CUBLAS_COMPUTE_real défini dans 'main.h'
- Indiquez l'algorithme CUBLAS_GEMM_DEFAULT_TENSOR_OP
- Testez
et validez votre implantation sur une matrice de 4096x4096 FLOAT.
- Vérifiez que vous obtenez les mêmes valeurs qu'avec le pseudo-kernel K6.
- Mesurez les performances obtenues sur les TensorCores avec une matrice de 4096x4096 FLOAT.
- Comparez-les performances obtenues maintenant avec le
pseudo-kernel K6 à celles obtenues précédemment
- Mesurez les performances obtenues sur les TensorCores avec une matrice de 4096x4096 DOUBLE
- Recompilez en double précision (-DDP dans le Makefile)
- Obtenez-vous les mêmes résultats que le pseudo-kernel K6 ?
- Comparez-les performances obtenues avec celles du
pseudo-kernel K6
5 - Utilisation des TensorCores avec des types de calculs en précision réduite
- Cette question ne sera valable que pour une compilation en Simple Précision
- Implantez le pseudo-kernel K8 dans le fichier 'gpu.cu' pour exploiter les TensorCores :
- Utilisez la routine cublasGemmEx(...)
- Utilisez le type de données : CUDA_R_32F
- Utilisez successivement les type de calcul en précision réduite :
- CUBLAS_COMPUTE_32F_FAST_TF32
- CUBLAS_COMPUTE_32F_FAST_16F
- CUBLAS_COMPUTE_32F_FAST_16BF
- Indiquez l'algorithme CUBLAS_GEMM_DEFAULT_TENSOR_OP
- Mesurez les performances obtenues sur les TensorCores avec une matrice de 4096x4096 FLOAT
- Pour les 3 types de calcul en précision réduite notez les
résultats obtenus (notamment les écarts avec la version CPU) et les
performances obtenues (exécutez chaque test plusieurs fois).
- Retenez le type de calcul le plus performant qui donne des résultats "encore acceptables"
- Comparez-les performances obtenues maintenant avec le
pseudo-kernel K6 à celles obtenues précédemment
6 - Evaluation du verouillage de la mémoire sur la vitesse des transferts de données CPU-GPU
- Commentez les transferts à partir ou vers des données CPU "vérouillées en mémoire" (pinned) dans gpuInit et gpuFinalize.
- Commentez le
verrouillage en mémoire des matrices A, B et C sur le CPU (ne changez
par leur allocation/déclaration) dans gpuInit.
- Commentez leur déverouillage dans gpuFinalize (qui a lieu après la fin des transferts et des calculs)
- Vérifiez que vos transferts mémoires ne déclenchent pas de messages d'erreur
- Vérifiez que vous obtenez toujours les mêmes résultats
(comparez-les à ceux d'un kernel CPU ou aux valeurs obtenues
précédemment sur GPU).
- Mesurez les performances obtenues sur une matrice de 4096x4096 FLOAT.
- Re-mesurez les performances obtenues avec les kernels 4, 5, 6, 7 et 8
- Remplissez le deuxième tableau de l'onglet "K4-8 CUBLAS" du fichier Excel
- Observez-vous des changements significatifs sur la Bw mesurée ? et sur la vitesse de calcul globale du produit de matrices ?
- Décommentez le verouillage et déverouillage de la mémoire
7 - Refaites vos tests sur la machine John8 ou John9 équipées d'un RTX3090
- Recompilez sur la machine John en Simple Précision
- Depuis chome faites un simple ssh john11
- Relancez chacun de vos kernels (notamment les kernels exploitant les Tensor Cores)
- Rassemblez les performances dans le troisième tableau de l'onglet "K4-8 CUBLAS" du fichier Excel
- Est-ce que les Tensor Cores apportent de la performance sur un RTX3090 (architecture "Ampere") ?
- Refaites une mesure de performances sans verrouillage de la mémoire CPU. Est-ce que la Bw est plus fortement impactée ?