openCL reduce、2d arrayを渡す

Question

ここでopenCLに変換したいループがあります。

for(n=0; n < LargeNumber; ++n) {  
    for (n2=0; n2< SmallNumber; ++n2) { 
     A[n]+=B[n2][n]; 
    }               
    Re+=A[n];  
} 

これまで私がこれまで行ってきたことはありますが、正しくないことが分かりました。

__kernel void openCL_Kernel(__global int *A, 
         __global int **B, 
         __global int *C, 
         __global _int64 Re, 
            int D) 
{ 

int i=get_global_id(0); 
int ii=get_global_id(1); 

A[i]+=B[ii][i]; 

//barrier(..); ? 

Re+=A[i]; 

} 

私はこのタイプのものの完全な初心者です。まず第一に、私はグローバルダブルポインタをopenCLカーネルに渡すことができないことを知っています。解決策を投稿する前に数日ほど待つことができますが、私は自分自身でこれを理解したいと思いますが、私が正しい方向に向けるように助けてくれれば幸いです。

Answer 1

doublepointersを渡す際の問題について：このような問題は通常、行列全体（または作業中のもの）を連続した1つのメモリブロックにコピーすることによって解決されます。個々の行のオフセットが含まれています（アクセスはB[index[ii]+i]のようになります）。

ここでは、Reに縮小しています：あなたが作業しているデバイスの種類について言及していないので、私はそのGPUを想定します。その場合、私は同じカーネルでの削減を避けるでしょう。あなたが投稿した方法と同じように遅くなるからです（何千ものスレッドに対してReへのアクセスをシリアル化しなければなりません。A[i]にもアクセスできます）。 代わりに私がA[i]にすべてB[*][i]を合計したカーネルを、欲しい書き込み、別のカーネルにReへAからの削減を入れて、いくつかのステップでそれを行うだろう、それはあなたがn要素上で動作し、何かにそれらを縮小する縮小カーネルを使用していますn/16（または他の任意の数）。あなたがあなたの結果である1つの要素になるまで繰り返しカーネルを呼び出す（あなたがあなた自身を考えていると言ったので、この記述を意図的に漠然としている）。

脇役として：元のコードには素晴らしいメモリアクセスパターンはありません。 Bが外側のインデックスを反復する内側ループを持つ比較的大きい（そして2番目の次元のためにAよりもはるかに大きい）と仮定すると、たくさんのカシミスが作成されます。コヒーレントメモリアクセスに非常に敏感であるGPUに移植するときにこれがさらに悪くなる

だから、大規模なパフォーマンスを向上させることがあり、このようにそれを並べ替え：あなたがあるコンパイラを持っている場合

for (n2=0; n2< SmallNumber; ++n2) 
    for(n=0; n < LargeNumber; ++n)  
    A[n]+=B[n2][n]; 
for(n=0; n < LargeNumber; ++n)             
    Re+=A[n];  

これが本当particularyでありますその構造をベクトル化することができるかもしれないので、自動ベクトル化では良いですが、元のコードではそうすることはできません（もしAとB[n2]が同じメモリを参照できないことが証明できなければ元のコードをこれに変換してください）。