__global__ または __device__ が付いている。__global__ は C++ から直接呼び出す用の関数で__device__ は __global__ の関数内から呼び出される関数だと思う。cudaMalloc で確保した GPUメモリに cudaMemcpy でCPU側からデータを移してcudaMemcpy でCPU側のメモリにデータを移す。<<<grid, threads>>> というテンプレート的な表記のものが付くのでわかりやすい。
Grid の中には複数の Block があり、Block の中には複数の thread がある。
画素の座標を知るには？
dim3 gridShape(a, b, c); : 3次元テンソルdim3 blockShape(d, e, f); : 3次元テンソル<<<gridShape, blockShape>>> で呼び出すとよい。例)
| サイズ | 実際の値 | インデックス名 | インデックス |
|---|---|---|---|
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 は次のように計算すれば良い。
int id = (gridDim.y*blockDim.x)*blockIdx.x + (blockDim.x)*blockIdx.y + threadIdx.x; // id = (131*256)*i + (256)*j + k ;
| 単位 | 説明 |
|---|---|
| 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 が無効化されている。
dim3cudaMalloccudaMemcpycudaMemcpyHostToDevicecudaMemcpyDeviceToHostcudaFreecudaGetDeviceCountcudaSetDevicecudaMemcpy はカーネル処理の終了がされたことを保証する。<<< >>> のこと。カーネル実行時の設定__global__ 関数をよぶものはすべて実行コンフィグを指定する必要がある。<<< Dg, Db, Ns, S >>>Dg ブロックの数。 Dg.x と Dg.y を指定。Db スレッドの数。 ブロックあたりのスレッド数と Db.x * Db.y * Db.z が同じになる必要があるNs シェアードメモリのバイトサイズ。デフォルト0のオプション引数S ストリーム。デフォルト0 のオプション引数。extern __shared__ として宣言しなければいけない。スレッドを複数まとめたもの。実行の同期がとられる一単位。 同じ Warp 内にあるスレッドは while 文など実行すると、他のスレッドの while が止まるまで待たなければいけない。