差分

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

--- tips:cuda:start [2021/09/16 12:53] – yuuho
+++ tips:cuda:start [2022/04/27 04:10] (現在) – yuuho
@@ 行 2: / 行 2: @@
-頻出データ構造/関数
+=== 基本事項 ===
+  * GPU 上で動く関数には ''%%_%%_global%%_%%_'' または ''%%_%%_device%%_%%_'' が付いている。
+  * ''%%_%%_global%%_%%_'' は C++ から直接呼び出す用の関数で
+  * ''%%_%%_device%%_%%_'' は ''%%_%%_global%%_%%_'' の関数内から呼び出される関数だと思う。
+  * ''cudaMalloc'' で確保した GPUメモリに ''cudaMemcpy'' でCPU側からデータを移して
+  * CUDAの関数を呼び出して実行されたら ''cudaMemcpy'' でCPU側のメモリにデータを移す。
+  * CUDAの関数(カーネル)は関数名の後に ''<%%<%%<grid, threads>%%>%%>'' というテンプレート的な表記のものが付くのでわかりやすい。
+=== CUDAにおける配列 ===
+Grid の中には複数の Block があり、Block の中には複数の thread がある。\\
+画素の座標を知るには？
+  * ''dim3 gridShape(a, b, c);'' : 3次元テンソル
+  * ''dim3 blockShape(d, e, f);'' : 3次元テンソル
+  * として、''<%%<%%<gridShape, blockShape>%%>%%>'' で呼び出すとよい。
+  * 全体としては 6次元テンソルのように見える。
+例)
+^ サイズ ^ 実際の値 ^ インデックス名 ^ インデックス ^
+| ''gridDim.x ''  | 131 | ''blockIdx.x '' | ''i'' |
+| ''gridDim.y ''  | 131 | ''blockIdx.y '' | ''j'' |
+| ''gridDim.z ''  | 1   | ''blockIdx.z '' |       |
+| ''blockDim.x''  | 256 | ''threadIdx.x'' | ''k'' |
+| ''blockDim.y''  | 1   | ''threadIdx.y'' |       |
+| ''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|最適なブロック数の決め方]] がある。
+=== 頻出データ構造/関数 ===
   * ''dim3''
   * ''cudaMalloc''
@@ 行 12: / 行 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 が止まるまで待たなければいけない。