差分

このページの2つのバージョン間の差分を表示します。

--- tips:cuda:start [2021/09/17 04:45] – yuuho
+++ tips:cuda:start [2022/04/27 04:10] (現在) – yuuho
@@ 行 16: / 行 16: @@
 画素の座標を知るには？
-  * ''gridShape'' : 3次元テンソル
+  * ''dim3 gridShape(a, b, c);'' : 3次元テンソル
-  * ''blockShape'' : 3次元テンソル
+  * ''dim3 blockShape(d, e, f);'' : 3次元テンソル
   * として、''<%%<%%<gridShape, blockShape>%%>%%>'' で呼び出すとよい。
+  * 全体としては 6次元テンソルのように見える。
 例)
 ^ サイズ ^ 実際の値 ^ インデックス名 ^ インデックス ^
-| gridDim.x   | 131 | blockIdx.x  | i |
+| ''gridDim.x ''  | 131 | ''blockIdx.x '' | ''i'' |
-| gridDim.y   | 131 | blockIdx.y  | j |
+| ''gridDim.y ''  | 131 | ''blockIdx.y '' | ''j'' |
-| gridDim.z   | 1   | blockIdx.z  |   |
+| ''gridDim.z ''  | 1   | ''blockIdx.z '' |       |
-| blockDim.x  | 256 | threadIdx.x | k |
+| ''blockDim.x''  | 256 | ''threadIdx.x'' | ''k'' |
-| blockDim.y  | 1   | threadIdx.y |   |
+| ''blockDim.y''  | 1   | ''threadIdx.y'' |       |
-| blockDim.z  | 1   | threadIdx.z |   |
+| ''blockDim.z''  | 1   | ''threadIdx.z'' |       |
 グローバルな ID は次のように計算すれば良い。
 <code c++>int id = (gridDim.y*blockDim.x)*blockIdx.x + (blockDim.x)*blockIdx.y + threadIdx.x;
 // id = (131*256)*i + (256)*j + k ;</code>
+=== 物理構造とデータモデル ===
+^ 単位      ^ 説明 ^
+| GPC       | Graphics Processing Cluster, has N x SM        |
+| SM        | Streaming Multi-Processor,   has N x CUDA Core |
+| CUDA Core | thread に対応(？)                              |
+一つの Block は一つの SM に割り当てられる。\\
+同じ Block 内のスレッドでは共有メモリを共有できるらしい。\\
+共有メモリはデバイスメモリよりも高速にアクセス可能。
+例) 1080ti は 6[GPC/GPU], 5[SM/GPC], 128[CUDACo./SM]\\
+ただし 30SM のうちの 2 つの SM が無効化されている。
+  * [[https://on-demand.gputechconf.com/gtc/2014/jp/sessions/4002.pdf|最適なブロック数の決め方]] がある。
 === 頻出データ構造/関数 ===
@@ 行 43: / 行 60: @@
   * ''cudaGetDeviceCount''
   * ''cudaSetDevice''
+=== 同期の保証 ===
+  * ''cudaMemcpy'' はカーネル処理の終了がされたことを保証する。
+  * カーネルを連続して実行した場合、thread内では順序が保証されていそう。
+=== 実行コンフィグレーション(execution configuration) ===
+  * 参考 [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#execution-configuration|doc]]
+  * ''%%<%%<%%<%% %%>%%>%%>%%'' のこと。カーネル実行時の設定
+  * ''%%_%%_global%%_%%_'' 関数をよぶものはすべて実行コンフィグを指定する必要がある。
+  * グリッドとブロックの次数を定義。
+  * 正式な文法
+    * ''%%<%%<%%<%% Dg, Db, Ns, S %%>%%>%%>%%''
+    * 省略された引数は 1 で初期化？
+    * ''Dg'' ブロックの数。 Dg.x と Dg.y を指定。\\  Dg.z は予約されているけど使われていない？(歴史的経緯で、昔は二次元のみだったので,環境によってうまく有効化されないことがある[[https://stackoverflow.com/questions/36843432/why-is-z-always-zero-in-cuda-kernel|stackoverflow]])
+    * ''Db'' スレッドの数。 ブロックあたりのスレッド数と Db.x * Db.y * Db.z が同じになる必要がある
+    * ''Ns'' シェアードメモリのバイトサイズ。デフォルト0のオプション引数
+    * ''S'' ストリーム。デフォルト0 のオプション引数。
+  * シェアードメモリのバイトサイズを指定するような状況とは、実行するまで必要なサイズがわからない場合。
+    * カーネル呼び出し側で指定したサイズが作成される。
+    * 動的な確保の場合はカーネル内で ''extern %%_%%_shared%%_%%_'' として宣言しなければいけない。
+=== Warp とは ===
+スレッドを複数まとめたもの。実行の同期がとられる一単位。
+同じ Warp 内にあるスレッドは while 文など実行すると、他のスレッドの while が止まるまで待たなければいけない。