cuda 基礎

関数修飾子

__global__ : hostがコール、deviceが実行。
__device__ : deviceがコール、deviceが実行。
__host__ : hostが実行、hostが実行。default。

変数修飾子

デバイス上で実行される関数内の変数種別。

指定なし: 組込み型はレジスタ。あふれるとローカルメモリ。非組込型はローカルメモリ
__shared__ : gpuオンチップの共有メモリ。同一スレッドブロック内で共有。ブロックが異なる場合はアクセスできない
__constant__ : デバイスの定数メモリ
__device__ : デバイスのグローバルメモリ

デバイスのレジスタとメモリ

f:id:qshino:20190622130004j:plain

デバイスのメモリ種別。下記に加えホストのメモリがある。

レジスタ、スレッド内。float4等の組込み型変数が配置される。
ローカルメモリ、スレッド内。非組込型、溢れた組込型。
共有メモリ、スレッドブロック内。
グローバルメモリ、全スレッド。
定数メモリ、全スレッド
テキスチャメモリ、全スレッド。

組み込みベクトル型

GPU、CPUコードのどちらでも使用可能

[u]char[1..4], [u]short[1..4], [u]int[1..4], [u]long[1..4], float[1..4], double[1..2]

構造体をx、y、z、wフィールドでアクセス

uint4 param;
int y = param.y;
dim3 uint3に基づく次元の指定に使用デフォルト値(1,1,1)

grid/block

kernel起動

kernel<<<dG,dB, smbytes,stream>>>(args,...)

dG : grid 2次元まで。グリッドのサイズ
dB : block 3次元まで。スレッドブロックのサイズ
smbytes : shared memory byte size
stream : stream, 同一ストリームは順次実行。デフォルト値は0

スレッドブロックの最大スレッド数は1024。

例

__global__ void k(int a){}

in main(){
  
  k<<<1,1>>>(1);

  dim3 g(3,4),b(2,3,4);
  k<<<g,b>>>(3);
}

スレッド内では、下記の変数が定義される。

gridDim
blockDim
blockIdx
threadIdx

index

blockは3次元まて指定可能。各次元のインデックスをidx,y,zとすると、下記により計算できる。

idx = blockDim.x * blockIdx.x + threadIdx.x

idy = blockDim.y * blockIdx.y + threadIdx.y

idy = blockDim.z * blockIdx.z + threadIdx.z

メモリ確保

cuda メモリ関数

cudaMalloc() : デバイスメモリ確保
cudaMallocHost() : pinned host mem
cudaHostAlloc() : new for cudaMallocHost() to avoid warning with g++
cudaMallocManaged() : for tegra uma
cudaFree() : デバイスメモリ解放
cudaFreeHost() : ホストメモリ解放

memory coherency

Host Pageable memory/H
Host Pinned memory/P
Device memory/D
Unified memory/U
CPU cached : H, U, P>=CC7.2
CPU uncashed : P<CC7.2
CPU n/a D
iGPU cached : D, U, P>=CC7.2
iGPU uncashed : P<CC7.2
iGPU n/a : H
dGPU cashed : D
dGPU uncashed : P
dGPU n/a : H, U
iGPU : tegra gpu
dGPU : pcie gpu
CC : Compute Capability
CC7.2 : Tegra on AGX Xavier, Volta 512core

cuda for tegra

https://docs.nvidia.com/cuda/cuda-for-tegra-appnote/index.html

複数gpu

cudaSetDevice()を使う。Streamでそれぞれで同期する事もできる。

http://www.e-em.co.jp/tutorial/chap10.htm

gpu間転送

gpu間転送前に

cudaDeviceEnablePeerAccess()

を実行する。

https://www.cc.kyushu-u.ac.jp/scp/support/faq/faq006.html

転送関数

cudaMemcpy( d, s,len, dir ) : 非pin可
cudaMemcpyAsync() : pinメモリ必須

memcpyはdmaで行われるため、非pinの場合は、pinメモリにコピー後に非pinメモリに転送される。つまり、2回コピーされる。

http://yusuke-ujitoko.hatenablog.com/entry/2016/02/07/222059

ref

http://cparch-mclearn.blogspot.com/2018/01/gpu-5.html?m=1

qshinoの日記

Powershell関係と徒然なこと

cuda basics