8

OpenCL を使用して、2 組の 3D ポイント間の最近傍を見つけています。

Nearest Neighbour: DataSet 内の各ポイント (x、y、z) について、モデル内で最も近いものを見つける必要があります。二乗距離 = (Ax-Bx)^2 + (Ay-By)^2 + (Az-Bz)^2

ここで私がこれまでに行ったこと：

struct point {
int x;
int y;
int z;
};

__kernel void 
nearest_neighbour(__global struct point *model,
__global struct point *dataset,
__global int *nearest,
const unsigned int model_size)
{
    int g_dataset_id = get_global_id(0);

    int dmin = -1;
    int d, dx, dy, dz;

    for (int i=0; i<model_size; ++i) {
        dx = model[i].x - dataset[g_dataset_id].x;
        dx = dx * dx;

        dy = model[i].y - dataset[g_dataset_id].y;
        dy = dy * dy;

        dz = model[i].z - dataset[g_dataset_id].z;
        dz = dz * dz;

        d = dx + dy + dz;

        if(dmin == -1 || d < dmin)
        {
            nearest[g_dataset_id] = i;
            dmin = d;
        }
    }
}

コードは機能しているように見えますが、最適化できると確信しています。ローカルメモリを利用して改善する方法を知りたいです。

ありがとう

PS kd-tree のように、最も近い隣人を見つけるための他の (より良い) 方法があることは知っていますが、今のところ、簡単な方法を実行したいと思います。

4

7 に答える 7

2

クイック プレフィルターを使用すると、速度が向上する可能性があります。二乗距離をすぐに計算する代わりに、3 つの座標すべての距離が dmin よりも近いかどうかを最初に確認できます。したがって、内側のループを次のように置き換えることができます

{
    dx = model[i].x - datum_x;
    if (abs(dx) >= dmin) continue;

    dy = model[i].y - datum_y;
    if (abs(dy) >= dmin) continue;

    dz = model[i].z - datum_z;
    if (abs(dz) >= dmin) continue;

    dx = dx * dx;    
    dy = dy * dy;
    dz = dz * dz;

    d = dx + dy + dz;

    if(d < dmin)
    {
        nearest[g_dataset_id] = i;
        dmin = d;
    }
}

abs()ポイントごとのおよびsへの余分な呼び出しがif十分なデータ ポイントを除外するかどうかはわかりませんが、試してみるには十分単純な変更だと思います。

于 2011-03-22T02:06:23.200 に答える
1

ヒースの提案は出力インデックスにも適用できます。変数nearest_idを維持し、ループ後に一度だけ書き込みます。

3 つのコンポーネントの構造体の代わりに、int4 ベクトルを試し、ベクトル演算を使用します。

于 2011-03-22T10:49:10.350 に答える
1

私が最初に思いついたのは、ヒースの提案でした。各作業項目はmodel[i]同時にメモリ項目にアクセスしています。コンパイラの最適化の程度によっては、各作業項目が配列の異なる要素にアクセスする方がよい場合があります。それをよろめかせる1つの方法は次のとおりです。

int datum_x, datum_y, datum_z;

datum_x = dataset[g_model_id].x;
datum_y = dataset[g_model_id].y;
datum_z = dataset[g_model_id].z;

for (int i=0; i<size_dataset; ++i) {
    j = (i + g_model_id) % size_dataset;  // i --> j by cyclic permutation

    dx = model[j].x - datum_x;
    dx = dx * x;

    dy = model[j].y - datum_y;
    dy = dy * dy;

    /* and so on... */
}

ただし、model[i]コード内の へのアクセスが「ブロードキャスト」として処理される可能性があります。その場合、コードの実行速度が遅くなります。

于 2011-03-21T21:17:07.470 に答える
1

Eric Bainville の提案の後、point struct を取り除こうとしました。示唆されているように、私はfloat4を使用しました。ここで私がやったこと:

__kernel void 
nearest_neighbour(__global float4 *model,
__global float4 *dataset,
__global unsigned int *nearest,
const unsigned int model_size)
{
    int g_dataset_id = get_global_id(0);

    float dmin = MAXFLOAT;
    float d;

    /* Ottimizzato per memoria locale */
    float4 local_xyz = dataset[g_dataset_id];
    float4 d_xyz;
    int imin;

    for (int i=0; i<model_size; ++i) {
        d_xyz = model[i] - local_xyz;

        d_xyz *= d_xyz;

        d = d_xyz.x + d_xyz.y + d_xyz.z;

        if(d < dmin)
        {
            imin = i;
            dmin = d;
        }
    }

    nearest[g_dataset_id] = imin; // Write only once in global memory
}

問題は、このバージョンが point 構造体に基づくバージョンよりも少し遅く実行されることです。おそらく、構造体でプレフィルターを使用したためです。

dx = model[i].x - local_x;
dx = dx * dx;
if (dx >= dmin) continue;

dy = model[i].y - local_y;
dy = dy * dy;
if (dy >= dmin) continue;

dz = model[i].z - local_z;
dz = dz * dz;
if (dz >= dmin) continue;

d = dx + dy + dz;

そのプレフィルター幅の float4 バージョンは使用できません。あなたの意見では、float4 バージョンでできる他の最適化はありますか?

皆様、貴重なご提案ありがとうございます

于 2011-03-22T20:54:19.107 に答える
1

現在の最小値を に書き込むのに多くの時間が費やされていると確信していnearest[g_dataset_id]ます。グローバル メモリへのアクセスは非常に遅いことが多いため、現在の最小値をレジスタに格納する方がよいでしょうdmin = d。

ちょうどこのような：

...
int dmin = MAX_INT;
int imin = 0;
...
for (...)
{
  ...
  if(d < dmin)
  {
    imin = i;
    dmin = d;
  }
}

nearest[g_dataset_id] = imin; //write to global memory only once
于 2011-03-22T11:14:05.923 に答える
1

「ローカルメモリを活用して改善する方法を知りたい」という具体的な質問に対して。

GPU のローカル メモリを使用するのは難しい場合があります。それに取り組む前に、SDK のコード サンプルとプログラミング ガイドで十分な時間を費やす必要があります。

基本的に、ローカル メモリを使用してグローバル データの一部のブロック (この場合は model[] 配列) をキャッシュし、グローバルから読み取るよりも高速に読み取ることができるようにします。試してみたい場合は、次の擬似コードのようになります。

For each block of the model array {
    1) Read data from __global and write it to __local
    2) Barrier
    3) For each model datum in the __local cache,
       Read it and process it.
    4) Barrier
}

ステップ 3 は基本的に現在のループですが、モデル データ全体ではなくチャンクのみを処理する点が異なります。

手順 2 と 4 は、ローカル メモリを使用する場合は常に不可欠です。ワークグループ内のすべてのスレッドを同期する必要があります。バリアは、すべての作業項目がバリアの後にコードを実行できるようになる前に、バリアの前にコードを完了するように強制します。これにより、他のスレッドによってローカル メモリにデータが書き込まれる前に、作業項目がローカル メモリからデータを読み取ることを防ぎます。バリア命令の構文は覚えていませんが、OpenCL のドキュメントに記載されています。

ステップ 1 各作業項目に、グローバルから異なるデータを読み取らせ、それをローカル キャッシュに書き込みます。

このようなもの（これは単純化しすぎてテストされていないことに注意してください！）：

__local float4 modelcache[CACHESIZE];
int me = get_local_id(0);

for (int j = 0; j < model_size; j += CACHESIZE) {
    modelcache[me] = dataset[j+me];
    barrier(CLK_LOCAL_MEM_FENCE);
    for (int i=0; i < CACHESIZE; ++i) {
        d_xyz = modelcache[i] - local_xyz;
        ... etc.
    }
    barrier(CLK_LOCAL_MEM_FENCE);
}

その場合の設計上の問題は、ローカル キャッシュをどのくらいの大きさにすべきかということです。作業グループのサイズは？

ローカル データ ストアは、ワーク グループ内のワーク アイテム間で共有されます。作業項目の ND 配列が多数の作業グループを並行して実行する場合、各作業グループには独自のモデル キャッシュのコピーがあります。

ローカル データ配列を小さくしすぎると、それらを使用してもメリットがほとんどまたはまったく得られません。それらを大きくしすぎると、GPU は多くのワーク グループを並行して実行できなくなり、実際にはかなり遅くなる可能性があります。

最後に、この特定のアルゴリズムは、ローカル メモリ キャッシュからあまりメリットを得られない可能性が高いと言わざるを得ません。あなたのプログラムでは、すべての作業項目が同時に同じモデル [i] の場所を読み取っています。ほとんどの GPU には、それを高速に実行するように特別に最適化されたハードウェアがあります。

于 2011-04-03T22:39:05.887 に答える