アクセラレータのデモとプログラミング手法

(1)

アクセラレータのデモと

プログラミング手法

会津大学中里直人

アクセラレータボードを使った高速化スクール 2009/12/07

(2)

アクセラレータとは(1)

• ホスト計算機を補佐して特定の計算を高速化

する計算機デバイス

– ホスト(CPU)で動作するプログラムを補佐

• アクセラレータの例

– Cell/PowerXCell8iブレード・ボード: 計算 – GPU ボード(NVIDIA, AMD, S3) : CGや計算 – FPGAボード : 計算や暗号処理

– ビデオ処理チップ : 映像ファイルの処理 – 他、暗号処理ボード等

(3)

アクセラレータとは(2)

• 今回のスクールでは具体的に

– GPU(AMD)とGRAPE-DRを指す

• アクセラレータの特徴

– チップに乗った並列計算機 • 大量(>100)の演算器が搭載されている – 単精度性能が非常に高速 ~ 1 Tflops • 倍精度の性能は単精度性能の0.1 – 0.5倍 – 現状では、どれも自律的には動作できない • ホスト計算機から制御される。 • 別のメモリ空間をもつのでデータ転送が必要 – 演算性能とメモリ性能のギャップが大きい • 1.2 Tflops vs. 100 GB s-1

(4)

アクセラレータとは(3)

• チップに乗った並列計算機なので

– 何らかの方法で並列プログラミングをする必要 • どのように複数の演算器を効率よく利用するか？ • そもそも、どのようにプログラミングするか？

• 二つのプログラム(コード)が必要

– ホストプログラム : CPUで動作するコード – kernelプログラム : GPUで動作するコード

• 二つのコードの橋渡しが必要

– データ転送をどのように実現するか？

(5)

アクセラレータのアーキテクチャ

• CPUからGPU/GRAPE-DRを操作する

– アプリケーションは、CPUで実行されるコードと、 GPUで実行されるkernelからなる

(6)

アクセラレータの困難点

• 複数レベルで並列化を考える必要あり

– CPUとGPUで並列動作 – GPU上の複数の演算器が並列動作 – 個々の演算器上で並列計算(SIMD演算) – (GPUクラスターでの並列化)

• メモリ空間が分離されている問題

– 効率の良いデータ転送の方法 – 効率の良いデータ構造

• GPU上でのメモリ帯域の問題

– AMD GPUはキャッシュ機構あり • ない場合：明示的なキャッシュ処理が必要

(7)

本スクールの目的

• GPUとGRAPE-DRをプログラムする手法を概説

• 1日目

– アクセラレータの利用方法を説明 – 各プログラミング環境によるデモとその利用法の説明

• 2-3日目

– 特にGPUにしぼって, より詳細なプログラミング手法を説明 – 各自の問題のテスト実装

(8)

デモと実習

• 問題設定

– N体計算: 常微分方程式を解く • N個の粒子が相互作用しながら進化する

• 複数のプログラミング環境の紹介

– それぞれの概要を説明します – 実習では、デモプログラムを修正・変更することで、アクセラレータを利用してもらいます









N j j i i

r

f

dt

v

d

1

)

(



(9)

プログラム環境について(1)

• Goose

– KFCR社開発によるプログラミング環境 – ディレクティブを利用したプログラミング – C言語+ディレクティブ

• OpenCL

– 業界標準として制定されている、アクセラレータのプログラミング環境 • 異なるハードウエアにひとつのソースで対応可能

– GPU(AMD, NVIDIA), CPU(SSE,multicore), Cell等

– GPU: 拡張されたC言語

(10)

プログラム環境について(2)

• LSUMP

– GPU(AMD)とGRAPE-DR用問題特化型コンパイラ – GPU: 独自言語 – CPU: APIをC言語やC++言語からよびだし – 特徴: 四倍精度に対応

• IL (AMD)

– GPUをアセンブリ言語でプログラミング – 詳しくは明日 – GRAPE-DRにも同等のアセンブリ言語があります

(11)

(12)

デモプログラムについて(1)

• 二重ループで加速度を計算

• 変数

– 位置ベクトル x[],y[],z[] – 質量 m[]

– 加速度 ax[], ay[], az[] – ε : softening parameter 2 / 3 2 2 ) | (| ) ( ) , (



    i j i j j j i r r r r m r r f  

(13)

デモプログラムについて(2)

• 二重ループの例

for(i = 0; i < n;i++) { ax[i] = 0.0; ay[i] = 0.0; az[i] = 0.0; for(j = 0; j < n;j++) { dx = x[j] - x[i]; dy = y[j] - y[i]; dz = z[j] - z[i]; r2 = dx*dx + dy*dy + dz*dz + eps2; rinv = 1.0/sqrt(r2); mrinv = m[j]*rinv; mr3inv = mrinv*rinv*rinv; ax[i] += mr3inv * dx;

ay[i] += mr3inv * dy; az[i] += mr3inv * dz; }

(14)

デモプログラムについて(3)

• 二重ループを各プログラミング手法で実行

– /usr/local/src/KEK.tar

– 展開した各ディレクトリにそれぞれのソースあり

• nbody_goose : Goose version • nbody_opencl : OpenCL version • nbody_il : IL version

• nbody_lsump_il : LSUMP for GPU version

– main.c以外はソース共通

• setup() : 初期化

(15)

デモプログラムについて(4)

• Goose version

– 元プログラムとほぼ同じ • #pragmaでディレクティブを追記 • rsqrt()関数を定義 – “make host”でCPU用コード生成 – “make”でgooseccによるGPU用コード生成 – 実行 “./run” • GPU用の場合コンソールで実行する必要あり

(16)

(17)

(18)

OpenCL/LSUMP/IL共通の説明(1)

• アクセラレータの計算モデル

– たくさんの演算プロセッサの集団

– 全プロセッサが同一のkernelを実行(SIMD動作) – ただし個々のプロセッサは別のデータを処理

(19)

OpenCL/LSUMP/IL共通の説明(2)

• 必要な処理

– kernelコードの生成 – アクセラレータの初期設定 – データ転送用メモリの設定 – データ転送: 書き込みと読み出し – カーネルの実行

• 以上の処理を実現するためのAPIがある

– OpenCLではユーザーがAPIを利用 – LSUMPではある程度自動生成

(20)

OpenCL/LSUMP/IL共通の説明(3)

• kernelコードの生成

1. kernelプログラムの記述  OpenCLなら拡張されたC言語  ILではアセンブリ言語  LSUMPではC言語ライクな独自言語  LSUMPによりIL/DRアセンブリ言語生成 2. kernelプログラムのコンパイルリンク • オンライン – APIによりプログラムからコンパイル処理する場合 • オフライン – あらかじめツールによりコンパイルリンク

(21)

OpenCL/LSUMP/IL共通の説明(4)

• アクセラレータの初期設定

– デバイスのオープン、初期化、保持など

• データ転送用メモリの設定

1. メモリの確保 2. kernel内の変数とホスト上のメモリを対応つける

• データ転送

• ホストメモリからGPUメモリへの書き込み • GPUメモリからホストメモリへの読み出し

• kernel実行

• 同期または非同期

(22)

(23)

AMD GPU用の環境設定(1)

• ATI Stream SDK

– AMD 社のGPU用のプログラミング環境 – 短縮URL http://bit.ly/4fxIZc • http://developer.amd.com/gpu/ATIStreamSDK/Pages/default.aspx – 現状で1.4betaと2.0betaがある – OpenCLには2.0betaより対応

• ATI Stream SDK v2.0 Beta ProgramよりDL

– 注意点

(24)

AMD GPU用の環境設定(2)

• ATI Stream SDKのインストール

– デバイスドライバのインストールと設定 • 公式サポート openSUSE 11.0, Ubuntu 9.04 • Ubuntu 8.04LTS, 9.04で1.4, 2.0の動作確認 – SDK自体のインストール

• インストール例 (1.4の場合)

– 短縮URL http://bit.ly/5u3oph – 2.0betaも基本的には同様にインストール可能 – インストール場所 • 1.4 : /usr/local/atical • 2.0 : /usr/local/opencl

(25)

AMD GPU用の環境設定(3)

• SDKのサンプルプログラム

– $(OPENCL) /samples/opencl/cl/app

• デモプログラムは”NBody”を改変

– 改変箇所 • 全てのエラーチェックを省略 • 利用変数を変更 • kernelプログラムを若干変更

• 以下OpenCLの詳しい説明はしません

(26)

OpenCLプログラムの流れ

1. GPUの確保と初期化 2. kernelプログラムのcompile/link 3. メモリ領域の確保 4. メモリとkernel変数の束縛 5. 入力データのsetup 6. kernelの実行 7. 結果の回収 8. 後処理 • 5から7を繰り返し実行。 – その他は一度おこなえばよい。

(27)

OpenCLプログラミング例 (1)

• KEK/nbody_opencl

– main.c : ホストプログラム • OpenCLデバイスの初期化 – 可視化用にOpenGLの初期化 • kernelプログラムのオンラインコンパイル • 数値積分と可視化 – source.cl : kernelプログラム • 加速度の計算

(28)

OpenCLプログラミング例 (2)

• kernelプログラム

__kernel void

nbody_sim(__global float4* pos ,

__global float4* acc_g, int n, float e2) {

unsigned int gid = get_global_id(0);

float4 myPos = pos[gid];

float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f); for(int j = 0; j < n; ++j) {

float4 r = pos[j] - myPos;

float distSqr = r.x * r.x + r.y * r.y + r.z * r.z; float invDist = 1.0f / sqrt(distSqr + e2);

float invDistCube = invDist * invDist * invDist; float s = pos[j].w * invDistCube;

acc += s * r; }

acc_g[gid] = acc; }

(29)

OpenCLプログラミング例 (3)

• kernelプログラムのキーワード

– __kernel : OpenCL kernel関数の定義

– __global : GPU上の配列変数 • get_global_id()関数により「自分のid」を得る(gid) • gidにより「自分の担当データ」読み出し書き込み – float4 : 単精度四要素の変数 • float4 v == v[4]と同等 – {v.x, v.y, v.z, v.w} == {v[0], v[1], v[2], v[3]}

(30)

OpenCLプログラミング例 (4)

• kernelプログラムの動作

– 元プログラムの外側のループは暗黙的に並列計算される • i = 0 → gid = 0のプロセッサ • i = 1 → gid = 1のプロセッサ • i = 2 → gid = 2のプロセッサ – ホストプログラム • kernelコードのNプロセッサでの並列計算を指令する

(31)

OpenCLプログラミング例 (5)

• ホストプログラム

– 24 – 36行 : デバイスの初期化 – 39, 40行 : ホスト側の配列確保 – 41, 42行 : GPU側の配列確保 – 45 – 72行 : kenerlプログラムの処理 – 74 – 82行: ホスト変数とGPU変数の束縛 • clSetKernelArg()の第2引数はkernelの引数の順番 – 関数 go() • clEnqeueWriteBuffer() : GPU側へ書き込み • clEnqeueReadBuffer() : GPUから読み出し • clEnqueueNDRangeKernel() : kernel実行

(32)

(33)

LSUMPの概要

• 環境設定について

– XXXよりDLできます(近日公開)

– ATI Stream SDKかGRAPE-DRの実行環境が必要 – KEK/bin/lsumpにバイナリ

• KEK/nbody_lsump_ilにデモプログラム

– “lsump –SCAL source.q”のように実行 – LSUMPは問題特化型コンパイラ • 以下の右辺の計算用に設計された









N j j i i

r

f

dt

v

d

1

)

(



(34)

プログラムの流れ

1. GPU/DRの確保と初期化 2. kernelプログラムのcompile/link 3. メモリ領域の確保 4. メモリとkernel変数の束縛 5. 入力データのsetup 6. kernelの実行 7. 結果の回収 8. 後処理 • 1-4までは自動生成したAPIを利用可能

(35)

LSUMPプログラミング例 (1)

• KEK/nbody_lsump_il

– main.c : ホストプログラム • デバイスの初期化 – 可視化用にOpenGLの初期化 • 数値積分と可視化 – source.q : kernelプログラム • 加速度の計算 – callib.c, callib.h • API用のソースコード

(36)

LSUMPプログラミング例 (2)

• kernelプログラム

VARI xi, yi, zi, e2; VARJ xj, yj, zj, mj; VARF ax, ay, az; dx = xj - xi;

dy = yj - yi; dz = zj - zi;

r1i = rsqrt(dx**2 + dy**2 + dz**2 + e2); af = mj*r1i**3;

ax += af*dx; ay += af*dy; az += af*dz;

(37)

LSUMPプログラミング例 (3)

• VARI, VARJ, VARFが変数宣言

– VARI : プロセッサローカルな入力変数 – VARF: プロセッサローカルな出力変数 • +=とするとループで積算される – VARJ : ループで逐次読み込まれる変数 • 内側のループは暗黙的に処理される • ホストプログラムでループ回数を指定

• 残りの部分 : kernel演算本体

– 関数f()のみを定義する – VARI変数とVARJ変数を入力として、最終的に VARF変数に結果を計算する

(38)

LSUMPプログラミング例 (4)

VARF

_{VARI VARJ}

– VARI,VARF

– プロセッサ固有の変数 – VARJ

(39)

LSUMPプログラミング例 (5)

• ホストプログラム

– OCLsetup1() : GPUの初期化やメモリ確保 • LSUMPによりkernelの変数宣言から自動生成 • ocal_lib.c – 関数 go() • OCLwritememory_float() : GPUへ書き込み • OCLreadmemory_float() : GPUから読み出し – 第2引数はkernelでの変数宣言に現れた順番 • OCLgo(mod, n) – kernelの実行 : ループ回数はn回

(40)

LSUMPの手法について

• ユーザーは以下をDSLで記述する

– 並列計算する部分

– 入力変数の性質の指定

• 変数宣言に意味を持たせることで最適化が可能 DSL : domain specific language

• 提案コンパイラは、指定された変数の性質に

基づいてアクセラレータ用コードを生成する

– 経験的な最適計算手法の適用

• これは問題に依存する：今回は総和演算 • さらにアクセラレータにも依存する

(41)

(42)

Example

(43)

Example : Feynman-loop integral (1)

VARI xx, yy, cnt4; VARJ x30_1, gw30; VARF res;

CONST tt, ramda, fme, fmf, s, one; zz = x30_1*cnt4;

d = -xx*yy*s-tt*zz*(one-xx-yy-zz)+(xx+yy)*ramda**2 +

(one-xx-yy-zz)*(one-xx-yy)*fme**2+zz*(one-xx-yy)*fmf**2; res += gw30/d**2;

(44)

Example : Feynman-loop integral (2)

• 元の三重ループ→内側の二重ループを計算

for(int i1 = 0; i1 < n; i1++) { xx = x30[i1]*cnt0+cnt1; by = 1.0-xx; cnt2 = (by-ay)*0.5; cnt3 = (by+ay)*0.5; for(int i2 = 0; i2 < n; i2++) { yy = x30[i2]*cnt2+cnt3; bz = 1.0-xx-yy; cnt4 = (bz-az)*0.5; cnt5 = (bz+az)*0.5; for (int i3 = 0;i3 < n; i3++) {

zz=x30_1[i3]*cnt4; d = -xx*yy*s-tt*zz*(one-xx-yy-zz)+ (xx+yy)*ramda*ramda+ (one-xx-yy-zz)*(one-xx-yy)*fme*fme+ zz*(one-xx-yy)*fmf*fmf ; w3=gw30[i3]/(d*d) ; sum3 = sum3 + w3; } } } 相互作用計算部分

(45)

LSUMPのオプション

• GPU用

– lsump –SCAL : 単精度コード – lsump –DCAL : 倍精度コード – lsump –DDCAL : 四倍精度コード

• DR用

– lsump –D : 倍精度コード – lsump –DD : 四倍精度

(46)

(47)

実習問題例

• CPUでの実行(nbody_opencl)

– main.cの”CL_DEVICE_TYPE_GPU”はGPUでの実行を指定しているので、”CL_DEVICE_TYPE_CPU”に変更するとCPUでの実行となる

• 相互作用のべき乗を変化

– source.cl/source.qを修正して実行 – 相互作用関数自体の変更

(48)

実習問題例

• 相互作用関数を変更

– Kernel変数を追加してみる

• OpenCLのサンプルプログラムをテスト

– エラー処理が多いので省略したプログラムを作成

• LSUMPによりGPUをプログラム

– 単精度、倍精度 – 異なる相互作用関数

(49)

実習問題例

• ILプログラミングのさわり

– KEK/nbody_il : ILによるデモプログラム – cal_source_f.ilがkernelプログラム • cal_source_f2.ilは別のkernelプログラム(冪変更) – main.c 24行

• OCLloadkernel(dev, mod, "cal_source_f.il", 1)

を以下に変更

• OCLloadkernel(dev, mod, "cal_source_f2.il", 1)

• レジスタ割りあてが必要なく、制御構造をサポートするので、案外容易に利用できます

(50)

アクセラレータのデモと プログラミング手法