4

Intel CPU と NVIDIA GPU で同じ OpenCL カーネル コードを実行していますが、最初の結果は間違っていますが、後者の結果は正しくありません。奇妙なことは、一見無関係な変更を行うと、両方の場合で出力が期待どおりに機能することです。

この関数の目的は、A (三角形) と B (正則) の間の行列の乗算を計算することです。ここで、演算における A の位置は、変数 の値によって決まりますleft。このバグは、leftが true で、for ループが少なくとも 2 回繰り返される場合にのみ発生します。

以下は、わかりやすくするために、影響を与えてはならないいくつかのビットを省略したコードの一部です。

__kernel void blas_strmm(int left, int upper, int nota, int unit, int row, int dim, int m, int n,
                         float alpha, __global const float *a, __global const float *b, __global float *c) {

  /* [...] */
  int ty = get_local_id(1);
  int y = ty + BLOCK_SIZE * get_group_id(1);
  int by = y;
  __local float Bs[BLOCK_SIZE][BLOCK_SIZE];
  /* [...] */

  for(int i=start; i<end; i+=BLOCK_SIZE) {
    if(left) {
      ay = i+ty;
      bx = i+tx;
    }   
    else {
      ax = i+tx;
      by = i+ty;
    }   

    barrier(CLK_LOCAL_MEM_FENCE);
    /* [...] (Load As) */
    if(bx >= m || by >= n)
      Bs[tx][ty] = 0;
    else
      Bs[tx][ty] = b[bx*n+by];
    barrier(CLK_LOCAL_MEM_FENCE);

    /* [...] (Calculate Csub) */
  }

  if(y < n && x < (left ? row : m)) // In bounds
    c[x*n+y] = alpha*Csub;
}

今、それは奇妙になります。

ご覧のとおり、が true の場合はby常に等しいです。私はチェックしました（いくつかの s を使用してください）、常に true であり、ループ内の else ブランチのコードは実行されません。それでも、そこの行を削除またはコメントアウトすると、コードは機能します。なんで？まだわかりませんが、期待値が割り当てられていないことに関連している可能性があります。yleftprintfleftby = i+tyby

私の考えでは、 と の間に矛盾があるかどうかを確認する必要がありましたby。yこれらは常に同じ値を持つ必要があるためです。チェックする行を追加しましたby != yが、期待どおり、その比較は常に false を返しました。byそれで、私は続けてforの外観を変更しましyた

if(bx >= m || by >= n)

に変換

if(bx >= m || y >= n)

by3行下の変数をまだ適切に使用しているにもかかわらず、再び機能しました。

心を開いて他のことを試してみたところ、ループ内に次の行を追加すると、最初の if/else の後、if 条件の前の任意の位置にある限り、コードが機能するようになりました。先ほど申し上げました。

if(y >= n) left = 1;

( ) 内のコードleft = 1は、何にでも置き換えることができます ( a printf、別の役に立たない代入など) が、条件はもう少し制限的です。コードが正しい値を出力する例を次に示します。

if(y >= n) left = 1;
if(y < n) left = 1;
if(y+1 < n+1) left = 1;
if(n > y) left = 1;

動作しないものもありますがm = n、私がテストしている特定の例では次のことに注意してください。

if(y >= n+1) left = 1;
if(y > n) left = 1;
if(y >= m) left = 1;
/* etc. */

それが私が今いるポイントです。プログラムにまったく影響を与えないはずの行を追加しましたが、それは機能します。この魔法の解決策は私にとって満足のいくものではなく、CPU 内で何が起こっているのか、そしてその理由を知りたいのです。

私が何も忘れていないことを確認するために、ここに完全な機能コードと例の入力と出力の要点があります。

どうもありがとうございました。

解決

DarkZeros と Sharpneli の両方のユーザーの仮定は正しかった: for ループ内のバリアが適切な回数ヒットされていなかったのだ。特に、各ローカル グループの最初の要素に関連するバグがあり、残りの反復よりも 1 反復少なく実行され、未定義の動作が引き起こされました。後知恵で見ると痛々しいほど明白でした。

回答と時間をありがとうございました。

4

2 に答える 2

1

私が最初に見たのは、それがひどく失敗する可能性があることです.forループ内でバリアを使用しているということです.

すべてのスレッドが同じ回数 for ループに入らない場合。その場合、結果は完全に未定義です。そして、forループが複数回実行された場合にのみ問題が発生すると明確に述べています。

この状態を確保していますか？

于 2013-11-04T12:12:00.213 に答える