Une introduction à CUDA


précédentsommairesuivant

V. Les mains dans le cambouis

 
Sélectionnez

Il n'y a pas de langage informatique dans lequel vous ne puissiez écrire de mauvais programme.
 
Sélectionnez

Si vous ne savez pas ce que votre programme est censé faire, vous feriez bien de ne pas commencer à l'écrire.

(Extraits de Les lois de Murphy).

V-A. Les kernels

Très simplement, un kernel est une fonction exécutée sur le GPU.

Il en existe différent types, qualifiés de :

  1. __global__ ;
  2. __device__ ;
  3. __host__.

Le premier correspond à un kernel exécuté sur le GPU mais appelé par le CPU ; le deuxième, à un kernel exécuté et appelé par le GPU ; le troisième, à une fonction exécutée et appelée par le CPU. Ce dernier n'est pas obligatoire : c'est le mode de fonctionnement par défaut.

Un kernel ne s'appelle pas de la même manière qu'une fonction. Voici un appel de fonction.

Appel de fonction
Sélectionnez
fonction(parametre, parametre);

Mais avant de vous parler de l'appel d'un kernel, il faut que vous compreniez bien le mode de fonctionnement d'un GPU.

Une grille représente la totalité de la tâche à effectuer. Chaque grille peut être divisée en un ou plusieurs blocs, chacun exécutant plusieurs threads.

Un thread sur un GPU n'a pas le même sens qu'un thread sur le CPU.
Sur un GPU, il s'agit de la plus petite subdivision de la tâche à effectuer.

Image non disponible

Un appel de kernel se fait en spécifiant 2 paramètres entre triples chevrons précédant les paramètres passés au kernel.

 
Sélectionnez
kernel <<< nBlocs, threadsParBloc >>> (arguments);

nBlocs est le nombre de subdivisions appliquées à la grille à calculer et est de type dim3 (le cast à partir d'un entier N initialise le dim3 à {N, 1, 1}).
threadsParBloc indique le nombre de threads à exécuter simultanément pour chaque bloc. Ici encore, cette valeur est de type dim3.

Les valeurs à appliquer dépendent simultanément du problème à résoudre (choix des dimensions des blocs) et du matériel utilisé (nombre de threads par bloc). Choisir un nombre de threads supérieurs à la quantité nativement supportée entraînera une perte de performances. Cette notation permet ainsi d'adapter dynamiquement le programme aux matériels passés, présents et futurs.

Chaque kernel dispose de variables implicites en lecture seule (toutes de type dim3).

  1. blockIdx : index du bloc dans la grille,
  2. threadIdx : index du thread dans le bloc,
  3. blockDim : nombre de threads par bloc (valeur de threadsParBloc du paramétrage du kernel).

La grille est ici considérée comme un seul et unique bloc à une seule dimension.

 
Sélectionnez
__global__ void vecAdd(float * A, float * B, float * C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    // utilisation du kernel
    vecAdd<<<1, N>>>(A, B, C);
        //     |-> vecteurs additionnés une seule fois
        //        |-> nombre de composante des vecteurs
}

Dans le cas où la grille est sous-divisée en N blocs (tous de 1 dimension), l'index pourrait être trouvé de la manière suivante.

 
Sélectionnez
__global__ void vecAdd(float * A, float * B, float * C)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    // utilisation du kernel
    const int nThreadsPerBlocks  = 4;
    const int nBlocks            = (arraySize / nThreadsPerBlocks) + ( (arraySize % nThreadsPerBlocks) == 0 ? 0 : 1);
    vecAdd<<<nBlocks, nThreadsPerBlocks>>>(A, B, C);
}

Les variables doivent être qualifiées, pour définir leur lieu de résidence : voyez la section qui y est réservée.

Les paramètres entre chevrons sont requis, car le kernel est de type __global__. S'il était d'un autre type, ils n'auraient pas dû être précisés !

V-B. Qualifieurs de kernels

V-B-1. __global__

  • Exécuté sur le périphérique,
  • Appelable de l'hôte.
  • Pas de récursion possible,
  • Pas de variables statiques,
  • Pas de liste de paramètres variable.
  • On ne peut demander leur adresse mémoire,
  • Incompatible avec __device__,
  • Ne peut rien retourner,
  • À l'exécution, on doit préciser la configuration,
  • Appel asynchrone (le kernel retourne avant d'avoir effectué les calculs),
  • Paramètres stockés dans la mémoire partagée, limités à 256 octets,
  • Dure aussi longtemps que le kernel.

V-B-2. __device__

  • Exécuté sur le périphérique,
  • Appelable du périphérique.
  • Pas de récursion possible,
  • Pas de variables statiques,
  • Pas de liste de paramètres variable.
  • On ne peut demander leur adresse mémoire,
  • Incompatible avec __global__,
  • Dure aussi longtemps que l'application.

V-B-3. __host__

  • Exécuté sur l'hôte,
  • Appelable de l'hôte.
  • Appliqué par défaut.
  • Compatible avec __device__ (dans ce cas, le kernel pourra être exécuté sur l'hôte et sur le périphérique),
  • Incompatible avec __global__,
  • Dure aussi longtemps que le kernel.

V-C. Configuration de l'exécution

Ceci n'est requis que pour les kernels __global__ ! Requis signifie bien que l'on ne peut s'en passer, sans quoi rien ne fonctionne (avec, à la clé, beaux plantages) !

Cette configuration doit être passée entre triples chevrons avant les paramètres.

 
Sélectionnez
//Définition du kernel
__global__ void func(float * parameter);
//Utilisation du kernel
func <<< Dg, Db, Ns, S >>> (parameter);

V-C-1. Dg

  • Type : dim3 ;
  • Utilité : spécifier la taille et la dimension de la grille (le produit des trois composantes est le nombre de blocs lancés) ;
  • Remarque : z n'est pas encore utilisé et doit valoir 1.

V-C-2. Db

  • Type : dim3 ;
  • Utilité : spécifier la taille et la dimension de chaque bloc (le produit des trois composantes est le nombre de threads par bloc)

V-C-3. Ns

  • Type : size_t ;
  • Utilité : spécifier le nombre d'octets en mémoire partagée alloués dynamiquement par bloc en plus de la mémoire allouée statiquement ;
  • Remarque : paramètre optionnel, valeur par défaut : 0.

V-C-4. S

  • Type : cudaStream_t ;
  • Utilité : spécifier le flux associé ;
  • Remarque : paramètre optionnel, valeur par défaut : 0 ;
  • La notion de flux sera abordée plus tard : sachez simplement qu'il s'agit d'une suite d'éléments de même type (comme une texture).

V-D. Qualificateurs de variables

V-D-1. __device__

Cette variable est et restera sur le périphérique. Elle ne vivra pas plus longtemps que l'application et est accessible à tous les threads de la grille et à l'hôte grâce au runtime.

Ce type peut se marier avec un des deux suivants.

V-D-2. __constant__

Ce type peut être utilisé avec __device__.

La variable restera en mémoire constante. Elle ne vivra pas plus longtemps que l'application et est accessible à tous les threads de la grille et à l'hôte par le runtime.

Ces variables ne peuvent être déclarées que de l'hôte, pas du périphérique !

V-D-3. __shared__

Ce type peut être utilisé avec __device__.

La variable résidera dans la mémoire partagée et ne survivra pas au bloc. Elle ne sera accessible qu'aux threads du bloc.

Avant que les modifications soient écrites dans la variable et visibles pour tous les autres threads, il faut appeler __syncthreads();. À noter que cet appel ne sert qu'à le garantir, il est possible que les modifications soient visibles avant.

Tableau externe
Sélectionnez
extern __shared__ float shared[];

Quand la variable est déclarée en tant que tableau externe, comme précédemment, sa taille sera fixée à l'exécution. Toutes les variables déclarées de cette manière ne sont pas contiguës : le premier bit de la première correspond au premier bit des autres, contrairement aux autres langages comme le C ou le C++.

Image non disponible

C'est pourquoi il faut préciser l'offset de début. Pour avoir l'équivalent de ce premier code, il faut écrire le contenu du second.

Code C++
Sélectionnez
short array0[128];
float array1[64];
int array2[256];
Équivalent CUDA
Sélectionnez
extern __shared__ char array[];
__device__ void func()        // kernel __device__ ou bien __global__
{
    short* array0 = (short*) array;
    float* array1 = (float*)&array0[128];
    int  * array2 = (int*)  &array1[64];
}

Ceci est le seul moyen d'utiliser le mot-clé extern sur des variables : tous les autres emplois sont interdits.

Ces variables ne peuvent pas être initialisées en même temps que leur déclaration !

Si nous avions utilisé le premier code dans CUDA, en écrivant une valeur dans le premier tableau, une partie de cette variable aurait été imputée au deuxième et au troisième tableau. Ce qui pourrait donner des résultats très aberrants.

V-D-4. Généralités

Ces paramètres ne sont pas permis sur des unions ou des structures.

En définissant une variable __shared__ ou __constant__, elle sera définie statique.

V-E. Compilation

NVIDIA, dans son immense bonté, nous fournit un compilateur prévu pour CUDA. Celui-ci dispose d'une interface en ligne de commande simple et comparable à celles que nous connaissons, cl, de Visual Studio, ou gcc, l'interface de GCC. Ce compilateur, nvcc, s'occupe de toutes les étapes de la compilation.

Pour pouvoir définir les portions de code spécifiques à ce compilateur, il définit la macro __CUDACC__.

Comme dit précédemment, il s'occupe de toutes les phases de la compilation : l'assemblage, la compilation, et l'édition des liens. Vous pouvez choisir ces parties grâce à la ligne de commande.

Ce compilateur fonctionne très bien avec les Makefiles, c'est d'ailleurs cette technique qui va être ici développée, compatible avec les chaînes de compilation GNU (make) et Microsoft (nmake).

 
Sélectionnez
# Précise le compilateur précis à utiliser
ifdef ON_WINDOWS
    export compiler-bindir := "a:/program files/microsoft visual studio 9.0/vc/bin"
endif

export NVCC := a:/cuda/bin/nvcc.exe

cpp.obj : cpp.cpp
    $(NVCC) -c cpp.cpp $(CFLAGS) -o cpp.obj

c.o : c.c
    $(NVCC) -c c.c $(CFLAGS) -o c.obj

cu.o : cu.cu
    $(NVCC) -c cu.cu $(CFLAGS) -o cu.obj

OBJECTS = cpp.obj c.obj cu.obj

all : $(OBJECTS)
    $(NVCC) $(OBJECTS) $(LDFLAGS) -o app.exe

clean :
    $(RM) $(OBJECTS)

Ce Makefile doit être utilisé après avoir appelé le script vsvars.bat s'il est utilisé avec Visual Studio !

Si vous utilisez un make d'origine GNU, vous pouvez utiliser ce Makefile

 
Sélectionnez
# Précise le compilateur précis à utiliser
ifdef ON_WINDOWS
    export compiler-bindir := "a:/program files/microsoft visual studio 9.0/vc/bin"
endif

export NVCC := a:/cuda/bin/nvcc.exe

%.o : %.cpp
    $(NVCC) -c %^ $(CFLAGS) -o $@
    $(NVCC) -M %^ $(CFLAGS) > $@.dep

%.o : %.c
    $(NVCC) -c %^ $(CFLAGS) -o $@
    $(NVCC) -M %^ $(CFLAGS) > $@.dep

%.o : %.cu
    $(NVCC) -c %^ $(CFLAGS) -o $@
    $(NVCC) -M %^ $(CFLAGS) > $@.dep

include $(wildcard *.dep) /dev/null

all : $(OBJECTS)
    $(NVCC) $(OBJECTS) $(LDFLAGS) -o app.exe

clean :
    $(RM) $(OBJECTS) *.dep

Vous pouvez aussi utiliser la ligne de commande directement :

 
Sélectionnez
nvcc -c cu.cu -o cu.obj
nvcc cu.obj -o app.exe

Vous pouvez aussi décider que le code CUDA sera exécuté sur le processeur, qui émulera alors un GPU. Il suffit d'ajouter emu=1 à la ligne de commandes, comme ceci.

 
Sélectionnez
make emu=1

précédentsommairesuivant

Vous avez aimé ce tutoriel ? Alors partagez-le en cliquant sur les boutons suivants : Viadeo Twitter Facebook Share on Google+   

CUDA et le GPGPU
Introduction à CUDA
CUDA approfondi
La FAQ GPGPU & CUDA
  

Copyright © 2009 Thibaut Cuvelier. Aucune reproduction, même partielle, ne peut être faite de ce site et de l'ensemble de son contenu : textes, documents, images, etc. sans l'autorisation expresse de l'auteur. Sinon vous encourez selon la loi jusqu'à trois ans de prison et jusqu'à 300 000 € de dommages et intérêts.