utilisant à la fois allouée dynamiquement et je statiquement mémoire partagée allouée

Question

Supposons que deux __device__ fonction CUDA, chacun ayant la variable locale suivante:

__shared__ int a[123]; 

et une autre fonction (dire qu'il est mon noyau, soit une fonction __global__) , avec:

extern __shared__ int b[]; 

Est-ce explicitement autorisé/interdit par nVIDIA? (Je ne le vois pas dans la section programming guide B.2.3 sur __shared__) Est-ce que toutes les tailles comptent ensemble vers la limite de mémoire partagée, ou est-ce le maximum possible en cours d'utilisation en même temps? Ou une autre règle?

_{Ceci peut être considéré comme une question complémentaire à this one.}

Answer 1

La mémoire partagée est divisée en deux parties: allouée statiquement et allouée dynamiquement. La première partie est calculée lors de la compilation, et chaque déclaration est une allocation réelle - activation ptxas information lors de la compilation illustre ici:

ptxas info : Used 22 registers, 384 bytes smem, 48 bytes cmem[0] 

Ici, nous avons 384 octets, ce qui est 3 tableaux de 32 ints. (voir l'exemple de corde ci-dessous).

Vous pouvez passer un pointeur vers la mémoire partagée depuis Kepler, vers une autre fonction permettant à une sous-fonction de périphérique d'accéder à une autre déclaration de mémoire partagée. Ensuite, vient la mémoire partagée allouée dynamiquement, dont la taille réservée est déclarée lors de l'appel du noyau.

Voici un exemple de diverses utilisations dans quelques fonctions. Notez la valeur de pointeur de chaque région de mémoire partagée.

__device__ void dev1() 
{ 
    __shared__ int a[32] ; 
    a[threadIdx.x] = threadIdx.x ; 

    if (threadIdx.x == 0) 
     printf ("dev1 : %x\n", a) ; 
} 

__device__ void dev2() 
{ 
    __shared__ int a[32] ; 
    a[threadIdx.x] = threadIdx.x * 5 ; 

    if (threadIdx.x == 0) 
     printf ("dev2 : %x\n", a) ; 
} 

__global__ void kernel(int* res, int* res2) 
{ 
    __shared__ int a[32] ; 
    extern __shared__ int b[]; 

    a[threadIdx.x] = 0 ; 
    b[threadIdx.x] = threadIdx.x * 3 ; 

    dev1(); 
    __syncthreads(); 
    dev2(); 
    __syncthreads(); 

    res[threadIdx.x] = a[threadIdx.x] ; 
    res2[threadIdx.x] = b[threadIdx.x] ; 

    if (threadIdx.x == 0) 
     printf ("global a : %x\n", a) ; 
    if (threadIdx.x == 0) 
     printf ("global b : %x\n", b) ; 
} 

int main() 
{ 
    int* dres ; 
    int* dres2 ; 

    cudaMalloc <> (&dres, 32*sizeof(int)) ; 
    cudaMalloc <> (&dres2, 32*sizeof(int)) ; 

    kernel<<<1,32,32*sizeof(float)>>> (dres, dres2); 

    int hres[32] ; 
    int hres2[32] ; 

    cudaMemcpy (hres, dres, 32 * sizeof(int), cudaMemcpyDeviceToHost) ; 
    cudaMemcpy (hres2, dres2, 32 * sizeof(int), cudaMemcpyDeviceToHost) ; 

    for (int k = 0 ; k < 32 ; ++k) 
    { 
     printf ("%d -- %d \n", hres[k], hres2[k]) ; 
    } 
    return 0 ; 
} 

Ce code renvoie l'info ptxas en utilisant 384 bytes smem, qui est un tableau pour tableau global a, une seconde pour le tableau méthode dev1 a, et un troisième pour le tableau de la méthode de dev2 a. Totaliser 3*32*sizeof(float)=384 bytes. Lors de l'exécution du noyau avec la mémoire partagée dynamique égale à 32*sizeof(float), le pointeur vers b commence juste après ces trois tableaux.

EDIT: Le fichier PTX généré par le code contient des déclarations de mémoire partagée statiquement définie,

.shared .align 4 .b8 _ZZ4dev1vE1a[128]; 
.shared .align 4 .b8 _ZZ4dev2vE1a[128]; 
.extern .shared .align 4 .b8 b[]; 

à l'exception du point d'entrée où elle est définie dans le corps de la méthode

// _ZZ6kernelPiS_E1a has been demoted 

l'espace commun de la mémoire est définie dans la documentation PTX here:

L'espace d'état partagé (.shared) est une zone de mémoire par CTA pour les threads dans un CTA pour partager des données. Une adresse dans la mémoire partagée peut être lue et écrite par n'importe quel thread dans un CTA. Utilisez ld.shared et st.shared pour accéder aux variables partagées.

Bien que sans détail sur le temps d'exécution. Il y a un mot dans le guide de programmation here sans autre détail sur le mélange des deux.

Lors de la compilation PTX, le compilateur peut connaître la quantité de mémoire partagée allouée de manière statique. Il pourrait y avoir de la magie supplémentaire. En regardant le SASS, les premières instructions utilisent la SR_LMEMHIOFF

1    IADD32I R1, R1, -0x8; 
2    S2R R0, SR_LMEMHIOFF; 
3    ISETP.GE.U32.AND P0, PT, R1, R0, PT; 

et les fonctions d'appel dans l'ordre inverse attribuent des valeurs différentes à la mémoire partagée statiquement allouée (ressemble beaucoup à une forme de stackalloc). Je crois que le compilateur ptxas calcule toute la mémoire partagée dont il pourrait avoir besoin dans le pire des cas quand toute la méthode peut être appelée (quand on n'utilise pas la méthode et en utilisant des pointeurs de fonction, l'adresse b ne change pas, et la non allouée la zone de mémoire partagée n'est jamais accessible). Enfin, comme le suggère einpoklum dans un commentaire, ceci est expérimental et ne fait pas partie d'une définition de norme/API.