c - CUDA コードのデータ収集部分が予期せず「0」を出力する

Question

編集
最初の投稿のコードスニペット (以下を参照) では、をに適切に送信していませんでしstructたdevice。これは修正されましたが、結果は同じです。私の完全なコードでは、この間違いはありませんでした。(私の最初の投稿では、そのコマンドに 2 つの間違いがありました。1 つは、構造がからコピーされていましHostToDeviceたが、実際には反転されており、コピーのサイズも間違っていました。お詫びします。両方のエラーは修正されましたが、再コンパイルされたコードはまだ残っています。私の完全なコードと同様に、以下に説明するゼロ現象を表示します。)

EDIT 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ませんが、精度が失われることが心配でした）。

小さい値に対して上記の直接アルゴリズムを使用すると、ゼロ以外の正しい値が得られますが、2 番目に中間値 ( 、など) を使用すると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」をGoogleで検索してください。

それがどうなるか教えてください。これを修正するには、おそらく数回の反復が必要です。

score 0 · Accepted Answer

ここで私が目にする最大の問題は、スコープの 1 つです。このコードの書き方から、C++ での変数のスコープが一般的にどのように機能するか、特に CUDA でデバイスとホストコードのスコープがどのように機能するかを理解していない可能性があると結論付けられます。いくつかの観察：

この種のことをコードで行う場合:
```
__device__ double *dev_arrA, *dev_arrB;
__global__ void TEST(int step, double dev_arrA[], double dev_arrB[], ....)
```
可変スコープの問題があります。dev_arrAコンパイル単位スコープと関数スコープの両方で宣言されています。2 つの宣言は同じ変数を参照していません。(カーネル内の) 関数ユニットスコープ宣言は、カーネル内のコンパイルユニットスコープ宣言よりも優先されます。その変数を変更すると、変数ではなくカーネルスコープ宣言が変更され__device__ます。これにより、あらゆる種類の微妙で予期しない動作が発生する可能性があります。複数のスコープで同じ変数を宣言することは避けたほうがよいでしょう。
指定子を使用して変数を宣言する場合、その変数はデバイスコンテキストシンボル__device__としてのみ使用されることを意図しており、デバイスコードで直接使用する必要があります。だから、このようなもの：
```
__device__ double *dev_arrA;
int main()
{
....
cudaMalloc( (void**)&dev_arrA, statCnt*sizeof(double) );
....
}
```
違法です。変数cudaMallocに対して直接API 関数を呼び出すことはできません。__device__コンパイルはできますが (ホストおよびデバイスコードの CUDA コンパイルの軌跡にハッカーが関与しているため)、そうするのは正しくありません。上記の例dev_arrAでは、デバイスシンボルです。API シンボル操作呼び出しを介して対話することができますが、それは技術的に合法なすべてです。コードでは、デバイスポインターを保持し、カーネル引数として渡されることを意図した変数 ( など) をスコープでdev_arrA宣言し、カーネルに値で渡す必要があります。main()

おそらく問題を引き起こしているのは、上記の2つの組み合わせです。

しかし、難点は、再現ケースとして大まかな 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