1
山崎研究室紹介
２０１２年１０月２２日山崎 勝弘
[email protected] １．研究室の目標 ２．育成したい人材像 ３．指導方針 ４．研究分野： 並列処理とハード/ソフト・コデザイン ５．研究テーマ ６．貴君らに提供できること 1 2１．研究室の目標
• ハードウェアとソフトウェアの両方分かる人
材の育成
• コミュニケーション能力、スケジューリング
能力、および知的体力の養成
• 社会人としての基本的素養をつけ、努力を
継続して、目標を達成できる人材の育成
• 並列処理とハード/ソフト・コデザインを融合
した高性能な問題解決システムの構築
2 3２．育成したい人材像
• あいさつ
• コミュニケーション能力
– 日本語で正しく表現、発表、議論
– 英語能力
• スケジューリング能力
– 立案、実行、チェック（Plan, do ,check）
• 知的体力
– 最後まであきらめずにやり遂げること
3 4３．指導方針
• 前向きに楽しく
： positive thinking
• 研究テーマの設定、研究環境の整備
• 社会人としての基本的素養
をつけて欲しい。
• 社会人になるための実力
をつけて欲しい。
• 英会話学習のきっかけ
をつかんで欲しい。
• 自分の夢
を将来にわたって実現して欲しい。
4 5４．研究分野
並列処理とハード/ソフト・コデザインを融合 した高性能な問題解決システムの構築 5研究経歴
（1976～現在）
FPGA ハードソフト アーキテクチャ 並列アプリ 知識工学 QA-1 MUNAP レイトレーシング ラジオシティ ハード/ソフト・コラー ニングシステム JPEG、MPEG N体、グリッド 電力系統 知的CAI 事例ベース 並列プログラミング ハッシュ関 数高速化 同期マルチ メディア 命令スケ ジューリング ガラス傷検出 OpenMP 動作合成 リアルタイム レイトレーシング モンテカルロ 囲碁 ルーティング トランスピュータ AP1000+ PCクラスタ SMPクラスタ KSR1 GPU SMPクラスタ ＧＰＵ FPGA プロセッサ４台の並列マシン マイクロプログラム制御 ハードウェア、システムソフト アプリケーション SMPクラスタ 情 報 工 学 科 情 報 学 科 電 子 情 報 デ ザ イ ン 学 科 電 子 情 報 工 学 科 1988 2004 1994 2012 MAP設計 データパス 可視化なぜハードとソフトか
• 人間： 強い肉体（体力）＋知恵（知力） • コンピュータ： 高い性能（ハード）＋使いやすさ（ソフト） • ハードのみ： 設計は難しい。組み立ては容易。 – 差別化が難しい。低価格化にさらされる。 • ソフトのみ： 技術者が多い。 – Cプログラミングだけなら理工系なら誰でもできる。 • ハードとソフト： できる人材が極めて少ない。ニーズは高い。 • Cプログラミング：アセンブリ言語やコンパイラのコード生成が 分かれば、コンパクトで速いプログラムができる • プロセッサ設計：命令の使われ方や使用頻度が分かれば、 良い命令セットを設計できる。 7 8ハード/ソフト・コデザインの目標
• プロセッサ設計 • 命令セットアーキテクチャ • アセンブリ言語、C言語、コンパイラ • ハードウェア設計言語 • ハードウェアとソフトウェアの最適バランス • FPGA：プログラム可能なLSI 目標 • 学生独自のプロセッサを設計してFPGA上で動かす。 • 画像処理、ガラス傷検出、囲碁などをFPGA上で動かす。 • FPGA上で動作させて、感激を味わう。 8 9並列処理の目標
• マルチコアプロセッサ • GPU(Graphics Processing Unit) • 超並列マルチスレッド処理 • SMPクラスタ • ハイブリッド並列処理 – 共有メモリ＋分散メモリ、OpenMP＋MPI 目標 • リアルタイムレイトレーシングや画像処理など、大規模問題 をSMPクラスタやGPU上で高速に解く。 • 実際に動作させて、高速化を実感し、感激を味わう。 9 10５．研究テーマ
５．１ マルチALUプロセッサの設計とFPGAボ
ード上での検証
５．２ FPGAを用いたガラス傷検出の高速化
５．３ GPUによるリアルタイムレイトレーシング
５．４ OpenMPハードウェア動作合成システム
５．５ プロセッサのデータパス可視化
５．６ モンテカルロ囲碁システム
５．７ 命令スケジューリング
10５．１ マルチALUプロセッサMAPの設計
• 複数ALUによる並列処理 ALU数：２，４，８，１６
• レジスタ数は32で、全ALUで共有
• 演算レベル並列処理
：ALUで
並列演算
と
連鎖演算
• HDLによるMAPの設計と論理シミュレーション
• FPGAボード上での実装・評価
• ハードとアセンブラによる並列演算と連鎖演算の検出
112ALUのMAPのデータパス
1213
MAPの命令セットアーキテクチャ
1命令３２ビット ４命令形式 MIPSサブセットALU並列処理
14ハードウェアによる並列性の検出
PPUで並列演算、連鎖演算、単一演算を判定
並列演算 連鎖演算MAPプログラミング例
整数同士の乗算 c=a×b
プログラマは並列演算や連鎖演算を考慮せずに 1命令ずつ記述する。 15MAP用アセンブラ
• 2命令を取り出し、各オペランドのデータ依存性を調べる • データ依存がなければ並列演算、あれば連鎖演算 16 並列 並列 連鎖 連鎖プロセッサデバッガ・モニタ
コマンドをデコード ボードコンピュータを制御 コマンドを入 力 17FPGAボード上への実装
• Spartan-3A Starter Kitと、論理合成ツールISE13.2を使用 • プロセッサデバッガを設計・開発中現状と今後の研究内容
• ２ALUのMAPをHDLで設計し、各モジュールをデバッグ • Booth乗算の論理シミュレーション OK • MAPアセンブラ、シミュレータを作成、MAPプログラミング • プロセッサデバッガのユーザインタフェースを設計 • プロセッサデバッガを作成中 今後の研究内容 • MAP本体をSpartan3A上で実動作させる • 並列演算、連鎖演算の有効性を検証する • ALU数：４，８のMAPを設計し、並列処理の有効性を評価する 195.2 FPGAを用いたガラス傷検出の高速化
• 液晶用ガラスの欠損検出を画像処理によって行い，それを ハードウェア化することによって高速に処理 • 本システムは雑音除去のためのTDI，欠損検出のためのラプ ラシアンフィルタ，2値化，ラベリングから構成される 画像撮影 画像データの 読み込み ノイズ除去 傷のエッジ抽出 閾値処理 画像内の オブジェクトの 認識 傷の検出 欠損（傷） 20欠損検出システムの構成
ラプラシアン &2値化 ラベリング 平均値計算 注目画素と 8近傍画素 の格納 ラプラシアン 処理 メモリに保存 2値化 注目画素と 左上,上,右上, 左画素の格納 仮ラベル 生成 注目ラベルと 8連結ラベル の格納 ラベル 補正 撮影した画像 の格納 TDI 画像出力 21①
②
③
TDI(Time Delay Integration)
• 対象物を縦方向に1ラインずつずらして複数回撮影 • 撮影画像の共通部分を重ね合わせて平均値を取る • ノイズの影響を減少 22 ・・・ 欠損（傷） 縦方向に1ラインずつずらして撮影 共通部分を重ね合わせて平均値を取るラベリング
• 画像中の連結成分にラベルをつける • 1回目のスキャンで仮ラベルを生成 • 2回目のスキャンでラベルを補正 仮ラベル生成 1 1 1 2 1 1 1 2 2 1 1 2 2 3 3 3 4 4 4 4 ラベル補正 1 1 1 1 1 1 1 1 1 1 1 1 1 3 3 3 4 4 4 4 23欠損検出システムの構成
CPU 画像ファイル TDI処理 ① BRAM ② Register File ラプラシアン ＆2値化 仮ラベル 生成 ラベル 補正 ③ ④ ⑤ ⑥ ⑦ ⑧ 画像出力 FPGA ⑨ 24ラプラシアン&2値化モジュール
①レジスタファイルにデータを保存 ②注目画素と周囲8画素を取り出 し，マスクパターンとかける ③ラプラシアン処理後のデータを 閾値処理 ④2値化後のデータをBRAMへ書き 込む ①～④を256*256分行う 画像中の傷は，はじめ黒色だが， ラプラシアンフィルタ，2値化処理後 は傷が白色に変化する 25 256 Register File BRAM 0 0 0 0 0 72 0 68 28 ラプラシアン フィルタ 2値化 0 256 3*3画素 ① ② ③ ④ 72ラベリングモジュール（仮ラベル生成）
256 Register File BRAM 仮ラベル 生成 1 256 ・白→flagを1に ・黒→そのまま 0 0 2 1 5画素 flag=1 label=0 flag=1 label=1 26 ①レジスタファイルにデータを保存 ②画素値が255ならflagに1を入れ て下位8bitをリセット ③flagが1なら周囲4画素（左上，上 ，右上，左の優先順位順）のラベル を参照しながら仮ラベルをつける ④仮ラベルをBRAMへ書き込む ①～④を256*256分行う ラベルの総数は入力画素と同じ 8bitまでとし，0から最大255まで ① ② ③ ④ラベリングモジュール（ラベル補正）
256 Register File BRAM 0 0 0 0 2 1 2 2 1 1 ラベル補正 256 3*3画素 27 ①レジスタファイルにデータを保存 ②注目画素と周囲8画素を取り出 し，ラベル補正を行う ③補正後のラベルをBRAMへ書き 込む ①～③を256*256分行う ③ ② ①実験結果
回路規模 処理時間 28 ラプラシアン 2値化 ラプラシアン 2値化 ラベリング LUT数 28867 31132 LUT使用率(%) 64.43 69.49 レジスタ数 76 93 レジスタ 使用率(%) 0.17 0.21 最大動作 周波数(MHz) 105.98 105.85 0 10 20 30 40 50 60 70 ラプラシアン 2値化 ラプラシアン 2値化 ラベリング CPU FPGA 実 行 時 間 （ ｍ ｓ ） 15 1.24 62 3.112.1倍
20倍
画像サイズは256*256 TDI処理 128回 Virtex5 ML507評価ボードCore 2 Quad CPU Q 9400 2.67GHz 実装メモリ4.00GB
5.3 GPUによるリアルタイムレイトレーシング
スクリーン，視点，物体，光源を作成 （１）スクリーン上の画素に向けて光線を発生 （２）すべての物体と光線の交差判定を行う （３）光線と接した物体のうち最も距離の近い 物体を抽出 （４）抽出した物体の輝度の計算をする （５）反射や屈折が起これば， それを新たな光線とみなして （２）～（４）を繰り返す 29 視点 スクリーン 光線 反射 光源 屈折 反射 屈折Geforce GTX 480
• NVIDIA社から2010年発表 • Graphics Processing Cluster(GPC)を4個搭載 • 各GPCはストリーミング・マル チプロセッサ（SM）を4個搭載 • GPU内にSMが16個 • SMが並列処理の単位 30 SM SM SM SM ラスターエンジン GPC ホストコンピュータ Giga スレッドスケジューラ GPC L2キャッシュ:768KB GPC GPU グローバルメモリ:1536MB GPCストリーミング・マルチプロセッサ(SM)
• スカラー・プロセッサ （SP）が32個 • 共有メモリ：64KB • レジスタファイル： 128KB • 32個のスレッド（ワー プ）による並列処理 • Single Instruction Multiple Thread （SIMT）型プロセッサ SP SP SP SP SP SP SP SFU SFU SFU SFU 共有メモリ/L1キャッシュ:64KB テクスチャ・キャッシュ 命令キャッシュ ワープ・スケジューラ ワープ・スケジューラ ディスパッチ・ユニット ディスパッチ・ユニット LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST SM RF RF RF RF RF RF RF RF SP SP SP SP SP SP SP SP RF RF RF RF RF RF RF RF SP SP SP SP SP SP SP SP RF RF RF RF RF RF RF RF SP SP SP SP SP SP SP SP RF RF RF RF RF RF RF RF SP 31SPDシーンデータ
• SPD：レイトレーシング用のベンチマークシーンデータ
• 視線，背景，光源，物体素材，物体形状，位置の定
義が含まれている
ｔeapot ｔetra mount
32
teapot tetra mount
ポリゴン数 2328 4096 8192 サイズ（KB) 370 249 617
画面のブロック分割
• スクリーンを複数の
ブロックに分割する
• 1個のブロックを1個
のSMで処理する
• 16ブロックを16個の
SMで並列処理
• SM内の32個のSPで
マルチスレッド処理
33 ブロックSM内でのマルチスレッド処理
• 1画素が1スレッドに
対応
• 各スレッドはSPで処
理される
• 32スレッド（ワープ）
の並列処理
• 64ピクセルで2ワー
プ処理
ブロック 1画素 1スレッドSP×32個
34GPU上での並列処理
SPDファイル情報 読み込み 画像生成 スクリーン 情報受信 CPU GPU YES NO SPD情報 配列変換 スクリーンを指定されたブロック数に分割 各ブロックをワープ単位でスレッド分割 スクリーンにレイトレーシング結果書き込み ・・・ ブロック数, スレッ ド数の指定 GPU側にSPD情報 転送 SM0 SP SP SP SP SP SP 32スレッド（1ワープ）で動作 SM1 SM15 ・・・ 光線の方向探索 交差判定 反射 輝度計算 35 マルチスレッド処理実験方法と実験環境
• CPUによる逐次処理とGPUによる並列処理 – 画像の解像度は512*512 • 画面分割はブロック数，スレッド数を変化させて実行 • CPU逐次処理の環境– Intel(R) Core i7 CPU 950 @3.07GHz 6.00GB
• GPU並列処理の環境
– Intel(R) Core i7 CPU 950 @3.07GHz 6.00GB
– GPU：Geforce GTX480 （コアプロセッサ数 512個） グラフィッククロック：700MHz，メモリクロック：1848MHz
– 並列処理環境：CUDA
ブロック数 256*256 128*128 64*64 32*32 16*16
実行時間と速度向上
ブロック数 スレッド数
teapot
tetra
mount
実行時間 (秒) 速度向上 (倍) 実行時間 (秒) 速度向上 (倍) 実行時間 (秒) 速度向上 (倍) 256*256 2*2 13.9 14.2 7.75 12.0 20.7 12.7 128*128 4*4 3.60 55.0 2.00 46.6 5.22 50.4 64*64 8*8 1.71 116 0.92 101 2.41 109 32*32 16*16 1.83 108 0.94 99.1 2.45 107 16*16 32*32 2.17 91.2 0.99 94.1 2.56 103 CPU 実行時間 １９８ ９３．２ ２６３
画面のサイクリック分割
• 画面を64＊64のブロックに分割、各ブロック8＊8
• 16個のSMがブロックをサイクリックに担当
• 負荷均衡を図る
適応型空間分割
• 各空間に含まれる物体 が最も多い空間を２つ に分割する • 上の処理を繰り返し，空 間を小さなブロックに分 割する • 分割されたブロックを各 SMに割り当てる 39 40 40５．４ OpenMPハードウェア動作合成システム
コード ジェネレータ OpenMP プログラム (動作記述) ハードウェ ア制約 SMP環境 (SMPクラスタ) シミュレーション 並列動作 ハードウェア OpenMP コンパイラ トランスレータ マルチスレッド プログラム 並列動作HW 中間表現 アルゴリズム評価 ハードウェア合成 並列アルゴリズム 並列効果 回路規模 性能評価 OpenMP構文を 利用した動作合成 早期に並列化手法の検討 検証時間の短縮 逐次プログラムから の段階的な設計CベースとOpenMPベースの比較
Cベース
41OpenMPベース
Cプログラム OpenMPプログラム 並列リージョン ユーザ記述 ハードウェア記述 自動並列化 並列リージョン からの変換 42データ並列の実行モデル
• データ並列性のある処理をノードで分担
• 100回の繰り返しを4ノードで分担する場合
#pragma omp parallel for for(i=0;i<100;i++) { 処理A(i); } #pragma omp parallel for
(スレッド生成 & fork)
処理Aスレッド (i=0～24) 処理Aスレッド (i=25～49) 処理Aスレッド (i=50～74) 処理Aスレッド (i=75～99)
join 終了時のリダクション演算 （結果の足し合わせなど) 各スレッドはforループ内の繰り返し処理を分担 sum= sum+i
43
データ並列のハードウェア構成
逐次データパス Memory Register 並列データパス Op処理A 処理A 処理A 処理A
Control Control 並列HWの起動・終了 Arbiter スレッド間共有データへ の同時アクセス制御 Memory Register 並列データパス 動作時は停止
コードジェネレータの改善
44State1,2 state3 State4,5 state6 state7 state8 false true true false a = a + 1 b = b + 2 state1 state2 state3 state4 state5 state6 state7 state8 false true true false a = a + 1 b = b + 2 ・Sジェネレータ（１演算）とMジェネレータ （複数演算） ・状態数と一時レジスタの削減による高 速化を図るコードジェネレータを作成中 ・マンデルブロー、エッジ抽出で速度が 25％向上 45
５．５ モンテカルロ囲碁 各種知的ゲームの比較
チェス 将棋 囲碁 局面数 １０１２０ １０２００ １０３００ コンピュータ チャンピオン 女流トッププロ の強さ に勝利 前名人以上 アマ四段程度 人間に勝つ １９９７年 ２０１０年 女流 ２０３０年？ のはいつ Kasparov ２０１５年？ 名人 トピック IBM 評価関数の モンテカルロ囲碁 DeepBlue 機械学習 人間との対戦 激指 清水上アマ名人に○ MoGo プロ8段に9子で○ 稲葉アマ準名人に× CrazyStone あから2010 清水女流王将に○ 青葉4段に8子で○ 2012年 ボンクラーズ 米長前名人に○ 武宮9段に6子で○ 46モンテカルロ囲碁とは
• 囲碁は陣地をたくさん取った方が勝ち
• ある局面から最後まで実際に打って（プレイアウト）、
最も勝率の高い手を次の手とする。
• ランダムプレイアウト：乱数を用いて、適当に最後まで
打つ（弱い）
• パターンプレイアウト：よく出てくるパターンに基づいて
最後まで打つ（上より強い）
• 次の候補手を５つ位選ぶ→プレイアウト→最も勝率の
高い手を次の手とする。
モンテカルロ囲碁の並列化
47 盤面認識 候補手生成（５つ） プレイアウト １ ２ ３ ４ ５ 0.4 0.5 0.3 0.7 0.6 勝率 １ ２ ３ ４ ５ ４を打つ プレイアウト 黒勝ち 白勝ち 0.8 0.66 0.6 0.5 大量のプレイアウトを高速に行う １万回/秒 ４プロセッサでは４万回/秒 48 Miss-prediction Bias (MPB) の発見と応用 予測ミスが一部の分 岐命令に集中する。 予測器LPHTを追加し て、再予測する。 分離式BTB (Branch Target Buffer)無条件分岐命令を従 来のBTBから分離し， 小さい特殊なBTBを追 加する。
５．６ 命令スケジューリング
条件分岐命令（ beq ）の分岐予測の精度を上げる。 ＢＴＢ：beq命令の分岐先の予測アドレスを格納する表 ＢＴＢを工夫、追加して精度を上げる。49
６．貴君らに提供できること
• 問題解決の仕方
– 卒論、進路、就職、‥
• スケジューリングの仕方
– Plan, do, check
• 研究発表の仕方
– 日本語文章、スライド作成、発表、‥
• 英会話勉強の仕方
– マンチェスター大学 客員研究員 1992年～1993年 – IEEE student branch カウンセラー 2006年から5年間 – 英語プレゼン大会 2004年から9回実施 49
