研究の最近のブログ記事

NVIDIA主催のOpenCLセミナー、OpenCL seminar〜GPUコンピューティングがもたらすもの〜に行ってきた。



その際の次世代GPU"Fermi"やOpenCLの概要のメモを公開
※個人的なメモなため、内容の正確性は保証できません

先ずは次世代GPU、コードネーム"Fermi"の話、

GT200と比較したFermiの仕様

開発コード名	GT200	Fermi
構成トランジスタ数	14億Tr	30億Tr
倍精度浮動小数点演算能力	30FMA Ops/clk	256FMA Ops/clk
単精度浮動小数点演算能力	240MAD Ops/clk	512MAD Ops/clk
Warpスケジューラ(SMあたり）	1	2
特別機能ユニット（SFU、SMあたり）	2	4
シェードメモリ（SMあたり）	16KB	48KB or 16KB
1次キャッシュ（SMあたり）	-	16KB or 48KB
2次キャッシュ（SMあたり）	-	768KB
ECCメモリサポート	-	○
コンカレントカーネル	-	16個まで
ロード/ストアアドレス幅	32bit	64bit

•   次世代CUDAアーキテクチャ”Fermi”

30億Tr、コア数倍増512(SMあたり32コア)


倍精度浮動小数点演算ピーク性能8倍(単精度の1/2)、従来は1/8


2個のThread Cchedular(コンテキストスイッチ最適化、複数カーネルの起動、Dual DMAエンジン)、より効率的なタスク並列とデータ並列が可能


IEEE754-2008準拠(最新CPUを凌駕)、非正規数サポート、Fused multiply add(FMAD)命令を倍精度・単精度でサポート、新設計の全命令32bitサポート整数ALU (64bitへの最適化、24bit divが消えた笑)


オンチップ・シェアードメモリともに本格的なキャッシュ階層構造、各SMにL1キャッシュ


GDDR5メモリインターフェイス、ECCサポート(レジスタ,L1,L2)、SECDECサポート


C/C++ポインタのフルサポート


NVIDIA Nexus IDE（統合化開発環境）, Nexus IDE 1.0 (Visual Studio 2008 SP1)

ちょっと赤字は困ったぞ、研究的な意味で(笑)

倍精度の性能が高くなれば、今私の単精度を用いた疑似倍精度の必要性なくなるし、ECCがサポートされれば東工大の某研究室のチームも笑うしかないだろう。


次にOpenCLの話、

• OpenCL概要
  

ヘテロジニアス並列コンピューティング環境のためのフレームワーク


Open CL C言語、Open CL Runtime API


ハイブリッドシステム　CPU+DSP, PPE+SPE(Cell/B.E.), CPU+GPU


Open CL Version 1.0 Revision 48

各種プロセッサに対応するOpen CLは今年中にリリース




統一された言語による記述、統一されたAPIによる演算デバイス制御が可能


CUDAとOpenCLは同等程度のパフォーマンス(?)

Porting/Optimization サービス、コンパイラは開発中

EclipseベースのOpenCLプラグイン




C99ベースの言語


拡張：ベクタ・ベクタ演算、組み込み関数、修飾子、Reinterpreting

厳密な定義：型のビット幅　演算子の意味

CUDAにはないDMA：グローバル（デバイスメモリ）ローカル（シェアード）間

制限：関数ポインタが使えない、可変長配列が使えない、再帰ができない

タスク並列とデータ並列は双方選択できる。CUDAと比較して高速になるかはものによる

個人的な感想として、CUDAとほとんど仕様同じじゃないかと(笑)。際立った違いはメモリ間のDMAだけど(cudaMallocに相当)、グローバル⇔ローカル間のメモリ転送って役に立つのかな・・・？　性能は書き方によるし、CUDA越せるとは限らないと言っていた。　

Open CLの規格統一で複数のプラットフォームに対応したことは良いことだけど、HPCの分野にはまだまだこれから伸びしろがあると感じた。

プログラミングセミナーはカメラが無くて記録できなかった(涙)
後日、資料くれたりしないかしら

• Openclの.clソースコードを一部抜粋
  

カーネル
__kernel void vecAdd4(__global float4 *A, __gloabal float4 *B, __gloabal float4 *C) { int n = get_global_id(0); C[n] = A[n] + B[n]; //n=1...128 }
ホストからの呼び出し
ret = clEnqueueNDRangeKernel1( comand_queue, kernel, //vecAddカーネル 1, //1次元 NULL, global_item_size, //128ワープアイテム local_item_size, //グループ内アイテム数128 0, NULL, NULL);

CEATEC JAPAN 2009の二日目である７日に行って参りました

様々の企業の展示ブースを回ったりしました。その時の写真の一部を載せます

それとアピール：

私の所属する早稲田大学　笠原・木村研究室は
日立・ルネサスと共同で展示デモを行なっております


RP2の周波数電圧制御と電源制御による低消費電力化実現のデモを行なっています

随分久しぶりのブログですが、生存報告を含めて更新

研究室ではGPUの計算精度に関しての研究をすることになりました

GPUってのは単精度浮動小数点演算が主であり、

倍精度性能は僅かしかないのが特徴です

よって単精度から倍精度の実現と高速化、

つまりは性能と速度のトレードオフを探していくことになります



手始めにCUDAプログラミングに慣れるために、blackscholes.cの手動書き換えを行なった
(ブラックショールズは金融に関する確率微分方程式のシミュレートプログラムである)

シミュレートする銘柄数を変化させてCPUとGPUを比較すると以下の結果が得られた

CPUはC言語@Core 2 Quad Q9300
GPUはCUDA@Geforce 9800GTX

blackscholes実行結果(単位：秒)

num	CPU	GPU	CPU/GPU
1000	0.104006	0.012000	8.7
10000	0.844052	0.012001	70.3
100000	8.164510	0.024001	408.2
1000000	81.709107	0.120008	680.9

この結果だけ見るとGPUすげー！！ってなるけれど

これだけじゃ終われないんだよね(涙)

CPUもGPUもどのように最適化したか。どのようにチューニングしたか
GPUならデータ転送の影響は？並列化の最適化
メモリにしたってグローバルメモリだけじゃなく他のキャッシュメモリを使用したか

色々と課題が残ってる

それと何よりもCPUとGPUの演算結果が微妙に誤差が生じている
これが一番の問題であり、制度保証を必要とするアプリでは致命的である

今後はGPUの精度について研究していきます！