====== CUDA ======


=== 基本事項 ===

  * 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''
  * ''cudaMemcpy''
    * ''cudaMemcpyHostToDevice''
    * ''cudaMemcpyDeviceToHost''
  * ''cudaFree''

  * ''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 が止まるまで待たなければいけない。