2010/7/26 高性能コンピューティング学論　講義資料 新しい並列処理「GPGPU」について 大島聡史　(東京大学 情報基盤センター)

Presentation on theme: "2010/7/26 高性能コンピューティング学論　講義資料 新しい並列処理「GPGPU」について 大島聡史　(東京大学 情報基盤センター)"— Presentation transcript:

1 2010/7/26 高性能コンピューティング学論　講義資料 新しい並列処理「GPGPU」について 大島聡史　(東京大学 情報基盤センター)

2 目次 GPGPUとは何か CUDAについて 今後のGPUとGPGPU GPGPUの概要、はじまりと発展について
2010/7/26 高性能コンピューティング学論　講義資料 目次 GPGPUとは何か GPGPUの概要、はじまりと発展について CUDAについて GPGPUの普及に大きな影響を与えたCUDA、そのアーキテクチャとプログラミング環境について 今後のGPUとGPGPU GPGPUの将来はどうなるのか、最新のGPU事情

3 自己紹介 大島聡史 経歴 2000年4月～2009年9月 電気通信大学 2009年10月～ 東京大学 情報基盤センター 助教
2010/7/26 高性能コンピューティング学論　講義資料 自己紹介 大島聡史 経歴 2000年4月～2009年9月　電気通信大学 2000年　入学　(情報工学科) 2003年　卒研時配属で弓場・本多研究室へ 2004年　博士前期課程進学　(並列処理学講座 弓場研究室) 2006年　博士後期課程進学　(並列処理学講座 本多研究室) 2009年　博士研究員　(高性能コンピューティング学講座 本多研究室) 2009年10月～　東京大学 情報基盤センター 助教

4 2010/7/26 高性能コンピューティング学論　講義資料 GPGPUとは何か GPGPUの概要、はじまりと発展について

5 「GPGPU」？ General-Purpose computation using GPUs GPU 「GPUを用いた汎用演算」
2010/7/26 高性能コンピューティング学論　講義資料 「GPGPU」？ General-Purpose computation using GPUs 「GPUを用いた汎用演算」 GPU Graphics Processing Unit いわゆるビデオカードや統合グラフィックスチップ（上のLSI） 代表的な製品：NVIDIA GeForce, ATI (AMD) Radeon, Intel GMA 用途＝画面出力全般 3Dグラフィックス処理 3Dゲーム、3DCAD、3DCG作成 エンコード・デコード支援 GPU上に専用チップを搭載していることが多い デスクトップ処理 Windows Aeroが比較的高性能なGPUを 要求したため普及

6 GPUの性能向上と並列処理 GPUに対する要求＝高速・高解像度・複雑な画像描画 GPUの高性能化におけるキーワード
2010/7/26 高性能コンピューティング学論　講義資料 GPUの性能向上と並列処理 GPUに対する要求＝高速・高解像度・複雑な画像描画 2000年以降、特に写実的な3DCGへの要求が高まる もちろん、家庭用ゲーム機にもGPUは積まれています GPUの高性能化におけるキーワード 並列化 プログラマブル化 今では携帯ゲームも3D！

7 一般的な3D描画の手順 頂点情報の設定 頂点処理：頂点が空間上のどこに対応するか（カメラ計算） ラスタライズ：画面上のピクセルとの対応付け
2010/7/26 高性能コンピューティング学論　講義資料 一般的な3D描画の手順 頂点情報の設定 頂点処理：頂点が空間上のどこに対応するか（カメラ計算） ラスタライズ：画面上のピクセルとの対応付け ピクセル処理：ピクセル色の決定、テクスチャ適用 (2, 2) (8, 3) (5, 7) 頂点ごと・ピクセルごとに独立計算（並列計算）可能なため、ハードウェアの並列化による性能向上が行われた

8 GPUのプログラマブル化 高度な描画処理手法の開発における問題 プログラマブルシェーダの登場 専用ハードウェアを用いる高度な手法を開発
2010/7/26 高性能コンピューティング学論　講義資料 GPUのプログラマブル化 高度な描画処理手法の開発における問題 専用ハードウェアを用いる高度な手法を開発 ハードウェアに実装されるまで時間が必要 対応ハードウェアを所有していないと使えない 次々に新たな手法が提案される ハードウェアと描画手法の開発サイクルがうまく回らない プログラマブルシェーダの登場 頂点処理やピクセル処理をソフトウェア処理で設定可能にした 新しい描画手法をすぐに利用可能になった シェーダのバージョン（世代）による実行可能不可能は存在 ハイエンドGPU＝多数の高速なシェーダユニット、高速大容量メモリ

9 描画処理以外にも使われ始めたGPU 初期のプログラマブルシェーダは限定的な処理のみ可能
2010/7/26 高性能コンピューティング学論　講義資料 描画処理以外にも使われ始めたGPU 初期のプログラマブルシェーダは限定的な処理のみ可能 GeForce6800(2004年)に搭載されたシェーダ(3.0)で大きく改善、様々なプログラムが記述可能に 複雑な照明処理や陰影処理が実装可能になり、リアルタイム3Dグラフィックス（主に3Dゲーム）の表現力が爆発的に向上 描画（照明・陰影）処理に物理計算（数値計算）が導入される 単精度浮動小数点演算への対応 左はGeForce6800のデモ。右はUnrealEngine（ゲーム開発用のフレームワーク）の進化過程。

10 GPUを用いた汎用演算(GPGPU) GPUによる処理を数値計算等に利用する研究が始まる GPGPUにおける入出力と演算の位置づけ
2010/7/26 高性能コンピューティング学論　講義資料 GPUを用いた汎用演算(GPGPU) GPUによる処理を数値計算等に利用する研究が始まる GPGPUにおける入出力と演算の位置づけ CPUからGPUへの転送（テクスチャへのデータセット） ＝入力 CPU GPU 描画 ＝演算 GPUからCPUへの書き戻し（描画結果の取得） ＝出力

11 描画処理を用いた汎用計算の例 配列加算 行列積 乗算描画結果の加算 tex1 a b c d tex2 w1 x1 y1 z1 w2 x2
2010/7/26 高性能コンピューティング学論　講義資料 描画処理を用いた汎用計算の例 配列加算 行列積 乗算描画結果の加算 tex1 a b c d tex2 w1 x1 y1 z1 w2 x2 y2 z2 w3 x3 z3 w4 x4 y4 z4 tex2 w x y z a*w b*x c*y d*z tex1 乗算描画 a1 b1 c1 d1 a2 b2 c2 d2 a3 b3 c3 d3 a4 b4 c4 d4 a1*w1 a1*x1 a1*y1 a1*z1 a2*w1 a2*x1 a2*y1 a2*z1 … 方法１：乗算ブレンド描画 方法２：tex1を描画し、 tex2を乗算描画

12 プログラマブルシェーダを用いた汎用計算の例
2010/7/26 高性能コンピューティング学論　講義資料 プログラマブルシェーダを用いた汎用計算の例 グラフィックスAPI（DirectX, OpenGL）による描画処理＋シェーダ言語（HLSL, GLSL）による演算 void gpumain(){ vec4 ColorA = vec4(0.0, 0.0, 0.0, 0.0); vec4 ColorB = vec4(0.0, 0.0, 0.0, 0.0); vec2 TexA = vec2(0.0, 0.0); vec2 TexB = vec2(0.0, 0.0); TexA.x = gl_FragCoord.x; TexA.y = gl_FragCoord.y; TexB.x = gl_FragCoord.x; TexB.y = gl_FragCoord.y; ColorA = texRECT( texUnit0, TexA ); ColorB = texRECT( texUnit1, TexB ); gl_FragColor = F_ALPHA*ColorA + F_BETA*ColorB; } GPUの処理 (GLSL) 各ピクセルに対して実行される シェーダ言語を用いた配列加算（vc=a*va + b*vb）の例 void main(){ glutInit( &argc, argv ); glutInitWindowSize(64,64);glutCreateWindow("GpgpuHelloWorld"); glGenFramebuffersEXT(1, &g_fb); glBindFramebufferEXT(GL_FRAMEBUFFER_EXT, g_fb); glGenTextures(4, g_nTexID); // create (reference to) a new texture glBindTexture(opt1, texid); glTexParameteri(opt1, GL_TEXTURE_MIN_FILTER, GL_NEAREST); glTexParameteri(......); glTexImage2D(opt1, 0, opt2, width, height, 0, GL_RGBA, GL_FLOAT, 0); ……（以下省略） CPUの処理 (OpenGL)

13 2010/7/26 高性能コンピューティング学論　講義資料 GPGPUのはじまり：まとめ 高性能な描画処理への要求 →GPUの高性能化（高並列、一定のプログラマビリティ） →高い並列浮動小数点演算性能の汎用計算への利用可能性 →数値計算や可視化などで成果 GPGPUの特徴（メリット） 高い浮動小数点理論演算性能 CPUが10GFLOPS程度の時代にGeForce6800は120GFLOPS 高並列（並列度の高い演算に有用） CPUがシングルコア～デュアルコアの時代に数十～百の並列度 GeForce6800は6(vp)+16(pp)並列 安価 CPUと変わらないレンジ（～10万円）、多くのPCに搭載 GPGPUの課題 直接的でなく扱いにくいプログラミング記述

14 余談 私が知る限り、一番はじめの頃は「GPGP」「GP2」「GPU computing」という呼称が多かった
2010/7/26 高性能コンピューティング学論　講義資料 余談 私が知る限り、一番はじめの頃は「GPGP」「GP2」「GPU computing」という呼称が多かった 「GPGPU」という名称が普及したが、最近また「GPU computing」という言葉も使われている（「コンピューティング」を強調したいらしい） そのうちまた呼称が変わるのだろうか？ GPU＝General-Purpose GPU、などの記述を見ると違和感を感じずにはいられない

15 CUDAについて GPGPUの普及に大きな影響を与えたCUDA、そのアーキテクチャとプログラミング環境について 2010/7/26
高性能コンピューティング学論　講義資料 CUDAについて GPGPUの普及に大きな影響を与えたCUDA、そのアーキテクチャとプログラミング環境について

16 黎明期GPGPUの課題 ハードウェア プログラミング ライブラリによる隠蔽 にわかに脚光を浴びたストリーミング言語 仕様がよくわからない
2010/7/26 高性能コンピューティング学論　講義資料 黎明期GPGPUの課題 ハードウェア 仕様がよくわからない 非公開（グラフィックスAPI＋シェーダ言語　レベルでしか見えない） 一応、共有メモリ型並列計算機のようなもの？ テクスチャとフレームバッファを共有メモリとして、描画ユニットが演算 プログラミング グラフィックスAPI＋シェーダ言語 画像処理プログラミングのスキルが必須 ライブラリによる隠蔽 にわかに脚光を浴びたストリーミング言語

17 mem PE PE PE PE PE PE PE PE PE PE PE PE 2010/7/26 高性能コンピューティング学論 講義資料
高性能コンピューティング学論　講義資料 mem PE PE PE PE PE PE PE PE 共有メモリ、分散共有メモリ、分散メモリのイメージ……はすでに習っているだろうか？ PE PE PE PE

18 ライブラリによる隠蔽 GPUプログラミングをライブラリで隠蔽し、既存のプログラミング言語から容易に使えるようにする
2010/7/26 高性能コンピューティング学論　講義資料 ライブラリによる隠蔽 GPUプログラミングをライブラリで隠蔽し、既存のプログラミング言語から容易に使えるようにする 良く用いられる処理を高速化 行列積やFFTなどの数値計算はCPU上でもライブラリがよく利用されている GPUを利用してくれるライブラリを作ればGPUの性能を誰でも容易に活用可能 ちなみに、私もこのタイプの研究を行った（行っている） CPUとGPUを両方使って（片方の方が良い場合は片方だけ使って）計算する行列計算ライブラリを作った

19 ストリーミング言語 GPU向けの言語やクラスライブラリを提供
2010/7/26 高性能コンピューティング学論　講義資料 ストリーミング言語 GPU向けの言語やクラスライブラリを提供 多数の入力データ（配列、ストリーム）に次々に演算操作（カーネル）を施すプログラミングモデル 入力配列と出力配列、配列に対する演算を明確に記述する 元々GPU用というわけではないが、GPGPUの普及をうけて注目された BrookGPU、Brook+、RapidMind、SPRATなど 一部の言語はマルチコアCPUやCellにも対応 演算操作（カーネル） 入力ストリーム 出力ストリーム

20 CUDAの登場 ライブラリもプログラミング言語も一般へ普及せず 2007年初頭、NVIDIA社がCUDAを公開
2010/7/26 高性能コンピューティング学論　講義資料 CUDAの登場 ライブラリもプログラミング言語も一般へ普及せず 理由：性能が出せない、使い勝手が良くない、etc. 2007年初頭、NVIDIA社がCUDAを公開 CUDA Unified Device Architecture アーキテクチャ＋プログラミング環境 GPUのアーキテクチャを以前より（とても）詳しく公開 Cベースのプログラミング環境とドキュメントを無償で公開 まともなGPU最適化プログラミングが可能に ユーザが爆発的に増加 パズル的になんとかすれば使える、を脱却

21 CUDAアーキテクチャ1 物理的な構成の概要 ※最新版ではキャッシュを搭載するなど変わってきているが、ここでは初期型について簡単に解説する
2010/7/26 高性能コンピューティング学論　講義資料 CUDAアーキテクチャ1 物理的な構成の概要 ※最新版ではキャッシュを搭載するなど変わってきているが、ここでは初期型について簡単に解説する HOST GPU Streaming Multiprocessor (MP) ×N PCI-Express MainMemory DeviceMemory ScalarProcessor (SP) SharedMemory TextureCache ConstantCache ×8 Register やっとまともなハードウェアアーキテクチャの話ができます。 今はこれ全部覚えなくても良い。CUDAを使うときに覚えれば良い。 Instruction Unit

22 ハードウェアの特徴 階層性のあるハードウェア構成 演算器の構成 メモリの構成 階層的な共有メモリ＋ローカルメモリ
2010/7/26 高性能コンピューティング学論　講義資料 ハードウェアの特徴 階層性のあるハードウェア構成 演算器の構成 階層性のある演算器配置（SP*8×MP*N） 同一MP内のSPは同時に同じ演算しか行えない（SIMD的な構成） NVIDIAはSIMTと呼んでいる CPUのコアとは趣が異なるので注意 メモリの構成 階層性と局所性のあるメモリ配置 GPU上に搭載された大容量でグローバルなDRAM：DeviceMemory MPローカルの小容量高速共有メモリ：SharedMemory 階層的な共有メモリ＋ローカルメモリ 全体で共有可能なメモリ＋部分的な共有メモリ＋ローカルメモリ （習ったハードウェアとはだいぶ違うかも）

23 CUDAアーキテクチャ2 実行モデル＋メモリ構成の概要 MPに対応 Grid ×α SPに対応 Host(CPU, MainMemory)
2010/7/26 高性能コンピューティング学論　講義資料 CUDAアーキテクチャ2 MPに対応 実行モデル＋メモリ構成の概要 Grid ×α SPに対応 Host(CPU, MainMemory) TextureMemory ConstantMemory GlobalMemory Block ×β Register SharedMemory Thread ×γ LocalMemory CPUのプロセスやスレッド同様に、Block・Threadは物理的な数を超えて生成可能 むしろThreadは物理的な数を超えて作成するべき（後述）

24 詳細なメモリ構成 特徴の異なる複数種類のメモリ 動かすだけならRegisterとGlobalMemoryだけで十分
2010/7/26 高性能コンピューティング学論　講義資料 詳細なメモリ構成 特徴の異なる複数種類のメモリ 動かすだけならRegisterとGlobalMemoryだけで十分 GlobalMemoryとSharedMemoryの最適化が重要 名称 Lifetime 共有範囲 速度 容量 GlobalMemory プログラム GPU全体 高速・高レイテンシ ～*GB ConstantMemory 高速・高レイテンシ ＋キャッシュ 64KB TextureMemory GlobalMemoryと共用 SharedMemory Grid MP単位 超高速・低レイテンシ 16～KB/MP Register SP単位 8～KB/MP LocalMemory ※ - ※実体はGlobalMemory、レジスタを使いすぎるとLocalMemoryに配置される

25 CUDAアーキテクチャ3 処理の流れ Grid ×α Host(CPU, MainMemory) TextureMemory
2010/7/26 高性能コンピューティング学論　講義資料 CUDAアーキテクチャ3 処理の流れ Grid ×α Host(CPU, MainMemory) TextureMemory ConstantMemory GlobalMemory Block ×β Register SharedMemory Thread ×γ LocalMemory

26 CUDAアーキテクチャ4 もう少し詳しい実行モデル解説 既存のCPU CUDA 命令H 命令G 命令F 命令E 命令D 命令C 命令B
2010/7/26 高性能コンピューティング学論　講義資料 CUDAアーキテクチャ4 もう少し詳しい実行モデル解説 既存のCPU CUDA 命令H 命令G 命令F 命令E 命令D 命令C 命令B 命令A ↓ 命令h 命令g 命令f 命令e 命令d 命令c 命令b 命令a ↓ 命令H 命令G 命令F 命令E 命令D 命令C 命令B 命令A ↓ core 0 core 1 SP 0 SP 1 SP 2 SP 3 … SP毎（Thread毎）に分岐させたい場合にはマスク処理 変数（Register）を活用することでSP毎に別のメモリを処理するのが一般的

27 CPUとGPUの比較 CPU (Intel Xeon W3520) GPU (NVIDIA GeForceGTX280) 演算器（コア）数
2010/7/26 高性能コンピューティング学論　講義資料 CPUとGPUの比較 CPU (Intel Xeon W3520) GPU (NVIDIA GeForceGTX280) 演算器（コア）数 4 240 クロック周波数 2.66GHz 602MHz(SP) 1296MHz(SP) 搭載メモリ容量 - 1GB メモリ帯域幅 25.6GB/s 141.7GB/s 最大消費電力 130W 236W 理論演算性能 42.56GFlops 933GFlops CPUとの接続 PCIe x16(Gen2) ターボモードとかHTとかは考えてません。 主な数字を出すとこんな感じだけど、中身を見るとさらに多くの違いがあることがわかったはず。 （主なカタログスペックだけでも多くの差があるが、中身を知るとそれ以上に違うことがわかるだろう）

28 プログラミング環境としてのCUDA CUDA C (RuntimeAPI)
2010/7/26 高性能コンピューティング学論　講義資料 プログラミング環境としてのCUDA CUDA C (RuntimeAPI) C/C++ベースの並列化プログラミング言語 基本的にC/C++ 接頭辞による実行対象ハードウェアおよびメモリ配置指定 実行対象指定 __global__ CPUから呼び出しGPU上で実行 __device__ GPUから呼び出しGPU上で実行 __host__ CPUから呼び出しCPU上で実行 配置指定 __device__ GlobalMemory __shared__ SharedMemory __constant__ ConstantMemory (TextureMemoryは専用のクラスを用いて扱う) バックエンドにGCCやVC++が必要 （より細かくCUDAを制御可能なDriverAPIもあるが、ここでは割愛）

29 CUDAプログラミング シンプルな配列加算の例 GPUの処理 grid,threadで指定された数だけ実行される
2010/7/26 高性能コンピューティング学論　講義資料 CUDAプログラミング シンプルな配列加算の例 __global__ void gpumain (float* vc, float* va, float* vb, int nSize, float a, float b){ for(int i=0; i<nSize; i++){ vc[i] = a*va[i] + b*vb[i]; } GPUの処理 grid,threadで指定された数だけ実行される （CPUからGPU上のプロセッサを個別に操作しない） CUDAを用いた配列加算（vc=a*va + b*vb）の例 void main(){ CUT_DEVICE_INIT(); cudaMalloc((void**)&d_va, n); cudaMalloc((void**)&d_vb, n); cudaMalloc((void**)&d_vc, n); cudaMemcpy(d_va, h_va, n, cudaMemcpyHostToDevice); cudaMemcpy(d_vb, h_vb, n, cudaMemcpyHostToDevice); gpumain<<< grid, threads >>>(d_vc, d_va, d_vb, nSize, alpha, beta); cudaMemcpy(h_vc, d_vc, n, cudaMemcpyDeviceToHost); } CUTは使ってるけどね 個別操作無し、まとめて実行されます CPUの処理

30 CUDAプログラミング シンプルな配列加算の例（並列計算版） GPUの処理 grid,threadで指定された数だけ実行される
2010/7/26 高性能コンピューティング学論　講義資料 CUDAプログラミング シンプルな配列加算の例（並列計算版） __global__ void gpumain (float* vc, float* va, float* vb, int nSize, float a, float b){ int begin = blockIdx.x*blockDim.x + threadIdx.x; int step = gridDim.x*blockDim.x; for(int i=begin; i<nSize; i+=step){ vc[i] = a*va[i] + b*vb[i]; } GPUの処理 grid,threadで指定された数だけ実行される （シェーダと比べて）非常にわかりやすくなった CUDAに関する理解は必要 CUDAを用いた配列加算（vc=a*va + b*vb）の例 void main(){ CUT_DEVICE_INIT(); cudaMalloc((void**)&d_va, n); cudaMalloc((void**)&d_vb, n); cudaMalloc((void**)&d_vc, n); cudaMemcpy(d_va, h_va, n, cudaMemcpyHostToDevice); cudaMemcpy(d_vb, h_vb, n, cudaMemcpyHostToDevice); gpumain<<< grid, threads >>>(d_vc, d_va, d_vb, nSize, alpha, beta); cudaMemcpy(h_vc, d_vc, n, cudaMemcpyDeviceToHost); } 並列処理の知識は必要、だが既存の並列化の知識がそのまま使えるわけではない。言語しかり、モデルしかり。 CPUの処理

31 CUDAプログラミング（実行方法） 通常のCプログラムと同様にコンパイル・実行が可能
2010/7/26 高性能コンピューティング学論　講義資料 CUDAプログラミング（実行方法） 通常のCプログラムと同様にコンパイル・実行が可能 CUDAコンパイラ(NVCC)がGPUカーネルを分離し、CPU部とGPU部それぞれをコンパイルし、単一の実行ファイルを生成 CPU部またはGPU部のみをコンパイルすることなども可能 中間表現ファイルを出力して解析することも可能 中間表現の仕様も公開されているが、実行時にさらに改変されることなどから最適化に使うのは現時点では困難 nvcc main.cu ./a.out

32 MPIやOpenMPとの実行モデル比較 × MPI OpenMP CUDA main thread PE0 PE1 PE2 PE3 CPU
2010/7/26 高性能コンピューティング学論　講義資料 MPIやOpenMPとの実行モデル比較 MPI OpenMP CUDA main thread PE0 PE1 PE2 PE3 CPU GPU Block, Thread (MP, SP) thread 0,1,2,3 × OpenMPに通信は要らないけど、CUDAではCPU-GPU間に必要。 分散メモリ（プロセス間通信） 共有メモリ CPU-GPU間分散メモリ GPU内階層型共有メモリ

33 CUDA最適化プログラミング 主な最適化技術の概要 （どのようなプログラムに対して高性能が得られるか） 大量のThreadを生成する
2010/7/26 高性能コンピューティング学論　講義資料 CUDA最適化プログラミング 主な最適化技術の概要 （どのようなプログラムに対して高性能が得られるか） 大量のThreadを生成する 理想的なBlockあたりThread数は128程度 GlobalMemoryアクセスのコアレスアクセス 複数SPに同時に連続したメモリをアクセスさせる →まとめられて性能向上 SharedMemoryのバンクコンフリクト回避 SharedMemoryにアクセスする際に別々のメモリバンクを叩かせる →アクセスが衝突せずに性能低下を回避 ループアンローリング 分岐回数が減るため重要 どのような最適化が行われているのか、逆にどのようなプログラムは速度が出ないのか 逆に言えば、コアレス化やバンクコンフリクト回避ができないと性能が出ない

34 大量のThreadを生成する SP(Thread)のコンテキスト切り替えコストは非常に小 大量のThreadを生成することで高性能
2010/7/26 高性能コンピューティング学論　講義資料 大量のThreadを生成する SP(Thread)のコンテキスト切り替えコストは非常に小 メモリアクセスを待つよりコンテキストを切り替えて別のThreadを処理した方が速い 大量のThreadを生成することで高性能 （むしろ、大量のThreadでGlobalMemoryメモリアクセスのレイテンシを隠蔽しないと高い性能が得られない） RegisterやSharedMemoryの使用量が多いと多数のThreadを保持できない GPUカーネルをシンプルにするべき

35 GlobalMemoryのコアレスアクセス
2010/7/26 高性能コンピューティング学論　講義資料 GlobalMemoryのコアレスアクセス 同一MP内の複数SPによるメモリアクセスを揃える（近づける）とまとめてアクセスできる 詳細な条件はGPUの世代によって異なる（緩和が進んでいる） アクセスが揃っていない場合 アクセスが揃っている場合 SP0 SP1 SP2 SP3 4回のメモリアクセスが行われる GlobalMemory アドレス境界とか、一部重なっている場合とかの細かい話はあるが…… SP0 SP1 SP2 SP3 1回のメモリアクセスに纏められる GlobalMemory

36 SharedMemoryのバンクコンフリクト回避
2010/7/26 高性能コンピューティング学論　講義資料 SharedMemoryのバンクコンフリクト回避 SharedMemoryは16個・32bitずつのバンクにより構成 同一バンクへのアクセスが集中すると性能低下 均等なアクセス＝性能低下しない アクセスが集中＝性能低下する SharedMemory 2-way バンクコンフリクトの例 SharedMemory

37 CUDA最適化プログラミングのまとめ どのようなプログラムで高い性能が得られるか 高い並列度を持っている メモリアクセスに規則性がある
2010/7/26 高性能コンピューティング学論　講義資料 CUDA最適化プログラミングのまとめ どのようなプログラムで高い性能が得られるか 高い並列度を持っている メモリアクセスに規則性がある GlobalMemoryのコアレスアクセスやSharedMemoryのバンクコンフリクト回避ができる テクスチャキャッシュに収まる範囲であれば多少のランダム性も容認 GlobalMemoryへのアクセスがランダムでも、SharedMemoryに落とせればOK CPU-GPU間の通信が頻繁に必要でない SP間の同期を繰り返す必要がない

38 GPGPUを試す方法（情報源） NVIDIAとAMDの開発者サイトに開発に必要なものと資料が揃っている NVIDIA AMD
2010/7/26 高性能コンピューティング学論　講義資料 GPGPUを試す方法（情報源） NVIDIAとAMDの開発者サイトに開発に必要なものと資料が揃っている ダウンロードしてサンプルの実行・改造をすると良い NVIDIA CUDA AMD ATI Stream SDK

39 CUDAプログラミングを試す方法 準備 ハードウェア CUDA関連ソフトウェア その他 普通のPC CUDA Driver
2010/7/26 高性能コンピューティング学論　講義資料 CUDAプログラミングを試す方法 準備 ハードウェア 普通のPC CUDA対応GPUが無くてもエミュレーションで試すことは可能だが、１万円未満のGPUでも良いので積んでおきたい CUDA関連ソフトウェア CUDA Driver CUDA Toolkit CUDA SDK code samples その他 Windows・Linux・MacOSXに対応 gccかVisualStudioが必要 個人的にはWindowsよりLinuxの方が環境を作りやすいが、VisualStudio統合デバッガが強力なのは捨てがたい 一部のメーカー製PCではドライバ導入が大変（最悪、不可能） 割と簡単に試せるので興味があれば試して欲しい

40 2010/7/26 高性能コンピューティング学論　講義資料 今後のGPUとGPGPU GPGPUの将来はどうなるのか、最新のGPU事情

41 現在のGPGPU ハードウェア プログラミング環境 NVIDIAが大きな影響力 OpenCLが普及し始めるとAMDも対等に戦えるか？
2010/7/26 高性能コンピューティング学論　講義資料 現在のGPGPU ハードウェア NVIDIAが大きな影響力 CUDAが大きな原動力 OpenCLが普及し始めるとAMDも対等に戦えるか？ AMDはほぼシェーダの開発環境のみ提供していた IntelがGPUを出すという噂は常にあるが…… Larrabee？ プログラミング環境 CUDAが中心 徐々に高まるOpenCLへの注目

42 OpenCL Khronosがまとめている並列計算フレームワーク OpenGLの並列計算版のようなもの Appleが提案
2010/7/26 高性能コンピューティング学論　講義資料 OpenCL Khronosがまとめている並列計算フレームワーク OpenGLの並列計算版のようなもの Appleが提案 Apple版CUDA……？ C言語(C99)の拡張言語 CUDAに似た接頭辞(__kernel、__globalなど) ベクトル型 マルチコアCPUやCell B/E、DSPなどの並列計算記述にも使えるように規格化されている

43 スパコンとGPU1 GPUスパコンの登場 TSUBAME1.2（東工大） Tianhe-1（中国） Nebulae（中国）
2010/7/26 高性能コンピューティング学論　講義資料 スパコンとGPU1 GPUスパコンの登場 TSUBAME1.2（東工大） TOP500( ) #29、Tesla S 台(x4) PeakFlops TFlops / GPU TFlops Tianhe-1（中国） TOP500( ) #5、RadeonHD4870x2 2560台(x2) PeakFlops TFlops / GPU * TFlops Nebulae（中国） TOP500(2010.6) #2、Tesla C 台(x1) PeakFlops CPU TFlops + GPU TFlops TSUBAME2.0（東工大） TOP500( ) 予定、Tesla M 台(x1) PeakFlops CPU TFlops + GPU TFlops Tianhe 天津、Nebulae 深セン

44 スパコンとGPU2 TOP500とGPU GPUスパコンの課題 スコアを大きく左右するのは行列積の性能（演算速度・メモリ速度）と通信性能
2010/7/26 高性能コンピューティング学論　講義資料 スパコンとGPU2 TOP500とGPU スコアを大きく左右するのは行列積の性能（演算速度・メモリ速度）と通信性能 GPUでスコアを出しやすい問題ではある GPUスパコンの課題 演算効率（性能が出せるか） Tianhe-1やNebulaeは50%も出せていない 稼働率（使われるか） ユーザがGPUプログラムを作成できるか→教育、ライブラリ GPUに適している問題ばかりではない 故障率（壊れないか） 年単位で動かした実績がほとんどない

45 最新GPUの動向 GPUが様々なアプリケーションに有用であることは広く知れ渡りつつある
2010/7/26 高性能コンピューティング学論　講義資料 最新GPUの動向 GPUが様々なアプリケーションに有用であることは広く知れ渡りつつある より様々な分野で利用され始めた プロダクトでの利用が始まった 動画編集ソフト、科学技術計算ライブラリ 倍精度浮動小数点演算性能やメモリの信頼性を求める声が大きくなってきた 倍精度浮動小数点演算性能が可能な（性能が高い）GPUの登場 Fermiアーキテクチャ、RadeonHD5xxx ECCサポートGPUの登場 Fermiアーキテクチャ版Tesla 研究ではなく製品

46 最新のCUDAアーキテクチャ Fermiアーキテクチャの主な更新ポイント チューニングに大きく影響
2010/7/26 高性能コンピューティング学論　講義資料 最新のCUDAアーキテクチャ Fermiアーキテクチャの主な更新ポイント 高速な倍精度浮動小数点演算に対応（単精度の半分の性能） 浮動小数点atomic演算に対応 BlockあたりThread数が増加（512->1024） MPあたりRegister数が増加（8K,16K->32K） MPあたりSharedMemory容量が増加（16KB->48KB） SharedMemoryのバンク数が増加（16->32） キャッシュを搭載（L1$新規搭載、L2$の汎用化） DirectXとCUDAの切り替えが高速化 チューニングに大きく影響 （アーキテクチャ更新が速いのは善し悪し？）

47 GPU vs メニーコアCPU 汎用化するGPU 増加するCPUコア 汎用化したGPUとメニーコアCPUは似たようなハードウェアになる？
2010/7/26 高性能コンピューティング学論　講義資料 GPU vs メニーコアCPU 汎用化するGPU よりプログラムを組みやすく（性能を出しやすく） 増加するCPUコア 2010年7月現在：最大6コア 微細化等によりさらなるコア数増加が予定されている 5年後には100コア規模に達するという予測も 課題：どうやって使う？何に使う？本当に100コアも必要か？（→CPUコアとGPUコアの統合） 汎用化したGPUとメニーコアCPUは似たようなハードウェアになる？ シンプルCPUコアを大量に並べるLarrabeeはCPUとGPUの両方を取り込む……予定だった 私見：「シンプル超多数コアのGPU」＋「複雑さを維持しコア数も抑えるCPU」

48 今後（近い将来）のGPU/GPGPU像 コア数・理論演算性能のさらなる増加 CPUとの統合が進む
2010/7/26 高性能コンピューティング学論　講義資料 今後（近い将来）のGPU/GPGPU像 コア数・理論演算性能のさらなる増加 CPUとの統合が進む IntelのCore iに統合されているGPUはGPGPUに使えないレベル 本命：AMDのFusion 疑問：メモリはどうなる？プログラミングモデルに影響は？ HPC向けGPUとグラフィックスGPUが分かれる可能性 そもそも要求が異なる 共通のベースアーキテクチャを持つ異なる者 プログラミング環境はどうなる CUDA=NVIDIA GPUアーキテクチャ 競争が起きる必要がある->OpenCL？ ※勝手な予想を含む

49 2010/7/26 高性能コンピューティング学論　講義資料 おまけ その他

50 SIMD型アクセラレータ GPU以外にもCPUと比べて浮動小数点演算を並列・高速に扱えるハードウェアは存在する（していた）
2010/7/26 高性能コンピューティング学論　講義資料 SIMD型アクセラレータ GPU以外にもCPUと比べて浮動小数点演算を並列・高速に扱えるハードウェアは存在する（していた） ClearSpeed（英ClearSpeed Technology社） TSUBAMEに搭載されているアクセラレータ 事実上GPUに駆逐された（高価だったのが最大の問題か） Cell Broadband Engine（SCE、ソニー、IBM、東芝） PS3に搭載されていることで有名なプロセッサ HPC向けの強化版や東芝のノートPC搭載SpursEngineなども 一時期GPGPUとともに盛り上がった（次の製品が出ないのが問題か） GRAPE 重力多体問題専用の計算機 現在もバージョンアップが続いており天文学の研究で活用されている 最近、TOP500のLittle Green500で一位になった(GRAPE-DR)

51 おわりに GPGPUについて、およびCUDAについて（広く浅く）紹介した 興味を持った場合は実際にプログラムを作成してみて欲しい
2010/7/26 高性能コンピューティング学論　講義資料 おわりに GPGPUについて、およびCUDAについて（広く浅く）紹介した 興味を持った場合は実際にプログラムを作成してみて欲しい GPUプログラミングコンテストも開催されているので是非挑戦して欲しい （恐らく今年度も開催する、はず） 非常に身近でHOTな（CPUと同等以上に発展著しい）ハードウェア・ソフトウェアなので、挑戦する価値あり 情報を追いかけるだけでも楽しいかもしれない？

52 大島聡史(東京大学 情報基盤センター 助教) ohshima@cc.u-tokyo.ac.jp
2010/7/26 高性能コンピューティング学論　講義資料 内容に関する質問・問い合わせ先 大島聡史(東京大学 情報基盤センター 助教)