c - CUDA 代码的数据采集部分意外输出“0”

Question

编辑
在最初发布的代码片段（见下文）中，我没有正确发送struct到device，这已得到修复，但结果仍然相同。在我的完整代码中，这个错误不存在。（在我最初的帖子中，该命令有两个错误 - 一个，结构是从复制的HostToDevice，但实际上是颠倒的，并且副本的大小也错误。道歉；两个错误都已修复，但重新编译的代码仍然显示下面描述的零现象，我的完整代码也是如此。）

编辑 2
在我对代码的去私有化重写的匆忙中，我犯了几个错误，dalekchef好心地向我指出（设备的副本是struct在设备上分配之前执行的，在我重写的代码和设备中cudaMalloc调用没有与sizeof(...)数组元素的类型相乘。我添加了这些修复，重新编译和重新测试，但它没有解决问题。还仔细检查了我的原始代码 - 它没有那些错误。再次道歉，对于困惑。

我正在尝试从大型模拟程序中转储统计数据。下面显示了一个类似的精简代码。两种代码都表现出相同的问题——它们在应该输出平均值时输出零。

#include "stdio.h"

struct __align__(8) DynamicVals 
{ 
   double a;
   double b;
   int n1;
   int n2;
   int perDump;
};

__device__ int *dev_arrN1, *dev_arrN2;
__device__ double *dev_arrA, *dev_arrB;
__device__ DynamicVals *dev_myVals;
__device__ int stepsA, stepsB;
__device__ double sumA, sumB;
__device__ int stepsN1, stepsN2;
__device__ int sumN1, sumN2;

__global__ void TEST
(int step, double dev_arrA[], double dev_arrB[],
 int dev_arrN1[], int dev_arrN2[],DynamicVals *dev_myVals)
{
   if (step % dev_myVals->perDump)
   {
      dev_arrN1[step/dev_myVals->perDump] = 0;
      dev_arrN2[step/dev_myVals->perDump] = 0;
      dev_arrA[step/dev_myVals->perDump] = 0.0;
      dev_arrB[step/dev_myVals->perDump] = 0.0;
      stepsA = 0;
      stepsB = 0;
      stepsN1 = 0;
      stepsN2 = 0;
      sumA = 0.0;
      sumB = 0.0;
      sumN1 = 0;
      sumN2 = 0;
   }

   sumA += dev_myVals->a;
   sumB += dev_myVals->b;
   sumN1 += dev_myVals->n1;
   sumN2 += dev_myVals->n2;
   stepsA++;
   stepsB++;
   stepsN1++;
   stepsN2++;

   if ( sumA > 100000000 )
   {
      dev_arrA[step/dev_myVals->perDump] +=
     sumA / stepsA;
      sumA = 0.0;
      stepsA = 0;
   }
   if ( sumB > 100000000 )
   {
      dev_arrB[step/dev_myVals->perDump] +=
     sumB / stepsB;
      sumB = 0.0;
      stepsB = 0;
   }
   if ( sumN1 > 1000000 )
   {
      dev_arrN1[step/dev_myVals->perDump] +=
     sumN1 / stepsN1;
      sumN1 = 0;
      stepsN1 = 0;
   }
   if ( sumN2 > 1000000 )
   {
      dev_arrN2[step/dev_myVals->perDump] +=
     sumN2 / stepsN2;
      sumN2 = 0;
      stepsN2 = 0;
   }

   if ((step+1) % dev_myVals->perDump)
   {
      dev_arrA[step/dev_myVals->perDump] +=
     sumA / stepsA;
      dev_arrB[step/dev_myVals->perDump] +=
     sumB / stepsB;
      dev_arrN1[step/dev_myVals->perDump] +=
     sumN1 / stepsN1;
      dev_arrN2[step/dev_myVals->perDump] +=
     sumN2 / stepsN2;
   }
}

int main() 
{
   const int TOTAL_STEPS = 10000000;
   DynamicVals vals;
   int *arrN1, *arrN2;
   double *arrA, *arrB;
   int statCnt;

   vals.perDump = TOTAL_STEPS/10;
   statCnt = TOTAL_STEPS/vals.perDump+1;
   vals.a = 30000.0;
   vals.b = 60000.0;
   vals.n1 = 10000;
   vals.n2 = 20000;

   cudaMalloc( (void**)&dev_arrA, statCnt*sizeof(double) );
   cudaMalloc( (void**)&dev_arrB, statCnt*sizeof(double) );
   cudaMalloc( (void**)&dev_arrN1, statCnt*sizeof(int) );
   cudaMalloc( (void**)&dev_arrN2, statCnt*sizeof(int) );
   cudaMalloc( (void**)&dev_myVals, sizeof(DynamicVals));
   cudaMemcpy(dev_myVals, &vals, sizeof(DynamicVals), 
          cudaMemcpyHostToDevice);

   arrA = (double *)malloc(statCnt * sizeof(double));
   arrB = (double *)malloc(statCnt * sizeof(double));
   arrN1 = (int *)malloc(statCnt * sizeof(int));
   arrN2 = (int *)malloc(statCnt * sizeof(int));

   for (int i=0; i< TOTAL_STEPS; i++)
      TEST<<<1,1>>>(i, dev_arrA,dev_arrB,dev_arrN1,dev_arrN2,dev_myVals);

   cudaMemcpy(arrA,dev_arrA,statCnt * sizeof(double),cudaMemcpyDeviceToHost);
   cudaMemcpy(arrB,dev_arrB,statCnt * sizeof(double),cudaMemcpyDeviceToHost);
   cudaMemcpy(arrN1,dev_arrN1,statCnt * sizeof(int),cudaMemcpyDeviceToHost);
   cudaMemcpy(arrN2,dev_arrN2,statCnt * sizeof(int),cudaMemcpyDeviceToHost);

   for (int i=0; i< statCnt; i++)
   {
      printf("Step: %d   ; A=%g  B=%g  N1=%d  N2=%d\n",
         i*vals.perDump,
         arrA[i], arrB[i], arrN1[i], arrN2[i]);
   }
}

输出：

Step: 0   ; A=0  B=0  N1=0  N2=0
Step: 1000000   ; A=0  B=0  N1=0  N2=0
Step: 2000000   ; A=0  B=0  N1=0  N2=0
Step: 3000000   ; A=0  B=0  N1=0  N2=0
Step: 4000000   ; A=0  B=0  N1=0  N2=0
Step: 5000000   ; A=0  B=0  N1=0  N2=0
Step: 6000000   ; A=0  B=0  N1=0  N2=0
Step: 7000000   ; A=0  B=0  N1=0  N2=0
Step: 8000000   ; A=0  B=0  N1=0  N2=0
Step: 9000000   ; A=0  B=0  N1=0  N2=0
Step: 10000000   ; A=0  B=0  N1=0  N2=0

现在，如果我要使用一小段时间进行转储，或者如果我的 #s 更小，我可以直接使用

添加
除以期间和期间结束

...算法，但我使用临时总和，否则我int会溢出（double不会溢出，但我担心它会失去精度）。

如果我对较小的值使用上述直接算法，我会得到正确的非零值，但第二次我使用中间值（例如stepsA、sumA等），值会变为零。我知道我在这里做一些愚蠢的事情......我错过了什么？

注意：
A.) 是的，我知道上面形式的代码不是并行的，并且本身不保证并行化。它是较长代码的一小部分统计信息收集部分的一部分。在该代码中，它被封装在线程索引特定的条件逻辑中以防止冲突（使其并行）并用作模拟程序的数据收集（保证并行化）。希望您能理解上述代码的来源，并避免对其缺乏线程安全性的讽刺评论。（此免责声明是根据过去的经验添加的，这些人不理解我发布的是摘录，而不是完整的代码，尽管我以不那么明确的术语写作。）

B.) 是的，我知道变量的名称不明确。这就是我想说的。我正在处理的代码是专有的，尽管它最终将是开源的。我之所以写这个，是因为我过去曾发布过类似的匿名代码，并且收到了关于我的命名约定的粗鲁评论。

C.) 是的，我已经多次阅读CUDA 手册，尽管我确实犯了错误，并且我承认有些功能我不理解。我在这里没有使用共享内存，但我在我的完整代码中使用了共享内存（当然）。

D.) 是的，上面的代码确实代表了与我的非工作代码的数据转储部分完全相同的功能，删除了与此特定问题无关的逻辑，并带有线程安全条件。变量名称已更改，但从算法上讲，它应该保持不变，并且通过完全相同的非工作输出（零）来验证。

E.) 我确实意识到struct上述代码段中的“动态”具有非动态值。我将结构命名为，因为在完整代码中，它struct包含模拟数据，并且是动态的。精简代码中的静态性质不应该使统计信息收集代码失败，它只是意味着每个转储的平均值应该是恒定的（并且非零）。

score 1 · Accepted Answer

有几件事：

在为它调用 cudaMalloc 之前，您似乎在为 dev_MyVals 调用 cudaMemcpy。这不是应该的样子。

另外：当您进行 cudaMalloc 调用时，您不会乘以 sizeof int。

你真的应该检查你所有的 CUDA 调用 cudaMalloc/cudaMemcpy 是否有错误代码。它们都应该返回错误或 CUDA_SUCCESS。我相信 CUDA 示例都展示了如何做到这一点。

此外，为了将来参考永远不要在 CUDA 中使用模运算符，它非常慢。只需谷歌搜索“Modulo CUDA”即可获得一些替代方案。

让我知道它是怎么回事，这可能需要几次迭代才能修复。

score 0 · Accepted Answer

我在这里看到的最大问题是范围之一。这段代码的编写方式使我得出结论，您可能不了解 C++ 中的变量作用域一般是如何工作的，尤其是设备和主机代码作用域在 CUDA 中是如何工作的。几点观察：

当您在代码中执行此类操作时：
```
__device__ double *dev_arrA, *dev_arrB;
__global__ void TEST(int step, double dev_arrA[], double dev_arrB[], ....)
```
你有一个变量范围问题。dev_arrA在编译单元范围和函数范围都声明。这两个声明不引用同一个变量——函数单元范围声明（在内核中）优先于内核中的编译单元范围声明。您修改该变量，您正在修改内核范围声明，而不是__device__变量。这可能导致各种微妙和未明确的行为。最好避免在多个范围内声明相同的变量。
当您使用说明符声明变量时__device__，它旨在专门用作设备上下文符号，并且只能在设备代码中直接使用。所以是这样的：
```
__device__ double *dev_arrA;
int main()
{
....
cudaMalloc( (void**)&dev_arrA, statCnt*sizeof(double) );
....
}
```
是非法的。您不能像cudaMalloc直接在__device__变量上那样调用 API 函数。即使它会编译（因为主机和设备代码的 CUDA 编译轨迹中涉及黑客），但这样做是不正确的。在上面的例子dev_arrA中是一个设备符号。您可以通过 API 符号操作调用与它进行交互，但这在技术上是合法的。在您的代码中，旨在保存设备指针并作为内核参数（如dev_arrA）传递的变量应在main()范围内声明，并按值传递给内核。

它是上述两件事的结合，可能会导致您的问题。

但困难在于您选择发布大约 150 行代码（其中很多是多余的）作为重现案例。我怀疑是否有人足够关心您的问题，以至于无法用细齿梳子检查那么多代码并查明确切的问题所在。此外，你习惯在你的问题中进行这些讨厌的“顶级编辑”，很快就会将可能合理编写的起点变成难以理解的伪变更日志，这些变更日志非常难以理解，而且不太可能对任何人有帮助。此外，温和的被动攻击性注释部分没有真正的目的——它没有为问题增加任何价值。

因此，我将为您提供您发布的代码的一个大大简化的版本，我认为它包含您正在尝试做的所有基本工作。我把它作为“读者练习”，把它变成你想要做的任何事情。

#include "stdio.h"

typedef float Real;
struct __align__(8) DynamicVals 
{ 
    Real a;
    int n1;
    int perDump;
};

__device__ int stepsA;
__device__ Real sumA;
__device__ int stepsN1;
__device__ int sumN1;

__global__ void TEST
(int step, Real dev_arrA[], int dev_arrN1[], DynamicVals *dev_myVals)
{
    if (step % dev_myVals->perDump)
    {
        dev_arrN1[step/dev_myVals->perDump] = 0;
        dev_arrA[step/dev_myVals->perDump] = 0.0;
        stepsA = 0;
        stepsN1 = 0;
        sumA = 0.0;
        sumN1 = 0;
    }

    sumA += dev_myVals->a;
    sumN1 += dev_myVals->n1;
    stepsA++;
    stepsN1++;

    dev_arrA[step/dev_myVals->perDump] += sumA / stepsA;
    dev_arrN1[step/dev_myVals->perDump] += sumN1 / stepsN1;
}

inline void gpuAssert(cudaError_t code, char *file, int line, 
                 bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code),
          file, line);
      if (abort) exit(code);
   }
}

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

int main() 
{
    const int TOTAL_STEPS = 1000;
    DynamicVals vals;
    int *arrN1;
    Real *arrA;
    int statCnt;

    vals.perDump = TOTAL_STEPS/10;
    statCnt = TOTAL_STEPS/vals.perDump;
    vals.a = 30000.0;
    vals.n1 = 10000;

    Real *dev_arrA;
    int *dev_arrN1;
    DynamicVals *dev_myVals;

    gpuErrchk( cudaMalloc( (void**)&dev_arrA, statCnt*sizeof(Real)) );
    gpuErrchk( cudaMalloc( (void**)&dev_arrN1, statCnt*sizeof(int)) );
    gpuErrchk( cudaMalloc( (void**)&dev_myVals, sizeof(DynamicVals)) );
    gpuErrchk( cudaMemcpy(dev_myVals, &vals, sizeof(DynamicVals), 
                cudaMemcpyHostToDevice) );

    arrA = (Real *)malloc(statCnt * sizeof(Real));
    arrN1 = (int *)malloc(statCnt * sizeof(int));

    for (int i=0; i< TOTAL_STEPS; i++) {
        TEST<<<1,1>>>(i, dev_arrA,dev_arrN1,dev_myVals);
        gpuErrchk( cudaPeekAtLastError() );
    }

    gpuErrchk( cudaMemcpy(arrA,dev_arrA,statCnt * sizeof(Real),
                cudaMemcpyDeviceToHost) );
    gpuErrchk( cudaMemcpy(arrN1,dev_arrN1,statCnt * sizeof(int),
                cudaMemcpyDeviceToHost) );

    for (int i=0; i< statCnt; i++)
    {
        printf("Step: %d   ; A=%g N1=%d\n",
                i*vals.perDump, arrA[i], arrN1[i] );
    }
}

c - CUDA 代码的数据采集部分意外输出“0”

2 回答 2

Related

Reference