enviando matriz 3d para o kernel CUDA

Tomei o código dado como uma resposta para Como posso sumr dois arrays 2D (inclinados) usando loops for nesteds? e tentou usá-lo para 3D em vez de 2D e mudou outras partes um pouco também, agora parece o seguinte:

__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]); } 

No código acima eu uso 2 como tamanhos para cada uma das dimensões de h_c, na implementação real terei esses tamanhos em números muito grandes e em diferentes para cada parte dos subarrays de “int ***” ou mais dimensões . Eu estou ficando problema com a parte após a chamada do kernel onde eu tento copiar os resultados para res array. Você pode me ajudar a consertar o problema? Plz você pode mostrar a solução da maneira que estou escrevendo acima. Obrigado!

Primeiro de tudo, eu acho que as garotas quando ele postou a resposta para a pergunta anterior que você mencionou, não pretendia que fosse representativo de uma boa codificação. Então, descobrir como estendê-lo para 3D pode não ser o melhor uso do seu tempo. Por exemplo, por que queremos escrever programas que usam exatamente um thread? Embora possa haver usos legítimos de tal kernel, este não é um deles. Seu kernel tem a possibilidade de fazer um monte de trabalho independente em paralelo , mas você está forçando tudo em um thread e serializando-o. A definição do trabalho paralelo é:

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

Vamos descobrir como lidar com isso em paralelo na GPU.

Outra observação introdutória que eu faria é que, uma vez que estamos lidando com problemas que têm tamanhos conhecidos antes, vamos usar o C para lidar com eles com o máximo de benefícios que pudermos obter da linguagem. Loops nesteds para fazer cudaMalloc podem ser necessários em alguns casos, mas não acho que seja um deles.

Aqui está um código que realiza o trabalho em 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 você pediu nos comentários, aqui está o menor número de alterações que eu poderia fazer no seu código para que ele funcione. Vamos também lembrar-nos de alguns comentários de garras da questão anterior que você mencionou:

"Por razões de complexidade de código e desempenho, você realmente não quer fazer isso, usando matrizes de pointers no código CUDA é mais difícil e mais lento que a alternativa usando memory linear."

"é uma idéia tão pobre se comparada ao uso de memory linear".

Eu tive que diagramar isso no papel para ter certeza de que a cópia do ponteiro estava correta.

 #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]); } 

Em suma, temos que fazer uma sequência sucessiva de:

  1. malloc uma multidimensional array de pointers (no host), uma dimensão menor que o tamanho do problema, com a última dimensão sendo um conjunto de pointers para as regiões cudaMalloc'ed no dispositivo em vez do host.
  2. crie outra multidimensional array de pointers, da mesma class criada na etapa anterior, mas uma dimensão menor que a criada na etapa anterior. este array também deve ter seus rankings finais cudaMalloc'ed no dispositivo.
  3. Copie o último conjunto de pointers do host do segundo passo anterior para a área cudaMalloced no dispositivo na etapa anterior.
  4. repita as etapas 2 a 3 até chegarmos a um único ponteiro (host) apontando para a multidimensional array de pointers, todos os quais agora residem no dispositivo.