发送3D数组到CUDA内核

我把给出的代码作为答案我怎样才能使用嵌套for循环加起来两个2d（pitched）数组？ 并试图使用它的3D而不是2D和稍微改变其他部分，现在看起来如下：

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

在上面的代码中，我使用2作为h_c的每个维度的大小，在实际的实现中，我将在“int ***”或更多维度的子数组的每个部分以非常大的数量和不同的大小。 内核调用后，我试图将结果复制回res数组。 你能帮我解决这个问题吗？ 你可以用我上面写的方式显示解决scheme。 谢谢！

Answer 1

首先，当他发表你对上一个问题的回应时，我认为他是个翘首以待的人，并不是要成为好的编码的代表。 因此，搞清楚如何将其扩展到3D可能不是您最好的时间。 例如，为什么我们要编写只使用一个线程的程序？ 虽然这样的内核可能有合法的用途，但这不是其中之一。 你的内核有可能并行地做一堆独立的工作，但是你可以把它全部强制到一个线程上，然后序列化它。 平行工作的定义是：

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

让我们弄清楚如何在GPU上并行处理。

我会做的另一个介绍性的观察是，由于我们正在处理提前知道的大小的问题，我们使用C来处理它们，并从语言中获得尽可能多的好处。 在某些情况下可能需要嵌套循环做cudaMalloc，但我不认为这是其中之一。

以下是并行完成工作的代码：

 #include <stdio.h> #include <stdlib.h> // 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<<<gridSize,blockSize>>>(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<nz; i++) for (unsigned j=0; j<ny; j++) for (unsigned k=0; k<nx; k++) if (c[i][j][k] != (i+j+k)) { printf("Mismatch at x= %d, y= %d, z= %d Host= %d, Device = %d\n", i, j, k, (i+j+k), c[i][j][k]); return 1; } printf("Results check!\n"); free(c); cudaFree(d_c); cudaCheckErrors("cudaFree fail"); return 0; } 

既然你已经在评论中提出了要求，下面是我为了让你的代码运行起来所能做的最less的改变。 让我们也提醒一下你前面提到的一些问题，

“对于代码复杂性和性能的原因，你真的不希望这样做，在CUDA代码中使用指针数组比使用线性内存的替代方法更难和更慢。

“与使用线性内存相比，这是一个糟糕的主意。”

我不得不在纸上画出来，以确保我的指针复制正确。

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

简而言之，我们必须做一个连续的序列：

1. malloc指针的multidimensional array（在主机上），一维小于问题的大小，最后一个维度是指向区域的一组指针，而不是主机。
2. 创build与上一步创build的类相同的类的另一个多维指针数组，但是一个维度小于上一步创build的维度。 这个数组还必须在设备上有最终的等级cudaMalloc。
3. 将上一个第二步中的最后一组主机指针复制到上一步中设备上的cudaMalloced区域。
4. 重复步骤2-3，直到我们得到一个指向多维指针数组的单个（主机）指针，所有指针都驻留在设备上。