6

私は OpenCL プログラミングの初心者です。私のアプリのために。(分子シミュレーション) レナード・ジョーンズ液体の分子間ポテンシャルを計算するためのカーネルを書きました。このカーネルでは、すべての粒子のポテンシャルの累積値を 1 で計算する必要があります。

__kernel void Molsim(__global const float* inmatrix, __global float* fi, const int c, const float r1, const float r2, const float r3, const float rc, const float epsilon, const float sigma, const float h1, const float h23)
{
   float fi0;
   float fi1;
   float d;

   unsigned int i = get_global_id(0); //number of particles (typically 2000)

   if(c!=i) {
      // potential before particle movement
      d=sqrt(pow((0.5*h1-fabs(0.5*h1-fabs(inmatrix[c*3]-inmatrix[i*3]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(inmatrix[c*3+1]-inmatrix[i*3+1]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(inmatrix[c*3+2]-inmatrix[i*3+2]))),2.0));
      if(d<rc) {
        fi0=4.0*epsilon*(pow(sigma/d,12.0)-pow(sigma/d,6.0));
      }
      else {
        fi0=0;
      }
      // potential after particle movement
      d=sqrt(pow((0.5*h1-fabs(0.5*h1-fabs(r1-inmatrix[i*3]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(r2-inmatrix[i*3+1]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(r3-inmatrix[i*3+2]))),2.0));
      if(d<rc) {
        fi1=4.0*epsilon*(pow(sigma/d,12.0)-pow(sigma/d,6.0));
      }
        else {
          fi1=0;
        }
      // cumulative difference of potentials
      // fi[0]+=fi1-fi0; changed to full size vector
      fi[get_global_id(0)]=fi1-fi0;
      }
}

私の問題は次の行にあります: fi[0]+=fi1-fi0;. 1 要素ベクトル fi[0] では、間違った結果になります。総和の縮約について読んだことがありますが、計算中にそれを行う方法がわかりません。

私の問題の簡単な解決策はありますか?

注意: ベクトル コンポーネントの合計のために次のカーネルを追加しようとしましたが (以下のコードを参照)、CPU を使用してベクトルを合計する場合よりもさらに遅くなりました。

__kernel void Arrsum(__global const float* inmatrix, __global float* outsum, const int inmatrixsize, __local float* resultScratch)
{
       // načtení indexu
      int gid = get_global_id(0);
      int wid = get_local_id(0);
      int wsize = get_local_size(0);
      int grid = get_group_id(0);
      int grcount = get_num_groups(0);

      int i;
      int workAmount = inmatrixsize/grcount;
      int startOffest = workAmount * grid + wid;
      int maxOffest = workAmount * (grid + 1);
      if(maxOffest > inmatrixsize){
        maxOffest = inmatrixsize;
    }

    resultScratch[wid] = 0.0;
    for(i=startOffest;i<maxOffest;i+=wsize){
            resultScratch[wid] += inmatrix[i];
    }
    barrier(CLK_LOCAL_MEM_FENCE);

    if(gid == 0){
            for(i=1;i<wsize;i++){
                    resultScratch[0] += resultScratch[i];
            }
            outsum[grid] = resultScratch[0];
    }
}
4

3 に答える 3

2

fi[0]+=fi1-fi0; のatomic_addアトミック関数が必要だと思います。

警告: アトミック関数を使用すると、パフォーマンスが低下します。

ここでは、インクリメント アトミック関数を使用した 2 つの例を示します。

アトミック関数と 2 つの作業項目のない例:

__kernel void inc(global int * num){
    num[0]++; //num[0] = 0
}
  1. 作業項目 1 は num[0]: 0 を読み取ります
  2. 作業項目 2 は num[0]: 0 を読み取ります
  3. ワークアイテム 1 インクリメント num[0]: 0 + 1
  4. ワークアイテム 2 インクリメント num[0]: 0 + 1
  5. 作業項目 1 書き込み num[0]: num[0] = 1
  6. 作業項目 2 書き込み num[0]: num[0] = 1

結果:数値[0] = 1

アトミック関数と 2 つのワークアイテムの例:

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable

__kernel void inc(global int * num){
    atom_inc(&num[0]);
}
  1. 作業項目 1 は num[0]: 0 を読み取ります
  2. ワークアイテム 1 インクリメント num[0]: 0 + 1
  3. 作業項目 1 書き込み num[0]: num[0] = 1
  4. 作業項目 2 は num[0]: 1 を読み取ります
  5. ワークアイテム 2 インクリメント num[0]: 1 + 1
  6. 作業項目 2 書き込み num[0]: num[0] = 2

結果: num[0] = 2

于 2012-10-05T07:32:26.420 に答える
0

アトミックな追加は 1 つのソリューションですが、アトミックな部分が作業項目をシリアル化するため、パフォーマンスの問題が発生する可能性があります。

より良い解決策は、すべての作業項目について、次のように独自の変数に書き込むことだと思います。

fi[get_global_id(0)] +=fi1-fi0;

次に、配列を CPU に転送してすべての要素を合計するか、並列に実行するアルゴリズムを使用して GPU で実行します。

于 2012-10-05T07:40:24.900 に答える
0

すべてのスレッドは「グループ」によって実行されます。get_local_id(dim) 関数を使用して、グループ内のスレッド ID を特定できます。各グループ内のスレッドは、共有メモリ (OpenCL では「ローカル メモリ」と呼ばれます) を使用して実行を同期できますが、異なるグループ内のスレッドは直接通信できません。

したがって、削減の典型的な解決策は次のとおりです。

  1. 一時配列 part_sum (グローバル) と tmp_reduce (ローカル) をカーネル引数に追加します。

    __kernel void Molsim(..., __global float *part_sum, __local float *tmp_reduce)

  2. カーネルのグループ数 (=global_size/local_size) に等しいサイズの float の配列を割り当て、パラメータ part_sum を設定します。

  3. パラメータ tmp_reduce をカーネルの「ローカル サイズ」x size_of(float) および NULL に設定します。

    clSetKernelArg(kernel,<par number>,sizeof(float)*<local_size>,NULL);

  4. カーネルで、コードを次のように置き換えます。

      int loc_id=get_local_id(0);

...

//      fi[0]+=fi1-fi0;
       tmp_reduce[loc_id]=fi1-fi0;
      }
   barrier(CLK_LOCAL_MEM_FENCE);
   if(loc_id==0) {
     int i;
     float s=tmp_reduce[0];
     for(i=1;i<get_local_size(0);i++)
       s+=tmp_reduce[i];
     part_sum[get_group_id(0)]=s;
   }
}

  5. カーネルの実行が終了したら、ホスト上で part_sum[array] の内容を合計します。これは、global_size よりもはるかに小さいです。

より複雑なアルゴリズムを使用して Log2(local_size) 操作を使用して tmp_reduce 配列を並列に合計できるため、これは完全な「並列削減」ではありませんが、これはアトミック操作よりもはるかに高速でなければなりません。

また、より良い並列削減方法については、 http: //developer.amd.com/Resources/documentation/articles/pages/OpenCL-Optimization-Case-Study-Simple-Reductions_2.aspx を参照してください。

于 2012-10-06T09:18:56.853 に答える