0

タイプの再帰方程式をシミュレートしようとしています

s i (t+1) = f[Σ j W ij s j (t) + v i *input(t)]

ここで、f(.) は何らかの非線形関数 (以下のコードでは、しきい値 th を持つ単なるステップ関数です) であり、s(t) は何らかの外部入力です。当然、x iごとに 1 つのワーカーを実装しました。すべての時間ステップで、すべてのワーカーが上記の方程式の結果を計算し、その後、この結果が他のすべてのワーカーと共有されます。したがって、すべてのワーカーは同じワークグループに属している必要があります。

現在の OpenCL カーネルは次のようになります

__kernel void part1(__global int* s, __global float* W, __global float* Th, __global float* V, __global float* input, int N, int T)
    {
        unsigned int i = get_global_id(0);

        float value = 0;
        float v = V[i];
        float th = Th[i];  

        for(int t = 0; t < T; t++){
            value = v*input[t];
            for(int j = 0; j < N; j++){
                value = value + W[i*N + j]*s[j];
            }
            barrier(CLK_GLOBAL_MEM_FENCE);
            if (value >= th){
                s[i] = 1;
            } else {
                s[i] = 0;
            }
            barrier(CLK_GLOBAL_MEM_FENCE);
        }
    }

残念ながら、このコードは実際には同等の C 実装よりも 3 倍遅くなります。また、ワーカーの数を変更しても大きな違いは生じないと予想していましたが (新しいワーカーは、他のワーカーと並行して実行される新しいスレッドに座っているため)、実際には処理時間はワーカーの数に比例して増加します。ボトルネックは、最初の障壁の後の書き込み操作のようです。この操作をなくすと (ただし、バリアはそのままにしておくと)、処理時間が 25 分の 1 に短縮され、線形依存がなくなります。

私は OpenCL にかなり慣れていないので、このコードを高速化するための助けをいただければ幸いです!

よろしくお願いします！Blue2script

4

1 に答える 1

0

コメントで既に述べたように、グローバル メモリへのアクセスは低速です。通常、ハードウェアは、同じ計算ユニットで実行されるスレッドのサブグループをいくつか持つことで、レイテンシーを隠します。私が言及しているサブグループは、NVIDIA 用語ではコール ワープ、AMD 用語ではウェーブフロントです。通常、ワークグループは複数のサブグループで構成されています。

そのため、1 つのサブグループがグローバル メモリからのデータの受信を待機している間に、必要なすべてのリソースが既にある別のサブグループを実行できます。グローバルメモリとの間でデータを読み書きする必要があるために実行中のものが停止すると、別のものが実行を開始できます。

ただし、あなたの場合、障壁のため、すべてのサブグループのすべてのワーカーは、計算を続行する前に、他のメンバーがメモリに書き込んだことを確認する必要があります (障壁はワークグループ レベルにあります)。したがって、レイテンシーはあなたを直撃します:)。

ここで、実装を改善する方法は、ローカル メモリを使用することです。今回は、ローカル メモリ レベルでバリアを使用します (フラグ CLK_LOCAL_MEM_FENCE を使用)。今説明したのと同じ原則が適用されますが、ローカル メモリへのアクセスははるかに高速です。

私があなたのコードを理解している限り (すべての微妙な点を理解できていない可能性が非常に高いです)、あなたの s 配列には N 個の要素があり、N 個のワーカーもあると思います。したがって、N 要素のローカル配列を作成し、次のようにします。

kernel void part1(global int* s, global float* W, global float* Th, global float* V, global float* input, local int* local_s, int N, int T)
    {
        unsigned int i = get_global_id(0);
        unsigned int local_i = get_local_id(0);

        float value = 0;
        float v = V[i];
        float th = Th[i];  
        //fetch from global to local and sync before computing
        local_s[local_i] = s[i];
        barrier(CLK_LOCAL_MEM_FENCE);

        for(int t = 0; t < T; t++){
            value = v*input[t];
            for(int j = 0; j < N; j++){
                value = value + W[i*N + j]*local_s[j];
            }
            barrier(CLK_LOCAL_MEM_FENCE);
            if (value >= th){
                local_s[i] = 1;
            } else {
                local_s[i] = 0;
            }
            barrier(CLK_LOCAL_MEM_FENCE);
        }
    //If necessary write some stuff to global (maybe the last s computed?)
    }

今、私はあなたに警告する必要があります:

  • 私はあなたのニーズを完全に誤解しているかもしれません:)
  • この回答を入力しているときにコードを編集したところなので、おそらくタイプミスなどがあります。
  • ローカル メモリを使用しても、バリアが非常に多いため、opencl バージョンは C バージョンよりも遅くなる可能性があります。

先頭の __ は不要であり、私の意見では読みやすいため、削除したことに注意してください。

編集: CLK_LOCAL_MEM_FENCE と CLK_GLOBAL_MEM_FENCE に関するコメントについて。バリアは常にワークグループ レベルで適用されるため、ワークグループ内のすべてのワーカーはそのバリアに到達する必要があります。パラメータとして指定されたフラグは、メモリ アクセスを参照します。フラグが CLK_GLOBAL_MEM_FENCE の場合、ワーカーが次のステートメントを実行し続ける前に、グローバル メモリに関するすべての読み取り/書き込み操作をすべてのワーカーが完了する必要があることを意味します。これは CLK_LOCAL_MEM_FENCE フラグとまったく同じですが、ローカル メモリの場合です。

于 2013-07-27T10:12:21.653 に答える