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 :

1

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

1
2
3

short array0[128];
float array1[64];
int array2[256];

Équivalent CUDA :

1
2
3
4
5
6
7

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.