Objectifs du TP :
Ce
TP a pour objectif de pratiquer la programmation d'un GPU à travers une
bibliothèque de type BLAS, au sein d'un noeud de calcul CPU+GPU : il
consiste à
implanter des appels CUBLAS de plus en plus efficaces.
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 "gpucs1_#i", où #i est une valeur entière
entre 1 et 16.
Depuis votre poste de travail :
- vous vous connecterez
par ssh sur la machine ghome.metz.supelec.fr en indiquant le login gpucs1_#i
- puis vous vous connecterez par ssh à la machine term2.grid.metz.supelec.fr,
- où vous réserverez UN nœud sur le cluster à l'aide d'une commande oarsub.
Ex: oarsub
-p "cluster='Tx'" -l nodes=1,walltime=4:00:00
-I
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 comparer les résultats numériques et les performances duCPU et du GPU
- 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.
- Squelette, ou bien allez le recopier sur votre compte de TP par la commande : cp ~vialle/tmp/MatrixProduct-CUBLAS-enonce.zip .
- Compilez ce squelette et testez son fonctionnement sur CPU (exécutez la commande ./MatrixProduct -h
pour voir les détails de fonctionnement de l'application).
- 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 4 sur Tx
- MatrixProduct -t CPU -cpu-k 1 -cpu-nt 6 sur Cameron
- Calculez le speedup GPU vs best-config-CPU.
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 - Améliorez la vitesse des transferts de données CPU-GPU
- Implantez des transferts à partir ou vers des données CPU "vérouillées en mémoire" (pinned).
- Verrouillez en mémoire les matrices A, B et C sur le CPU (ne changez par leur allocation/déclaration).
- Puis déverouillez-les après avoir fini les transferts et les calculs.
- Testez les codes de retour de vos opérations CUDA, pour vérifier qu'elles retournent bien cudaSuccess
- 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 0, 1, 4, 5 et 6
- Observez-vous des gains significatifs ?
- Mettez-à-jour l'onglet "K4-7 CUBLAS" du fichier Excel
5 - Utilisation des TensorCores
- Implantez le pseudo-kernel K7 dans le fichier 'gpu.cu' pour utilisez les TensorCores :
- Utilisez la routine cublasGemmEx(...)
- et précisez le type T_CUBLAS_real défini dans 'main.h'
- et précisez l'algorithme CUBLAS_GEMM_DEFAULT_TENSOR_OP
- Testez
et validez votre implantation sur une matrice de 4096x4096 DOUBLE
- 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
-
Comparez-les aux performances de la meilleure configuration sur le CPU
- 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
- Vous pouvez expérimenter d'autres algorithmes :
- De CUBLAS_GEMM_ALGO0_TENSOR_OP à CUBLAS_GEMM_ALGO15_TENSOR_OP
- Relevez les performances et l'écart relatif maximum avec le calcul en FLOAT sur CPU (affichés par le programme)
- L'algorithme le plus performant affiche-t-il une erreur relative plus forte que les autres ?
Document à rendre :
Vous rendrez un SEUL document par binôme : un fichier Word (.doc) ou PDF :
- il sera nommé avec vos noms précédés
de "TP2-CUDA". Ex : "TP2-CUDA-Pignon.pdf"
- il comportera en haut de la première page :
- vos nom et prénoms, et le nom de votre option de 3A,
- la date de remise du rapport,
- le titre ("TP2 de programmation en CUBLAS").
- il contiendra en général :
- une description de la parallélisation réalisée (pourquoi ? comment ?),
- les
codes sources des parties importantes du code (ex : les routines des
kernels, des définition et lancement de grilles de blocs, et de
transferts CPU/GPU),
- vos
mesures de performances (Texec et/ou GigaFlops, SpeedUp...) sur GPU, sous forme de tableaux ou de figures,
- votre analyse des performances obtenues.
- Taille maximale de vos rapports : <voir taille maximale annoncée en cours>
- Date limite de remise de vos rapports : <voir date annoncée en cours>
- Remise par e-mail à Stephane Vialle (Stephane.Vialle@centralesupelec.fr)
