C/JAVAの最近のブログ記事

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);

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

研究室では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の精度について研究していきます！

Microsoftの寄付講座『リアルタイム3Dグラフィックスプログラミング』で分かったことのメモ。
今回はアンチエイリアシングのお話。いろんな種類があるのね。

• 点と2点を結ぶ直線の符号付き距離

• アンチエイリアシング手法の一例

1*1　真中の色を使用する






1*2　上下の位置の平均値の色を使用する





2*2　4か所の平均値の色を使用する





2*2　4か所の平均値の色を使用する(RGSS)





4*4　16か所の平均値の色を使用する





5点　中央と4隅を使用する。ただし、中央の重みは4倍

いそっちノートのひとは実際Flexで実装してるね。すごいｗ

Microsoftの寄付講座『リアルタイム3Dグラフィックスプログラミング』で分かったことのメモ。
今回は光源のお話。イメージだけはつかめたかも。

• 平行光源


• アンピエント光源


• 点光源


• スポットライト


• マテリアル
  

Microsoftの寄付講座『リアルタイム3Dグラフィックスプログラミング』で分かったことのメモその2。

• クォータニオン

• 線形球面補間

以下、図とプログラム例