GPUとそのプログラム

天野

アクセラレータとは？

• 特定の性質のプログラムを高速化するプロセッサ
• 典型的なアクセラレータ
• GPU(Graphics Processing Unit)
• Xeon Phi
• FPGA(Field Programmable Gate Array)
• 最近出て来たDeep Learning用ニューロチップなどDomain Specific Architecture

AI用DSA（Domain Specific Architecture)

クラウド用

エッジ用

巨大IT企業が直接DSAを作る時代になった！

①GPGPU：General Purpose Computing with GPUグラフィックプロセッサをアクセラレータとして使う

• TSUBAME2.0(Xeon+Tesla,Top500 2010/11 4^th )
• 天河一号(Xeon+FireStream,2009/11 5^th )

​

​

​

​

​

※()内は開発環境

PBSM

PBSM

Thread Processors

PBSM

PBSM

Thread Processors

PBSM

PBSM

Thread Processors

PBSM

PBSM

Thread Processors

PBSM

PBSM

Thread Processors

…

Thread Execution Manager

Input Assembler

Host

Load/Store

Global Memory

GeForce

GTX280

240 cores

GPU　(NVIDIA’s GTX580)

​

512 GPU cores ( 128 X 4 )

768 KB L2 cache

40nm CMOS 550 mm^2

128個のコアは

SIMD動作をする

４つのグループは

独立動作をする

​

もちろん、このチップを

たくさん使う

NVIDIAのGPUの名前が訳が分からん問題

• 目的用途別の名前とアーキテクチャの名前が混乱しがち
• 目的別製品シリーズの名前
• デスクトップ用、ゲーム用：GeForce（ジーフォース）
• GeForce GTX＞GeForce GT＞GeForceで高性能
• TITAN Xというグラフィック用のカードがあるがこれはPascalアーキテクチャを使っている。
• コスト性能比が高い
• プロ用：Quadro
• 使ったことがないので良く分からないが凄そう
• モバイル用：Tegra
• 車載などの用途のための低電力
• Tegra X1:Maxwell アーキテクチャを使っている
• Tegra K1:Keplarアーキテクチャを使っている
• Tegra 3,2はGPUが付いていないARMだけ
• 高性能用（AI用）：Tesla
• 以前はGPGPU用のをTeslaと呼んでいたが最近は大きくAI用にシフトした
• Tesla P100：Pascalアーキテクチャ
• Tesla V100： Voltaアーキテクチャ
• アーキテクチャの名前
• Fermi, Maxwell, Kepler, Pascal, Volta
• プロセッサの構造を示す
• どんどん新しいのが出てきて追従できない。。。

​

CoolChips　2019のKeynoteのスライドより引用

CoolChips　2019のKeynoteのスライドより引用

CUDA/OpenCL

• CUDA はNVIDEAのGPUプログラム用の言語
• ホストプログラムとデバイス（GPU）側のプログラムに分離
• データに3次元的なスレッドを割り当てる
• 32スレッド＝Warp
• SIMDプログラミング
• プログラマがメモリのレベルを考える
• OpenCLは、ベンダに依存しない標準言語
• 考え方はCUDAに似ている
• FPGAでも使える

​

なんといっても本家を見よう

• https://http.download.nvidia.com/developer/cuda/jp/CUDA_Programming_Basics_PartI_jp.pdf
• https://http.download.nvidia.com/developer/cuda/jp/CUDA_Programming_Basics_PartII_jp.pdf

​

アクセラレータのプログラム

…

ホストのプログラム

Device

…

ホストのプログラム

Device

CPU:Serial Code

Parallel Kernel

KernelA(args);

アクセラレータ

CPU:Serial Code

Parallel Kernel

KernelB(args);

アクセラレータ

ホストのプログラムが準備してアクセラレータのプログラムにデータを渡す

処理が終わったら回収

CUDA、OpenCLはこの考え方を取る

スレッドとスレッドブロック

0

1

2

3

4

5

6

7

Thread Block 0

threadID

…

float x =

input[threadID];

float y=func(x);

output[threadID]=y;

…

0

1

2

3

4

5

6

7

Thread Block 1

…

float x =

input[threadID];

float y=func(x);

output[threadID]=y;

…

0

1

2

3

4

5

6

7

Thread Block N-1

…

float x =

input[threadID];

float y=func(x);

output[threadID]=y;

…

…

各スレッドは同じコードを実行

同一スレッドブロック内のスレッドはバリア同期

_syncthreads();

スレッドブロック間では同期されない。

CUDA threadはスレッドIDを使って各データへ割り付ける

メモリ階層

Thread

Per-thread

Local memory

Block

Per-block

Shared

Memory

…

…

Per-device

Global

Memory

Kernel 0

Kernel 1

Kernelは

順番に実行

ホストのメモリとの間では

cudaMemcpy();

を用いて転送

ログインとサンプルプログラムの実行

まずITCにログイン

• ssh login_name@XXXX.educ.cc.keio.ac.jp

https://keio.box.com/s/zd3i2rhx0254bkqms7ky5lq5jbojy2yu

次にGPUが使えるマシンにログイン

• ssh ex20XX@comparc01.am.ics.keio.ac.jp
• ログインアカウントはメールで配布された　はず
• ファイルの転送
• wget http://www.am.ics.keio.ac.jp/arc/cuda_ex1.tar
• tar xvf cuda_ex1.tar
• cd ex1
• make sample1
• nvcc sample1.cu sample1_kernel.cu –o sample1
• ./sample1
• 時間計測付きはmake sample1_timeを実行

​

​

今回使うGPU：NVIDIA Pascal

• グラフィックスプロセッサ NVIDIA® GP100
• CUDAコアプロセッサ数 3584 コア
• ベースクロック 1189 MHz
• ブーストクロック 1328 MHz
• 半精度浮動小数点性能 18.7 TFLOPS(最大ブースト)
• 単精度浮動小数点性能 9.3 TFLOPS(最大ブースト)
• 倍精度浮動小数点性能 4.7 TFLOPS(最大ブースト)
• メモリ 16 GB HBM2(バンド帯域幅 720 GB/s) or 12GB HBM2 (バンド帯域幅 540GB/s)

サンプルプログラム (sample1.cu, sample_kernel1.cu)

• 浮動小数の二つの配列の和を求める
• プログラムの流れ:
• ホストでの前処理
1. デバイス（GPU)でのメモリ割り付け
2. ホストからデータ転送
• Kernel 呼び出し→ここでGPUで実行
• ホストでの後処理
• デバイスからデータ転送
• ホストでの処理
• デバイスのメモリの解放

​

​

#include <stdio.h>

#include <stdlib.h>

#include "header.h" // Library files

int main(int argc, char **argv) {

float *h_A, *h_B, *h_C; // variables in the host

float *d_A, *d_B, *d_C; // variables in the device

float result = 0.0f; // results

dim3 dim_grid(LENGTH/BLOCK_SIZE, 1); // For kernel call

dim3 dim_block(BLOCK_SIZE, 1, 1); //

// Allocation in the host memory and Generation of array

h_A = (float *)malloc(sizeof(float) * LENGTH);

h_B = (float *)malloc(sizeof(float) * LENGTH);

h_C = (float *)malloc(sizeof(float) * LENGTH);

for (int i = 0; i < LENGTH; ++i) {

h_A[i] = 1.0f; h_B[i] = 2.0f; h_C[i] = 0.0f; }

host: sample.cu

ホストでの初期化

dim3 dim_grid(LENGTH/BLOCK_SIZE, 1); // For kernel call

ブロックによるグリッドの次元　（2次元(3次元の定義もOK)）

ブロック数　dim_grid.x * dim_grid.y

dim3 dim_block(BLOCK_SIZE, 1, 1); //

スレッドによるブロックの次元　（3次元）

スレッド数 dim_block.x * dim_block.y* dim_block.z

…..

Sample1Kernel<<<dim_grid, dim_block>>>(d_A, d_B, d_C);

dim3: 組み込みデバイス変数

​

<<<… >>> がCUDA独特の記法

​

​

Kernel 呼び出し (host: sample1.cu)

//デバイスのメモリ割り当て

cudaMalloc((void **)&d_A, sizeof(float) * LENGTH);

cudaMalloc((void **)&d_B, sizeof(float) * LENGTH);

cudaMalloc((void **)&d_C, sizeof(float) * LENGTH);

// デバイスへのデータコピー

cudaMemcpy(d_A, h_A, sizeof(float) * LENGTH,

cudaMemcpyHostToDevice);

cudaMemcpy(d_B, h_B, sizeof(float) * LENGTH,

cudaMemcpyHostToDevice);

Sample1Kernel<<<dim_grid, dim_block>>>(d_A, d_B, d_C);

​

​

メモリ割り当てとデータのコピー

カーネル呼び出しの例

blockIdx.x=0

blockDim.x=4

threadIdx.x=0,1,2,3

idx=0,1,2,3

blockIdx.x=1

blockDim.x=4

threadIdx.x=0,1,2,3

idx=4,5,6,7

blockIdx.x=2

blockDim.x=4

threadIdx.x=0,1,2,3

idx=8,9,10,11

blockIdx.x=3

blockDim.x=4

threadIdx.x=0,1,2,3

idx=12,13,14,15

LENGTH=16, BLOCK_SIZE=4の場合

int idx = blockDim.x * blockId.x + threadldx.x;

により、ローカルindexであるthreadldxをグローバルなidxにマップしている

blockDim は実際のコードでは 32 以上でないとまずい

多い分はメモリアクセス遅延の隠蔽に使われ、有効な場合もある

実行モデル

Kernel

1

Block

(0,0)

Block

(1,0)

Block

(2,0)

Block

(0,1)

Block

(1,1)

Block

(2,1)

Grid1

Device

Host

Kernel

2

Block

(0,0)

Block

(1,0)

Block

(2,0)

Block

(0,1)

Block

(1,1)

Block

(2,1)

Grid2

Thread

(0,0)

…

Thread

(31,0)

Thread

(32,0)

…

Thread

(63,0)

Warp 0

Warp 1

Thread

(0,1)

…

Thread

(31,1)

Thread

(32,1)

…

Thread

(63,1)

Warp 2

Warp 3

Thread

(0,2)

…

Thread

(31,2)

Thread

(32,2)

…

Thread

(63,2)

Warp 4

Warp 5

Block (1,1)

Block内の

32スレッドはWarpという単位で

並列実行される

Kernel: sample1_kernel.cu

__global__ void Sample1Kernel(float *d_A, float *d_B,

float *d_C) {

// Getting its thread id

int thread_id = blockDim.x * blockIdx.x + threadIdx.x;

// Compute sum of array

d_C[thread_id] = d_A[thread_id] + d_B[thread_id];

}

thread_idを使うことで、一重分ループを並列実行することができる

// 結果のホストへのコピー

cudaMemcpy(h_C, d_C, sizeof(float) * LENGTH, cudaMemcpyDeviceToHost);

// デバイスメモリの解放

cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);

// 結果のプリント

for (int i = 0; i < LENGTH; ++i) result += h_C[i];

result /= (float)LENGTH;

printf("result = %f\n", result);

// 終了

free(h_A); free(h_B); free(h_C);

return 0;

}

Post processing (host: sample1.cu)

演習 ex1

• A[i] 、 B[i] はサイズ65536（256×256）の配列
• 以下のコードを実行するカーネル

ex1_kernel.cu　を記述せよ。

for (i=0; i<LENGTH; i++)

C[i] = 0.0;

for(j=0; j<LENGTH; j++)

C[i] += (A[i]-B[j])*(A[i]-B[j]);

ex1.cuのex1Kernelのコメントをはずして実行

CPUとGPUの答が一致するはず

提出：ex1_kernel.cuをkeio.jpに提出

​

最終レポートと授業の成績

• 1回、2回、4回の演習の結果　5点×３
• 3回、5回、6回の実習の結果　10点×3
• 最終レポート　15点
• 60点を4つに分けてA,B,C,Dを付ける。Sは特に凄いのがあれば例外的に付ける
• 最終レポート
• OpenMP,MPI,Cudaのプログラムモデルを比較し、対象とするアーキテクチャとの関連を論ぜよ。A4　2ページ以内とする。
• 全ての結果の締め切りは6月2日