‘un access mémoire illégal’ lors d’une tentative d’écriture dans un tableau 2D alloué à l’aide de cudaMalloc3D

J’essaie d’allouer et de copier la mémoire d’un tableau 2D aplati sur le périphérique à l’aide de cudaMalloc3D pour tester les performances de cudaMalloc3D. Mais lorsque j’essaie d’écrire dans le tableau à partir du kernel, une exception «Un access mémoire illégal a été rencontré» est lancée. Le programme fonctionne correctement si je ne fais que lire dans le tableau, mais lorsque j’essaie d’y écrire, une erreur se produit. Toute aide à ce sujet sera grandement appréciée. Vous trouverez ci-dessous mon code et la syntaxe permettant de le comstackr.

Comstackr en utilisant

nvcc -O2 -arch sm_20 test.cu 

Code: test.cu

 #include  #include  #include  #define PI 3.14159265 #define NX 8192 /* includes boundary points on both end */ #define NY 4096 /* includes boundary points on both end */ #define NZ 1 /* needed for cudaMalloc3D */ #define N_THREADS_X 16 #define N_THREADS_Y 16 #define N_BLOCKS_X NX/N_THREADS_X #define N_BLOCKS_Y NY/N_THREADS_Y #define LX 4.0 /* length of the domain in x-direction */ #define LY 2.0 /* length of the domain in x-direction */ #define dx (REAL) ( LX/( (REAL) (NX) ) ) #define cSqrd 5.0 #define dt (REAL) ( 0.4 * dx / sqrt(cSqrd) ) #define FACTOR ( cSqrd * (dt*dt)/(dx*dx) ) #define IC (i + j*NX) /* (i,j) */ #define IM1 (i + j*NX - 1) /* (i-1,j) */ #define IP1 (i + j*NX + 1) /* (i+1,j) */ #define JM1 (i + (j-1)*NX) /* (i,j-1) */ #define JP1 (i + (j+1)*NX) /* (i,j+1) */ // Macro for checking CUDA errors following a CUDA launch or API call #define cudaCheckError() {\ cudaError_t e = cudaGetLastError();\ if( e != cudaSuccess ) {\ printf("\nCuda failure %s:%d: '%s'\n",__FILE__,__LINE__,cudaGetErrorSsortingng(e));\ exit(EXIT_FAILURE);\ }\ } typedef double REAL; typedef int INT; void meshGrid ( REAL *x, REAL *y ) { INT i,j; REAL a; for (j=0; j<NY; j++) { a = dx * ( (REAL) j ); for (i=0; i<NX; i++) { x[IC] = dx * ( (REAL) i ); y[IC] = a; } } } void initWave ( REAL *u, REAL *uold, REAL *x, REAL *y ) { INT i,j; for (j=1; j<NY-1; j++) { for (i=1; i<NX-1; i++) { u[IC] = 0.1 * (4.0*x[IC]-x[IC]*x[IC]) * ( 2.0*y[IC] - y[IC]*y[IC] ); } } for (j=1; j<NY-1; j++) { for (i=1; i0 && i 0 && j < (NY-1) ) { char *unewPtr = (char *) unew.ptr; REAL *unew_row = (REAL *) (unewPtr + i * unew.pitch); REAL tmp = unew_row[j]; // no error on this line unew_row[j] = 1.2; // this is where I get the error } } INT main(INT argc, char *argv[]) { INT nTimeSteps = 10; // pointers for the host side REAL *unew, *u, *uold, *uFinal, *x, *y; // allocate memory on the host unew = (REAL *)calloc(NX*NY,sizeof(REAL)); u = (REAL *)calloc(NX*NY,sizeof(REAL)); uold = (REAL *)calloc(NX*NY,sizeof(REAL)); uFinal = (REAL *)calloc(NX*NY,sizeof(REAL)); x = (REAL *)calloc(NX*NY,sizeof(REAL)); y = (REAL *)calloc(NX*NY,sizeof(REAL)); // pointer for the device side size_t pitch = NX * sizeof(REAL); cudaPitchedPtr d_u, d_uold, d_unew, d_tmp; cudaExtent myExtent = make_cudaExtent(pitch, NY, NZ); // allocate 3D memory on the device cudaMalloc3D( &d_u, myExtent ); cudaCheckError(); cudaMalloc3D( &d_uold, myExtent ); cudaCheckError(); cudaMalloc3D( &d_unew, myExtent ); cudaCheckError(); // initialize grid and wave meshGrid( x, y ); initWave( u, uold, x, y ); // copy host memory to 3D device memory cudaMemcpy3DParms cpy3D = { 0 }; cpy3D.kind = cudaMemcpyHostToDevice; // copying u to d_u cpy3D.srcPtr = make_cudaPitchedPtr(u, pitch, NX, NY); cpy3D.dstPtr = d_u; cpy3D.extent = myExtent; cudaMemcpy3D( &cpy3D ); cudaCheckError(); // copying uold to d_uold cpy3D.srcPtr = make_cudaPitchedPtr(uold, pitch, NX, NY); cpy3D.dstPtr = d_uold; cpy3D.extent = myExtent; cudaMemcpy3D( &cpy3D ); cudaCheckError(); // set up the GPU grid/block model dim3 dimGrid ( N_BLOCKS_X , N_BLOCKS_Y ); dim3 dimBlock ( N_THREADS_X, N_THREADS_Y ); for ( INT n = 1; n < nTimeSteps + 1; n++ ) { solveWaveGPU <<>> ( d_uold, d_u, d_unew ); cudaThreadSynchronize(); cudaCheckError(); d_tmp = d_uold; d_uold = d_u; d_u = d_unew; d_unew = d_tmp; } // copy the memory back to host cpy3D.kind = cudaMemcpyDeviceToHost; // copying d_unew to uFinal cpy3D.srcPtr = d_unew; cpy3D.dstPtr = make_cudaPitchedPtr(uFinal, pitch, NX, NY); cpy3D.extent = myExtent; cudaMemcpy3D( &cpy3D ); cudaCheckError(); free(u); cudaFree(d_u.ptr); free(unew); cudaFree(d_unew.ptr); free(uold); cudaFree(d_uold.ptr); free(uFinal); free(x); free(y); return EXIT_SUCCESS; } 

La raison pour laquelle l’erreur ne se produit pas sur cette ligne:

 REAL tmp = unew_row[j]; // no error on this line 

est parce que le compilateur optimise cette ligne. Cela ne sert à rien, et le compilateur l’élimine complètement. L’avertissement du compilateur:

 xxx.cu(87): warning: variable "tmp" was declared but never referenced 

est un indice à cet effet.

Votre code est presque presque correct. Le problème est ici:

 REAL *unew_row = (REAL *) (unewPtr + i * unew.pitch); 

CA devrait etre:

 REAL *unew_row = (REAL *) (unewPtr + j * unew.pitch); 

La variable i dans votre kernel est la dimension largeur (c’est-à-dire X). La variable j est la dimension de hauteur (ie Y).

La hauteur est celle qui fait référence à la rangée sur laquelle vous vous trouvez. Par conséquent, le pas de la rangée doit être multiplié par le paramètre height, c’est-à-dire j et non pas i .

De même, bien que ce ne soit pas la source de l’échec spécifique pour vos dimensions particulières, ce code peut ne pas être ce que vous vouliez soit:

 REAL tmp = unew_row[j]; // no error on this line unew_row[j] = 1.2; // this is where I get the error 

Si, par exemple, vous aviez l’intention de calculer le décalage par rapport à la ligne, puis d’indexer dans la ligne (par exemple, pour définir tous les éléments de l’allocation), je pense que vous voudriez utiliser i j comme index final:

 REAL tmp = unew_row[i]; // no error on this line unew_row[i] = 1.2; // this is where I get the error 

Toutefois, pour cet exemple particulier, il ne s’agit pas de la source réelle de l’access illégal à la mémoire.