6

私は特殊な行列関数のOpencLコードの一部に取り組んでいます：Dx1ベクトルvの場合、2つのDxD行列AとB定数の場合、ベクトルcを返します。1xDrr[i] = c * sum_over_j (v[j] * A[i][j] * B[i][j])

以下は私がこれまでに持っているものですが、それは異常に遅いです。行列を返す合計なしのバージョンDxDは、約10倍高速です。違いがある場合は、PyOpenCLから呼び出されます。

何か間違ったことはありますか？最適化できますか？

#define D 1000
...

   __kernel void element_mult(
      __global float *result,
      __global const float *vector,
      __global const float *matrix,
      __global const float *matrix2,
        const float factor)
      {
         int y = get_global_id(1);
         float sum = 0;
         for(int k = 0; k < D; k++)
         {
            sum += vector[k] * matrix[(y*D) + k]
            * matrix2[(y*D) + k ];
         }
         result[y] = sum * factor;
      }

乾杯！

4

1 に答える 1

6

最適化 #1: ベクトルを __local にする。

これでの私の最初のパスは、パフォーマンスがかなり改善されました。各 vector[k] が合計 D 回読み取られることに気付いたので、__local にコピーしました。これが可能なのは、D が十分に小さいためです。上記のカーネルは、5870 と 6970 gpu の両方で 0.08 というひどい ALU:fetch 比率に悩まされています。低速の gpu でさえ、メモリ アクセスを待機しています。

   #define D 1000
    __kernel void element_mult(
    __global float *result,
    __global const float *vector,
    __global const float *matrix,
    __global const float *matrix2,
    const float factor)
    {
        int y = get_global_id(0);
        float sum = 0;

        __local float vectCopy[D];
        int ls = get_local_size(0);
        int lid = get_local_id(0);
        for(int i=0;i<D;i+=ls){
            vectCopy[i+lid] = vector[i+lid];
        }
        mem_fence(CLK_LOCAL_MEM_FENCE);

        for(int k = 0; k < D; k++)
        {
            sum += vectCopy[k] * matrix[(y*D) + k] * matrix2[(y*D) + k ];
        }
        result[y] = sum * factor;
    }

この変更により、APP プロファイラーは、5870 および 6970 gpu の新しい ALU:fetch 比率 0.20 を示しています。同じカードで平均時間が 1513 から 1034、1261 から 861 に変更されました。ローエンドの GPU は、フェッチではなく ALU によってバインドされるようになりました。(4:1 以上の比率)

最適化 #2: ワーク グループ全体を使用して各結果 [y] を計算します。

これを行う必要があります id D ははるかに大きかった (100k+)。アイデアは、ワーク グループを使用して一度に結果の 1 つの要素を計算することにより、最適なメモリ アクセス パターンを取得することです。ここでは ls (ローカル サイズ) を 64 に定義しました。これは、ほとんどのベンダーと同様に、私のハードウェアでも機能するためです。ホスト側から使用するワークグループ サイズは、その定義を変更しない限り 64 にする必要があります。sum[ls] ストレージを __local として作成するように定義する必要があり、可変サイズの __local 変数をカーネルに渡すのは好きではありません。

結果: 5870 ALU:fetch=0.59:1, avg=708. 6970 ALU:フェッチ=0.72、平均=590。APP プロファイラーによると、これは元のリストの約 2 倍の速さです。

#define D 1000
#define ls 64
__kernel void element_mult(
__global float *result,
__global const float *vector,
__global const float *matrix,
__global const float *matrix2,
const float factor)
{
    __local float vectCopy[D];
    int lid = get_local_id(0);
    for(int i=0;i<D;i+=ls){
        vectCopy[i+lid] = vector[i+lid];
    }
    mem_fence(CLK_LOCAL_MEM_FENCE);

    int ng = get_num_groups(0);
    int gid = get_group_id(0);
    int y, k;
    __local float sum[ls];
    for(y = gid; y < D; y+=ng){
        for(k = lid; k < D; k+=ls)
        {
            sum[lid] += vectCopy[k] * matrix[(y*D) + k] * matrix2[(y*D) + k ];
        }
        if(lid==0){
            result[y] = sum[0];
            for(k=1;k<ls;k++){
                result[y] += sum[k];
            }
            result[y] *= factor;
        }
        mem_fence(CLK_LOCAL_MEM_FENCE);
    }
}

編集: APP プロファイラー = AMD APP KernelAnalyzer

于 2012-02-23T16:24:42.450 に答える