7

これが私が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カーネルに渡すことができないことを知っています。可能であれば、解決策を投稿する前に数日ほど待ってください。私はこれを自分で理解したいと思いますが、正しい方向に私を向けるのを手伝っていただければ幸いです。

4

1 に答える 1

12

ダブルポインターの受け渡しに関する問題について：この種の問題は通常、行列全体（または作業中のもの）をメモリの1つの連続ブロックにコピーし、ブロックの長さが異なる場合は、次のオフセットを含む別の配列を渡すことで解決されます。個々の行（アクセスは次のようになりますB[index[ii]+i]）。

さて、次のように減らしますRe。作業しているデバイスの種類について言及していなかったので、GPUを想定します。その場合、同じカーネルでの削減は避けます。投稿した方法と同じくらい遅くなるからですRe（数千を超えるスレッドへのアクセス（およびアクセスも）をシリアル化する必要がありA[i]ます）。代わりに私はwant kernelを記述します。これは、すべてB[*][i]を合計して別のカーネルA[i]にリダクションを入れ、Aいくつかのステップで実行します。つまり、要素を操作して次のようなものReにリダクションするリダクションカーネルを使用します。nn / 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];

これは、自動ベクトル化が得意なコンパイラーを使用している場合に特に当てはまります。これは、その構成をベクトル化できる可能性があるためですが、元のコードではそうすることができない可能性が非常に高いです（それを証明できAずB[n2]、元のコードをこれに変換できない同じメモリを参照しないでください）。

于 2012-01-07T17:27:43.720 に答える