将 3d 数组发送到 CUDA 内核
Posted
技术标签:
【中文标题】将 3d 数组发送到 CUDA 内核【英文标题】：sending 3d array to CUDA kernel 【发布时间】：2012-10-07 02:39:58 【问题描述】：我将给出的代码作为How can I add up two 2d (pitched) arrays using nested for loops? 的答案，并尝试将其用于 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 数组的内核调用后遇到问题。你能帮我解决这个问题吗？请您按照我上面写的方式显示解决方案。谢谢！
【问题讨论】：
【参考方案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;
既然您已在 cmets 中要求它，那么我可以对您的代码进行最少的更改以使其正常工作。让我们也提醒自己您提到的上一个问题中的一些 talonmies cmets：
“出于代码复杂性和性能原因，您真的不想这样做，在 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]);
简而言之，我们必须做一个连续的序列：
-
malloc 一个多维指针数组（在主机上），比问题大小小一维，最后一维是一组指向 cudaMalloc'ed 到设备而不是主机上的区域的指针。
创建另一个多维指针数组，与上一步中创建的类相同，但比上一步中创建的少一维。此数组还必须在设备上具有最终排名 cudaMalloc。
将上一步中的最后一组主机指针复制到上一步设备上的 cudaMalloced 区域中。
重复步骤 2-3，直到我们最终得到一个指向多维指针数组的（主机）指针，所有这些指针现在都驻留在设备上。
【讨论】：
谢谢，请您告诉我如何以我目前的方式解决它。非常感谢！ 您能否提供一个完整的、可编译的示例来说明您正在尝试做什么？对于那些试图帮助你的人来说，这是一个方便的问题。 对不起，如果我没有正确询问。上面的例子是完整的。如果我要改变任何东西，那将是“h_c”的大小，它可能会变成“int****”，其中每个子数组都将具有非常大且不同的大小。谢谢！如果我问错了，再次抱歉。非常感谢您的帮助。 您提供的示例不完整或无法编译。例如，一个原因是您没有在任何地方定义 GPUerrchk。是的，我可以转到上一个问题，并通过将该示例的一些部分与您的结合来拼凑您的意图，但这很不方便，特别是因为您可能拥有完整的代码并正在编译它。无论如何，我已经用回复更新了我的答案，向您展示如何以您目前的方式解决它。 我知道这可能太多了，但是是否可以只用 h_c 做所有事情而没有 h_c1 的部分？谢谢！！！以上是关于将 3d 数组发送到 CUDA 内核的主要内容，如果未能解决你的问题，请参考以下文章