このOpenCLコードを最適化できますか？

Question

私は、特殊な行列関数のOpenCLのコードの一部に取り組んでいます：Dx1ベクトルv、2つのDxD行列AとBと定数cのために、以下r[i] = c * sum_over_j (v[j] * A[i][j] * B[i][j])

は、私は持っているものである1xDベクトルrを返します遠いですが、それは狂ってゆっくりと走ります。 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; 
     } 

乾杯！

Answer 1

最適化＃1：ベクトル__localを作成します。

これで私の最初のパスは、パフォーマンスのまともな改善を得ました。私は、各ベクトル[k]が合計D回読み取られていることに気づいたので、それを__localにコピーしました。これは、Dがこれを可能にするのに十分小さいためにのみ可能です。上記のようなカーネルは、5870と6970gpusの両方で0.08のひどいALU：フェッチ比に苦しんでいます。遅いgpusでさえメモリアクセスを待っています。この変更に伴い

#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プロファイラは、新たなALUを示している：5870と6970のGPUのための0.20の比率を取得。同じカードで平均時間が1513 - > 1034、および1261 - > 861に変更されました。ローエンドのgpusはフェッチではなくALUでバインドされるようになりました。 （4：1より大きい）

オペレーション＃2：作業グループ全体を使用して各結果[y]を計算する。

あなたはこのIDをもっと大きくする必要があります（100k +）。考え方は、ワークグループを使用して一度に結果の単一要素を計算することによって、最良のメモリアクセスパターンを得ることです。私はls（ローカルサイズ）をここでは64に定義しました。なぜなら、それは私のハードウェアだけでなく、ほとんどのベンダーにとっても機能するからです。その定義を変更しない限り、ホスト側から使用するワークグループのサイズは64でなければなりません。これは、__localとして合計[ls] storageを作成するために定義する必要があります。また、可変サイズの__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); 
    } 
} 

EDIT：APPプロファイラ= AMD APP KernelAnalyzer