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

Question

编辑
在最初发布的代码 sn-p（见下文）中，我没有正确地将 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 更小，我可以直接使用

1. 添加
2. 除以期和期末

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

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

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

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

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

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

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

Answer 1

有几点：

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

另外：在进行 cudaMalloc 调用时，不要乘以 sizeof int。

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

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

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

Answer 2

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

1. 当您在代码中执行此类操作时： <pre>__device__ double *dev_arrA, *dev_arrB; __global__ void TEST(int step, double dev_arrA[], double dev_arrB[], ....)

   你有一个变量范围问题。 dev_arrA 在编译单元范围和函数范围都声明。这两个声明不引用同一个变量——函数单元范围声明（在内核中）优先于内核中的编译单元范围声明。您修改该变量，您正在修改内核范围声明，而不是__device__variable。这可能导致各种微妙和未明确的行为。最好避免永远在多个范围内声明相同的变量。

2. 当您使用__device__ 说明符声明变量时，它旨在专门用作设备上下文符号，并且只能在设备代码中直接使用。所以是这样的： <pre>__device__ double *dev_arrA; int main() { .... cudaMalloc( (void**)&dev_arrA, statCnt*sizeof(double) ); .... }

   是非法的。您不能直接在 __device__ 变量上调用像 cudaMalloc 这样的 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] );
    }
}