enviando matriz 3D al kernel de CUDA

Tomé el código proporcionado como respuesta para ¿Cómo puedo agregar dos matrices 2d (lanzadas) usando bucles nesteds? e intenté usarlo para 3D en lugar de 2D y también cambié otras partes ligeramente, ahora se ve de la siguiente manera:

__global__ void doSmth(int*** a) { for(int i=0; i<2; i++) for(int j=0; j<2; j++) for(int k=0; k<2; k++) a[i][j][k]=i+j+k; } int main() { int*** h_c = (int***) malloc(2*sizeof(int**)); for(int i=0; i<2; i++) { h_c[i] = (int**) malloc(2*sizeof(int*)); for(int j=0; j<2; j++) GPUerrchk(cudaMalloc((void**)&h_c[i][j],2*sizeof(int))); } int*** d_c; GPUerrchk(cudaMalloc((void****)&d_c,2*sizeof(int**))); GPUerrchk(cudaMemcpy(d_c,h_c,2*sizeof(int**),cudaMemcpyHostToDevice)); doSmth<<>>(d_c); GPUerrchk(cudaPeekAtLastError()); int res[2][2][2]; for(int i=0; i<2; i++) for(int j=0; j<2; j++) GPUerrchk(cudaMemcpy(&res[i][j][0], h_c[i][j],2*sizeof(int),cudaMemcpyDeviceToHost)); for(int i=0; i<2; i++) for(int j=0; j<2; j++) for(int k=0; k<2; k++) printf("[%d][%d][%d]=%d\n",i,j,k,res[i][j][k]); } 

En el código anterior utilizo 2 como tamaños para cada dimensión de h_c, en la implementación real tendré estos tamaños en grandes cantidades y en diferentes para cada parte de los subconjuntos de “int ***” o más dimensiones . Tengo problemas con la parte después de la llamada al kernel, donde trato de copiar los resultados a res array. ¿Puedes ayudarme a solucionar el problema? Por favor, pueden mostrar la solución en la forma en que la escribo arriba. ¡Gracias!

En primer lugar, creo que talonmies cuando publicó la respuesta a la pregunta anterior que mencionas, no tenía la intención de ser representativo de una buena encoding. Así que descubrir cómo extenderlo a 3D podría no ser el mejor uso de su tiempo. Por ejemplo, ¿por qué queremos escribir progtwigs que usan exactamente un hilo? Si bien puede haber usos legítimos de dicho kernel, este no es uno de ellos. Su kernel tiene la posibilidad de hacer un montón de trabajo independiente en paralelo , pero en cambio lo está forzando todo en un hilo y serialándolo. La definición del trabajo paralelo es:

 a[i][j][k]=i+j+k; 

Veamos cómo manejar eso en paralelo en la GPU.

Otra observación introductoria que haría es que, dado que estamos lidiando con problemas que tienen tamaños que se conocen de antemano, usemos C para abordarlos con el mayor beneficio que podamos obtener del lenguaje. Es posible que se necesiten ciclos nesteds para hacer cudaMalloc en algunos casos, pero no creo que este sea uno de ellos.

Aquí hay un código que realiza el trabajo en paralelo:

 #include  #include  // set a 3D volume // To compile it with nvcc execute: nvcc -O2 -o set3d set3d.cu //define the data set size (cubic volume) #define DATAXSIZE 100 #define DATAYSIZE 100 #define DATAZSIZE 20 //define the chunk sizes that each threadblock will work on #define BLKXSIZE 32 #define BLKYSIZE 4 #define BLKZSIZE 4 // for cuda error checking #define cudaCheckErrors(msg) \ do { \ cudaError_t __err = cudaGetLastError(); \ if (__err != cudaSuccess) { \ fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ msg, cudaGetErrorString(__err), \ __FILE__, __LINE__); \ fprintf(stderr, "*** FAILED - ABORTING\n"); \ return 1; \ } \ } while (0) // device function to set the 3D volume __global__ void set(int a[][DATAYSIZE][DATAXSIZE]) { unsigned idx = blockIdx.x*blockDim.x + threadIdx.x; unsigned idy = blockIdx.y*blockDim.y + threadIdx.y; unsigned idz = blockIdx.z*blockDim.z + threadIdx.z; if ((idx < (DATAXSIZE)) && (idy < (DATAYSIZE)) && (idz < (DATAZSIZE))){ a[idz][idy][idx] = idz+idy+idx; } } int main(int argc, char *argv[]) { typedef int nRarray[DATAYSIZE][DATAXSIZE]; const dim3 blockSize(BLKXSIZE, BLKYSIZE, BLKZSIZE); const dim3 gridSize(((DATAXSIZE+BLKXSIZE-1)/BLKXSIZE), ((DATAYSIZE+BLKYSIZE-1)/BLKYSIZE), ((DATAZSIZE+BLKZSIZE-1)/BLKZSIZE)); // overall data set sizes const int nx = DATAXSIZE; const int ny = DATAYSIZE; const int nz = DATAZSIZE; // pointers for data set storage via malloc nRarray *c; // storage for result stored on host nRarray *d_c; // storage for result computed on device // allocate storage for data set if ((c = (nRarray *)malloc((nx*ny*nz)*sizeof(int))) == 0) {fprintf(stderr,"malloc1 Fail \n"); return 1;} // allocate GPU device buffers cudaMalloc((void **) &d_c, (nx*ny*nz)*sizeof(int)); cudaCheckErrors("Failed to allocate device buffer"); // compute result set<<>>(d_c); cudaCheckErrors("Kernel launch failure"); // copy output data back to host cudaMemcpy(c, d_c, ((nx*ny*nz)*sizeof(int)), cudaMemcpyDeviceToHost); cudaCheckErrors("CUDA memcpy failure"); // and check for accuracy for (unsigned i=0; i 

Como lo solicitó en los comentarios, esta es la cantidad más pequeña de cambios que puedo hacer en su código para que funcione. También recordemos algunos de los comentarios de la pregunta anterior que hace referencia a los talonmies:

"Por razones de complejidad y rendimiento del código, realmente no desea hacer eso, el uso de matrices de punteros en el código CUDA es más difícil y más lento que la alternativa que utiliza la memoria lineal".

"Es una idea tan pobre en comparación con el uso de la memoria lineal".

Tuve que diagtwigr esto en papel para asegurarme de que obtuve toda la copia de mi puntero correcta.

 #include  inline void GPUassert(cudaError_t code, char * file, int line, bool Abort=true) { if (code != 0) { fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code),file,line); if (Abort) exit(code); } } #define GPUerrchk(ans) { GPUassert((ans), __FILE__, __LINE__); } __global__ void doSmth(int*** a) { for(int i=0; i<2; i++) for(int j=0; j<2; j++) for(int k=0; k<2; k++) a[i][j][k]=i+j+k; } int main() { int*** h_c = (int***) malloc(2*sizeof(int**)); for(int i=0; i<2; i++) { h_c[i] = (int**) malloc(2*sizeof(int*)); for(int j=0; j<2; j++) GPUerrchk(cudaMalloc((void**)&h_c[i][j],2*sizeof(int))); } int ***h_c1 = (int ***) malloc(2*sizeof(int **)); for (int i=0; i<2; i++){ GPUerrchk(cudaMalloc((void***)&(h_c1[i]), 2*sizeof(int*))); GPUerrchk(cudaMemcpy(h_c1[i], h_c[i], 2*sizeof(int*), cudaMemcpyHostToDevice)); } int*** d_c; GPUerrchk(cudaMalloc((void****)&d_c,2*sizeof(int**))); GPUerrchk(cudaMemcpy(d_c,h_c1,2*sizeof(int**),cudaMemcpyHostToDevice)); doSmth<<<1,1>>>(d_c); GPUerrchk(cudaPeekAtLastError()); int res[2][2][2]; for(int i=0; i<2; i++) for(int j=0; j<2; j++) GPUerrchk(cudaMemcpy(&res[i][j][0], h_c[i][j],2*sizeof(int),cudaMemcpyDeviceToHost)); for(int i=0; i<2; i++) for(int j=0; j<2; j++) for(int k=0; k<2; k++) printf("[%d][%d][%d]=%d\n",i,j,k,res[i][j][k]); } 

En pocas palabras, tenemos que hacer una secuencia sucesiva de:

  1. malloc una matriz multidimensional de punteros (en el host), una dimensión menor que el tamaño del problema, con la última dimensión es un conjunto de punteros a las regiones cudaMalloc'ed en el dispositivo en lugar del host.
  2. cree otra matriz multidimensional de punteros, de la misma clase que la creada en el paso anterior, pero una dimensión menor que la creada en el paso anterior. esta matriz también debe tener sus rangos finales cudaMalloc'ed en el dispositivo.
  3. copie el último conjunto de punteros de host del segundo paso anterior en el área cudaMalloced en el dispositivo en el paso anterior.
  4. repita los pasos 2-3 hasta que terminemos con un puntero único (host) apuntando a la matriz multidimensional de punteros, todos los cuales ahora están en el dispositivo.