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

标签 c cuda statistics simulation

编辑
在最初发布的代码段（请参见下文）中，我没有正确地将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


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


加
除以期和期末


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

如果我对较小的值使用上述直接算法，则会得到正确的非零值，但是第二个我使用中间值（例如stepsA，sumA等），则值将变为零。
我知道我在这里做傻事...我想念什么？

笔记：
A.）是的，我知道上面形式的代码不是并行的，并且本身不保证并行化。它是更长代码的少量统计信息收集的一部分。在该代码中，将其封装在特定于线程索引的条件逻辑中，以防止发生冲突（使其并行），并用作将数据收集到模拟程序（保证并行化）中。希望您能理解以上代码的起源，并避免就其缺乏线程安全性发表sn贬不一的评论。 （此免责声明是根据过去的经验而获得的，这些经验来自那些不理解我是在摘录而不是完整代码的人提供的毫无用处的评论，尽管我这样写得不太明确。）
B.）是的，我知道变量的名称不明确。这就是我想说的。我正在处理的代码是专有的，尽管它最终将是开源的。我之所以写这篇文章，是因为我过去曾张贴过类似的匿名代码，并收到了有关我的命名约定的粗鲁评论。
C.）是的，尽管我确实犯了错误并且承认有些功能我不理解，但我已经读过几次CUDA manual了。我不是在这里使用共享内存，而是在完整代码中使用了共享内存（OF COURSE）。
D.）是的，上面的代码确实代表了与我的非工作代码的数据转储部分完全相同的功能，并且删除了与该特定问题无关的逻辑，并且使用了线程安全性作为条件。变量名称已更改，但是在算法上应保持不变，并且通过完全相同的非工作输出（零）进行验证。
E.）我确实意识到上述代码段中的“动态” struct具有非动态值。之所以命名该结构，是因为在完整的代码中，该struct包含模拟数据，并且是动态的。简化后的代码中的静态性质不应使统计信息收集代码失败，这仅意味着每个转储的平均值应为常数（且不为零）。

最佳答案

我在这里看到的最大问题是范围之一。编写此代码的方式使我得出一个结论，即您可能不了解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”，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/10356842/

上一篇：c - 体系结构 x86_64 的 undefined symbol "_min"

下一篇：cURL 获取响应

相关文章：

r - 在 R 中计算多个数据帧的 z 分数

ElasticSearch 中的随机文档

c - TAG 替换 C 字符串

关于C中指针递减的困惑

c - 如何在c中获取相对路径

c - 我编写了矩阵乘法 CUDA 程序，但与我的 CPU 结果相比，它总是得到错误的答案？

unit-testing - 是否可以模拟 GPU 以进行 CUDA/OpenCL 单元测试？

python - python 中 Weibull 分布的拟合优度检验

c - 尝试包含头文件时出现语法错误

cuda - cuda内核运行时访问cuda设备内存

©2024 IT工具网  联系我们