logo标签列表

将3维数组发送到CUDA内核

ccudagpugpgpu
6

我使用一个回答中提供的代码来实现二维数组的加法，该回答来源于如何使用嵌套for循环相加两个二维（pitched）数组？ 现在我想将其应用于三维数组而非二维数组，并做了一些细微的修改，代码如下：

 __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个回答
11

首先，我认为talmonmies在回答你提到的上一个问题时，并没有打算让那个代码代表良好的编程。因此，想要将其扩展到三维可能不是你时间的最佳利用方式。例如，为什么我们要编写只使用一个线程的程序？虽然这种内核可能有合法的用途，但这不是其中之一。你的内核有可能进行大量独立的工作并行处理，但你却强制将它们全部放在一个线程上，从而使其串行化。并行工作的定义如下：

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;
}

由于您在评论中要求，这里是我能够使您的代码运行所做的最少更改数量。让我们也回想一下talonmies在您参考的上一个问题中的一些评论：

"出于代码复杂性和性能原因，您真的不希望这样做，在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. 在主机上分配一个多维指针数组，比问题规模少一维，最后一维是指向cudaMalloc分配在设备上的区域的指针集合。
  2. 创建另一个多维指针数组，与前一步所创建的类别相同，但比前一步少一维。该数组必须也将其最终排名cudaMalloc分配在设备上。
  3. 将第二步中最后一组主机指针复制到前一步中cudaMalloc在设备上的区域中。
  4. 重复执行2-3步，直到最终得到一个指向多维指针数组的单个（主机）指针，这些指针现在都驻留在设备上。
谢谢，请问您可以向我展示如何解决目前我正在做的这种方式吗？非常感谢！ - starter
你能提供一个完整的、可编译的示例来说明你想要做什么吗？这对于那些试图帮助你的人来说是很方便的。 - Robert Crovella
我知道这可能太多了，但是只使用h_c是否有可能完成所有操作而不需要h_c1部分？谢谢！ - starter
没发生过这样的事情。但是你可能会找到一种方法。 - Robert Crovella
@RobertCrovella，谢谢您！我曾考虑过在这里与主题发起者做类似的事情，现在我知道这只是错误的方式，以及如何以正确的方式做到它... - itsid
显示剩余2条评论
网页内容由stack overflow 提供, 点击上面的
可以查看英文原文，
原文链接