スキップしてメイン コンテンツに移動

GPUを支える技術読み始めた 第5章[前半]

GPUを支える技術読み始めた 第5章[前半]

新年の休みもとうに終わり、皆さんどうお過ごしだろうか。
ちなみに主は年末読む予定であった本や論文を全く消化できずに今に至ってしまった。
年末年始は時間があったので、ひたすら読んでいたはずなのになぜだ。。＿|￣|○
読むのが遅いのがいけないのか？翻訳するのが遅いのか？？それともそれをまとめるのに時間がかかっているのか。
おそらく全て当てはまるが、時間がかかるものは仕方がないので、少しずつ慣れていくしかないなー。。（-_-;）

余談だが、年末にブログ執筆環境を再構築した。
今まではwebエディタのClasseurを利用していたが、やはり純粋なWebアプリケーションなのでレンダリングやアップデートに問題があった。そこで、今までも利用していたAtomを少しカスタマイズして試してみたが、どうにも動作が重くかつ、vim-modeがいい感じにならずに残念に思っていたところ、vscodeのことを思い出した。どちらもエレクトロンベースだが、vscodeはAtomより全然軽く、markdownプラグインも豊富にあるので、すぐにmarkdown+mathjax環境を構築することができた。いやはや、世の中は便利になったものだ（＾ω＾）
せっかくなので、下に執筆環境のスクリーンショットを自慢げに貼ってみようと思う。(markdownは頻繁に見る必要はないので、普段はプレビューは別のタブで開いている)

https://gyazo.com/33452309e25481bedacaec053ea7ddd7

もしかして、こういうことしてるから時間がかかるのかな？？／(^o^)＼

それでは、本題に戻ろう。
今回はGPU支える技術の第5章だ。4章と同じく楽しみにしていた章なのでじっくり読んでいきたい。
なお今回も長い章なので、前半と後半に分けてまとめと感想を書いていく。

第5章 GPUプログラミングの基本[前半]

GPUの超並列プロセッサでプログラムを実行するには、超並列で実行できようなプログラムの記述に適したプログラミング言語が求められる。また、これらはGPU側メモリへのデータの送信や計算後のデータ読み出し、同期処理などの計算以外の処理も必要とする。
上記のようなGPUプログラミング言語はいくつか存在する。ハードウェアに近いレベルの命令記述(PTX, SPIL-V)やよく知られているCUDAとOpenCL、C言語をディレクティブで拡張できるOpenACCやOpenMP等だ。

GPUの互換性の考え方

ハードウェアを開発する場合、従来の命令への互換性を保ったまま、新しい命令を追加する「上位互換」という考え方が一般的に用いれられる。

ハードウェアの互換性、機械語命令レベルの互換性

CPUの世界ではIntelとAMDのCPUは命令互換なので同じ機械命令を用いたプログラムであれば、どちらの上でも同様に動作する(もちろん各ベンダ固有の命令もある)。もちろん、IBM POWERやARMプロセッサの上ではこれらのプログラムは動作しないが、中にはAndroidのような仮想マシン技術を利用して、異なるハードウェアを用いた互換性を実現しているものもある。
一方でGPUは各社でまちまちであり、世代が変わっただけでプログラムの互換性が亡くなってしまう。

NVIDIAの抽象化アセンブラPTX

NVIDIAはその実際の機械語命令を公開する代わりに、PTXという命令セットを公開している。これはGPUドライバによって機械語に変換される。PTXはこれまでの全ての命令を保存し、新しい機能の追加に関しては上位互換で拡張を行う。しかし、GPUプログラミングに於いて直接PTXが利用されることは殆どなく、CUDAが用いられる。一方でAMDは機械語命令を公開している。

GPU言語レベルの互換性

CUDAはNVDIA GPUを前提に作られているため、他のGPUでは利用できない。従って、Khronous Gruupを中心にOpenCL言語が開発された。OpenCLはNVIDIA GPUでもサポートされており、業界標準のプログラミング言語となっている(現在の最新バージョンは2.2)。

CUDA

1章で説明したように、NVIDIAがG80 GPUと同時に公開したC言語の拡張。ちなみに「Compute Unified Device Architecture」の略。

C言語拡張

CPTとGPUのメモリが別れていることは既に説明したが、CUDAではこの指定を修飾子で行うことができる。これらは、__global__, __host__, __device__の3種類がある。

  • global
    デバイスで実行されるが、ホストCPUから直接呼び出される関数

  • device
    デバイスで実行され、呼び出しもデバイスから行われる。

  • host
    ホストCPUで実行される。呼び出しはホストからのみ可能。


メモリ構造

GPUのメモリは以下の5つの種類に分けられて利用される。

メモリ名 説明 搭載メモリ
グローバルメモリ 全てのスレッドで共有される デバイスメモリ
コンスタントメモリ 全てのスレッドで共有される定数を格納する デバイスメモリ
テクスチャメモリ テクスチャパターンを記憶する 数デバイスメモリ
シェアードメモリ スレッドブロック単位で割り当てられる 高速SRAM
プライベートローカルメモリ スレッドブロック単位で割り当てられ、他のスレッドからはアクセスできない デバイスメモリ

enter image description here
出展: 「GPUを支える技術」，p176，Hisa Ando，2017


メモリ領域の修飾子により、そのデータがどのメモリに配置されるかが決まる。例えば、__device__ で指定されたらデバイスメモリとなる。__global__を付ける、または指定がない場合はグローバルメモリに置かれる。__device__ __constant__ とするとデバイスメモリの定数領域に置かれ、__device__、__shared__ とすると、シェアードメモリに置かれる。ここに置かれたデータは異なるSM間では利用できない。また、各スレッドが必要とするautomatic変数は通常はGPU内のレジスタにおかれ、その容量が大きい場合等はローカルメモリに置かれる。

変数

C言語と同様に、char, short, int, long, longlong, float, double がある。また、それぞれ符号なしの uchar, ushort, uint, ulong, ulonglong がサポートされている。
int3などどするとベクトル変数となる。これは3つの連続した4バイト整数のことである。なお、インデックスは1〜4、長さは16バイト以下に制限され、2, 4, 8, 16バイトの境界でアラインされる。uint3の特殊化でdim3という指定も存在する。
上記の他に、スレッドが自信の位置を知るための組み込み変数として、 gridDim, blockIdx, blockDim, threadIdx等が定義されている。

データ転送

計算に必要なデータをGPUに転送したり、受け取ったりするために、ホストプログラムは以下のような処理ステップで記述されます。

  1. ホスト側にカーネルへのデータの入出力となる領域の確保: malloc()
  2. デバイス側に入力用のデバイスメモリの確保: cudaMalloc()
  3. デバイス側に出力用のデバイスメモリの確保: cudaMalloc()
  4. ホスト側からデバイス側に入力データをコピー: cudaMemcpy()
  5. カーネルを起動して計算を実行
  6. 計算結果をデバイス側からホスト側にコピー
  7. 使い終わったメモリを開放: free() & cudaFree()

具体的には次のようなソースコードで表される。

// Device code
__global__ void VecAdd(float* A, float* B, float* C, int N)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N)
        C[i] = A[i] + B[i];
}

// Host code
int main()
{
    int N = ...;
    size_t size = N * sizeof(float);

    // Allocate input vectors h_A and h_B in host memory
    float* h_A = (float*)malloc(size);
    float* h_B = (float*)malloc(size);

    // Initialize input vectors
    ...

    // Allocate vectors in device memory
    float* d_A;
    cudaMalloc(&d_A, size);
    float* d_B;
    cudaMalloc(&d_B, size);
    float* d_C;
    cudaMalloc(&d_C, size);

    // Copy vectors from host memory to device memory
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // Invoke kernel
    int threadsPerBlock = 256;
    int blocksPerGrid =
            (N + threadsPerBlock - 1) / threadsPerBlock;
    VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // Copy result from device memory to host memory
    // h_C contains the result in host memory
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // Free device memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // Free host memory
    ...
}

出展: http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-memory


行列積を計算するCUDAプログラム

N行N列の３つの行列A, B, Cがある時、その行列積は以下のCUDAコードで計算できる。以下の場合、1つのカーネルはC[i][j]という1つのデータだけを計算する。このサンプルコードはN × Nのスレッドを含んだスレッドブロックを1つ起動させ、それらの行列要素の計算を行う。

__global__ void MatMul(float A[N][N], float B[N][N], float C[N][N])
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    int j = blockDim.y * blockIdx.y + threadIdx.y;
    for(int k=0; k<N; k++)
	    C[i][j] += A[i][k] + B[k][j];
}

int main()
{
...
    // N * N * 1スレッドの1つのブロックを起動
    int numBlocks = 1;
    dim3 threadsPerBlock (N, N);
    MatMul<<<numBlocks, threadsPerBlock>>>(A, B, C);
...
}

出展: 「GPUを支える技術」，p183，Hisa Ando，2017


4章で説明したように、1つのスレッドブロック内のスレッドは全て同じSMで実行されるため、上記のコードは1つのSMのみしか利用することができない。また、1つのスレッドブロックには最大1024個のスレッドしかNには入れられないという制限もある。従って、本書では行列演算をブロック化して複数スレッドブロックで動かすことで効率化している。ブログには掲載しないが、気になる人は実際に本を参照してほしい。

CUDA数学ライブラリ

CUDAは標準の数学ライブラリの他にIntrinsic Functionsというライブラリを持っている。これは、IEEE準拠の標準ライブラリと比べると少し誤差が大きいが、高速に演算を行うことができる。NVIDIAはその他にも以下のようなライブラリを揃えている。

  • cuBLAS: 行列、ベクトル演算
  • cuSPARSE: 疎行列演算
  • cuFFT: FFTを行う
  • cuDNN: ディープラーニング

GPUのコンピュート能力

NVIDIAはGPUの持つSMの性能をコンピュート能力という指標で表している。これは各数値によって、SMを構成するプロパティが分かるというものである。各GPUのコンピュート能力を知りたい方はこちらを参照されたい。

CUDAプログラムの実行制御

CUDAにはデータ転送と計算を並行して行えるようにするstreamという危機構も提供されている。このため複数のカーネルを同時に実行する事ができるGPUの場合、そのオーバーラップにより全体的な計算時間を短縮する事ができる。
また、CUDAのプログラムはインオーダ形式で実行されるが、異なるワープ(32スレッド)間では実行の順番は決まっていない。これらはタイミングによって結果が変わってしまう計算(メモリアクセスを利用する関数など)を制御するのに重要であるため別途用意されている。
以下の図のようにA, B, Cという3つのワープが同じ命令を実行しており、それぞれがRead, Writeで同一のアドレスアクセスを必要としている場合、その実行順序によって結果が変わってしまうことが予測される。従って、CUDAでは threadfence_block()という関数を使って、同一スレッドブロック内の全スレッドに於いて、この関数の実行(フェンス)以前に発行された全てのReadとWriteが終了していることを保証することができる。

enter image descripotion here
出展: 「GPUを支える技術」，p188，Hisa Ando，2017


threadfence_block()と似たような機能を持つ関数として、__threadfence()や__threadfence_system()がある。

CUDAの関数実行を同期させるsyncthreads関数

GPU上では1つのワープ（32スレッド）が並行で実行されるが、それらが32を超えた場合、同時には終了しない。また、その実行時間はスレッドにおけるメモリアクセス時間(L2キャッシュ、DDRアクセス、リプレイ回数など）によってよって大きく異なってくる。
このため、CUDAには同じスレッドブロックに含まれる全てのワープの実行が終わるのを待ち合わせる関数__syncthreads()が用意されている。また、似た他の命令として__syncthreads_count(int predicates)、__syncthreads_and()、__syncthreads_or()がある。

ストリーム実行を動悸させるcudaDeviceSynchronize関数

cudaDeviceSynchronize()関数は、ストリーム内の全てのコマンドが終了するのを待ち合わせる。また、cudaStreamSynchronize()は1つのストリームのみを対象として、同様に動作する。

ユニファイドメモリ

4章で詳しく説明した、ユニファイドメモリの具体的な利用方法が記載してある。まず、そのマネージド領域を獲得するためにcudamallocManaged()関数はまたは、__managed__ float A[100]のように指示子を宣言部分に付加する。これらはホストとデバイスの両方からアクセス可能なポインタとなり、明示的なデータコピーを記述する必要がない。これは非常に便利であるが、そのメモリ転送はデータが必要となった時点で行われるので、計算と転送をオーバラップさせることはできない。従って、ストリームなどを用いたカーネル実行と並列に転送を実行するプログラムのほうが高速になる可能性が高い。

複数GPUシステム制御

計算能力向上のため、1つのCPUに複数のGPUを接続する構成が一般的になってきている。CUDAは1台のホストでGPUを最大8台まで制御でき、以下のようなコードで記述される。

size_t size = 1024 * sizeof(float);
cudaSetDevice(0);            // Set device 0 as current
float* p0;
cudaMalloc(&p0, size);       // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1);            // Set device 1 as current
float* p1;
cudaMalloc(&p1, size);       // Allocate memory on device 1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1

出展: http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#multi-device-system


CUDAはある程度は高級化されているけれど、依然としてアプリケーションに対して最適な実装を行うには同期処理、ストリーム、デバイス選択などを細かくコントロールする必要があるように見える。(゜.゜)
CUDAはGPUが広く用いられていることもあって結構進んでいる方だと思うけど、
アルゴリズムと計算構造の分離や自動化はまだまだ今後の課題のようだ。


OpenCL

OpenCLはKhronos Groupが中心となって作成した業界標準のGPUプログラミング環境である。その仕様にはAMD, Apple, ARM, IBM, Intel, Qualcom, NVIDIAなどから多数のエンジニアが参加しているため、各社のGPUで動作するという大きなメリットがある。OpenCLはCUDAと同様にCPU、GPU、DSP等のヘテロジニアスな構成をサポートしている。
OpenCLはCUDAを参考にして作られたため、メモリモデルの思想などが似通っているが、その呼び名は異なる。以下にCUDAとの類似点および用語の対応表をまとめる。

類似点

  • デバイスメモリにglobal, constant, local, privateのメモリ領域がある
  • グリッドとスレッドブロックと共通した考え方で多数のカーネルスレッドを同時に実行する方式

用語対応表

CUDA OpenCL
シェアードメモリ ローカルメモリ
ローカルメモリ プライベートメモリ
スレッド ワークアイテム
スレッドブロック ワークグループ
グリッド NDレンジ
ユニファイドメモリ SVM(Shared Virtual Memory)

また、OpenCLは各社のGPUをサポートするためにより柔軟に構成されている。例えば、CUDAはそのグリッドに含まれるブロックが最大3次元であるが、OpenCLは最大値は任意である（最小値は3次元）。OpenCLからはSPIL-Vという中間表現データが作られるので、各社はそこからGPUの機械語への変換系を用意するだけで対応することができる。OpenCL 2.0ではC++ベース、2.2ではC++14ベースのカーネル記述がサポートされている。

変数

整数型のchar, short, int, longとそれぞれの符号なし型とhalf, float, doubleの浮動小数点型に対応している。また、CUDA同様にそれぞれのベクトル型を用いることができる。

実行環境

OpenCLは1つのCPUと複数GPUで構成されるプラットフォームを想定しているため、実行プログラムはホストで実行する部分とデバイスで実行する部分に分けられる。CUDAと同様に、デバイスで実行されるプログラムをカーネルと呼ぶ。
初めにOpenCLはカーネルオブジェクトをキューに入れ、その実行条件を満たしたら実行を開始する。その後、実行が終了したカーネルをキューから取り除く。キューはインオーダとアウトオブオーダの2つの実行タイプをもち、異なるキュー同士は独立して動作する。このため、複数のキューを用いるとCUDAのストリームのような構成となる。

カーネルの実行

上記の類似点にも記載したように、OpenCLもCUDAと同様に、その最低実行単位はワークアイテム(スレッド)であり、それらをまとめたワークグループ(スレッドブロック)をGPUのCUに割り当てて実行する（以下に図を掲載）。

https://gyazo.com/d562da5f950c56b6a68578f921dd6d28
出展: 「GPUを支える技術」，p198，Hisa Ando，2017


上記の表にある$(sx,sy)$はワークグループ内のワークアイテムの位置であり、組み込み関数get_local_ID()で得ることができる。同様に、そのワークグループのサイズはget_local_size()、グローバルIDはget_global_ID()で取得できる。また、OpenCL2.0からはCUDAのダイナミックパラレリズムに類似した実装も行うことができるようになった。

メモリ構造

OpenCLではホストCPUのメモリとGPUなどのデバイスのメモリを区別して利用することができる。デバイスメモリには1つのワークグループ内のワークアイテムで共有されるローカルメモリと1つのワークアイテムのみが利用できるプライベートメモリがある。以下にKhronosのOpenCL2.1の仕様から抜粋した図を示す。

https://gyazo.com/94bf940eac757183230b8fed5c2cfa15
出展: https://www.khronos.org/registry/OpenCL/specs/opencl-2.1.pdf


CUDAと同様に、ローカルメモリはオンチップSRAM、プライベートメモリはレジスタファイルもしくはデバイスメモリ、グローバルメモリとコンスタントメモリはデバイスメモリを利用する。
メモリ領域はBufferと呼ばれ、その獲得はclCreateBuffer()関数で行われる。また、clEnqueueReadBuffer()関数でバッファのデータをホストのメモリ領域に読み出し、clEnqueueWriteBuffer()関数でホストのメモリ領域からバッフェに書き込み、clEnqueueCopyBuffer()関数でバッファ間のデータコピーを行うことができる。
また、SVMを用いるとCUDAのユニファイドメモリのようにホストとデバイスの両方からアクセスできる領域を作ることができる。OpenCLでは粗粒度のSVMサポートは必須であり、細粒度のSVMサポートはオプションである。どちらにせよ仕様であるので、その性能は実装に依存する。ソースコードのサンプルはAMD APP SDK OpenCL User Guideにのっているので、参考にすると良いだろう。
なお、OpenCLは業界標準なので各社のGPUで動作するが、NVIDIAのサポート状況は芳しくなく、未だに1.2の段階で止まっている。また、CUDAで書いたほうが速度が出やすいという利点もある。しかし、NVIDIA以外のGPUではOpenCLを用いる場合が多い。

CUDAとOpenCLについて触りだけ理解することができた。(^^)
次回はGPUプログラムの最適化からスタートする。

≡

コメント

コメントを投稿

このブログの人気の投稿

GPUを支える技術読み始めた 第2章

2日連続で第2回目の投稿である。 そもそも書き溜めをしてあったり、本ブログを初めたのが土日で比較的時間をとれたのは大きい。 前回でBloggerのエディタではMarkdownやLatexが使えず不便だったので、どうにかならないかと探していたところ、 このブログ を見つけ Classeur というWebエディタを使い始めた。 普段はメモ用として使っているQuiverやAtomと言ったオープンソースの高性能エディタを使うこともできたが、やはり変更のたびに投稿を行わなければいけないのはめんどくさい。ClasseurはMarkdownが使え、かつBloggerと連携して１クリックで更新ができる。いつもながら有益な情報を残してくれた先人に感謝である。 物事を継続するのコツとして、いかにメイン作業以外の手間を減らすが重要だと感じているので、ここでもそれに従うことにする。 もちろん環境の導入やそもそもの事前調査に時間は使ったが、長く続ければ十分にもとがとれるので、初めのうちにやっておいた( ･´ｰ･｀) 上記以外にもLatexを使えるようにするためにMathjaxをBloggerのテンプレートに導入した。 準備は万端、早速書き始めよう。 ２章 計算処理の変革 ゲームと画面描画の内容が多く、少し前回と同様に流し読みぎみ。 初期グラフィックボードである1983 年に Intel が発表した iSBX 275というボードは、256 x 256ピクセル解像度で8種類の色を利用可能だった。 続く1996年にha3dfx Interactive社がVoodooシリーズを、1999年にはNVIDIAはGeForde 256を発表した。 GeForce 256を含む初期GPUの主な目的は、3Dグラフィックスに必要なT & L(Transpose & Lighting)処理を高速に行うためであったが、当初の性能はハイエンドCPUに負ける程度であった。その後、ムーアの法則に従って搭載できるトランジスタの量が増えるにつれて、CPUの１０倍以上の性能を出せるようになっていく。 当初はTransposeとLightingの処理は別々のパイプラインとなっていたが、 リソースを柔軟に使いまわしすことで、効率を上げるUnified shaderが用いられるよう
コメントを投稿
続きを読む

GPUを支える技術読み始めた 第4章 [前半]

今年もあと少しになってきたが、なんとか目標だったもう一本を投稿することができて良かった。 私情だが、先日の社内年末パーティでは年間MVPに選出していただいた。＼(^o^)／ 非常に嬉しく思うと同時に、いろんな面でサポートをしてくれたHWチームメンバやバックオフィスに感謝したい。 また、来年は社内だけでなくて社外にも影響を与えられるよう頑張っていきたい。 少し気が早いが、来年度の本ブログの方針として「基礎と応用」というコンセプトで書いていきたと考えている。古典的名著と最新の論文の要約などができたら上出来だろうか。（＾ω＾） 直近だと、DeepLearning × HWに関するの新しめの論文や並列処理技法系の本のまとめを計画している。もしかすると、年内にまだいけるかもしれない。 良い報告もできたところで、早速続きを初めていこうと思う。 4章は個人的には一番楽しみな章でじっくり読んでいる。内容が多いので前半と後半に分割して投稿していく。 4章 GPUの超並列処理 [前半] GPUの並列処理方式 先の章で並列処理方式について以下のように説明した。 SIMD: 1つの計算を幾つかのデータに対して並列に実行する SIMT: 1つの計算を別々の演算機で並列に実行する 4章では上記2つについてもう少し詳しく解説している。 SIMD方式 以下2つのベクトルXと行列Aがあるとする。 \[ \begin{align} \bf{X} &= (a, b, c) \\ \bf{A} &= \left( \begin{array}{ccc} a00 & a01 \\ a10 & a12 \\ a21 & a21 \end{array} \right) \\ \bf{Y} &= \bf{X} \cdot \bf{A} \end{align} \ \] Yを計算する時、SIMDでは先にXの列要素(a)を各演算機にブロードキャストし、Aの行要素(a00, a10, a20)と計算する。この動作をXの列要素分繰り返すことで計算を完了する。 仮に、Xの要素がシェアードメモリ(後述)など、レイテンシのあるメモリに格納されている場合、各ブロードキャストでサイク
コメントを投稿
続きを読む