ストリーミングmultiprocessors、ブロック、スレッド(CUDA)

Question

の関係をどう考えていけばよいか、CUDAコアは、ストリーミングマルチプロセッサのCUDAモデルのブロック、スレッド?

何がマップされるいは並列化といかがですか？何より効率の最大化をブロックまたはスレッドの数?

---

現在の私の理解では、ある8cudaコア／マルチプロセッサ.一cudaコアを実行することができるようにな一cudaブロックです。すべてのスレッドがブロックの実行を直列にそのコアです。

すが、いかがでしょうか？

Answer 1

スレッド /ブロックレイアウトは、 CUDAプログラミングガイド. 。特に、第4章では次のように述べています。

CUDAアーキテクチャは、マルチスレッドストリーミングマルチプロセッサ（SMS）のスケーラブルな配列を中心に構築されています。ホストCPUのCUDAプログラムがカーネルグリッドを呼び出すと、グリッドのブロックが列挙され、利用可能な実行能力を備えたマルチプロセッサに分布します。スレッドブロックのスレッドは、1つのマルチプロセッサで同時に実行され、複数のスレッドブロックは1つのマルチプロセッサで同時に実行できます。スレッドブロックが終了すると、空いているマルチプロセッサで新しいブロックが起動されます。

各SMには8つのCUDAコアが含まれており、いつでも32スレッドの単一のワープを実行しているため、ワープ全体に1つの命令を発行するには4クロックサイクルが必要です。特定のワープのスレッドはロックステップで実行されると仮定できますが、ワープ全体で同期するには、使用する必要があります __syncthreads().

Answer 2

のGTX970が13ストリーミングMultiprocessors(SM)128Cudaコアです。Cudaコアとも呼ばれるストリームプロセッサ数(SPと略).

を定義することができるグリッドマップブロックを用いています。

を定義することができるブロック図スレッドがストリームプロセッサ（128Cudaコア／SM).

一warpは、常に形成される32スレッドは、すべてのスレッドの反りを行simulaneously.

完全には利用可能な電力のGPUコンピューティングに必要なものスレッド一SMのSMはSPs.各計算能力が一定数のスレッドできる一つのSMます。すべてのブロック定義するのはキューに、お待ちSMの資源（SPs無料であることを検証することが読み込まれます。SM開始を実行し反り.今、経糸だけで32スレッドおよびSMは、例えば128SPs、SM実行することができ4の反りをすることができるものとします。この場合、スレッドなメモリアクセスのスレッドまでブロックメモリ要求を満足しています。数字:演算の計算には、SPはのアイデンティティに18-22サイクルが非キャッシュされたグローバルメモリアクセスにつきましては300-400サイクルです。この場合、スレッドの反りお待ちしてデータのみサブセットの128SPsいます。そのスケジュールのスイッチを実行う経糸が可能です。このワープブロックで実行し、次います。このコンセプトと呼び遅延隠されているのです。の反りのブロックサイズを決めるの稼働率（からどのように多くの反りのSMを選んで実行).た場合の稼働率が高いのではありませんのSPs.

お算書るcudaコアの実行をブロック時には間違っています。場合についてお話しいただけまストリーミングMultiprocessorsて実行することができ反りからすべてのスレッドに入居する、SM.場合にブロックサイズの最大256スレッド、お持ちのGPU allowes2048スレッドの住民たSM各SMい8ブロックに居住するSM選択できる反りをを実行することを示しています。すべてのスレッドの実行の反りを行う。

す表示件数を増やす表示件数を減らの異なる計算能力とGPU築。https://en.wikipedia.org/wiki/CUDA#Limitations

ダウンロードできる宿泊人数計算シートからNvidia 稼働率の計算シート(Nvidia).

Answer 3

Compute Work Distributorは、SMがスレッドブロック（共有メモリ、ワープ、レジスタ、バリアなど）に十分なリソースがある場合にのみ、SMでスレッドブロック（CTA）をスケジュールします。スレッドブロックレベルのリソースは、そのような共有メモリが割り当てられます。割り当ては、スレッドブロック内のすべてのスレッドに対して十分なワープを作成します。リソースマネージャーは、使用を使用してワープを割り当てます ラウンドロビン SMサブパーティションへ。各SMサブパティションには、ワープスケジューラ、レジスタファイル、および実行ユニットが含まれています。ワープがサブパーティションに割り当てられると、コンテキストスイッチ（Pascal Architecture）が完了または先制されるまで、サブパティションに残ります。コンテキストスイッチの復元で、ワープは同じSM同じWarp-IDに復元されます。

ワープのすべてのスレッドが完了したら、ワープスケジューラがWarpによって発行されたすべての未解決の指示が完了するまで待機し、リソースマネージャーはWarp-IDと登録ファイルを含むWarpレベルのリソースをリリースします。

スレッドブロック内のすべてのワープが完了すると、ブロックレベルのリソースがリリースされ、SMはブロックが完了したことを計算作業ディストリビューターに通知します。

ワープがサブパティションに割り当てられ、すべてのリソースが割り当てられると、ワープスケジューラがワープの状態を積極的に追跡していることを意味します。各サイクルで、ワープスケジューラは、どのアクティブワープが失速し、どのアクティブワープが命令を発行する資格があるかを決定します。ワープスケジューラは、最優先の適格なワープを選択し、ワープから1〜2の連続した指示を発行します。二重問題のルールは、各アーキテクチャに固有です。ワープがメモリ負荷を発行すると、依存命令に達するまで独立した命令を実行し続けることができます。ワープは、負荷が完了するまで失速した報告を行います。依存する数学の指示にも同じことが言えます。 SMアーキテクチャは、ワープ間でサイクルごとに切り替えることにより、ALUとメモリの両方のレイテンシを隠すように設計されています。

この答えは、Cuda Coreという用語を使用しません。これは、誤ったメンタルモデルを導入するためです。 CUDAコアは、パイプ化された単一精度の浮動点/整数実行ユニットです。発行率と依存関係の遅延は、各アーキテクチャに固有です。各SMサブパティションとSMには、ロード/ストアユニット、二重精度の浮動点ユニット、半精度の浮動点ユニット、ブランチユニットなどを含む他の実行ユニットがあります。

パフォーマンスを最大化するために、開発者はブロック対ワープ対レジスタ/スレッドのトレードオフを理解する必要があります。

占有という用語は、SMのアクティブワープと最大ワープの比率です。 Kepler -Pascal Architecture（GP100を除く）には、SMごとに4つのワープスケジューラーがあります。 SMあたりのワープ数の最小数は、少なくともワープスケジューラの数に等しくなければなりません。アーキテクチャに6サイクル（MaxwellとPascal）の依存実行レイテンシがある場合、レイテンシをカバーするためにSMあたり24（24/64 = 37.5％の占有率）であるスケジューラごとに少なくとも6ワープが必要になります。スレッドに命令レベルの並列性がある場合、これは減少する可能性があります。ほとんどすべてのカーネルは、80〜1000サイクルかかる可能性のあるメモリ負荷などの可変レイテンシ命令を発行します。これには、レイテンシを隠すために、ワープスケジューラごとによりアクティブなワープが必要です。各カーネルには、ワープの数と共有メモリやレジスタなどの他のリソースの間にトレードオフポイントがあるため、他の犠牲がおそらく行われる可能性が高いため、100％の占有率を最適化することはお勧めしません。 CUDAプロファイラーは、開発者がそのバランスを決定するのを支援するために、指導の発行率、占有率、および失速の理由を特定するのに役立ちます。

スレッドブロックのサイズは、パフォーマンスに影響を与える可能性があります。カーネルに大きなブロックがあり、同期障壁を使用している場合、障壁は失速の理由になる可能性があります。これは、スレッドブロックごとのワープを減らすことで軽減できます。

Answer 4

1つのデバイスに複数のストリーミングマルチプロセッサがあります。
SMには複数のブロックが含まれる場合があります。各ブロックにはいくつかのスレッドが含まれる場合があります。
SMには複数のCUDAコアがあります（開発者として、Warpによって抽象化されているため、これを気にしないでください）。 SMは常にスレッドのワープに取り組んでいます（常に32）。ワープは、同じブロックのスレッドでのみ作業します。
SMとブロックの両方に、スレッド数、レジスタの数、共有メモリの制限があります。