汎用の並列ハードウェア
汎用マルチコアプロセッサ
や
インテル汎用CPUのマルチコアの例
• 2005/4 Pentium エクストリームエディション
– ２コア×ハイパースレデッィング ⇒ ４スレッド
• 2006/6 デュアルコアXeon
– サーバー用Xeonでのデュアルコア
• 2006/11 クワッドコアXeon
Xeonでの４コア
• 2007/1 Core2 Quad
• 2010/3 Xeon 7500 8コア×HT⇒16スレッド
• 2011/4 Xeon E7 10コア、2～8ソケット可
• 2013/9 Xeon E5-2600 12コア
• 2014/2 Xeon E7-v2 15コア
• 2014/9 Core i7-5960X 8コアHaswell
• 2014/9 Xeon E5-2600v3 18コア
2 インテル・ミュージアム マイクロプロセッサーの歴史 http://japan.intel.com/contents/museum/hof.html コア自体の向上が目覚 ましいので、単純な コア数の比較は、意味 がないことに注意インテル汎用CPUのマルチコアの構造
3 コア コア コア コア L1 L1 L1 L1 L3キャッシュ L2 キャッ シュ L2 キャッ シュ L2 キャッ シュ L2 キャッ シュ CPUチップ コア コア コア コア L1 L1 L1 L1 L3キャッシュ L2 キャッ シュ L2 キャッ シュ L2 キャッ シュ L2 キャッ シュ CPUチップ…
メモリ
マザーボード 例えば 256KB 例えば 8MB 例えば 32KBx2・端末用のCoreではこの他にGPUがチップ上に載っている
・共有メモリのキャッシュ統一性の問題
GPGPU
4
General Purpose Graphic Processing Unit
（汎用のグラフィックプロセッサ）
従来
の
GPU
グラフィック演算を高速化する
グラフィック演算は各点・各オブジェクトに対する
同じ処理から成っている
⇒ パイプラインを用いて高速化できる
GP
GPU
パイプラインをメモリ共有の多コアとして構成
グラフィック以外の処理も （←「汎用」）
メモリ共有多コア上で並列実行可能にする
プログラミング環境も提供し、プログラマに開放
ＧＰＧＰＵ（１）
• GPUは一般的に画像処理を専門とする演算装置であり、
CPUの制御の下で用いられる動画信号生成専用の補助演算用ICである。
• 動画像の実生成は演算能力が要求されるが、その多くが定式化された単純
な演算の繰り返しであるためハードウェア化に向いており、高速なメモリ
・インターフェース機能と高い画像演算能力を備えたIC製品のシリーズが
いくつも製造・販売されている。
• 90年代中盤以降は3D描画性能が劇的に向上し、それに伴い行列演算を中
心としたSIMD演算機の色彩が強くなってきた。
2000年代に入ると、表現力の向上を求めて固定機能シェーダからプログ
ラマブルシェーダへの移行が進み、演算の自由度が飛躍的に増した。
これをグラフィック・レンダリングのみならず、他の数値演算にも利用す
るのがGPGPUのコンセプトである。
• GPUの浮動小数点演算能力は2010年頃に2000GFLOPSをオーバーした一
方で、CPUの浮動小数点演算能力は頭打ち状態にあり、2011年時点で数
十GFLOPS台に留まっている。GPUは構成が単純であるために浮動小数点
演算での効率がよく、GPU専用にローカル接続されたメモリICとのメモリ
バンド幅を広く備えるために、CPUと比べて性能比で安価かつ成長の伸び
率が高い。
http://ja.wikipedia.org/wiki/GPGPUＧＰＧＰＵ（２）
• GPUはシェーダプロセッサと呼ばれる多数の演算ユニットを持ち、複数
のシェーダプロセッサをまとめてクラスタとしている。
これらの演算器に命令を与えるインストラクション・ユニットはクラス
タごとに1台しか無く、クラスタを構成するシェーダプロセッサはそれ
ぞれ異なるデータを与えられ、そのデータに対して同じ命令内容を一度
に実行する。
• このようなSIMD型データ処理は3次元演算やマルチメディア処理に効
果を発揮する一方で、命令中に条件分岐による分岐が入るとオーバーヘ
ッドがかさみ、途端に効率を落としてしまう。
• また、シェーダプロセッサ間でデータをやりとりする場合、遠くのデー
タバスを経由することになり、それがボトルネックとなってしまう。
http://ja.wikipedia.org/wiki/GPGPUＧＰＧＰＵ（３）
• GPUは倍精度の浮動小数点演算が不得意である。GPUが扱う多くの画
像演算では、整数演算や単精度の浮動小数点演算で足りてしまうため
に浮動小数点演算器は仮数部が24ビット程度とそれほど広くなく、倍
精度の浮動小数点演算を行うには分割して幾度も演算器を使う必要が
ある。
• メモリ環境についても、演算入力は少数の格子点データと幾分大きな
テクスチャ・データだけであり、演算出力は画像1枚程度の大きさのピ
クセルごとに3色のデータを保持しながら順次それらを送り出すだけで
済むため、相応に大きな外部の半導体メモリとかなり広いメモリバン
ド幅による接続で十分に対応しており、演算対象データの局所性が高
いのでローカルメモリと内部キャッシュによってデータの読み書き性
能が向上すると同時に演算も途切れずに順次行える傾向が強い。
http://ja.wikipedia.org/wiki/GPGPUＧＰＧＰＵ（４）
• 基本的にGPUは、配列構造の単純なデータを単精度程度の浮動小数
点演算によって順番に処理することで2次元の動画像データを実時
間内に生成することに特化しているため、それ以外の用途ではあま
り高い性能は期待できない。画像処理専用ICの流用では、科学技術
計算でも倍精度以上の浮動小数点演算を必要としたり、演算の局所
性が低いものではそれほど高い性能は得られない。画像処理専用で
はなく、GPUから派生して新たに開発されたGPGPU用のICでは、
倍精度浮動小数点演算やより広いメモリ空間に対応したものがあり
、これらは広範な科学技術計算への利用が期待される。
http://ja.wikipedia.org/wiki/GPGPUCUDAの概略
• CUDA = NVIDIA GPGPUプログラム開発環境
• ヘテロジニアス (非均一) なハード環境を対象
– ホスト側CPU vs GPGPU側PU(コア)
• 命令体系も違う、置き場所(下記)も違う
– グローバルメモリ vs ローカルメモリ
• 転送が必要になる ⇒ 転送時間がかかる
– でもなるべく一様に (少なくともC/C++言語)
• モジュール (ファイル) 分け、陽にデータ転送指示
– イメージ：転送⇒fork⇒処理⇒join⇒結果転送
• GPGPUハードに密着したランタイムを含む
– スレッドのスケジューリングなど
• GPGPUハードの細かい知識が前提
– 結局パフォーマンス追及したいので、必須
9CUDAの概要（NVIDIA資料）
CUDA はハードウェアとソフトウェアを組み合わせたプラットフォームで、C、C++ 、Fortran などさまざまな言語で書かれたプログラムを NVIDIA GPU で実行すること ができます。 CUDA プログラムはカーネルという名前の並列関数を呼び出します。各カーネルは、 複数の並列スレッドによる並列実行となります。 図1に示すように、スレッドをまとめたものをスレッドブロック、スレッドブロックを まとめたものをグリッドと呼び、プログラムやコンパイラではこれらを単位として取 り扱います。スレッドブロックを構成するスレッド1本1本がそれぞれカーネルのイン スタンス一つを実行するのです 各スレッドは自分が属するスレッドブロックおよびグリッドにおけるスレッド ID と ブロック ID を持つほか、プログラムカウンター、レジスタ、スレッド単位のローカ ルメモリ、入力、出力結果を持ちます。 同時並行で処理を行うスレッドのセットがスレッドブロックです。同じスレッドブロ ックに属するスレッドはバリア同期と共有メモリーにより協調して動作します。スレ ッドブロックも自分が属するグリッドにおけるブロック ID を持ちます。 スレッドブロックを行列としてまとめたものがグリッドです。グリッドは全体で一つ のカーネルを実行し、グローバルメモリーからの入力データの読み取り、グローバル メモリーへの出力データの書き出し、依存関係にあるカーネルコールの同期といった 処理を行います。 http://www.nvidia.co.jp/content/apac/pdf/tesla/nvidia-kepler-gk110-architecture-whitepaper-jp.pdf
CUDAの概要（NVIDIA資料）
CUDA 並列プログラミングモデルでは、レジスタスピルや関数呼び出し、C の自動配列変数など に使うローカルメモリー空間がスレッドごとに確保されます。 スレッドブロックのレベルでは、並列アルゴリズムに必要なスレッド間の通信、データの共有、 結果の共有に使う共有メモリー空間がブロックごとに確保されます。 複数スレッドブロックで構成されるグリッドは、カーネル全体をカバーするグローバルな同期を 行ったあと、グローバルメモリー空間で結果を共有します。 http://www.nvidia.co.jp/content/apac/pdf/tesla/nvidia-kepler-gk110-architecture-whitepaper-jp.pdf#define STRIPS 65536*256 /* Number of total strips */ __global__ void Kernel(FLOAT *Result);
FLOAT h_result[THREAD*BLOCK]; int main(int argc, char** argv){ unsigned int timer;
FLOAT pi;
CUT_DEVICE_INIT(argc, argv); FLOAT *d_result;
cudaMalloc((void**) &d_result, sizeof(FLOAT)*THREAD*BLOCK); cudaMemset(d_result, 0, sizeof(FLOAT)*THREAD*BLOCK);
cudaMemcpy(d_result, h_result, sizeof(FLOAT)*THREAD*BLOCK, cudaMemcpyHostToDevice)；
dim3 grid(BLOCK, 1, 1); dim3 threads(THREAD, 1, 1);
Kernel<<< grid, threads >>>(d_result);
cudaMemcpy(h_result, d_result, sizeof(FLOAT)*THREAD*BLOCK, cudaMemcpyDeviceToHost);
pi=0.0;
for (int k=0; k<THREAD*BLOCK; k++) { pi += h_result[k];
}
printf(“GPU結果 =%15.12f¥n”,
(pi - (((FLOAT)0.5)/((FLOAT)STRIPS))) * (FLOAT)4.0); cudaFree(d_result);
}
__global__ void Kernel(FLOAT *Result) {
//GPUでの処理
int ix = blockIdx.x*blockDim.x + threadIdx.x; FLOAT tmp = 0.0; for(int k=0; k<(STRIPS/BLOCK/THREAD); k++){ FLOAT x = ((FLOAT)((ix*(STRIPS/BLOCK/THREAD))+k)) / ((FLOAT)STRIPS); tmp += ((FLOAT)sqrt(1.0-x*x)) / ((FLOAT)STRIPS); } d_result[ix] = tmp; }
CUDAのイメージ
区分求積でπ計算の例
ホスト側プログラム GPU側プログラム グローバルメモリ上に結果の 書込み用領域h_resultを用意 GPU側メモリ上に結果の 書込み用領域d_resultを用意し グローバルメモリのh_result をd_resultに転送コピー GPU側プログラムを fork ⇒ 実行 ⇒ join GPU側メモリ上の結果の値 d_resultをグローバルメモリ 上のh_resultに転送コピー ホスト側でそれぞれの GPUの結果を集計 GPU側では、自分の割当て範囲 のストリップの高さを集計し、 d_resultに書き込む 割当て範囲はホストからもらう のではなく、自分のIDと総スト リップ数から計算しているNVIDIA Titan
http://www.nvidia.co.jp/object/geforce-gtx-titan-jp.html
