• GPU 上で動く関数には __global__ または __device__ が付いている。
• __global__ は C++ から直接呼び出す用の関数で
• __device__ は __global__ の関数内から呼び出される関数だと思う。
• cudaMalloc で確保した GPUメモリに cudaMemcpy でCPU側からデータを移して
• CUDAの関数を呼び出して実行されたら cudaMemcpy でCPU側のメモリにデータを移す。
• CUDAの関数(カーネル)は関数名の後に <<<grid, threads>>> というテンプレート的な表記のものが付くのでわかりやすい。

Grid の中には複数の Block があり、Block の中には複数の thread がある。
画素の座標を知るには？

• dim3 gridShape(a, b, c); : 3次元テンソル
• dim3 blockShape(d, e, f); : 3次元テンソル
• として、<<<gridShape, blockShape>>> で呼び出すとよい。
• 全体としては 6次元テンソルのように見える。

例)

グローバルな ID は次のように計算すれば良い。

int id = (gridDim.y*blockDim.x)*blockIdx.x + (blockDim.x)*blockIdx.y + threadIdx.x;
// id = (131*256)*i + (256)*j + k ;

一つの Block は一つの SM に割り当てられる。
同じ Block 内のスレッドでは共有メモリを共有できるらしい。
共有メモリはデバイスメモリよりも高速にアクセス可能。

例) 1080ti は 6[GPC/GPU], 5[SM/GPC], 128[CUDACo./SM]
ただし 30SM のうちの 2 つの SM が無効化されている。

• 最適なブロック数の決め方 がある。

• dim3
• cudaMalloc
• cudaMemcpy
  • cudaMemcpyHostToDevice
  • cudaMemcpyDeviceToHost
• cudaFree

• cudaGetDeviceCount
• cudaSetDevice

• cudaMemcpy はカーネル処理の終了がされたことを保証する。
• カーネルを連続して実行した場合、thread内では順序が保証されていそう。

• 参考 doc
• <<< >>> のこと。カーネル実行時の設定
• __global__ 関数をよぶものはすべて実行コンフィグを指定する必要がある。
• グリッドとブロックの次数を定義。
• 正式な文法
  • <<< Dg, Db, Ns, S >>>
  • 省略された引数は 1 で初期化？
  • Dg ブロックの数。 Dg.x と Dg.y を指定。
    Dg.z は予約されているけど使われていない？(歴史的経緯で、昔は二次元のみだったので,環境によってうまく有効化されないことがあるstackoverflow)
  • Db スレッドの数。 ブロックあたりのスレッド数と Db.x * Db.y * Db.z が同じになる必要がある
  • Ns シェアードメモリのバイトサイズ。デフォルト0のオプション引数
  • S ストリーム。デフォルト0 のオプション引数。
• シェアードメモリのバイトサイズを指定するような状況とは、実行するまで必要なサイズがわからない場合。
  • カーネル呼び出し側で指定したサイズが作成される。
  • 動的な確保の場合はカーネル内で extern __shared__ として宣言しなければいけない。

スレッドを複数まとめたもの。実行の同期がとられる一単位。 同じ Warp 内にあるスレッドは while 文など実行すると、他のスレッドの while が止まるまで待たなければいけない。