CUDAでの占有管理

Question

おはようございます。

私はプログラミングの勉強を始めています。私はパフォーマンスを上回ります。私は他に、我々は考慮して4つの物事を取るべき良好な性能を持っていることをCUDAのウェブサイトで読む：

http://docs.nvidia.com/gameworks/content/developertools/desktop/analysis/report/cudaexperiments/kernellevel/achievedoccupancy.htm

-warpsあたりSM（システムマルチプロセッサ）SM あたり -blocks SM -Sharedあたり - レジスタメモリごとに

私は最初のものに行きます.GPUに応じて、SMあたりの最大ワープ数とSMあたりのブロック数に応じて、カーネルの寸法を定義しました。私の仕事は、どの方法がうまくいくかを測定するために10億の合計を実行しています。

私が行うことは、各繰り返しで占有を最大化するカーネルを起動するforループです。たとえば、私が読んNVidiaの1080 GPUのために：

int max_blocks = 32; //maximum number of active blocks per SM int max_threads_per_Block = 64; //maximum number of active threads per SM int max_threads = 2048;

これはSMあたりの合計2048件のスレッドで与え、最大収容人数を保証します。このGPUには32スレッドのアクティブワーピングが64個あります。このGPUでは、1つのアクティブなブロックに2つのワープがあり、これは各ブロックが同時に64のアクティブなスレッドを持つことができることを意味します。次のようにこれで私はカーネルを起動します。

dim3 threadsPerBlock(max_threads_per_Block); dim3 numBlocks(max_blocks); VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,max_threads);

私は意外に気づくことである私のような直接カーネルを起動した場合：

int N = total_ops; //in this case one thousand millions dim3 threadsPerBlock(256); dim3 numBlocks(2*N/threadsPerBlock.x); VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,);

パフォーマンスは良いです（tイム消費）。同じ実験を5回繰り返して、異常値を避ける。私の質問です：占有率を管理する方法はありますかコンパイラと実行時APIより良い結果を持っている？私がしようとしている最適化は、GPUによって何らかの形ですでに管理されていることを理解しています。私は、良いパフォーマンスを達成するために、ソフトウェアをどのように立ち上げるべきかを説明した文書（上記のリンク先）があれば、これを制御する方法でなければならないことを理解しています。あなたの最初の例で

おかげ

Answer 1

、

int max_blocks = 32;   //maximum number of active blocks per SM 
int max_threads_per_Block = 64; //maximum number of active threads per SM 
int max_threads = 2048; 

dim3 threadsPerBlock(max_threads_per_Block); 
dim3 numBlocks(max_blocks); 
VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,max_threads); 

完全 SMをロードするために、必要に応じて、あなたはブロックごとのように多くのブロックとスレッドを立ち上げています。ただし、GTX 1080には SMがあるため、占有率はわずか1/20 = 5％です。第2の例で

、あなたは100％の占有率に到達するために必要なGPUは並列に限り多くを実行することを可能にする多数のブロックを、起動している

int N = total_ops;    //in this case one thousand millions 
dim3 threadsPerBlock(256); 
dim3 numBlocks(2*N/threadsPerBlock.x); 
VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,); 

（リソースはなく、許可します簡単なベクトル加算の場合には問題になる）。したがって、より良いパフォーマンス。

最初の例ではブロック数を20倍に増やして2番目の例と同じパフォーマンスにすることができますが、2番目の例のパターンは、使用されるGPUの特定の設定。したがって、コードは広範囲のGPUのどれかを完全にロードします。

メモでは、メモリ結合アルゴリズムとしてのベクトル加算は、占有の影響を実証するのに特に適していません。しかし、メモリーサブシステムを完全にロードするには、飛行中の特定の最小メモリートランザクション数（メモリー帯域幅とメモリーアクセスのレイテンシーを掛け合わせたもの）が必要であり、5％占有率の例がこの最小値に満たないため、 。