yuuho.wiki

カオスの欠片を集めて知恵の泉を作る

ユーザ用ツール

サイト用ツール

現在位置: start » tips » cuda
トレース: • cuda • hack
tips:cuda:start

差分

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

この比較画面へのリンク

両方とも前のリビジョン前のリビジョン
次のリビジョン		前のリビジョン
tips:cuda:start [2021/09/17 04:48] – yuuhotips:cuda:start [2022/04/27 04:10] (現在) – yuuho
行 33: 行 33:
 <code c++>int id = (gridDim.y*blockDim.x)*blockIdx.x + (blockDim.x)*blockIdx.y + threadIdx.x; <code c++>int id = (gridDim.y*blockDim.x)*blockIdx.x + (blockDim.x)*blockIdx.y + threadIdx.x;
 // id = (131*256)*i + (256)*j + k ;</code> // 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|最適なブロック数の決め方]] がある。
  
 === 頻出データ構造/関数 === === 頻出データ構造/関数 ===
行 44: 行 60:
   * ''cudaGetDeviceCount''   * ''cudaGetDeviceCount''
   * ''cudaSetDevice''   * ''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 が止まるまで待たなければいけない。
 +
  
  
tips/cuda/start.1631854116.txt.gz · 最終更新: 2021/09/17 04:48 by yuuho