-1

カーネルが基本的に次のことを行う CUDA プログラムがあります。

  • デカルト座標 (x_i,y_i) の n 個の点のリストを次元 dim_x * dim_y の平面に提供します。それに応じてカーネルを呼び出します。
  • この平面 (x_p,y_p) 上のすべての点について、n 個の点のそれぞれがそこに到達するのにかかる時間を数式で計算します。これらの n 個の点が特定の速度で移動しているとします。
  • これらの時間を t_0、t_1、...t_n の昇順で並べます。t_i の精度は 1 に設定されます。つまり、t'_i=2.3453 の場合、t_i=2.3 のみを使用します。
  • 時間が正規分布から生成されると仮定して、最も速い 3 つの時間をシミュレートして、それらの 3 つのポイントに最も早く到達した時間の割合を見つけます。したがって、ランダムな実験により、prob_0 = 0.76、prob_1 = 0.20、prob_2 = 0.04 と仮定します。t_0 は 3 つの中で最初に到達するため、ポイントの元のインデックス (時間の並べ替え前) も返します。idx_0 = 5 (整数) とします。
  • したがって、この平面上のすべての点について、ペア (prob,idx) を取得します。

これらのポイントの n/2 が 1 つの種類で、残りが他のポイントであるとします。生成されたサンプル画像は次のようになります。

最適化されていない

特に、時間の精度が 1 に設定されている場合、時間の一意の 3 つのタプル (t_0、t_1、t_2) の数は、全データ ポイント、つまり平面上のポイントの数のわずか 2.5% であることに気付きました。これは、ほとんどの場合、カーネルが以前のシミュレーションの値を使用するだけで無駄にシミュレートしていたことを意味していました。したがって、キーを時間の 3 タプル、値をインデックスと確率として持つ辞書を使用できます。私が知る限り、カーネル内で STL にアクセスできないため、サイズ 201000000 の float の配列を作成しました。この選択は、上位 3 回のいずれも 20 秒を超えなかったため、実験によるものでした。したがって、t_0 は {0.0,0.1,0.2,...,20.0} から任意の値を取ることができるため、201 の選択肢があります。次のような辞書のキーを作成できます

  • キー = t_o * 10^6 + t_1 * 10^3 + t_2

値に関する限り、(prob+idx) のようにすることができます。idx は整数であり、0.0<=prob<=1.0 であるため、後でこれらの値の両方を取得できます。

  • prob=dict[key]-floor(dict[key])
  • idx = フロア (dict[キー])

だから今、私のカーネルは次のようになります

__global__ my_kernel(float* points,float* dict,float *p,float *i,size_t w,...){
  unsigned int col = blockIdx.y*blockDim.y + threadIdx.y;
  unsigned int row = blockIdx.x*blockDim.x + threadIdx.x;
  //Calculate time taken for each of the points to reach a particular point on the plane
  //Order the times in increasing order t_0,t_1,...,t_n
  //Calculate Key = t_o * 10^6 + t_1 * 10^3 + t_2
  if(dict[key]>0.0){
    prob=dict[key]-floor(dict[key])
    idx = floor(dict[key])
  }
  else{
    //Simulate and find prob and idx

    dict[key]=(prob+idx)
  }
  p[row*width+col]=prob;
  i[row*width+col]=idx;
}

結果は、ほとんどの点で元のプログラムと非常に似ていますが、一部は間違っています。

最適化されているが正しくない

これは競合状態によるものだと確信しています。dict がすべてゼロで初期化されていることに注意してください。基本的な考え方は、dict の特定の場所でデータ構造を「複数回書き込みを一度に読み取る」ようにすることです。

大量のメモリを割り当てるよりも、この問題を解決するためのより最適化された方法があるかもしれないことは承知しています。その場合はお知らせください。しかし、この特定のソリューションが失敗する理由を本当に理解したいと思います。特に、この設定でatomicAddを使用する方法を知りたいです。私はそれを使用するのに失敗しました。

4

1 に答える 1

0

ブランチでのシミュレーションelseが非常に長い (数百の浮動小数点演算) でない限り、グローバル メモリ内のルックアップ テーブルは、計算の実行よりも遅くなる可能性があります。グローバルメモリアクセスは非常に高価です!

いずれにせよ、条件分岐を利用して「作業をスキップ」して時間を節約する方法はありません。GPU の単一命令、複数スレッド アーキテクチャは、ブロック内のすべてのスレッドが同じ分岐をたどらない限り、分岐の両側の命令が連続して実行されることを意味します。

編集：

条件付き分岐を導入した結果、パフォーマンスが向上し、デッドロックの問題が発生していないという事実は、各ブロック内のすべてのスレッドが常に同じ分岐を取っていることを示唆しています。dict人口が増え始めると、パフォーマンスの向上はなくなると思います。

おそらく私は何かを誤解していますが、イベントの確率を計算したい場合x、正規分布を仮定し、平均muと標準偏差sigmaが与えられた場合、乱数の負荷を生成してガウス曲線を近似する必要はありません. 確率を直接計算できます。

p = exp(-((x - mu) * (x - mu) / (2.0f * sigma * sigma))) / 
    (sigma * sqrt(2.0f * M_PI));
于 2013-11-27T11:20:43.293 に答える