IdentifiantMot de passe
Loading...
Mot de passe oublié ?Je m'inscris ! (gratuit)

Une introduction à CUDA

Une introduction CUDA


prcdentsommairesuivant

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

Trs simplement, un kernel est une fonction excute sur le GPU.

Il en existe diffrent types, qualifis de :

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

Le premier correspond un kernel excut sur le GPU mais appel par le CPU ; le deuxime, un kernel excut et appel par le GPU ; le troisime, une fonction excute et appele par le CPU. Ce dernier n'est pas obligatoire : c'est le mode de fonctionnement par dfaut.

Un kernel ne s'appelle pas de la mme manire 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 reprsente la totalit de la tche effectuer. Chaque grille peut tre divise en un ou plusieurs blocs, chacun excutant plusieurs threads.

Un thread sur un GPU n'a pas le mme sens qu'un thread sur le CPU.
Sur un GPU, il s'agit de la plus petite subdivision de la tche effectuer.

Image non disponible

Un appel de kernel se fait en spcifiant 2 paramtres entre triples chevrons prcdant les paramtres passs au kernel.

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

nBlocs est le nombre de subdivisions appliques 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 excuter simultanment pour chaque bloc. Ici encore, cette valeur est de type dim3.

Les valeurs appliquer dpendent simultanment du problme rsoudre (choix des dimensions des blocs) et du matriel utilis (nombre de threads par bloc). Choisir un nombre de threads suprieurs la quantit nativement supporte entranera une perte de performances. Cette notation permet ainsi d'adapter dynamiquement le programme aux matriels passs, prsents 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 paramtrage du kernel).

La grille est ici considre 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 additionns une seule fois
        //        |-> nombre de composante des vecteurs
}

Dans le cas o la grille est sous-divise en N blocs (tous de 1 dimension), l'index pourrait tre trouv de la manire 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 qualifies, pour dfinir leur lieu de rsidence : voyez la section qui y est rserve.

Les paramtres entre chevrons sont requis, car le kernel est de type __global__. S'il tait d'un autre type, ils n'auraient pas d tre prciss !

V-B. Qualifieurs de kernels

V-B-1. __global__

  • Excut sur le priphrique,
  • Appelable de l'hte.
  • Pas de rcursion possible,
  • Pas de variables statiques,
  • Pas de liste de paramtres variable.
  • On ne peut demander leur adresse mmoire,
  • Incompatible avec __device__,
  • Ne peut rien retourner,
  • l'excution, on doit prciser la configuration,
  • Appel asynchrone (le kernel retourne avant d'avoir effectu les calculs),
  • Paramtres stocks dans la mmoire partage, limits 256 octets,
  • Dure aussi longtemps que le kernel.

V-B-2. __device__

  • Excut sur le priphrique,
  • Appelable du priphrique.
  • Pas de rcursion possible,
  • Pas de variables statiques,
  • Pas de liste de paramtres variable.
  • On ne peut demander leur adresse mmoire,
  • Incompatible avec __global__,
  • Dure aussi longtemps que l'application.

V-B-3. __host__

  • Excut sur l'hte,
  • Appelable de l'hte.
  • Appliqu par dfaut.
  • Compatible avec __device__ (dans ce cas, le kernel pourra tre excut sur l'hte et sur le priphrique),
  • Incompatible avec __global__,
  • Dure aussi longtemps que le kernel.

V-C. Configuration de l'excution

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 passe entre triples chevrons avant les paramtres.

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

V-C-1. Dg

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

V-C-2. Db

  • Type : dim3 ;
  • Utilit : spcifier 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 : spcifier le nombre d'octets en mmoire partage allous dynamiquement par bloc en plus de la mmoire alloue statiquement ;
  • Remarque : paramtre optionnel, valeur par dfaut : 0.

V-C-4. S

  • Type : cudaStream_t ;
  • Utilit : spcifier le flux associ ;
  • Remarque : paramtre optionnel, valeur par dfaut : 0 ;
  • La notion de flux sera aborde plus tard : sachez simplement qu'il s'agit d'une suite d'lments de mme type (comme une texture).

V-D. Qualificateurs de variables

V-D-1. __device__

Cette variable est et restera sur le priphrique. Elle ne vivra pas plus longtemps que l'application et est accessible tous les threads de la grille et l'hte grce 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 mmoire constante. Elle ne vivra pas plus longtemps que l'application et est accessible tous les threads de la grille et l'hte par le runtime.

Ces variables ne peuvent tre dclares que de l'hte, pas du priphrique !

V-D-3. __shared__

Ce type peut tre utilis avec __device__.

La variable rsidera dans la mmoire partage 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 dclare en tant que tableau externe, comme prcdemment, sa taille sera fixe l'excution. Toutes les variables dclares de cette manire ne sont pas contigus : le premier bit de la premire correspond au premier bit des autres, contrairement aux autres langages comme le C ou le C++.

Image non disponible

C'est pourquoi il faut prciser l'offset de dbut. 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 initialises en mme temps que leur dclaration !

Si nous avions utilis le premier code dans CUDA, en crivant une valeur dans le premier tableau, une partie de cette variable aurait t impute au deuxime et au troisime tableau. Ce qui pourrait donner des rsultats trs aberrants.

V-D-4. Gnralits

Ces paramtres ne sont pas permis sur des unions ou des structures.

En dfinissant une variable __shared__ ou __constant__, elle sera dfinie statique.

V-E. Compilation

NVIDIA, dans son immense bont, nous fournit un compilateur prvu 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 dfinir les portions de code spcifiques ce compilateur, il dfinit la macro __CUDACC__.

Comme dit prcdemment, il s'occupe de toutes les phases de la compilation : l'assemblage, la compilation, et l'dition des liens. Vous pouvez choisir ces parties grce la ligne de commande.

Ce compilateur fonctionne trs bien avec les Makefiles, c'est d'ailleurs cette technique qui va tre ici dveloppe, compatible avec les chanes de compilation GNU (make) et Microsoft (nmake).

 
Sélectionnez
# Prcise le compilateur prcis  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 aprs 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
# Prcise le compilateur prcis  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 dcider que le code CUDA sera excut 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

prcdentsommairesuivant

Copyright © 2009 Thibaut Cuvelier. Aucune reproduction, même partielle, ne peut être faite de ce site ni 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.