2016-05-09 4 views
1

radeon HD 7970 Tahiti XTでOpenCLを使用した2段階合計削減(このAMD linkから)に関するランタイムベンチマークを実行したいと考えています。OpenCLベンチマーク - 変更するパラメータについてのアドバイス

最初に、サイズの入力配列からサイズNworkItemsの出力配列への縮小を実行する最初のループを使用しないコードの最初のバージョンを使用しました。ここでカーネルのコードに、この最初のループは次のとおり

int global_index = get_global_id(0); 
    float accumulator = 0; 
    // Loop sequentially over chunks of input vector 
    while (global_index < length) { 
    float element = buffer[global_index]; 
    accumulator += element; 
    global_index += get_global_size(0); 
    } 

したがって、この最初のバージョンで、Iは、入力配列のサイズの関数として実行時に測定(スレッドの合計数に等しい)とワークの異なるサイズを持っていますグループ。ここでの結果です:

enter image description here

は今、私は上記のこの最初のループを使用してベンチマークを行いたいと思います。しかし、どのパラメータを変えなければならないのか分かりません。

this linkから、AMDはWorkGroupのサイズ(NVIDIAでは32)に64の倍数を推奨しています。

さらに、this other linkの最後のコメントから、WorkGroup size = (Number of total threads)/(Compute Units)のように作業グループのサイズを設定することをお勧めします。 私のGPUカードには、32個の計算ユニットがあります。

この2番目のバージョン(最初のリダクションループ)でランタイムを比較するために、どのパラメータを変更するのが面白いかを知るためのアドバイスを受けたいと思います。例えば、私は、(上記式を参照)比(N size of input array)/(total NworkItems)WorkGroup sizeための固定値に対して異なる値をとる

又は逆に行うことができ、すなわちIはWorkGroup sizeの値を変化させると、比(N size of input array)/(total NworkItems)を修正する必要がありますか?すべてのアイデアは大歓迎です

、おかげ

答えて

2

あなたは、ローカルデータの代わりに広がるデータを合計しなければならないメモリ転送を支援するためには、(データアクセスを合体しました)。したがって、代わりにこれを使用してください:

int chunk_size = length/get_global_size(0)+(length%get_global_size(0) > 0); //Will give how many items each work item needs to process 
    int global_index = get_group_id(0)*get_local_size(0)*chunk_size + get_local_id(0); //Start at this address for this work item 
    float accumulator = 0; 

    for(int i=0; i<chunk_size; i++) 
    // Loop sequentially over chunks of input vector 
    if (global_index < length) { 
     float element = buffer[global_index]; 
     accumulator += element; 
     global_index += get_local_size(0); 
    } 
    } 

また、キャッシングを助けるために、2のべき乗のサイズを使用する必要があります。

+0

ありがとうございました。私がうまくいけば、入力配列の最初の「長さ」と「get_local_size(0)」=「各ワークグループのサイズ」の2つのパラメータだけを変更する必要があります。だから私は私の最初のベンチマークの上の図を生成するために使用したのと同じパラメータを変更することができます。 – youpilat13

+0

あなたのコードの問題は、item0が 'global_size()'で区切られた 'buffer []'から散在したデータを読み込んでいることです。作業項目のサイズによっては、メモリパターンが非常に悪くなる可能性があります。各作業グループがバッファの小さな部分で動作すると、キャッシュがよりうまく機能し、グループ内のすべての作業項目でもキャッシュを共有できるようになると良いでしょう。 – DarkZeros

+0

を「ワークグループサイズ」に設定している場合は、私のベンチマークで「GPUカードのCompute Units = 32」と「Compute Units = 8」の「合計スレッド数=(ワークグループサイズ)*(Compute Units)私のCPU)? – youpilat13

関連する問題