CUDA(GPU)について

Question

#CUDA(GPU)について

CUDAについていくつか質問があります
どなたかお答えいただけませんか？
全てではなくとも一部だけでも答えていただけると非常に助かります.

1. threadは32thread(1warp)単位で並列に動作すると思うのですがこれはカーネル内のすべての処理が終了してから次のwarpの処理へと移行するのでしょうか.それともカーネル内の処理の途中途中で別のワープの処理に切り替わり,平均的に同じペースで進みながらカーネルの処理を終えるのでしょうか.
2. 例えばSM(ストリーミングマルチプロセッサー)が80, 各SMの最大スレッド数?(maxThreadsPerMultiProcessor)が2048のGPUがあるとしたとき,1blockのthread数を1024とすると160blockまでは同時に動かすことができるのでしょうか?またその場合,二つのblockで一つのSMのシェアードメモリを共有するのでしょうか?

Answer 1

1について

warpは非同期で進行します（warp内の32 threadは同期的に動作します）
それを利用して条件分岐が生じる処理を高速化することができます（warp divergenceの削減といいます）
具体的に，サイズが64の配列の奇数番目には1を，偶数番目には0をいれるというタスクを考えます
通常であれば，スレッドIDにインデックスを割り当てると思います
疑似コードは以下

idx = threadIdx.x
if idx is odd: array[idx] = 1
else: array[idx] = 0

この場合，奇数番目のスレッドが同時に奇数を処理し，その間他のスレッドは待機しています。その後，偶数番目のスレッドが処理を行います。つまりSM数を考慮しなければ2ステップ必要になります
しかしwarpは非同期で動作するので，カーネル内で， 条件分岐を以下に書き換えることで，1ステップで処理が完了します

warpIdx = idx / 32
laneIdx = idx % 32

if warpIdx is 0: array[2*laneIdx+1] = 1
else: array[2*laneIdx] = 0

やっていることは前半のwarpに奇数の処理を，後半のwarpに偶数の処理を割り当てています。warp間は非同期で進行するので，thread間でのifの待機がなくなり，結果１つ目より早く進行します
ちなみに，32というマジックナンバーがどうしても気に入らなければ，warpSizeという予約変数がありますので，書き換えてもよいですが，今のところGPUによって変わらないので32で問題なさそうです

2について

理論上そうなります。できるだけこの処理リソースを生かす方が処理は早くなります（Occupancyを100%に近くする）
しかし実際にはそうはなりません。というのも，最近のCUDAの実行速度のほとんどがメモリアクセス時間であるため，少し凝った処理をしようとすると，1 blockに対し1024 threadではすぐにレジスタとL1キャッシュを使い切り，オフチップであるL2キャッシュにローカル変数が確保されてしまい，結局遅くなります
なので，目安としては1 block に対し， 256 thread か，512 threadと言われています
シェアドメモリは実装上ではブロックでのみ共有されますが，物理的にはSMで共有されているので，二つのblockで一つのSMのシェアードメモリを共有しています

Comments

[pypy7]

とても丁寧な回答ありがとうございます
1番について、理解しました.warp処理の切り替わるタイミングというのはメモリにアクセスした瞬間などという認識で間違いないですか？
2番についてですが"1 blockに対し1024 threadではすぐにレジスタやキャッシュを使い切り,"というのはどういうことでしょうか?
私の認識は,SMごとにレジスタやキャッシュの上限値が決められていて,それを超えるようなblock数はSMに割り当てられない,という認識です.
SMのthread数の上限が2048の場合,1block1028threadにしていると,例えば1024threadが上限の55%の資源を使うとするとSMに1block(1024thread)しか割り当てられないですが,1block512threadにしていると3block(1536thread)割り当てることができるので同時に動くthread数が増えるということでしょうか？
また少し話が飛ぶのですが1threadのレジスタの使用数は,カーネル内で宣言された変数のバイトの合計という認識であっていますか？

[DaichiIshida]

すみません、warp処理が切り替わるタイミングというのは、全SMに含まれるwarp数を超えたwarpIdxの処理に移行するタイミングという意味でしょうか？これに関しては時間方向に次の処理が予約されていきます（どのように割り当てられているかはコンパイラに詳しくないため申し訳ないですがわかりません）
2番目について、すみません、少し補足します。
カーネル内で宣言される変数はレジスタとキャッシュに確保されます（SMの保有するリソースを超えた場合はお書きの通りです）
ここで使い切ると問題なのはレジスタとL1キャッシュで、これを使い切るとL2キャッシュに割り当てられます。L2キャッシュはキャッシュという名前ながらも、その実オフチップメモリのため、グローバルメモリとアクセスの時間はほぼ変わりません。したがって、スレッド数が多いと例えSMのリソースの範囲内であっても遅くなります
ソース: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#local-memory
レジスタの使用数はその認識で間違いありませんが、レジスタを超えた時にどの変数がレジスタに割り当てられるのかはコンパイラが決定するので注意が必要です

[pypy7]

返信ありがとうございます.
遅くなる原因については理解しました.
SMのキャッシュのリソース範囲内であってもL1キャッシュに入りきらない場合はローカルメモリのL2キャッシュに割り当てられ遅くなるということですね.
しかしスレッド数が多いとL1キャッシュを使い切りやすくなり？遅くなるということだと思うのですがなぜそうなるのでしょうか? スレッド数が1024でも512でもSMに入る上限のスレッド数は2048だとするとブロック数はそれぞれ2と4になり結局のところ同じ量のキャッシュを使うと思ったのですが見当違いですか？
たびたびの質問申し訳ございません.

[DaichiIshida]

回答が遅くなりましてすみません
確かに考えてみればその通りですね
先の回答までは自分の方でも検証して確認していますので，ここまでは事実ととらえてもらって大丈夫ですが，おっしゃっていることも理にかなっていると思います。ただ，なぜそうならないかについては把握していないので自分でも少し調べてみたいと思います
はっきり回答できなくて申し訳ありません