2017年7月 : sonots:blog

CUDA C プロフェッショナル第１章から第３章までのメモ。

この本はとても良かった。丁寧に詳細まで解説してある。中級〜上級者向け。Impressのサイトにソースコードがおいてあるので、落として手を動かしながら進めていくのが吉。

CUDA C プロフェッショナルプログラミング (impress top gear)

posted with amazlet at 17.07.12

John Cheng Max Grossman Ty McKercher
インプレス
売り上げランキング: 42,483

Amazon.co.jpで詳細を見る

第１章

nvcc -arch sm_35 hello.cu -o hello

CC(Compute Capability) 3.5 以降 (Kepler 世代以降) に最適なバイナリを生成する。それ以前では動かない。

nvcc -gencode arch=compute_20,code=sm_20 -gencode arch=compute_35,code=sm_35

gencode オプションを使うと複数世代に最適なコードを生成し、ワンバイナリに含めることができる

第２章

１つのカーネル呼び出しによって生成されたすべてのスレッドを、まとめてグリッドと呼ぶ
ブロックとは、互いに協調して動作できるスレッドのグループであり、以下の機能を利用できる
- ブロックに属するスレッド間の同期 (ブロックをまたいだ同期は難しい)
- ブロック内で共有されるシェアドメモリ
kernel<<nBlocks, nThreadsPerBlock>> のように呼び出す
- nBlocks 1グリッドあたりのブロック数
- nThreadsPerBlock 1ブロックあたりのスレッド数
CUDAでは32本のスレッド(１ Warp)に対して１つのインストラクションを１度に実行する。
- nThreadsPerBlock は 32の倍数にしておかないと、損をする
- インストラクションは Warp ごとだが、レジスタとシェアドメモリと同期はブロックごと

Fermi 以降なら、カーネルで printf が使える

CUDA API の実行失敗をチェックするマクロ CHECK()。例えば
CHECK(cudaDeviceSynchronize()); とすればそこまでの処理全部に対して失敗しているものがいないか確認できる。

#define CHECK(call)                                                  \
{                                                                    \
    const cudaError_t error = call;                                  \
    if (error != cudaSuccess)                                        \
    {                                                                \
        printf("Error: %s:%d, ", __FILE__, __LINE__);                \
        printf("code:%d, reason: %s\n", error,                       \
                cudaGetErrorString(error));                          \
        exit(1);                                                     \
    }                                                                \
}

kernel を送り込むところで CPU で時間がかかる。なので、非同期に送り込みたくなるし、実際 kernel 実行は非同期になっている

Fermi の場合、
グリッドあたりの各次元(x,y,z,w)ごとの最大ブロック数は65536
ブロックあたりのスレッドの数は最大 1024

なお、現行の chainer(cupy) の場合は以下のように x しか使っていない。

134     cpdef linear_launch(self, size_t size, args, size_t shared_mem=0,
135                         size_t block_max_size=128, stream=None):
137         gridx = size // block_max_size + 1
138         if gridx > 65536:
139             gridx = 65536
140         if size > block_max_size:
141             size = block_max_size

スレッド数 128、ブロック数 = 要素数 // スレッド数(128)。
(65536 を超えるようだったら、y を使うようにしても良いかもしれないが、ソフトウェア的にブロック数を増やした所で、ハードウェアのスペックをすでに超えてるから意味ないだろうなぁ)

第３章

SIMDとSIMT

SIMD: 複数データにたいして１つのインストラクションを１度に実行する。全体で１つのインストラクション。命令ユニットが少なくてすむ。
SIMT: NVIDIA独自(というか提案?)のアーキテクチャ。SIMDの改善。32本のスレッド(１ Warp)に対して１つのインストラクションを１度に実行する。全体としては同時に複数のインストラクションを発行可能(別Streamを指定して発行)。

ワープダイバージェンス

同じワープ内のスレッドが異なる命令を実行することをワープダイバージェンスと呼ぶ
ワープ内のスレッドが異なる分岐パスを選択した場合、ワープは各分岐パスを逐次的に実行する(分岐 true を実行してから、分岐 false を実行する)
- SIMDの場合、そもそも分岐処理できないので、それの改良ではある。
- ただし、遅くはなる

ダイナミックパラレリズム

kepler 以降
kernel から kernel 呼び出しができる

nvidia-smi -a -q -d CLOCK | fgrep -A 3 "Max Clocks" | fgrep "Memory

メモリ転送速度を帯域幅じゃなくて周波数でみることができる
帯域幅 / 周波数 = bytes / サイクル

3.3 並列性の確保

nvprof --metrics arhieved_occupancy

占有率 = アクティブワープの数 / ワープの最大数を見る。GPUを有効に使えているかどうかの指標
ワープの最大数: SMあたりの最大スレッド数(M2070ではmaxThreadsPerMultiProcessor == 1,536) / 32
アクティブワープの数: 有効に使っているワープの数
占有率をあげるには、基本的には nBlocks の数をあげれば良い
- といって、nThreadsPerBlock を 128 より小さくして nBlocks を大きくしようとすると占有率はあがるものの、１ブロック(ハードウェア的にはSM)あたりのスレッド数が減って並列処理できなくなるので、遅くなる。<= Q. そこもあわせて見る指標はないのか？

nvprof --metrics gld_throughput

メモリ読み取りスループットのチェック
補足: cudaMalloc とか cudaMemcpy とかではなくて、GPUプロセッサとGPUメモリ間の転送
(CPUでこれを見ようと思うと難しいので簡単に見れるのは嬉しい)
Q. ブロック数とスレッド数の変更で、メモリ読み取りスループットが変わる理屈がよくわからない。そもそも説明が書いてない。

nvprof --metrics gld_efficiency

メモリ読み取り効率のチェック
(さらに%で見れるとは嬉しい)

グリッドのサイズ(nBlock)とブロックのサイズ(nThreadsPerBlock)の決定に関するガイドライン

nThreadsPerBlock をワープサイズ（32）の倍数に保つ
nThreadsPerBlock は 128 または 256 からチューニングを始める
カーネルのリソース要件に従ってブロックサイズを調整する
nBlock がSMの数をはるかに上回るようにする (k80では13)
- 元々のデータサイズが小さいと、SMの数を上回れなくなるので、データをまとめて一気に処理するようにするとか?

3.4, 3.5. リダクション

http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf

(日本語) http://gpu-computing.gsic.titech.ac.jp/Japanese/Lecture/2010-06-28/reduction.pdf

うーむ、難しい。あとでまた読む

1章 [入門]プロセッサとGPU

SMが独立の命令列を実行できる単位

CPUがアクセスするメモリとGPUがアクセスするメモリは別になっているのが現状は一般的

GPUメモリは高バンド幅メモリが必要。CPUメモリはバンド幅はそこまで必要ないが大容量であって欲しい
特質が違って統一させるともったいない
Skylakeぐらいになってくると、同じぐらいのメモリでも有効に使える、とかなんとか

グラフィックスでは1ピクセルちょっと違っても気にならないので、エラー検知をしない

GPGPU向け Fermi では特別に誤り訂正ユニットを載せた

SIMDとSIMT

SIMD: 複数データにたいして１つのインストラクションを１度に実行する。全体で１つのインストラクション。命令ユニットが少なくてすむ。
SIMT: NVIDIA独自(というか提案?)のアーキテクチャ。SIMDの改善。32本のスレッド(１ Warp)に対して１つのインストラクションを１度に実行する。全体としては同時に複数のインストラクションを発行可能(別Streamを指定して発行)。SIMT アーキテクチャでないと実現不可能なのかがいまいちわからないが、Predicate という機構が入っていて、条件分岐処理を書くことができるようになっている。ただし、MIMDのように if の TRUE 側も FALSE 側も並列実行できるわけではなくて、どちらか一方の path のみが同時実行になる。
ref. このサイトが詳しかった http://yosefk.com/blog/simd-simt-smt-parallelism-in-nvidia-gpus.html

4章 [詳説] GPUの超並列処理

kepler 以降で、gpu カーネルから gpu カーネルを起動するダイナミックパラレリズムという機能が入っている

float16を使うと32に比べて2倍になると思いきや、メモリバンド幅のほうがネックになっていて、4倍ぐらいになる。とにかくバンド幅の進化が、GPUの進化においついていない

5章 GPUプログラミングの基本

PTX アセンブラっぽいやつ

nvidia は機械語を公開していない。互換性をばんばん壊してる
が、ドライバでPTXを機械語に実行時に変換していて、PTXレイヤで互換性を保っている

同期

__syncthreads(): 同一ブロック内の全スレッド待ち合わせ
__threadfence_block(): メモリアクセス(グローバルメモリとシェアードメモリ)全てがデバイス内(シェアードメモリの場合はブロック内)の全てのスレッドで確認できるようになるまで待機
- 同一ブロックじゃなくてもいいのかな？ => そんなかんじっぽい Shared Memory and Synchronizations
cudaStreamSynchronize(): 指定したストリームの処理完了を待つ
cudaDeviceSynchronize(): 指定したデバイスの処理完了を待つ(全ストリームの処理完了を待つ)
- ref. When to When to call cudaDeviceSynchronize?

一つのワープに含まれる32のアクセスのアドレスが128バイト境界にアラインされた128バイトの範囲内(補足: 128の倍数でメモリ領域が分割されている)にまとまっていれば１回のメモリアクセスでまとめて処理(コアレス)できる。

１回で処理できず、もう一度アクセスしにいくことをリプレイという
リプレイを避けることが高速化の鍵

シェアドメモリは on chip なので、グローバルメモリより速い

列方向で飛び飛びアクセスになる場合は、シェアドメモリに転送しておくと改善できるかも(転送するには一回アクセスはするんだよな？何度もアクセスする場合は一回ですむ、ということかな)
シェアドメモリも同じバンクからアクセスすると遅くなるので、ズラしたほうが良いらしい。そこまでしなくても十分速いけど

ダブルバッファ

ストリームを使って、非同期にメモリー転送すると、転送中に計算させられるから良い。
計算中のメモリ領域は上書きできないので、半分しか使えない。ダブルバッファというテクニック

DMA転送する場合に、スワップアウトが起きると困る。起きないようにするのが、pinned memory

感想

単著の割に、内容が重複していることが多いと感じた。

例えば、SIMDとSIMTの話は、何度も出てくる。

内容は網羅的だが、詳細までは説明されていないと感じた。

重要なキーワードは漏れなく載っている。
しかし、わかりにくいものは、わかりにくいままで、結局ウェブ検索した。
例えば、SIMT がわかりにくかった。

sonots:blog

2017年07月

CUDA C プロフェッショナル第一章〜第三章

第１章

第２章

第３章

3.3 並列性の確保

3.4, 3.5. リダクション

GPUを支える技術読了

1章 [入門]プロセッサとGPU

4章 [詳説] GPUの超並列処理

5章 GPUプログラミングの基本

感想

sonots:blog

2017年07月

CUDA C プロフェッショナル 第一章 〜 第三章

第１章

第２章

第３章

3.3 並列性の確保

3.4, 3.5. リダクション

GPUを支える技術読了

1章 [入門]プロセッサとGPU

4章 [詳説] GPUの超並列処理

5章 GPUプログラミングの基本

感想

CUDA C プロフェッショナル第一章〜第三章