CUDA approfondi

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.

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.

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.

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-1. Mémoire linéaire et tableaux▲

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

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

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.

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

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

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

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.

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.

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.

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

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

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.

// 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;
}