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.

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.

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.

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.

__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.

__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-1. __global__▲

V-B-2. __device__▲

V-B-3. __host__▲

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.

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

V-C-1. Dg▲

V-C-2. Db▲

V-C-3. Ns▲

V-C-4. S▲

V-D-1. __device__▲

V-D-2. __constant__▲

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++.

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▲

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).

# 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

# 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 :

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.

make emu=1