GPU のアーキテクチャとプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓

GPUのアーキテクチャと

プログラム構造

今回の内容

2015/04/22 GPGPU実践プログラミング 58



_{GPUのアーキテクチャ}



_CUDA



_{CUDAによるプログラミング}

GPU（Graphics Processing Unit）とは



画像処理専用のハードウェア



具体的には画像処理用のチップ



チップ単体では販売されていない



PCI‐Exカードで販売（チップ単体と区別せずにGPUと呼ぶことも多い）



ノートPCに搭載



_{PCI‐Exカードとして販売されるGPUには，ビデオメモリと呼ばれ}

るRAMが搭載

GPU（Graphics Processing Unit）の役割



グラフィックスを表示するために様々な処理を行い，処

理の結果をディスプレイに出力



_{3次元グラフィックスの発展に伴って役割が大きく変化}

3次元座標変換 ポリゴンとピクセルの 対応付け ピクセル色計算 テクスチャ参照 フレームバッファ（ビデ オメモリ）への書き込み ディスプレイ出力 CPU GPU 3次元座標変換 ポリゴンとピクセルの 対応付け ピクセル色計算 テクスチャ参照 フレームバッファ（ビデ オメモリ）への書き込み ディスプレイ出力

現在

過去

 _{CPU が 3D 描 画} の演算を実行  _GPUが出力 描画情報 画面出力  _{GPUが演算から} 出力までの全て を担当  _{CPUは描画情報} の生成やGPUへ の情報の引き渡 し ， GPU の 制 御 を行う 描画情報 画面出力 2015/04/22 GPGPU実践プログラミング 60 ディスプレイコントローラ

GPUの描画の流れ

1.

CPUからGPUへ描画情報を送信

2.

頂点処理（頂点シェーダ）



座標変換



画面上での頂点やポリゴンの位置・大きさの決定



頂点単位での照明の計算

3.

頂点やポリゴンからピクセルを生成

（ラスタライザ）

4.

ピクセル処理（ピクセルシェーダ）



画面上のピクセルの色



テクスチャの模様

5.

画面出力



ピクセルの色情報をフレームバッファに書き込み

2. 3. 4.

ビデオカードから

GPUへ



_{CGの多様化と共に固定機能の実装が困難に}



頂点処理とピクセル処理をユーザが書き換えられるプロ

グラマブルシェーダの実装

頂点処理用回路 ピクセル処理用回路

グラフィックスカード

頂点シェーダユニット ピクセルシェーダユニット

GPU

2015/04/22 GPGPU実践プログラミング 62

レンダリングパイプライン処理

頂点情報 光源情報 視野変換 陰影計算 材質情報 投影変換 クリッピング ビューポート変換 走査変換 合成 テクスチャ 出力画像 投影像を画素 へ変換 整数演算とメモリ アクセス 形状データの画面 への投影像 実数演算

レンダリングパイプライン処理

2015/04/22 GPGPU実践プログラミング 64 頂点情報 光源情報 視野変換 陰影計算 材質情報 投影変換 クリッピング ビューポート変換 走査変換 合成 テクスチャ 出力画像 形状データの画面 への投影像 ハードウェアで処理 （固定機能） 実数演算 実数演算を行うハード ウェアは高価だった

レンダリングパイプライン処理

頂点情報 光源情報 視野変換 陰影計算 材質情報 投影変換 クリッピング ビューポート変換 走査変換 合成 テクスチャ 出力画像 ハードウェアで処理 （固定機能） ハードウェアで処理 （固定機能）

レンダリングパイプライン処理

2015/04/22 GPGPU実践プログラミング 66 頂点情報 光源情報 視野変換 陰影計算 材質情報 投影変換 クリッピング ビューポート変換 走査変換 合成 テクスチャ 出力画像 ピクセルシェーダ 頂点シェーダ

ビデオカードから

GPUへ



頂点処理とピクセル処理をユーザが書き換えられるプロ

グラマブルシェーダの実装



処理によっては利用効率に差が生じる

GPU 頂点シェーダユニット ピクセルシェーダユニット

頂点処理重視の処理

GPU 頂点シェーダユニット ピクセルシェーダユニット

ピクセル処理重視の処理

空きユニット 空きユニット

ビデオカードから

GPUへ



頂点シェーダとピクセルシェーダを統合したユニファイド

シェーダへの進化



頂点処理とピクセル処理を切り替えることで利用率を高める

GPU ユニファイドシェーダユニット

頂点処理重視の処理

ピクセル処理重視の処理

GPU ユニファイドシェーダユニット 2015/04/22 GPGPU実践プログラミング 68

ビデオカードから

GPUへ



各ピクセルに対して処理を並列に実行



単純な処理を行う演算器(Streaming Processor, 

SP)を大量に搭載



演算器は現在CUDA Coreという名称に変更



高い並列度で処理を行う

Teslaアーキテクチャの構造

2015/04/22 GPGPU実践プログラミング 70



_{Tesla C1060の仕様}



_SM数

₃₀



_{CUDA Core数}

_{240(=8 Core/SM×30 SM)}



キャッシュを搭載せず

Teslaアーキテクチャの構造



_{Tesla C1060の仕様}

CUDAコア数（単精度） 240 Cores CUDAコアクロック周波数 1,296 MHz 単精度演算ピーク性能 _622*1 _(933*2_) GFLOPS 倍精度演算ユニット数 _30*3 _Units 倍精度演算ピーク性能 _78 GFLOPS メモリクロック周波数 _800 MHz メモリバス幅 _512 bit 最大メモリバンド幅*4 _102 GB/s *1_{単精度演算ピーク性能 = コアクロック周波数×コア数×命令の同時発行数（2）} *2_{CUDA CoreとSFUが同時に命令を発行できれば1296 MHz×240×3} *3_{一つのSMに倍精度演算器が一つ搭載} *4_{最大メモリバンド幅=メモリクロック周波数×メモリバス幅/8×2(Double Data Rate)}

Fermiアーキテクチャの構造

2015/04/22 GPGPU実践プログラミング 72



_{Tesla M2050の仕様}



_SM数

₁₄



_{CUDA Core数}

_{448(=32 Core/SM×14 SM)}



_{L1/L2 キャッシュを搭載}



_{ECC（誤り訂正機能）を搭載}

Fermiアーキテクチャの構造



_{Tesla M2050の仕様}

CUDAコア数（単精度） 448 Cores CUDAコアクロック周波数 1,150 MHz 単精度演算ピーク性能 _1.03 TFLOPS 倍精度演算ユニット数 _0*1 _Unit 倍精度演算ピーク性能 _515 GFLOPS メモリクロック周波数 _1.55 GHz メモリバス幅 _384 bit 最大メモリバンド幅 _148 GB/s *1_{単精度CUDA Coreを2基使って倍精度演算を実行}

Keplerアーキテクチャの構造

2015/04/22 GPGPU実践プログラミング 74



_{Tesla K20c/mの仕様}



_SMX数

₁₃



_{Streaming Multiprocessor eXtreme (?)}



_{CUDA Core数}

_{2,496(=192 Core/SM×13 SMX)}

Keplerアーキテクチャの構造



_{Tesla K20c/mの仕様}

CUDAコア数（単精度） 2,496 Cores CUDAコアクロック周波数 706 MHz 単精度演算ピーク性能 _3.52 TFLOPS 倍精度演算ユニット数 _832*1 _Units 倍精度演算ピーク性能 _1.17 TFLOPS メモリクロック周波数 _2.6 GHz メモリバス幅 _320 bit 最大メモリバンド幅 _208 GB/s *1_{64基/SMX×13基}

Maxwellアーキテクチャ



_{GeForce GTX TITAN Xの仕様}



_SM数

₂₄



_{CUDA Core数}

_{3,072(=128 Core/SM×24 SM)}

2015/04/22 GPGPU実践プログラミング 76

Maxwellアーキテクチャ



_{GeForce GTX TITAN Xの仕様}

* CUDAコア数（単精度） 3,072 Cores CUDAコアクロック周波数 1,002 MHz 単精度演算ピーク性能 _{6.14 TFLOPS} 倍精度演算ユニット数 _0*1 _Units 倍精度演算ピーク性能 _192 GFLOPS*2 メモリクロック周波数 _3.5 GHz*3 メモリバス幅 _384 bit 最大メモリバンド幅 _336.5 GB/s *1_{http://www.4gamer.net/games/121/G012181/20141225075/} *2_{倍精度演算は単精度演算の性能の1/32 (1/16 Flop/Core/clock)} *3_{DDR(Double Data Rate) 7GHz相当と書かれている場合もある} http://http://www.geforce.com/hardware/desk top‐gpus/geforce‐gtx‐titan‐x/specifications *http://ja.wikipedia.org/wiki/FLOPS

Pascalアーキテクチャ



_{2016年にリリース予定}



倍精度演算器を搭載予定



_NVLink



_{GPU同士やGPUとCPUを接続する独自の方式}



通信（CPU ↔ メモリ ↔ PCI Express ↔ メモリ ↔ GPU）の

ボトルネックを解消（PCI Express3.0の5~12倍）



複数のGPUを使って大規模な計算が可能



_{3Dメモリ（High Bandwidth Memory, HBM）*}



_{3次元積層技術を利用し，メモリの容量と帯域を大幅に増加}



最大32GB，メモリ帯域1TB/s

2015/04/22 GPGPU実践プログラミング 78 *http://pc.watch.impress.co.jp/docs/column/kaigai/20150421_698806.html

Voltaアーキテクチャ



_{Pascalの後継}



詳しい情報は不明



アメリカの次世代スーパーコンピュータへ採用予定



オークリッジ国立研究所

_SUMMIT

_{150~300PFLOPS}



ローレンス・リバモア研究所 SIERRA

_{100PFLOPS以上}



地球シミュレータと同等の演算性能を1ノードで実現



現在Top500 2位のスーパーコンピュータと同じ電力で5~10

倍高速，サイズは1/5

CUDA 

Core

CUDA 

_Core

CUDA 

_Core

CUDA 

Core

CUDA 

_Core

CUDA 

Core

CUDA 

_Core

CUDA 

_Core

Streaming  Multiprocessor

GPUの模式図

2015/04/22 GPGPU実践プログラミング 80

GPU

Streaming  Multiprocessor L2キャッシュ ローカル メモリ コンスタントメモリ テクスチャメモリ GPU Chip グローバルメモリ ローカル メモリ SM SM SM ・・・ SM ・・・ SM SM SM ・・・ SM _レジ スタ レジ スタ レジ スタ レジ スタ CUDA  Core CUDA  Core CUDA  Core CUDA  Core L1キャッ シュ 共有 メモリ Streaming Multiprocessor SM SM SM ・・・ SM

GPUの並列化の階層



グリッド－ブロック－スレッドの3階層



グリッド（Grid）



並列に実行する処理



_{GPUが処理を担当する領域全体}



スレッド（Thread）



_{GPUの処理の基本単位}



_{CPUのスレッドと同じ}



ブロック（Block）もしくはスレッドブロック



スレッドの集まり

GPUの並列化の階層



_{GPUのハードウェアの構成に対応させて並列性を管理}

GPU

Streaming  Multiprocessor CUDA  Core ハードウェア構成

並列に実行する

処理

スレッドの集

まり

スレッド

並列化の階層

Grid

Block

Thread

CUDA 2015/04/22 GPGPU実践プログラミング 82

CUDA



_{Compute Unified Device Architecture}



_{NVIDIA社製GPU向け開発環境(Windows,Linux,Mac OS X)}



_{2007年頃発表}



_{C/C++言語＋独自のGPU向け拡張}



専用コンパイラ（nvcc）とランタイムライブラリ



いくつかの数値計算ライブラリ（線形代数計算，FFTなど）



_{CUDA登場以前}



グラフィックスプログラミングを利用



足し算を行うために，色を混ぜる処理を実行



汎用計算のためには多大な労力が必要

CUDAによるプログラミング



_{CPUをホスト（Host），GPUをデバイス（Device）と表現}



ホスト（CPU）



処理の流れやGPUを利用するための手続きを記述



プログラムの書き方は従来のC言語と同じ



利用するGPUの決定，GPUへのデータ転送，GPUで実行する関

数の呼び出し等

2015/04/22 GPGPU実践プログラミング 84

CUDAによるプログラミング



_{CPUをホスト（Host），GPUをデバイス（Device）と表現}



デバイス（GPU）



処理する内容を関数として記述



引数は利用可能，返値は利用不可（常にvoid）



関数はkernelと呼ばれる



関数呼び出しはlaunch, invokeなどと呼ばれる

Hello World



何を確認するか



最小構成のプログラムの作り方



ファイル命名規則（拡張子は.c/.cpp）



コンパイルの方法(gcc, cl等を使用)

#include<stdio.h>

int main(void){

printf("hello world¥n");

return 0;

}

2015/04/22 GPGPU実践プログラミング 86 helloworld.c

CUDAでHello World



何を確認するか



最小構成のプログラムの作り方



ファイル命名規則（拡張子は

_.cu

）



コンパイルの方法（

_nvcc

を使用

）

#include<stdio.h>

int main(void){

printf("hello world¥n");

return 0;

}

#include<stdio.h>

int main(void){

printf("hello world¥n");

return 0;

}

違いは拡張子だけ？

CUDAプログラムのコンパイル



ソースファイルの拡張子は

_.cu



_{nvccを用いてコンパイル}



_{CPUが処理する箇所はgcc等がコンパイル}



_{GPUで処理する箇所をnvccがコンパイル}



_{helloworld.cuにはCPUで処理する箇所しかない}

2015/04/22 GPGPU実践プログラミング 88

CUDAでHello World



_{CUDA専用の処理を追加}

#include<stdio.h>

__global__ void kernel(){}

int main(void){

kernel<<<1,1>>>();

printf("hello world¥n");

return 0;

}

GPUで実行される関数（カーネル） __global__が追加されている ・・・ 通常の関数呼出とは異なり， <<<>>>が追加されている ・・・ helloworld_kernel.cu

CUDAプログラムの実行



実行時の流れ

（CPU視点）



利用するGPUの初期化やデータの転送などを実行



_{GPUで実行する関数を呼び出し}



_{GPUから結果を取得}

初期化の指示

初期化

カーネルの実行指示

カーネルを実行

結果の取得

実行結果をコピー

time

CPUとGPUは非同期 CPUは別の処理を実行可能 2015/04/22 GPGPU実践プログラミング 90

必要なデータのコピー

メモリに書込

CPU

GPU

GPUの構造とカーネルの書き方



プログラムからGPUで実行する関数を呼出



GPUで実行する関数という目印が必要



_{GPUはPCI‐Exバスを経由してホストと接続}



_{GPUはホストと別に独立したメモリを持つ}



関数の実行に必要なデータはGPUのメモリに置く



_{GPUはマルチスレッド（メニースレッド）で並列処理}



関数には1スレッドが実行する処理を書く



関数を実行する際に並列処理の度合いを指定

GPUの構造とカーネルの書き方



_{GPUで実行する関数（カーネル）という目印}



修飾子__global__を付ける



_{GPUはPCI‐Exバスを経由してホストと接続}



_{GPUはホストと別に独立したメモリを持つ}



カーネルの返値をvoidにする



_{GPUはマルチスレッド}

（メニースレッド）

で並列処理



カーネルには1スレッドが実行する処理を書く



カーネル名と引数の間に<<<1,1>>>を付ける

2015/04/22 GPGPU実践プログラミング 92

Hello Thread（Fermi世代以降）



_{GPUの各スレッドが画面表示}

#include<stdio.h>

__global__ void hello(){

printf("Hello Thread¥n");

}

int main(void){

hello<<<1,1>>>();

cudaThreadSynchronize();

return 0;

}

画面表示（Fermi世代以降で可能） コンパイル時にオプションが必要 ‐arch=sm_20以降 ・・・ カーネル実行 ・・・ ホストとデバイスの同期をとる CPUとGPUは原則同期しないので， 同期しないとカーネルを実行した 直後にプログラムが終了 ・・・ hellothread.cu

Hello Thread（Fermi世代以降）



_{<<< >>>内の数字で並列度が変わることの確認}

2015/04/22 GPGPU実践プログラミング 94

#include<stdio.h>

__global__ void hello(){

printf("Hello Thread¥n");

}

int main(void){

hello<<<

?

,

?

>>>();

cudaThreadSynchronize();

return 0;

}

<<<>>>内の数字を変えると画面 表示される行数が変わる <<<1,8>>>, <<<8,1>>>,  <<<4,2>>>等 ・・・ hellothread.cu

CPUとGPUのやりとり



_{GPUの想定される使い方}



ホスト(CPU)からデータを送り，デバイス(GPU)で計算し，

結果を受け取る



_{CPUとGPUのデータのやり取りが必要}



_{GPUは原則データを返さない}



_{PCI‐Ex経由で描画情報を受け取り，画面に出力}



カーネルの返値がvoidの理由

CPUとGPUのやりとり

2015/04/22 GPGPU実践プログラミング 96



_{CUDA独自の命令とC言語のポインタを利用}



_{GPUのメモリ上に計算に必要なサイズを確保}



確保したメモリのアドレスをC言語のポインタで格納



ポインタの情報を基にデータを送受信

CPUとGPUのやり取り（単純な加算）



_{int型の変数2個を引数として受け取り，2個の和を返す}



_{C言語らしい書き方}

#include<stdio.h>

int add(int a, int b){ return a + b; } int main(void){ int c; c = add(6, 7); printf("6 + 7 = %d¥n", c); return 0; } 引数で渡された変数の和を返す ・・・ 関数呼び出し ・・・ add_naive.c

CPUとGPUのやり取り（単純な加算）

2015/04/22 GPGPU実践プログラミング 98



関数の返値をvoidに変更し，メモリの動的確保を使用

#include<stdio.h> #include<stdlib.h>

void add(int a, int b, int *c){ *c = a + b;

}

int main(void){ int c;

int *addr_c;

addr_c = (int *)malloc(sizeof(int)); add(6, 7, addr_c); c = *addr_c; printf("6 + 7 = %d¥n", c); return 0; } 引数で渡された変数の和を，cが指す アドレスに書き込み ・・・ 引数にアドレスを追加 ・・・ アドレスを基に結果を参照 ・・・ add.c

CPUプログラム（メモリの動的確保）



_malloc



指定したバイト数分のメモリを確保



_{stdlib.hをインクルードする必要がある}



_sizeof



データ型1個のサイズ（バイト数）を求める

#include<stdlib.h> int *a; a = (int *)malloc( sizeof(int)*100 ); printf("%d, %d¥n", sizeof(float), sizeof(double)); 実行すると4,8と表示される

CPUとGPUのやり取り（単純な加算）

2015/04/22 GPGPU実践プログラミング 100



_{add.cの処理の一部をGPUの処理に置き換え}

#include<stdio.h>

__global__ void add(int a, int b, int *c){ *c = a + b; } int main(void){ int c; int *dev_c; cudaMalloc( (void **)&dev_c, sizeof(int) ); add<<<1, 1>>>(6, 7, dev_c); cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost); printf("6 + 7 = %d¥n", c); cudaFree(dev_c); return 0; } __global__を追加 ・・・ GPU上のメモリに確保される変数のアドレス ・・・ GPU上にint型変数 一個分のメモリを確保 ・・・ ↑GPUから結果をコピー メモリを解放 ・・・ add.cu

CUDAでカーネルを作成するときの制限



カーネルの引数



値を渡すことができる



_{GPUのメモリを指すアドレス}



_{CPUのメモリを指すアドレスも渡すことは可能}



そのアドレスを基にホスト側のメモリを参照することは不可能



_{printfなどの画面出力}



_{Fermi世代以降のGPUで，コンパイルオプションを付与}



_{‐arch={sm_20|sm_21|sm_30|sm_32|sm_35|sm_50|sm_52}}



エミュレーションモード



新しいCUDA（4.0以降）では消滅

CPUプログラムの超簡単移植法



とりあえずGPUで実行すればいいのなら･･･



拡張子を.cuに変更



_{GPUの都合を反映}



関数の返値をvoidにし，__global__を付ける



関数名と引数の間に<<<1,1>>>を付ける



_{GPUで使うメモリをcudaMallocで確保}



_{mallocでメモリを確保していればそれをcudaMallocに置き換え}



_{GPUからデータを受け取るためにcudaMemcpyを追加}



最適化は追々考えればいい

2015/04/22 GPGPU実践プログラミング 102 カーネルの完成

Hello Thread（Fermi世代以降）



_{<<< >>>内の数字で並列度が変わる}



この情報を利用すれば並列処理が可能

#include<stdio.h> __global__ void hello(){ printf("Hello Thread¥n"); } int main(void){ hello<<<?,?>>>(); cudaThreadSynchronize(); return 0; } <<<>>>内の数字を変えると画面表示 される行数が変わる <<<1,8>>>, <<<8,1>>>,  <<<4,2>>>等 ・・・ hellothread.cu

GPUの並列化の階層



_{GPUのハードウェアの構成に対応させて並列性を管理}



並列化の各階層における情報を利用

GPU

Streaming  Multiprocessor CUDA  Core ハードウェア構成

並列に実行する

処理

スレッドの集

まり

スレッド

並列化の階層

Grid

Block

Thread

CUDA 2015/04/22 GPGPU実践プログラミング 104

GPUの並列化の階層



グリッド－ブロック－スレッドの3階層



各階層の情報を参照できる変数



_{x,y,zをメンバにもつdim3型構造体}



グリッド（Grid）



_gridDim

グリッド内にあるブロックの数



ブロック（Block）



_blockIdx

ブロックに割り当てられた番号



_blockDim

ブロック内にあるスレッドの数



スレッド（Thread）



_threadIdx

スレッドに割り当てられた番号

Hello Threads（Fermi世代以降）



_{<<< >>>内の数字で表示される内容が変化}

2015/04/22 GPGPU実践プログラミング 106 #include<stdio.h> __global__ void hello(){ printf("gridDim.x=%d, blockIdx.x=%d, blockDim.x=%d, threadIdx.x=%d¥n", gridDim.x, blockIdx.x, blockDim.x, threadIdx.x); } int main(void){ hello<<<?,?>>>(); cudaThreadSynchronize(); return 0; } <<<>>>内の数字を変えると画面表示 される内容が変わる <<<>>>内の数字とどのパラメータが 対応しているかを確認 ・・・ hellothreads.cu

GPUの構造とカーネルの書き方



_{GPUはマルチスレッド}

（メニースレッド）

で並列処理



関数には1スレッドが実行する処理を書く



関数を実行する際に並列処理の度合いを指定



カーネルと引数の間に追加した<<<,>>>で並列処理の

度合を指定



_{<<<グリッド内にあるブロックの数,1ブロックあたりのスレッド}

の数>>>

プログラム実習



以下のプログラムをコンパイルし，正しく実行できること

を確認せよ



_helloworld.c



_{helloworld.cu}



_{hellothread.cu}



_{hellothreads.cu}



_{hellothreads.cuについては，<<<>>>内の数字を変}

更し，実行結果がどのように変わるか確認せよ

GPGPU実践プログラミング 108 2015/04/22

レポート課題

1（提出期限は1学期末）



二つの値を交換する関数swapをGPUに移植せよ



並列化する必要はなく，1スレッドで実行すればよい

#include<stdio.h>

void swap(int *addr_a, int *addr_b){

int c;       //cは値を一時的に保持するための変数 c = *addr_a; //*は間接参照演算子

*addr_a = *addr_b; //メモリアドレス(=addr_a,addr_bの値)にある変数の値を参照 *addr_b =  c; } int main(void){ int a=1,b=2; printf("a = %d, b = %d¥n", a, b); swap(&a, &b); //変数a, bのメモリアドレスを渡す．&はアドレス演算子 printf("a = %d, b = %d¥n", a, b); return 0; } swap.c

レポート課題

1（提出期限は1学期末）

GPGPU実践プログラミング 110 2015/04/22 #include<stdio.h> カーネルという目印 _{void swap(int *a, int *b){} } int main(void){ int a=1,b=2; GPUで使う変数を宣言 printf("a = %d, b = %d¥n", a, b); GPU上のメモリを確保（aの分） GPU上のメモリを確保（bの分） CPUからGPUにメモリの内容をコピー（aの分） CPUからGPUにメモリの内容をコピー（bの分）

swap実行時の並列度の指定(GPUで使う変数, GPUで使う変数);

GPUのメモリの内容をCPUにコピー（aの分） GPUのメモリの内容をCPUにコピー（bの分） printf("a = %d, b = %d¥n", a, b); return 0; } ?