matrice équation de la chaleur dans CUDA - adresse illégale erreur

Question

Après this question en référence au shared memory example dans le guide officiel, je suis en train de construire la matrice de l'équation de la chaleur, ce qui est tout comme dans cette image mal dessiné j'ai fait

Voici ce que je l'ai fait jusqu'à présent, par exemple un minimum

#define N 32 
#define BLOCK_SIZE 16 
#define NUM_BLOCKS ((N + BLOCK_SIZE - 1)/ BLOCK_SIZE) 

__global__ void heat_matrix(int* A) 
{ 
    const unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x; 
    __shared__ int temp_sm_A[N*N]; 
    int* temp_A = &temp_sm_A[0]; memset(temp_A, 0, N*N*sizeof(int)); 

    if (tid < N) //(*) 
    { 
     #pragma unroll 
     for (unsigned int m = 0; m < NUM_BLOCKS; ++m) 
     {   
      #pragma unroll 
      for (unsigned int e = 0; e < BLOCK_SIZE ; ++e) 
      { 
       if ((tid == 0 && e == 0) || (tid == (N-1) && e == (BLOCK_SIZE-1))) 
       { 
        temp_A[tid + (e + BLOCK_SIZE * m) * N] = -2; 
        temp_A[tid + (e + BLOCK_SIZE * m) * N + (tid==0 ? 1 : -1)] = 1; 
       } 
       if (tid == e) 
       { 
        temp_A[tid + (e + BLOCK_SIZE * m) * N - 1] = 1; 
        //printf("temp_A[%d] = 1;\n", (tid + (e + BLOCK_SIZE * m) * N -1)); 
        temp_A[tid + (e + BLOCK_SIZE * m) * N] = -2; 
        //printf("temp_A[%d] = -2;\n", (tid + (e + BLOCK_SIZE * m) * N)); 
        temp_A[tid + (e + BLOCK_SIZE * m) * N + 1] = 1; 
        //printf("temp_A[%d] = 1;\n", (tid + (e + BLOCK_SIZE * m) * N +1)); 
       } 
      } 
     } 
     __syncthreads(); //(**) 
     memcpy(A, temp_A, N*N*sizeof(int)); 
    } 
} 
int main(){ 
    int* h_A = (int*)malloc(N*N*sizeof(int)); memset(h_A, 0, N*N*sizeof(int)); 
    int* d_A; 
    checkCudaErrors(cudaMalloc((void**)&d_A, N*N*sizeof(int))); 
    checkCudaErrors(cudaMemcpy(d_A, h_A, N*N*sizeof(int), cudaMemcpyHostToDevice)); 
    dim3 dim_grid((N/2 + BLOCK_SIZE -1)/ BLOCK_SIZE); 
    dim3 dim_block(BLOCK_SIZE); 

    heat_matrix <<< dim_grid, dim_block >>> (d_A); 
    checkCudaErrors(cudaMemcpy(h_A, d_A, N*N*sizeof(int), cudaMemcpyDeviceToHost)); 
... 
} 

Le code est conçu pour convenir à un grand N (supérieur à 32). J'ai profité de la division en blocs. Lors de l'exécution nvcc donne

CUDA error at matrix.cu:102 code=77(cudaErrorIllegalAddress) "cudaMemcpy(h_A, d_A, N*N*sizeof(int), cudaMemcpyDeviceToHost)" 

Et cuda-memcheck ne fournit qu'une seule erreur (en fait il y a un autre, mais il vient de cudasuccess=checkCudaErrors(cudaDeviceReset()); ...)

========= CUDA-MEMCHECK 
========= Invalid __shared__ write of size 4 
=========  at 0x00000cd0 in heat_matrix(int*) 
=========  by thread (0,0,0) in block (0,0,0) 
=========  Address 0xfffffffc is out of bounds 
... 

Je ne vois pas où je fait de mal dans le code. Comment le fil 0 dans le premier bloc peut-il provoquer un accès illégal? Il y a même le cas spécifique if à traiter, et la ligne du code dans lequel l'erreur s'est produite n'est pas signalée.

De plus, y a-t-il un moyen plus efficace pour mon code que de traiter tous ces if s? Bien sûr il y a, mais je ne pouvais pas trouver une meilleure expression parallèle pour diviser les cas dans le second for.

---

Sur une note de côté, pour moi le (*) semble inutile; au lieu (**) est nécessaire si je veux suivre avec d'autres appels de fonction GPU. Ai-je raison?

Answer 1

1. Voir cette ligne:

     temp_A[tid + (e + BLOCK_SIZE * m) * N - 1] = 1; 
   

   Pour le fil avec tid égal à zéro au cours de la première itération, tid + (e + BLOCK_SIZE * m) * N - 1 évalue à un indice de -1. C'est exactement ce à quoi la sortie de cuda-memcheck se plaint (avec l'adresse qui s'est enroulée en raison d'un dépassement de capacité).

2. Un accès similaire hors des limites du terrain se produira plus tard pour la ligne

     temp_A[tid + (e + BLOCK_SIZE * m) * N + 1] = 1; 
   

   quand tid, e et m tous assument leur valeur maximale.

3. Plusieurs threads écrivent dans le même emplacement de mémoire. Chaque thread doit écrire exactement dans un élément de tableau par itération de boucle interne. Il n'est pas nécessaire d'écrire les éléments voisins car ils sont déjà couverts par leurs propres threads.

4. Vous avez une condition de concurrence entre l'initialisation memset() et les magasins à l'intérieur des boucles principales. Mettez un syncthreads() après le memset().

5. Les appels à memset() et memcpy() conduiront à chaque thread faisant un ensemble/copie complète, en faisant les opérations N fois au lieu d'une seule fois.
   La manière courante de gérer ceci est d'écrire explicitement l'opération, en divisant le travail entre les threads du bloc.
   Cependant ...

6. Il n'y a aucun avantage à créer d'abord la matrice dans la mémoire partagée, puis à la copier plus tard dans la mémoire globale. Ecrire directement à A dans la mémoire globale élimine le besoin de memset(), memcpy() et syncthreads() tout à fait. L'utilisation d'une taille de bloc de seulement 16 threads laisse la moitié des ressources inutilisées, car les blocs de thread sont alloués en unités de 32 threads (une chaîne).

Vous pouvez relire la section sur le Thread Hierarchy dans le Guide de programmation CUDA C.

Answer 2

Dans votre noyau, temp_A est un pointeur local au début de votre matrice de mémoire partagée. Considérant:

N = 32;

BLOCK_SIZE = 16;

m (0,1);

e (0, BLOCK_SIZE)

Accède comme temp_A[tid + (e + BLOCK_SIZE * m) * N] peut facilement sortir des limites de 1024 éléments à long tableau.