質問をすることでしか得られない、回答やアドバイスがある。

15分調べてもわからないことは、質問しよう!

ただいまの
回答率

87.34%

CUDA(GPU)について

受付中

回答 1

投稿

  • 評価
  • クリップ 0
  • VIEW 779

score 15

CUDA(GPU)について

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

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

    クリップした質問は、後からいつでもマイページで確認できます。

    またクリップした質問に回答があった際、通知やメールを受け取ることができます。

    クリップを取り消します

  • 良い質問の評価を上げる

    以下のような質問は評価を上げましょう

    • 質問内容が明確
    • 自分も答えを知りたい
    • 質問者以外のユーザにも役立つ

    評価が高い質問は、TOPページの「注目」タブのフィードに表示されやすくなります。

    質問の評価を上げたことを取り消します

  • 評価を下げられる数の上限に達しました

    評価を下げることができません

    • 1日5回まで評価を下げられます
    • 1日に1ユーザに対して2回まで評価を下げられます

    質問の評価を下げる

    teratailでは下記のような質問を「具体的に困っていることがない質問」、「サイトポリシーに違反する質問」と定義し、推奨していません。

    • プログラミングに関係のない質問
    • やってほしいことだけを記載した丸投げの質問
    • 問題・課題が含まれていない質問
    • 意図的に内容が抹消された質問
    • 過去に投稿した質問と同じ内容の質問
    • 広告と受け取られるような投稿

    評価が下がると、TOPページの「アクティブ」「注目」タブのフィードに表示されにくくなります。

    質問の評価を下げたことを取り消します

    この機能は開放されていません

    評価を下げる条件を満たしてません

    評価を下げる理由を選択してください

    詳細な説明はこちら

    上記に当てはまらず、質問内容が明確になっていない質問には「情報の追加・修正依頼」機能からコメントをしてください。

    質問の評価を下げる機能の利用条件

    この機能を利用するためには、以下の事項を行う必要があります。

回答 1

+2

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の待機がなくなり,結果1つ目より早く進行します
ちなみに,32というマジックナンバーがどうしても気に入らなければ,warpSizeという予約変数がありますので,書き換えてもよいですが,今のところGPUによって変わらないので32で問題なさそうです

2について

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

投稿

編集

  • 回答の評価を上げる

    以下のような回答は評価を上げましょう

    • 正しい回答
    • わかりやすい回答
    • ためになる回答

    評価が高い回答ほどページの上位に表示されます。

  • 回答の評価を下げる

    下記のような回答は推奨されていません。

    • 間違っている回答
    • 質問の回答になっていない投稿
    • スパムや攻撃的な表現を用いた投稿

    評価を下げる際はその理由を明確に伝え、適切な回答に修正してもらいましょう。

  • 2019/09/20 12:44

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

    キャンセル

  • 2019/09/20 15:01 編集

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

    キャンセル

  • 2019/09/21 17:13

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

    キャンセル

  • 2019/09/24 02:07

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

    キャンセル

15分調べてもわからないことは、teratailで質問しよう!

  • ただいまの回答率 87.34%
  • 質問をまとめることで、思考を整理して素早く解決
  • テンプレート機能で、簡単に質問をまとめられる

関連した質問

同じタグがついた質問を見る