CUDA approfondi


précédentsommairesuivant

II. Variables intégrées

Ces variables sont disponibles dans tous les kernels et indiquent les options de configuration lors du lancement du kernel ainsi que la position du thread dans la grille.

Nom de la variable Type Utilité
gridDim dim3 Dimensions de la grille
blockIdx uint3 Index du bloc dans la grille
blockDim dim3 Dimensions du bloc
threadIdx uint3 Index du thread dans le bloc
warpSize int Taille du warp


Cependant, il y a quelques restrictions.

  • On ne peut demander l'adresse de ces variables ;
  • On ne peut leur affecter une nouvelle valeur.

Voici un exemple d'utilisation.

 
Sélectionnez
printf("Dimensions de la grille : ( x = %d ; y = %d ; z= %d )", gridDim.x, gridDim.y, gridDim.z);
printf("Dimensions du bloc : ( x = %d ; y = %d ; z= %d )", blockDim.x, blockDim.y, blockDim.z);
printf("Identifiant de ce bloc : ( x = %d ; y = %d ; z = %d )", blockIdx.x, blockIdx.y, blockIdx.z);
printf("Identifiant de ce thread : ( x = %d ; y = %d ; z = %d )", threadIdx.x, threadIdx.y, threadIdx.z);
printf("Taille du warp : %d", warpSize);

III. Types de vecteurs intégrés

Voici la liste exhaustive de ces types.

  • char1
  • uchar1
  • char2
  • uchar2
  • char3
  • uchar3
  • char4
  • uchar4
  • short1
  • ushort1
  • short2
  • ushort2
  • short3
  • ushort3
  • short4
  • ushort4
  • int1
  • uint1
  • int2
  • uint2
  • int3
  • uint3
  • int4
  • uint4,
  • long1
  • ulong1
  • long2
  • ulong2
  • long3
  • ulong3
  • long4
  • ulong4
  • float1
  • float2
  • float3
  • float4
  • double2

Ces types sont dérivés des types du C char, unsigned char, short, unsigned short, int, unsigned int, long, unsigned long, float et double, comme indiqué par leur nom.

Il s'agit de structures. Ils représentent des vecteurs de une à quatre dimensions, accessibles via X, Y, Z et W, dans l'ordre.

On les initialise à l'aide des fonctions TYPE make_TYPE (TYPE x, TYPE y, TYPE z, TYPE w), comme le montre ce court exemple.

 
Sélectionnez
float1 floatVectorOne    =    make_float1(1.0);
float2 floatVectorTwo    =    make_float2(1.0, 4.2);
float3 floatVectorThree  =    make_float3(1.0, 4.2, 1475.41742);
float4 floatVectorFour   =    make_float4(1.0, 4.2, 1475.41742, 0.000000000004105);

On peut accéder aux composantes de ces vecteurs très simplement.

 
Sélectionnez
floatVectorOne.x;

floatVectorTwo.x;
floatVectorTwo.y;

floatVectorThree.x;
floatVectorThree.y;
floatVectorThree.z;

floatVectorFour.x;
floatVectorFour.y;
floatVectorFour.z;
floatVectorFour.w;

Il existe aussi le type dim3, qui sert à spécifier les dimensions (d'une grille, d'un bloc ...). Il s'agit d'un vecteur à trois dimensions d'entiers. Toutes les composantes non initialisées valent 1.

Nous avons appris dans l'article précédent à les utiliser et à les déclarer. Ce petit bout de code vous rappellera les méthodes à employer.

 
Sélectionnez
dim3 dim     (5);
dim3 dimgrid (1, 1);
dim3 dimblock(5, 24, 240);

IV. Textures

Les textures ne sont pas des types semblables aux autres. Elles se déclarent à la manière d'un kernel, mais, avant d'entrer dans ces subtilités, il faut déjà savoir ce qu'est une texture, dans le langage CUDA.

CUDA supporte une partie du matériel de texturing des GPU, utilisé pour les opérations sur la mémoire des textures. Cette mémoire est lue par le kernel quand une fonction de texture fetching (littéralement : aller chercher une texture) est appelée.

Dans cette section, quand je parlerai de tableaux, il s'agira de tableaux CUDA. Vous verrez l'importance de ceci à la prochaine section, concernant la mémoire.

IV-A. Références de texture

Le premier paramètre passé à une telle fonction s'appelle une référence de texture. Une référence de texture précise quels endroits de la mémoire seront utilisés pour cette texture. On doit la lier à une région de la mémoire, la texture, avant de pouvoir l'utiliser. Plusieurs références peuvent pointer sur une même texture, ou sur des textures qui se superposent.

Une référence possède plusieurs attributs. Parmi ceux-ci, la dimensionnalité, qui spécifie la manière d'accéder à la texture : via un tableau à une dimension (avec une seule coordonnée de texture) ; à deux dimensions (avec deux coordonnées de texture) ; ou à trois dimensions (avec trois coordonnées de texture). Les éléments d'une texture sont appelés texels, des éléments de texture (texture elements).

D'autres attributs sont les types des données d'entrée et de sortie, la manière d'interpréter les coordonnées d'entrée, entre autres.

IV-B. Déclaration & attributs à la compilation

Quelques attributs doivent être connus à la compilation et ne peuvent donc pas être changés plus tard. Ces attributs sont définis dès la déclaration de la référence de texture.

 
Sélectionnez
texture <Type, Dim, ReadMode> refDeTex;

Type y représente le type de données retourné lors du fetch. Cet attribut est limité aux entiers, aux flottants de simple précision et à tous les types de vecteurs, décrits plus haut.

Dim est la dimensionnalité, le nombre de dimensions, de la texture, qui peut valoir 1, 2 ou 3. Ce paramètre est optionnel et vaut par défaut 1.

ReadMode ne peut prendre que deux valeurs : cudaReadModeNormalizedFloat ou cudaReadModeElementType.

Si Type est un entier de 16 ou de 8 bits et que ReadMode vaut cudaReadModeNormalizedFloat, alors la valeur est, en vérité, retournée en tant que flottant. Toute la plage de valeurs de l'entier est reportée dans l'intervalle [ 0.0 ; 1.0 ] pour un non-signé, dans [ - 1.0 ; + 1.0 ] pour un signé. Par exemple, 0xFF sera lu comme 1.

Si ReadMode vaut cudaReadModeElementType, aucune conversion n'est effectuée.

Il s'agit d'un paramètre optionnel, dont la valeur par défaut est cudaReadModeElementType.

IV-C. Référence des attributs déclarés à l'exécution

Les autres attributs d'une référence sont mutables et peuvent sans problème être changés à l'exécution. Ils spécifient si les coordonnées de la texture sont normalisées ou non, le mode d'adressage et le filtrage.

Par défaut, les textures sont référencées avec des coordonnées flottantes dans l'intervalle [ 0 ; N ], où N est la taille de la texture dans la dimension correspondant à la coordonnée.

Par exemple, une texture de taille 64 x 32 sera référencée avec des coordonnées dans les intervalles [ 0 ; 63 ] et [ 0 ; 31 ] pour les coordonnées x et y.

Des coordonnées normalisées reportent ces intervalles dans l'intervalle [ 0.0 ; 1.0 ]. Cette normalisation convient parfaitement à certaines applications, s'il est nécessaire que les coordonnées soient indépendantes de la taille de la texture, par exemple.

Le mode d'adressage définit ce qui arrive lorsqu'un élément hors dimensions est demandé. Les valeurs en dessous de 0 sont mises à 0 et les variables plus grandes que N sont mises à N-1, dans le cas de coordonnées non normalisées. Dans le cas de coordonnées normalisées, les coordonnées sont ramenées dans l'intervalle [ 0.0 ; 1.0 ]. Dans ce dernier cas, il existe aussi un mode wrap, qui n'utilise que la partie fractionnaire de la coordonnée. 1.25 devient 0.25 ; -1.25 devient 0.75.

Le filtrage de textures linéaires ne peut être effectué que sur des textures qui retournent des flottants. Il effectue des interpolations de basse précision entre les texels proches. Quand ce mode est activé, les texels proches de la cellule recherchée sont lus et la valeur retournée est interpolée, sur base de l'espace entre la valeur recherchée et les données. Une interpolation simple est effectuée pour des textures à une dimension, une interpolation bilinéaire est effectuée pour des textures à deux dimensions.

IV-D. Comparaison des mémoires des textures

IV-D-1. Mémoire linéaire et tableaux

Une texture peut être stockée en mémoire linéaire ou sous forme de tableaux. La première méthode possède quelques désavantages.

  • Dimensionnalité forcée à 1 ;
  • Pas de support de filtrage ;
  • Adressage uniquement non normalisé ;
  • Pas de mode d'adressage : toute valeur hors intervalle est ramenée à 0.

Le matériel met en vigueur une politique d'alignement sur les adresses des textures. Pour rendre ceci plus transparent pour les programmeurs, les fonctions de bind des références renvoient un offset à appliquer aux recherches pour lire la mémoire désirée. Les pointeurs de base retournés par les fonctions d'allocation de CUDA se conforment à cette contrainte, ce qui fait qu'il n'est pas obligatoire de passer l'offset lorsqu'on les utilise.

IV-D-2. Mémoire constante et mémoire globale

La lecture sur le périphérique avec le principe des textures présente bien des avantages en comparaison des mémoires globale ou constante.

  • Elles sont en cache (ce qui améliore fortement les performances si les données sont disponibles) ;
  • Elles ne sont pas sujettes aux contraintes sur l'accès à la mémoire que les deux autres doivent respecter pour de bonnes performances ;
  • La latence est mieux cachée, améliorant les performances des applications qui lisent aléatoirement la mémoire ;
  • Les données d'entrée en entiers sur 8 ou 16 bits peuvent être converties en flottants sur 32 bits de l'intervalle [ 0.0 ; 1.0 ] ou [ -1.0 ; 1.0 ].

Si, en plus, la texture est stockée sous forme de tableau, le matériel fournit d'autres capacités, qui peuvent être utiles pour certaines applications, spécialement dans le domaine du traitement d'images.

Capacité Utilité Problème
Filtrage Interpolation rapide, mais peu précise, de texels Valide uniquement si la référence de texture renvoie un flottant
Textures à coordonnées normalisées Codage indépendant de la résolution  
Modes d'adressage Gestion automatique des cas de bordure Ne peut être utilisé qu'avec des coordonnées normalisées

Cependant, dans le même kernel, le cache des textures n'est pas gardé cohérent en fonction des écritures sur la mémoire globale. Ainsi, toute recherche sur une adresse qui a subi une écriture globale dans le même appel de kernel retourne des données indéfinies.

En d'autres termes, un thread peut lire en toute quiétude un endroit de la mémoire si et seulement si cet endroit a été mis à jour par un autre kernel ou une copie mémoire, mais pas si elle l'a été par le même thread ou un autre thread du même appel de kernel.

Ceci n'a d'utilité que lors de la recherche depuis la mémoire linéaire, comme un kernel ne peut pas écrire dans un tableau.

IV-E. Fonctions de fetching

Dans le cas de la mémoire linéaire, nous utilisons les fonctions de la famille tex1Dfetch().

 
Sélectionnez
float tex1Dfetch
(
    texture < unsigned char, 1, cudaReadModeNormalizedFloat > refDeTex,
    int x
);

x est l'ordonnée du flottant dont on recherche la valeur dans la texture dont une référence est refDeTex.

Le filtrage et les modes d'adressage ne sont pas supportés, contrairement aux couples et aux quadruplets.

 
Sélectionnez
float4 tex1Dfetch
(
    texture < uchar4, 1, cudaReadModeNormalizedFloat > texRef,
    int x
);

Par contre, dans le cas de tableaux, nous utiliserons les familles tex1D(), tex2D() et tex3D().

 
Sélectionnez
float tex1D
(
    texture < float, 1, cudaReadModeNormalizedFloat > texRef,
    float x
);

float tex2D
(
    texture < float, 2, cudaReadModeNormalizedFloat > texRef,
    float x, float y
);

float tex3D
(
    texture < float, 3, cudaReadModeNormalizedFloat > texRef,
    float x, float y, float z
);

IV-F. Fonctions de configuration

Le type texture, tel que défini par l'API haut niveau, est un dérivé du type textureReference, lui défini par l'API bas niveau.

 
Sélectionnez
struct textureReference
{
    int                            normalized     ;
    enum    cudaTextureFilterMode  filterMode     ;
    enum    cudaTextureAddressMode addressMode[3] ;
    struct  cudaChannelFormatDesc  channelDesc    ;
}

normalized non-nul signifie que la texture est normalisée ; toute autre valeur signifie la non-normalisation.

filterMode spécifie le mode de filtrage de la texture. S'il s'agit de cudaFilterModePoint, le texel le plus proche sera retourné ; s'il s'agit de cudaFilterModeLinear, une interpolation sera retournée.

addressMode précise le mode d'adressage : cudaAddressModeWrap impose que les valeurs incorrectes soient remises dans l'ensemble des correctes en n'utilisant que la partie fractionnaire ; cudaAddressModeClamp utilise l'autre mode.
Chacun des éléments du tableau correspond à une dimension à laquelle est appliqué le mode d'adressage.

channelDesc désigne le format de la valeur retournée à la lecture de la texture.

 
Sélectionnez
struct cudaChannelFormatDesc
{
    int x;
    int y;
    int z;
    int w;
    enum cudaChannelFormatKind f;
};

Chacun des entiers donne le nombre de bits de chaque composante retournée, f, leur type.

  • cudaChannelFormatKindSigned : entier signé ;
  • cudaChannelFormatKindUnsigned : entier non signé ;
  • cudaChannelFormatKindFloat : flottant.

Tous ces paramètres peuvent donc être spécifiés à l'exécution depuis l'hôte. Ils ne s'appliquent, comme dit précédemment, qu'aux textures liées à un tableau.

Avant qu'un kernel ne puisse utiliser une référence de texture pour la lire, la référence doit être liée à une texture, avec les fonctions cudaBindTexture() ou cudaBindTextureToArray().

Ce code lie une référence à un espace en mémoire linéaire, pointée par devPtr.

 
Sélectionnez
//Avec l'API bas niveau
texture <float, 1, cudaReadModeElementType> texRef;
textureReference * texRefPtr;
cudaGetTextureReference(&texRefPtr, "texRef");
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
cudaBindTexture(0, texRefPtr, devPtr, &channelDesc, size);

//Avec l'API haut niveau
texture<float, 1, cudaReadModeElementType> texRef;
cudaBindTexture(0, texRef, devPtr, size);

Ce code lie une référence à un tableau, cuArray.

 
Sélectionnez
//Avec l'API bas niveau
texture <float, 2, cudaReadModeElementType> texRef;
textureReference * texRefPtr;
cudaGetTextureReference(& texRefPtr, "texRef");
cudaChannelFormatDesc channelDesc;
cudaGetChannelDesc(& channelDesc, cuArray);
cudaBindTextureToArray(texRef, cuArray, & channelDesc);

//Avec l'API haut niveau
texture <float, 2, cudaReadModeElementType> texRef;
cudaBindTextureToArray(texRef, cuArray);

Le format précisé lors de la liaison d'une texture à une référence doit correspondre aux paramètres de la déclaration de référence. Sinon, les résultats sont indéfinis.

La fonction cudaUnbindTexture() sert à délier une référence d'une texture.

IV-G. Remarque

Vous ne pouvez pas déclarer une texture dans un .cu et l'utiliser dans un autre .cu, sous peine de problèmes à l'édition des liens.

Cependant, ceci peut être contourné. Il vous suffit d'utiliser des fonctions comme celles-ci pour les récupérer.

 
Sélectionnez
// Declaration d'une texture 2D
texture<float, 2, cudaReadModeElementType> dataTest_tex;

// Accesseur depuis une fonction __host__
texture<float, 2, cudaReadModeElementType>
& getTexture()
{
    return dataTest_tex;
}

// Accesseur depuis une fonction __device__ ou __global__
static __inline__ __device__
texture<float, 2, cudaReadModeElementType> 
& getDeviceTexture()
{
    return dataTest_tex;
}

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.