カーネルレジスタが理論上の最大占有率を %50 に制限している CUDA プログラムがありました。そこで、ブロック スレッド間で一定であり、カーネルの実行中はほとんど読み取り専用であった変数に対して、レジスタの代わりに共有メモリを使用することにしました。ここでソース コードを提供することはできません。私がしたことは、概念的には次のようなものでした：
私の最初のプログラム:
__global__ void GPU_Kernel (...) {
__shared__ int sharedData[N]; //N:maximum amount that doesn't limit maximum occupancy
int r_1 = A; //except for this first initialization, these registers don't change anymore
int r_2 = B;
...
int r_m = Y;
... //rest of kernel;
}
上記のプログラムを次のように変更しました。
__global__ void GPU_Kernel (...) {
__shared__ int sharedData[N-m];
__shared__ int r_1, r_2, ..., r_m;
if ( threadIdx.x == 0 ) {
r_1 = A;
r_2 = B;
...
r_m = Y; //last of them
}
__syncthreads();
... //rest of kernel
}
ブロック内のワープのスレッドがブロードキャスト読み取りを実行して、新しく作成された共有メモリ変数にアクセスするようになりました。同時に、スレッドは達成された占有を制限するためにあまり多くのレジスターを使用しません。
2 番目のプログラムでは、理論上の最大占有率が %100 に等しくなります。実際の実行では、最初のプログラムで達成された平均占有率は ~%48 で、2 回目のプログラムでは約 ~%80 でした。しかし、問題は正味速度の向上が約 %5 から %10 であり、占有率の向上を考慮して予想していたものよりもはるかに少ないことです。この相関関係が線形でないのはなぜですか?
Nvidia のホワイトペーパーの下の画像を考えると、私が考えていたのは、占有率が %50 に達した場合、たとえば、SMX (新しいアーキテクチャの場合) コアの半分が一度にアイドル状態になるということでした。これは、他のコアによって過剰に要求されたリソースがそれらを停止するためです。アクティブ。私の理解は間違っていますか？それとも、上記の現象を説明するのは不完全ですか? それとも、追加され__syncthreads();て共有メモリアクセスのコストがかかりますか?
