TD 4. Introduction à CUDA
par Eric Goubault et Sylvie Putot
Utiliser CUDA dans les salles machines de l'X
Quelles cartes NVIDIA et ou sont-elles?
- En salles 30, 35 et 36 (liste des machines)
- Cartes GeForce GTX 260:
- Total amount of global memory: 938803200 bytes
- Number of multiprocessors: 24
- Number of cores: 192
- Total amount of constant memory: 65536 bytes
- Total amount of shared memory per block: 16384 bytes
- Total number of registers available per block: 16384
- Warp size: 32
- Maximum number of threads per block: 512
- Maximum sizes of each dimension of a block: 512 x 512 x 64
- Maximum sizes of each dimension of a grid: 65535 x 65535 x 1
- Maximum memory pitch: 262144 bytes
- Texture alignment: 256 bytes
- Clock rate: 1.24 GHz
Ecrire, compiler et exécuter un programme CUDA
- La première fois, il faut positionner les variables d'environnement. Créer ou ouvrir le fichier ~/.bashrc
et y ajouter:
PATH=$PATH:/usr/local/cuda/bin:/usr/local/cuda/open64/bin
LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/cuda/lib
export PATH
export LD_LIBRARY_PATH
- Récupérer l'arborescence d'exemples (dans laquelle par la suite vous mettrez vos propres programmes) :
cp -r /usr/local/cuda/sdk .
- Les exemples sont dans sdk/C/src:
- il y a une doc explicative dans le sous-répertoire doc de certains exemples
- compiler par make dans le répertoire sdk/C (on peut ensuite ne recompiler qu'un exemple dans le sous-répertoire correspondant)
- l'exécutable est dans sdk/bin/linux/release/ (par exemple sdk/C/bin/linux/release/scan pour exécuter l'exemple "scan")
- Un template se trouve dans le répertoire sdk/C/src/template, vous pouvez par exemple recopier ce répertoire
sous un autre nom mais au même niveau (par exemple sdk/C/src/MonEssai) pour démarrer rapidement votre propre code (il suffira
de changer les noms de fichiers dans le(s) Makefile)
- Pour débugger, compiler par make dbg=1 dans le répertoire sdk/C , on peut ensuite
lancer cuda-gdb sur l'exécutable présent dans sdk/C/bin/linux/debug, mais seulement en
mode non graphique (voir aussi la doc de cuda-gdb sur CudaZone).
- Une autre alternative est de compiler par make emu=1 (exécutables dans sdk/C/bin/linux/emurelease) pour pouvoir émuler sur le CPU ce qui se passe
sur le device et faire des affichages.
- Pour compiler vos fichiers indépendamment (en dehors donc) de la SDK, vous pouvez, par exemple pour le répertoire template, partir du Makefile suivant
- Vous pouvez aussi tout mettre dans un seul fichier mon_essai.cu et compiler puis exécuter
par
nvcc mon_essai.cu -I/usr/local/cuda/include/ -L/usr/local/cuda/lib -o mon_essai
./mon_essai
Si vous voulez utiliser la librairie cutil de la sdk (divers utilitaires, dont les timers et tout ce qui est en cut), il faut ajouter des include et des librairies, par exemple:
nvcc mon_essai.cu -I/usr/local/cuda/include/ -I/usr/local/cuda/sdk/C/common/inc/ -L/usr/local/cuda/lib/
-L/users/profs/info/goubaul1/sdk/C/lib/ -lcutil -o mon_essai
- Ne pas hésiter à consulter Cuda Zone
Mise en oeuvre : multiplication de matrices
- Commencer par la méthode naïve avec tout en mémoire globale, et une case du tableau destination par threads.
- Tester avec différentes répartitions de threads et blocs, puis essayer de faire des versions améliorées,
en chargeant des sous-blocs de la matrice en mémoire partagée.
- Enfin, comparer par rapport à l'implémentation
de cette multiplication de la librairie CUBLAS (fonction cublasSgemm)
Quelques outils pour vous aider:
Visual profiler
Occupancy calculator
Mode d'emploi inclus dans le fichier excel.
CUBLAS
CUBLAS est une implémentation de la librairie d'algèbre linéaire BLAS pour CUDA. Elle est incluse dans
la distribution de CUDA: documentation (d'une version légèrement antérieure), également une courte présentation
Pour utiliser cette bibliothèque, il suffit d'ajouter en tête de vos fichiers .cu
#include "cublas.h"
et d'ajouter -lcublas à la ligne de compilation.Vous avez un exemple de multiplication utilisant CUBLAS dans sdk/C/src/simpleCUBLAS.
Attention, CUBLAS utilise un stockage en colonnes des matrices ("column-major order" dans Wikipedia), contrairement à celui du C qui est en lignes. Il faut initialiser
(cublasInit()) et arrêter (cublasShutdown()) la librairie avant et après usage. L'allocation et les transferts
de mémoire entre CPU et GPU se font de manière assez semblable à CUDA (utilisation de cublasAlloc(...), cublasFree(...),
cublasSetVector(...) et cublasGetVector(...).)
JCuda
Des bindings Java pour Cuda
- Disponibles sur jcuda.org, également installés en salles machines,
- Le code exécuté sur les cartes doit toujours être écrit en C (à moins d'utiliser les packages JCublas
ou JCufft ou JCudpp),
- Mais une possibilité est par exemple pour vos projets de distribuer les calculs sur plusieurs cartes graphiques en utilisant Java/RMI.
L'utiliser en salles machines:
- Variables d'environnement à positionner dans le .bash_profile:
JCUDA=/usr/local/JCuda
LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$JCUDA
CLASSPATH=.:$JCUDA/jcublas-0.2.3.jar:$JCUDA/jcuda-0.2.3.jar:$JCUDA/jcudpp-0.2.3.jar:$JCUDA/jcufft-0.2.3.jar
export JCUDA
export LD_LIBRARY_PATH
export CLASSPATH
- Compiler puis exécuter par:
javac -classpath $CLASSPATH mon_example.java
java -classpath $CLASSPATH mon_example