c++ - Cuda "invalid argument"二维数组 - 元胞自动机

Question

我正在尝试使用 Cuda 计算二维元胞自动机重新分布。我对此完全陌生，所以我不知道我做错了什么。我已经尝试了许多我在这里看到的解决方案，但当我调用内核时，所有解决方案都给出了“无效参数”。

这是内核的简化版本:

//kernel definition
__global__ void stepCalc(float B[51][51], int L, int flag, float m, float en)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    int j = blockDim.y * blockIdx.y + threadIdx.y;

    float g=B[i][j]-0.25*(B[i+1][j]+B[i-1][j]+B[i][j+1]+B[i][j-1]);
    flag = 0;

    if (i < L-2 && j < L-2 && i>2 && j>2 && abs(g)>m)
    {
          flag = 1;
          en+=-16*g*g+8*B[i][j]*abs(g);
          B[i][j]+=-4*f*g;
          B[i+1][j]+=f*g;
          B[i-1][j]+=f*g;
          B[i][j+1]+=f*g;
          B[i][j-1]+=f*g;
    }
}

主要功能如下:

#define L 50
float B[L+1][L+1];
//initialize B[i][j]

float g=0;
int flag = 1;
float m=0.1;
float en = 0;
while (flag==1)
{
    float (*dB)[L+1];
    int *dFlag=NULL;
    float *dEn=NULL;

    cudaMalloc((void **)&dFlag,sizeof(int));
    cudaMalloc((void **)&dEn,sizeof(float));
    cudaMalloc((void **)&dB, ((L+1)*(L+1))*sizeof(float));

    cudaMemcpy(dB, B, sizeB, cudaMemcpyHostToDevice);
    cudaMemcpy(dFlag, &flag, sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(dEn, &en, sizeof(float), cudaMemcpyDeviceToHost);

    dim3 threadsPerBlock(16,16);
    dim3 numBlocks((L+1)/threadsPerBlock.x,(L+1)/threadsPerBlock.y);

    stepCalc<<<numBlocks, threadsPerBlock>>>(dB, L, dflag, m, dEn);
    GPUerrchk(cudaPeekAtLastError());  //gives "invalid argument" at this line

    cudaMemcpy(B, (dB), sizeB, cudaMemcpyDeviceToHost);
    cudaMemcpy(&flag, dFlag, sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(&en, dEn, sizeof(float), cudaMemcpyDeviceToHost);

    cudaFree(dB);
    cudaFree(dFlag);
    cudaFree(dEn);
}

我需要提取所有线程的新数组 B、标志值和总和“en”。我是否已经接近解决方案的样子了？有可能吗？我还尝试将主机数组 B 设置为 float** B，但没有成功。

Answer 1

您的代码存在各种问题。

1. 您可能忽略了向内核传递值和传递指针之间的区别:

   __global__ void stepCalc(float B[51][51], int L, int flag, float m, float en)
                                  ^                      ^
                                  |                      |
                                a pointer              a value
   

   我们稍后会回到 B，但对于像 flag 和 en 这样的值，请按值传递这些 到内核与按值传递到 C 函数具有类似的含义。这是一种单向通信路径。由于从您的代码中可以明显看出您希望稍后在主机代码中使用内核修改的这些值，因此您将需要传递指针。在某些情况下，您已经为此目的分配了指针，因此您会遇到另一种类型的错误，因为在某些情况下(dFlag)您传递的是指针，而内核定义需要一个值。

2. 关于 B，由于深复制问题，将 2D 数组从主机传递到设备可能比您最初预期的更加困难。在此处不涵盖所有内容的情况下，搜索本页右上角的“CUDA 2D array”，您将获得大量有关它的信息以及处理它的各种方法。由于您似乎愿意考虑固定宽度的数组(在编译时已知)，因此我们可以通过利用编译器帮助我们处理特定的 typedef 来简化 2D 数组的处理。

3. 当您在使用 cuda 代码时遇到问题时，最好在整个代码中(而不是仅在一个地方)进行严格的 CUDA 错误检查。原因之一是在特定位置发生的 CUDA 错误通常会在代码中的任何后续位置返回。如果您不检查每个 CUDA API 调用，这会让人感到困惑，因为特定的“无效参数”错误可能不是由内核本身引起的，而是由之前发生的某些 API 调用引起的。

4. 您通常不希望在数据处理 while 循环中进行 cudaMalloc 操作。这些通常是您在代码开头执行一次的操作。在 while 循环的每次迭代中执行 cudaMalloc 会带来几个负面问题，其中之一是您将耗尽内存(尽管您有 cudaFree 语句，所以也许不是)，最终，您在每次迭代中实际上都会丢弃数据。此外，它还会对您的表现产生负面影响。

5. 您的一些 cudaMemcpy 传输方向错误，如下所示:

   cudaMemcpy(dFlag, &flag, sizeof(int), cudaMemcpyDeviceToHost);
   

6. 在内核代码中将 flag 设置为零将会出现问题。 Warp 可以按任意顺序执行，并且在某些 Warp 稍后在内核中将 flag 设置为 1 后，其他 warp 可以开始执行并再次将 flag 设置为零。这可能不是您想要的。一种可能的修复方法是在执行内核之前将 flag 设置为零(即在主机代码中，并将其复制到设备)。

7. 您的内核将在此处生成越界索引:

   float g=B[i][j]-0.25*(B[i+1][j]+B[i-1][j]+B[i][j+1]+B[i][j-1]);
   

   (问问自己当 i=0 且 j=0 时会发生什么)。解决这个问题的方法是将这行代码移到 if-check 中，以便在其之后进行边界检查。

8. 你的内核使用了一个变量f，它的定义在我看不到的地方，例如这里:

     B[i+1][j]+=f*g;
   

以下代码是我尝试修改您的代码、创建一个完整的示例并消除上述问题。它没有做任何有用的事情，但它编译没有错误，运行对我来说也没有错误。我没有提供任何数据，所以目前这只是一个概念验证。我确信它仍然包含数据处理错误。

#include <stdio.h>

#define my_L 50

typedef float farray[my_L+1];

#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"); \
            exit(1); \
        } \
    } while (0)

//kernel definition
__global__ void stepCalc(farray B[], int L, int *flag, float m, float *en)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    int j = blockDim.y * blockIdx.y + threadIdx.y;

    //float g=B[i][j]-0.25*(B[i+1][j]+B[i-1][j]+B[i][j+1]+B[i][j-1]);
    // flag = 0;
    float f = 1.0f;
    if (i < L-2 && j < L-2 && i>2 && j>2){
      float g=B[i][j]-0.25*(B[i+1][j]+B[i-1][j]+B[i][j+1]+B[i][j-1]);
      if (abs(g)>m)
        {
          *flag = 1;
          *en+=-16*g*g+8*B[i][j]*abs(g);
          B[i][j]+=-4*f*g;
          B[i+1][j]+=f*g;
          B[i-1][j]+=f*g;
          B[i][j+1]+=f*g;
          B[i][j-1]+=f*g;
        }
      }
}


int main(){

  farray B[my_L+1];
//initialize B[i][j]

  farray *dB;
  int flag = 1;
  float m=0.1;
  float en = 0;
  int *dFlag=NULL;
  float *dEn=NULL;

  cudaMalloc((void **)&dFlag,sizeof(int));
  cudaCheckErrors("1");
  cudaMalloc((void **)&dEn,sizeof(float));
  cudaCheckErrors("2");
  size_t sizeB = (my_L+1)*sizeof(farray);
  cudaMalloc((void **)&dB, sizeB);
  cudaCheckErrors("3");
  cudaMemcpy(dB, B, sizeB, cudaMemcpyHostToDevice);
  cudaCheckErrors("4");
  cudaMemcpy(dEn, &en, sizeof(float), cudaMemcpyHostToDevice);
  cudaCheckErrors("5");

  dim3 threadsPerBlock(16,16);
  dim3 numBlocks((my_L+1)/threadsPerBlock.x,(my_L+1)/threadsPerBlock.y);
  while (flag==1)
  {
    flag = 0;
    cudaMemcpy(dFlag, &flag, sizeof(int), cudaMemcpyHostToDevice);
    cudaCheckErrors("6");
    stepCalc<<<numBlocks, threadsPerBlock>>>(dB, my_L, dFlag, m, dEn);
    cudaDeviceSynchronize();
    cudaCheckErrors("7");
    cudaMemcpy(&flag, dFlag, sizeof(int), cudaMemcpyDeviceToHost);
    cudaCheckErrors("8");
  }
  cudaMemcpy(B, (dB), sizeB, cudaMemcpyDeviceToHost);
  cudaCheckErrors("9");
  cudaMemcpy(&en, dEn, sizeof(float), cudaMemcpyDeviceToHost);
  cudaCheckErrors("10");
// process B
  cudaFree(dB);
  cudaFree(dFlag);
  cudaFree(dEn);
}