大規模な行列計算のためにOpenCLのローカルワーカーを使用する

Question

私はVisual Studio C＃でOpenCL（Op​​enCL.NETライブラリ）を使用する初心者で、現在は大きな3Dマトリックスを計算するアプリケーションで作業しています。行列の各ピクセルでは、192の一意の値が計算され、合計されてそのピクセルの最終値が得られます。だから、機能的には、（X 161 X 161 161）、4-Dマトリックスのようなものです今私はこのような私のホスト・コードからカーネルを呼んでいる192

を×：

//C# host code 
... 
float[] BigMatrix = new float[161*161*161]; //1-D result array 
CLCalc.Program.Variable dev_BigMatrix = new CLCalc.Program.Variable(BigMatrix); 
CLCalc.Program.Variable dev_OtherArray = new CLCalc.Program.Variable(otherArray); 
//...load some other variables here too. 
CLCalc.Program.Variable[] args = new CLCalc.Program.Variable[7] {//stuff...} 

//Here, I execute the kernel, with a 2-dimensional worker pool: 
BigMatrixCalc.Execute(args, new int[2]{N*N*N,192}); 
dev_BigMatrix.ReadFromDeviceTo(BigMatrix); 

サンプルカーネルコードは以下に掲載されています。

__kernel void MyKernel(
__global float * BigMatrix 
__global float * otherArray 
//various other variables... 
) 
{ 
    int N = 161; //Size of matrix edges 
    int pixel_id = get_global_id(0); //The location of the pixel in the 1D array 
    int array_id = get_global_id(1); //The location within the otherArray 


    //Finding the x,y,z values of the pixel_id. 
    float3 p; 
    p.x = pixel_id % N;  
    p.y = ((pixel_id % (N*N))-p.x)/N; 
    p.z = (pixel_id - p.x - p.y*N)/(N*N); 

    float result; 

    //... 
    //Some long calculation for 'result' involving otherArray and p... 
    //... 

    BigMatrix[pixel_id] += result; 
} 

私のコードは、現在、しかし、私は、このアプリケーションのためにスピードを探しています、と私は私の労働者/グループのセットアップが寸法のための最善のアプローチ（すなわち161 * 161 * 161と192があるかどうかわからないんだけど作品

ワーカープールの）。

グローバルワーカープールをローカルワーカーグループに編成して効率を上げるという他の例を見てきましたが、OpenCL.NETでその方法を実装する方法があまりよく分かりません。また、ワーカー・プールに別の次元を作成するだけとはどう違うかもわかりません。

私の質問は次のとおりです：ここでローカルグループを使用することはできますか？もしそうなら、私はそれらをどのように整理するのですか？一般的に、n次元のワーカープールを呼び出すだけではなく、どのようにローカルグループを使用するのですか？ （つまり、Execute（args、new int [] {（N * N * N）、192}）のローカルワークグループのサイズは192ですか？

ありがとうございました！

Answer 1

メモリアクセスを待っていると多くのパフォーマンスが失われていると思います。私は答えましたsimilar SO question.私の投稿があなたを助けてくれることを願っています。ご質問があればお尋ねください。

最適化：

1. カーネルの私のバージョンで大きな後押しをローカルメモリにotherArray読んでから来ています。
2. 各作業項目は、BigMatrixで4つの値を計算します。つまり、同じキャッシュライン上に同時に書き込むことができます。実行する作業項目が依然として1M以上あるため、並列処理の損失は最小限に抑えられます。

...

#define N 161 
#define Nsqr N*N 
#define Ncub N*N*N 
#define otherSize 192 

__kernel void MyKernel(__global float * BigMatrix, __global float * otherArray) 
{ 
    //using 1 quarter of the total size of the matrix 
    //this work item will be responsible for computing 4 consecutive values in BigMatrix 
    //also reduces global size to (N^3)/4 ~= 1043000 for N=161 

    int global_id = get_global_id(0) * 4; //The location of the first pixel in the 1D array 
    int pixel_id; 
    //array_id won't be used anymore. work items will process BigMatrix[pixel_id] entirely 

    int local_id = get_local_id(0); //work item id within the group 
    int local_size = get_local_size(0); //size of group 


    float result[4]; //result cached for 4 global values 
    int i, j; 
    float3 p; 

    //cache the values in otherArray to local memory 
    //now each work item in the group will be able to read the values efficently 
    //each element in otherArray will be read a total of N^3 times, so this is important 
    //opencl specifies at least 16kb of local memory, so up to 4k floats will work fine 
    __local float otherValues[otherSize]; 
    for(i=local_id; i<otherSize; i+= local_size){ 
     otherValues[i] = otherArray[i]; 
    } 
    mem_fence(CLK_LOCAL_MEM_FENCE); 

    //now this work item can compute the complete result for pixel_id 
    for(j=0;j<4;j++){ 
     result[j] = 0; 
     pixel_id = global_id + j; 

     //Finding the x,y,z values of the pixel_id. 
     //TODO: optimize the calculation of p.y and p.z 
     //they will be the same most of the time for a given work item 
     p.x = pixel_id % N;  
     p.y = ((pixel_id % Nsqr)-p.x)/N; 
     p.z = (pixel_id - p.x - p.y*N)/Nsqr; 

     for(i=0;i<otherSize;i++){ 
      //... 
      //Some long calculation for 'result' involving otherValues[i] and p... 
      //... 
      //result[j] += ... 
     } 
    } 
    //4 consecutive writes to BigMatrix will fall in the same cacheline (faster) 
    BigMatrix[global_id] += result[0]; 
    BigMatrix[global_id + 1] += result[1]; 
    BigMatrix[global_id + 2] += result[2]; 
    BigMatrix[global_id + 3] += result[3]; 
} 

注：

1. グローバルワークサイズは4の倍数である必要があります。理想的には、4 *のワークグループサイズの倍数。これは、各pixel_idが0..N^3-1の範囲内に入るかどうかを確認するエラーチェックがないためです。未処理の要素は、カーネルが実行されるのを待つ間に、CPUによってクランチされます。
2. ワークグループのサイズはかなり大きいはずです。これにより、キャッシュされた値がより強く使用され、LDSのデータをキャッシュする利点が拡大します。

3. 多すぎる高価な除算やモジュロ演算を避けるために、p.x/y/zの計算をさらに最適化する必要があります。以下のコードを参照してください。

   __kernel void MyKernel(__global float * BigMatrix, __global float * otherArray) { 
   int global_id = get_global_id(0) * 4; //The location of the first pixel in the 1D array 
   int pixel_id = global_id; 
   
   int local_id = get_local_id(0); //work item id within the group 
   int local_size = get_local_size(0); //size of group 
   
   
   float result[4]; //result cached for 4 global values 
   int i, j; 
   float3 p; 
   //Finding the initial x,y,z values of the pixel_id. 
   p.x = pixel_id % N;  
   p.y = ((pixel_id % Nsqr)-p.x)/N; 
   p.z = (pixel_id - p.x - p.y*N)/Nsqr; 
   
   //cache the values here. same as above... 
   
   //now this work item can compute the complete result for pixel_id 
   for(j=0;j<4;j++){ 
       result[j] = 0; 
   //increment the x,y,and z values instead of computing them all from scratch 
       p.x += 1; 
       if(p.x >= N){ 
        p.x = 0; 
        p.y += 1; 
        if(p.y >= N){ 
         p.y = 0; 
         p.z += 1; 
        } 
       } 
   
       for(i=0;i<otherSize;i++){ 
        //same i loop as above... 
       } 
   } 
   

Answer 2

私はあなたのためにいくつかの提案を持っている：

1. 私はあなたのコードは、競合状態を持っていると思います。あなたの最後のコード行は、複数の異なる作業項目によって変更されているBigMatrixと同じ要素を持っています。
2. あなたのマトリックスが本当に161x161x161である場合、これらの寸法を唯一の寸法として使用するための多くの作業項目があります。すでに4百万以上の作業項目があります。これは、マシンに十分な並列処理が必要です。あなたはそれを192回も必要としません。さらに、個々のピクセルの計算を複数の作業項目に分割しない場合は、最後の追加を同期する必要はありません。
3. グローバルワークサイズが2の偉大な倍数でない場合は、それを1になるようにパッドしてみてください。あなたのローカルワークサイズとしてNULLを渡しても、いくつかのOpenCL実装は、分割しないグローバルサイズに対して非効率なローカルサイズを選択します。
4. アルゴリズムにローカルメモリや障壁が必要ない場合は、ローカルワークグループをほとんどスキップできます。

希望します。