5

OpenCLを使用してn次元の点間のユークリッド距離を計算しています。n次元ポイントの2つのリストを取得し、最初のテーブルのすべてのポイントから2番目のテーブルのすべてのポイントまでの距離だけを含む配列を返す必要があります。

私のアプローチは、通常の二重ループを実行することです(Table1のすべてのポイント{Table2のすべてのポイント{...}}次に、並列のポイントのすべてのペアに対して計算を実行します。

次に、ユークリッド距離は3つの部分に分割されます。1。ポイントの各次元間の差を取ります2.その差を2乗します(まだすべての次元に対して)3。2で得られたすべての値を合計します。4。の平方根を取ります3.で取得した値(この例では、このステップは省略されています)。

すべての違いの合計を累積しようとするまで、すべてが魅力のように機能します(つまり、上記の手順のステップ3、以下のコードの49行目を実行します)。

テストデータとして、それぞれ2ポイントのDescriptorListsを使用しています。DescriptorList1:001,002,003、...、127,128; (p1)129,130​​,131、...、255,256; (p2)

DescriptorList2:000,001,002、...、126,127; (p1)128,129,130​​、...、254,255; (p2)

したがって、結果のベクトルの値は128、2064512、2130048、128になります。現在、実行ごとに変化する乱数を取得しています。

私が間違っていることについての助けやリードに感謝します。うまくいけば、私が取り組んでいるシナリオについてすべてが明確になっています。

#define BLOCK_SIZE 128

typedef struct
{
    //How large each point is
    int length;
    //How many points in every list
    int num_elements;
    //Pointer to the elements of the descriptor (stored as a raw array)
    __global float *elements;
} DescriptorList;

__kernel void CompareDescriptors_deb(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float As[BLOCK_SIZE])
{

    int gpidA = get_global_id(0);

    int featA = get_local_id(0);

    //temporary array  to store the difference between each dimension of 2 points
    float dif_acum[BLOCK_SIZE];

    //counter to track the iterations of the inner loop
    int loop = 0;

    //loop over all descriptors in A
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){

        //take the i-th descriptor. Returns a DescriptorList with just the i-th
        //descriptor in DescriptorList A
        DescriptorList tmpA = GetDescriptor(A, i);

        //copy the current descriptor to local memory.
        //returns one element of the only descriptor in DescriptorList tmpA
        //and index featA
        As[featA] = GetElement(tmpA, 0, featA);
        //wait for all the threads to finish copying before continuing
        barrier(CLK_LOCAL_MEM_FENCE);

        //loop over all the descriptors in B
        for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){
            //take the difference of both current points
            dif_acum[featA] = As[featA]-B.elements[k*BLOCK_SIZE + featA];
            //wait again
            barrier(CLK_LOCAL_MEM_FENCE);
            //square value of the difference in dif_acum and store in C
            //which is where the results should be stored at the end.
            C[loop] = 0;
            C[loop] += dif_acum[featA]*dif_acum[featA];
            loop += 1;
            barrier(CLK_LOCAL_MEM_FENCE);
        }
    }
}
4

2 に答える 2

7

問題は次のコード行にあります。

C[loop] = 0;
C[loop] += dif_acum[featA]*dif_acum[featA];

ワークグループ内のすべてのスレッド(実際にはすべてのスレッドですが、後で説明します)は、同期をまったく行わずに、この配列の位置を同時に変更しようとしています。いくつかの要因がこれを本当に問題にします:

  1. ワークグループが完全に並行して動作することは保証されていません。つまり、一部のスレッドでは、他のスレッドが次の行を実行した後にC [loop]=0を呼び出すことができます。
  2. 並行して実行されるものはすべて、C [loop]から同じ値を読み取り、増分で変更して、同じアドレスに書き戻そうとします。その書き戻しの結果が何であるかは完全にはわかりませんが(スレッドの1つは書き戻しに成功し、他のスレッドは失敗すると思いますが、完全にはわかりません)、どちらにしても間違っています。

これを修正しましょう:アトミックを使用してグローバルメモリでこれを機能させることができるかもしれませんが、高速ではないので、ローカルメモリに蓄積しましょう:

local float* accum;
...
accum[featA] = dif_acum[featA]*dif_acum[featA];
barrier(CLK_LOCAL_MEM_FENCE);
for(unsigned int i = 1; i < BLOCKSIZE; i *= 2)
{
    if ((featA % (2*i)) == 0)
        accum[featA] += accum[featA + i];
    barrier(CLK_LOCAL_MEM_FENCE);
}
if(featA == 0)
    C[loop] = accum[0];

もちろん、これには他のローカルバッファを再利用できますが、要点は明らかだと思います(ところで、dif_acumはローカルメモリに作成されると思います。これはローカルメモリに配置されないことをどこかで読んだと思いますが、これにより、Aをローカルメモリにプリロードすることは無意味になります)。

このコードに関するその他のポイント:

  1. あなたのコードはワークグループでのみ使用するように調整されているようです(どのアイテムを処理するかを確認するためにgroupidもglobal idも使用していません)。最適なパフォーマンスを得るには、それ以上を使用することをお勧めします。
  2. 個人的な好みかもしれませんが、私にget_local_size(0)は、Defineを使用するよりもworkgroupsizeに使用する方が良いようです(openclコードをに変更する必要があることに気付かずにホストコードで変更する可能性があるため)
  3. 別のスレッドによって書き込まれたローカルメモリ内の要素にアクセスするスレッドはないため、コード内のバリアはすべて不要です。したがって、これにローカルメモリを使用する必要はありません。

あなたが簡単にできる最後の弾丸を考えると:

float As = GetElement(tmpA, 0, featA);
...
float dif_acum = As-B.elements[k*BLOCK_SIZE + featA];

これにより、コードが作成されます(最初の2つの箇条書きは考慮されません)。

__kernel void CompareDescriptors_deb(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float accum[BLOCK_SIZE])
{
   int gpidA = get_global_id(0);
   int featA = get_local_id(0);
   int loop = 0;
   for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){
       DescriptorList tmpA = GetDescriptor(A, i);
       float As = GetElement(tmpA, 0, featA);
       for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){
           float dif_acum = As-B.elements[k*BLOCK_SIZE + featA];

           accum[featA] = dif_acum[featA]*dif_acum[featA];
           barrier(CLK_LOCAL_MEM_FENCE);
           for(unsigned int i = 1; i < BLOCKSIZE; i *= 2)
           {
              if ((featA % (2*i)) == 0)
                 accum[featA] += accum[featA + i];
              barrier(CLK_LOCAL_MEM_FENCE);
           }
           if(featA == 0)
              C[loop] = accum[0];
           barrier(CLK_LOCAL_MEM_FENCE);

           loop += 1;
        }
    }
}
于 2010-09-23T02:43:31.147 に答える
3

Grizzlyのおかげで、カーネルが機能するようになりました。Grizzlyの回答に基づいて変更する必要があるいくつかのこと:

ルーチンの先頭にIFステートメントを追加して、使用している配列内の有効な位置を参照しないすべてのスレッドを破棄しました。

if(featA > BLOCK_SIZE){return;}

最初の記述子をローカル(共有)メモリ(igからBs)にコピーする場合、関数GetElementは呼び出しごとに1つの要素のみを返すため、インデックスを指定する必要があります(質問ではスキップしました)。

Bs[featA] = GetElement(tmpA, 0, featA);

次に、各反復後にバッファが上書きされ、どのスレッドが最初にデータにアクセスするかを制御できないため、SCANループを少し調整する必要がありました。そのため、dif_acumバッファーを「リサイクル」して部分的な結果を保存し、そのようにして、そのループ全体での不整合を防ぎます。

dif_acum[featA] = accum[featA];

SCANループには、一緒に追加する用語を確実に決定するための境界制御もあります。

if (featA >= j && next_addend >= 0 && next_addend < BLOCK_SIZE){

最後に、ループ変数のインクリメントを最後のIFステートメントに含めて、1つのスレッドだけがそれを変更するのが理にかなっていると思いました。

if(featA == 0){
    C[loop] = accum[BLOCK_SIZE-1];
    loop += 1;
}

それでおしまい。group_sizeを使用してそのBLOCK_SIZE定義を削除するにはどうすればよいのでしょうか。また、スレッドの使用に関してより適切なポリシーを採用できるかどうかは疑問です。

したがって、コードは最終的に次のようになります。

__kernel void CompareDescriptors(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float accum[BLOCK_SIZE], __local float Bs[BLOCK_SIZE])
{

    int gpidA = get_global_id(0);
    int featA = get_local_id(0);

    //global counter to store final differences
    int loop = 0;

    //auxiliary buffer to store temporary data
    local float dif_acum[BLOCK_SIZE];

    //discard the threads that are not going to be used.
    if(featA > BLOCK_SIZE){
        return;
    }

    //loop over all descriptors in A
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){

        //take the gpidA-th descriptor
        DescriptorList tmpA = GetDescriptor(A, i);

        //copy the current descriptor to local memory
        Bs[featA] = GetElement(tmpA, 0, featA);

        //loop over all the descriptors in B
        for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){
            //take the difference of both current descriptors
            dif_acum[featA] = Bs[featA]-B.elements[k*BLOCK_SIZE + featA];

            //square the values in dif_acum
            accum[featA] = dif_acum[featA]*dif_acum[featA];
            barrier(CLK_LOCAL_MEM_FENCE);

            //copy the values of accum to keep consistency once the scan procedure starts. Mostly important for the first element. Two buffers are necesarry because the scan procedure would override values that are then further read if one buffer is being used instead.
            dif_acum[featA] = accum[featA];

            //Compute the accumulated sum (a.k.a. scan)
            for(int j = 1; j < BLOCK_SIZE; j *= 2){
                int next_addend = featA-(j/2);
                if (featA >= j && next_addend >= 0 && next_addend < BLOCK_SIZE){
                    dif_acum[featA] = accum[featA] + accum[next_addend];
                }
                barrier(CLK_LOCAL_MEM_FENCE);

                //copy As to accum
                accum[featA] = GetElementArray(dif_acum, BLOCK_SIZE, featA); 
                barrier(CLK_LOCAL_MEM_FENCE);
            }

            //tell one of the threads to write the result of the scan in the array containing the results.
            if(featA == 0){
                C[loop] = accum[BLOCK_SIZE-1];
                loop += 1;
            }
            barrier(CLK_LOCAL_MEM_FENCE);

        }
    }
}
于 2010-09-24T14:52:54.277 に答える