2つのほぼ同じOpenCLカーネル間のパフォーマンスの差

Question

GFLOPSで性能を計算したいのは、ほぼ同じ2つのOpenCLカーネルです。カーネル＃1は、次のとおりです。

__kernel void Test41(__global float *data, __global float *rands, int index, int rand_max){ 

    float16 temp; 
    int gid = get_global_id(0); 

    temp = data[gid]; 
    temp = (float) rands[1] * temp; 
    temp = (float) rands[2] * temp; 
    temp = (float) rands[3] * temp; 
    temp = (float) rands[4] * temp; 
    . 
    . 
    . 
    temp = (float) rands[497] * temp; 
    temp = (float) rands[498] * temp; 
    temp = (float) rands[499] * temp; 
    data[gid] = temp.s0; 

} 

二カーネルは次のとおりです。

__kernel void Test42(__global float *data, __global float *rands, int index, int rand_max){ 

    float16 temp[500]; 
    int gid = get_global_id(0); 

    temp[0] = data[gid]; 
    temp[1] = (float) rands[1] * temp[0]; 
    temp[2] = (float) rands[2] * temp[1]; 
    temp[3] = (float) rands[3] * temp[2]; 
    temp[4] = (float) rands[4] * temp[3]; 
    . 
    . 
    . 
    temp[497] = (float) rands[497] * temp[496]; 
    temp[498] = (float) rands[498] * temp[497]; 
    temp[499] = (float) rands[499] * temp[498]; 
    data[gid] = temp[index].s0; 

} 

あなたは、私はすべてのカーネルは、オペレーション、それぞれの500行を持っている16のストリームサイズを使用しているコードで見ることができるようにそのうちの1つだけが浮動小数点演算を行います。私は合計で約1048576のカーネルをデプロイするので、並列に実行するために約1048576の作業項目が必要になります。私はプを計算するために

：私の周り1.4 TFLOPSを得るが、2番目のカーネルのため、私は38 GFLOPSを取得する最初のカーネルのため残念ながら

flops = #numWorkItems(1048576) * (500) * StreamSize(16)/timeTaken; 

。私はこの大きなギャップを説明することができませんでした。単一の温度の代わりにtempのベクトルを使用することは大きな問題になります。また、実際のアプリケーションはほとんどが第2カーネルに似ているようです。最初のカーネルは実際のアプリケーションでは単純すぎます。

ここで何が起こっているのか、そして第2カーネルのパフォーマンスが最初のものに到達する方法を理解する助けになれますか？一般的に、デバイスのベンチマークを行う場合、理論値に近いパフォーマンスが期待できますか？

P.S.私はランドを__localメモリにコピーする必要があることを理解していますが、今のところそれをスキップしましょう。

Answer 1

2可能な問題があります。

• は、あなたはそれが非常に高いアクセスレイテンシでglobalメモリ空間に割り当てられます{のOpenCLのデフォルトである}と、最も可能性の高い__privateようfloat16一時バッファを宣言しました。デバイスのローカルメモリに合う場合は、__local float16と宣言してみてください。 tempバッファを追加する
• は...オリジナルのコードは、（たとえばインテル）いくつかのGPUアーキテクチャ上で簡単にベクトル化可能であるコンパイラのためのいくつかの問題を作成して、あなたは、私が実際に提出したいstore + load

を追加することにより、人工的な依存関係を追加しましたそのためのコンパイラに関するレポートを発行してください。コンパイラが依存関係を把握し、最適化を行い、さらにtempのバッファを取り除くだけでも簡単にできるはずです。

Answer 2

@pmdjがコメントに示唆しているように、2番目のカーネルの主な問題はレジスタ圧です。ハードウェアレジスタを多数使用しているため、同時に実行するワークグループの数が減ります。 一般に、大きなプライベートアレイはOpenCL/CUDAではお勧めできません。 この場合、パフォーマンスを最適化するコンパイラはほとんどありません。 アレイにローカルメモリを使用できますが、アクセスには適切な同期を追加する必要があります。