データストリーム処理における
GPU タスク並列を用いたスケーラブルな異
常検知
上野 晃司（東京工業大学）
鈴村 豊太郎（東京工業大学／ IBM 東京基礎研究所）
データストリーム処理における GPU タスク並列
を用いたスケーラブルな異常検知
入力・工場のセンサーデータ
・ e.t.c.
リアルタイムに解析
データストリーム処理系に実装
GPGP
U
膨大な数のセンサーデータ
異常？異常？
センサーデータ センサーデータ
複数センサーの
データを同時に計算
GPU タスク並列
複数センサーの
データを同時に計算
GPU タスク並列
高速化
2/38
目次
研究背景
既存手法とその問題点
GPU タスク並列
実装と評価
関連研究
まとめと今後の展望
3/38
目次
研究背景
既存手法とその問題点
GPU タスク並列
実装と評価
関連研究
まとめと今後の展望
4/38
研究背景 データストリーム処理とは
データを蓄積せず、オンメモリ上でリアルタイムに
処理を行う
従来は、データを蓄積してから計算するバッチ処理が主流
大規模なデータの処理やリアルタイムアプリケーションの実
現に有効
データストリーム処理系の例
IBM : System S
Yahoo! : S4
意味のある必要なデー タ
時々刻々と流れてくる 大規模なデジタルデー タ
時々刻々と流れてくる 大規模なデジタルデー タ
リアルタイムアプリ ケーションの実現 リアルタイムアプリ ケーションの実現 リアルタイムに
オンメモリ上で処理 リアルタイムに オンメモリ上で処理
データストリーム 処理系
5/38
本研究で利用 本研究で利用
研究背景 データストリーム処理系 System S
IBM Research が開発したデータストリーム処理系
製品名 : InfoSphere Streams
高い柔軟性と多様な最適化機構
容易に並列分散処理が可能
X86 Box
X86 Blade Cell
Blade X86 Blade
FPGA Blade X86
Blade
X86 Blade X86
Blade X86 Blade
X86 Blade
Operating System Transport
System S Data Fabric Processing
Element Container
Processing Element Container
Processing Element Container
Processing Element Container
Processing Element Container
最適スケジューラー：資源管理など自動化 最適スケジューラー：資源管理など自動化
Linuxマシンなどコモディティクラスタ上で稼
働
Linuxマシンなどコモディティクラスタ上で稼
働
6
研究背景 変化点検知アルゴリズム SST
*
変化点検知
時系列データの特徴が変化したこ とを検知すること
異常検知に使用される
SST
Singular Spectrum Transformation
特異スペクトル変換
優れている点
少ないパラメータで広範囲の異常 検知に適用可能
欠点
計算量が多い
SVD( 特異値分解 ) を使っているた
* Tsuyoshi Ide, et al. Knowledge Discovery from Heterogeneous Dynamic め
Systems using Change-Point Correlations. 2005. SIAM International Conference on Data Mining (SDM 05).
SVD
Ul = [u1; u2; …; ul] β(t)
z(t) = β(t)
T U l Tβ(t)
||Ul Tβ(t)||
SST 計算アルゴリズム SST 計算アルゴリズム
H(t)T G(t)T
time ウィンドウサイズ
7/38
研究背景 リアルタイム異常検知における問題
膨大なセンサーデータのリアルタイム処理
工場のセンサーの数＝ 数万以上
5 万個のセンサーがあるとする
4node(16 コア ) で分散処理しても、 SST で各センサーからの入
力データを 1 つ処理するのに、 10 分かかる
IKA-SST
*1*2(SST の近似を高速に求めるアルゴリズム ) で計算し
ても、 3.2 秒かかる
データの間隔 ( 例えば 0.1 秒 ) によっては計算が間に合わない
高速化に GPU を利用する
CPU:AMD Phenom X4 9850(4 コ ア )Software: LAPACK, ATLAS
行列サイズ : 256 の場合
*1 Implicit Krylov Approximation based Singular Spectrum Transformation
*2 Tsuyoshi Ide, Koji Tsuda. Change-point detection using Krylov subspace
learning. SDM2007, 2007. 8/38
目次
研究背景
既存手法とその問題点
GPU タスク並列
実装と評価
関連研究
まとめと今後の展望
9/38
先行研究
10/38
2010 年 6 月
森田康介、鈴村豊太郎「データストリーム処理を用いた変化点検知アル ゴリズム SST の GPU による性能最適化」電子情報通信学会技術研究報告
既存ライブラリを使った高速化
2010 年 10 月
上野晃司 , 鈴村豊太郎 . データストリーム処理における GPU タスク並 列を用いたスケーラブルな異常検知機構の実現 . インターネットコン ファレンス 2010.
GPU タスク並列を採用
特異値分解の前半のみ（＝ SST の半分 ）しか実装していなかっ
た
本発表
SST をすべて実装し、さらに IKA-SST も実装
既存ライブラリを利用した場合とその問題点
CULA を利用した場合の問題
目的
SVD を GPU で高速化
CULA
GPU を利用した線形代数計算ライ
ブラリ
提供元 : EM Photonics
小さい行列には最適化されてい
ない
SST は実用的には 500 以下の小さ
な行列を使用する場合が多い
複数センサーに対応できない
１秒 / １個 → 100 秒 /100 個
0 10 20 30 40 50 60
GPU
ウィンドウサイズ（行列サイズ）
実行時間(秒)
0 1 2 3 4 5
ウィンドウサイズ（行列サイズ）
実行時間(秒)
SST 実行時間
（データ１つあたりの時間） SST 実行時間
（データ１つあたりの時間）
小さい行列サイズ では CPU より遅
い
↓ 高速
better
参考：森田康介、鈴村豊太郎「データストリーム処理 を用いた変化点検知アルゴリズム SST の GPU による
性能最適化」電子情報通信学会技術研究報告 . 11/38
既存ライブラリを利用した場合とその問題点
行列の大きさと性能
小さい行列は並列度が小さくコアを使い切れない
データ並列だけでは並列度が足りない
大きい行列
大きい行列 小さい行列小さい行列
いないコア使われて いないコア使われて 利用率
GPU GPU
目次
研究背景
既存手法とその問題点
GPU タスク並列
実装と評価
関連研究
まとめと今後の展望
13/38
GPU タスク並列
小さい行列でも並列化すれば速くなる
タスク並列性も使う
小さい行列 小さい行列 小さい行列
小さい行列 小さい行列小さい行列小さい行列小さい行列 小さい行列 小さい行列
小さい行列 小さい行列
従来の計算方法 GPU タスク並列
（提案手法）
複数センサーの
データ ( 行列 ) を
同時に計算
複数センサーの
データ ( 行列 ) を
同時に計算
GPU タスク並列の実現方法
CUDA によるプログラム
Kernel1 ()
Kernel2 ()
CPUCPU
GPUGPU
CPUCPU
GPUGPU Serial code
Serial code
ホスト側のコードは
シーケンシャル
並列に計算できない
部分は CPU で計算する
（ SVD では QR 法による 特異値の計算が該当）
並列に計算できない
部分は CPU で計算する
（ SVD では QR 法による 特異値の計算が該当）
GPU コードは カーネル立ち上げ (≒ 関数呼び出し ) によって実行される
GPU コードは カーネル立ち上げ (≒ 関数呼び出し ) によって実行される
15/38
GPU タスク並列の実現方法
ホストスレッドのマルチスレッド化
タスクの数だけホストスレッドを作成
スレッド 1
(Task1) スレッド 2(Task2) スレッド 3(Task3)
OS のスケジューリング による並列化
OS のスケジューリング による並列化
CPUCPU
GPUGPU GPUGPU GPUGPU CPUCPU CPUCPU
CPUCPU CPUCPU CPUCPU
コンカレントカーネル による並列化
コンカレントカーネル による並列化
GPU タスク並列の実現方法
GPU 部分の並列化方法
ただし、
コンカレントカーネルによる並列化
複数カーネルの同時実行機能
Fermi 以降の GPU でサポートされる
最大同時実行数は 16 まで
小さい行列では性能が低下する
そこで、
スレッドブロックごとにタスクを割り当てて並列化
１つのカーネルで複数タスクを並列処理
スレッドブロックの独立性を利用
CUDA では warp(32 スレッド ) 毎に異なるコードを実行可能
17/38
GPU タスク並列の実現方法
解決方法
計算スレッドと GPU 管理 ( ホスト ) スレッドを分け
る
GPU 部分の並列実行が可能になる
CPUCPU
GPUGPU GPUGPU
GPUGPU CPUCPU CPUCPU
CPUCPU CPUCPU CPUCPU
CPU
CPU スレッド 1
(Task1) スレッド 2(Task2) スレッド 3(Task3)
GPU ホストスレッド
GPUGPU GPUGPU
GPUGPU GPU 処理を委譲
CPU 処理を再開
カーネルコード例
GPU タスク並列の実現方法
スレッドブロックごとにタスクを割り当てる
struct Task {
int m, n;
float* a;
};
/* taskArray は Task へのポインタの配列 */
__global__ void kernel_func(struct Task* taskArray[]) {
Task* task = taskArray[blockIdx.y];
int m = task ->m, n = task ->n;
float* a = task ->a;
/* 何かを実行する */
}
１つのタスクを 表す構造体 １つのタスクを
表す構造体
ブロック ID を 使ってタスクを
識別
ブロック ID を 使ってタスクを
識別
複数のタスクを
並列処理する
複数のタスクを
並列処理する
19/38
GPU タスク並列の実現方法
CPU スレッド数過多による性能低下
タスクの数だけ CPU スレッドが作成される
問題：スレッド数が多すぎて性能が低下する
e.g. タスク数 100 ～ 1000 の場合
CPUCPU
GPUGPU GPUGPU
GPUGPU CPUCPU CPUCPU CPU
・・・・
……
タスクの数だけ CPU スレッドが
作成される タスクの数だけ CPU スレッドが
作成される スレッド 1
(Task1)
スレッド 2
(Task2) スレッド N(TaskN)
GPU ホストスレッド
GPU タスク並列の実現方法
解決方法（完成）
CPU カーネルスレッドの数を減らす
スレッドプールを利用
CPU コードをイベントドリブン方式に書き換え
GPUGPU GPUGPU
GPUGPU CPU
…… CPUCPU
スレッドプール CPUCPU
CPUCPU CPUCPU
… …
GPU ホストスレッド
21/38
目次
研究背景
既存手法とその問題点
GPU タスク並列
実装と評価
関連研究
まとめと今後の展望
GPU タスク並列の実装と評価
実装したアルゴリズム
変化点検知アルゴリズム SST
少ないパラメータで広範囲の異常検知に適用可能
SVD （特異値分解）を使った変化点検知アルゴリズム
SVD 演算の遅さが問題
変化点検知アルゴリズム IKA-SST
SST の近似を高速に求めるアルゴリズム
SST と同じように行列計算で求める
GPU での高速化が見込める
精度は SST より少し劣る
近似する関係で必要なパラメータが増える
2 つのアルゴリズムを GPU タスク並列で実装
23/38
SST の GPU タスク並列実装
SVD （特異値分解）演算の処理フロー
CPU 通信 GPU
① 行列データ
② 二重対角化
左右の行列も同時に計算
③ 二重対角要素
④ 特異値を計算
(∑) ⑤ 計算データ
⑥ 特異ベクトル を計算 (U, V)
⑦ 特異ベクトル ハウスホルダー
変換による 二重対角化
QR 法による 対角化
左右の行列
特異ベクトルの 中間データ
シーケンシャルな 計算なので CPU で
実行 GPU では
複数の行列を タスク並列で計算
SVD を GPU タスク並列で高速化
参考Sheetal Lahabar, et al. Singular value decomposition on GPU using CUDA. IPDPS. 2009.
32.0 96.0 160.0 224.0 288.0 352.0 416.0 480.0 GPU/CPU 1 core (task=64)
GPU/CPU 1 core (task=256)
Matrix Size (N x N)
Speedup
SST の GPU タスク並列実装
SVD の性能（ CPU 実装と比較）
CPU 実装は LAPACK,ATLAS を使用 , 単精度
タスク数は 64, 256 で計測 , データ転送時間を含む
CPU:AMD Phenom X4 9850(4 コア ), MEM:4GB OS:Cent OS 5.4 GPU:Tesla C1060, Software:CUDA 3.2, GCC 4.1.2
CPU1 コア と比較
CPU4 コア と比較
CPU
計算部分が
速度に影響
3~4 倍
に高速化25/38
SST の GPU タスク並列実装
変化点検知アルゴリズム SST に適用
SVD の GPU タスク並列実装を使って SST を実装
GPU CPU 参考 :CULA
0 0.2 0.4 0.6 0.8 1 1.2
0.05
0.37
実行時間(秒) 1.06
CPU と比較して 約
8
倍高速化 CPU と比較して約
8
倍高速化データ１つあたりの SST の実行時間
（ CPU １コア）
データ１つあたりの SST の実行時間
（ CPU １コア）
better
CPU:AMD Phenom X4 9850, MEM:4GB OS:Cent OS 5.4 GPU:Tesla C1060, Software:CUDA 3.2, GCC 4.1.2
行列サイズ：
256
行列サイズ：
256
26/38
SST の GPU タスク並列実装
System S 実装のデータフロー図
IBM のデータストリーム処理系 System S 上に実装
センサーデータを入力
計算結果を集約
4 台のマシンで
SST を計算
4 台のマシンで
SST を計算
CPU:AMD Phenom X4 9850(4 コア ) GPU:GeForce 8800 GTS 512
OS:Cent OS 5.4
Linux クラスタ
System S による分散処
理
マシン 1 台に 4 台
の GPU を搭載
27/38
0 4 8 12 16 0
50 100 150 200 250 300 350
Number of GPUs/CPUs
Throughput (# of scores/sec)
SST の GPU タスク並列実装
評価
最大４台マシン (16GPU,16CPU コア ) で分散処理
最大スループットを測定
CPU:AMD Phenom X4 9850(4 コア ), MEM:4GB OS:Cent OS 5.4 GPU:GeForce 8800 GTS 512, Software:CUDA 3.2, GCC 4.1.2
CPU 実装と比較して
約 7 倍
に高速化行列サイズ：
320
行列サイズ：
320
28/38
*1 Implicit Krylov Approximation based Singular Spectrum Transformation
*2 Tsuyoshi Ide, Koji Tsuda. Change-point detection using Krylov subspace learning. SDM2007, 2007.
変化点検知アルゴリズム IKA-SST
*1*2
IKA-SST にも GPU タス
ク並列を適用して高速
化
H(t)T G(t)T
time
IKA-SST 計算アルゴリズム IKA-SST 計算アルゴリズム
ウィンドウサイズ
F/B to (t+1) べき乗法べき乗法 Lanczos 反復
Lanczos 反復
の固有ベクトル の固有ベクトル
ウィンドウサイズと
同じ次元での計算
圧縮された次元での計算
（ 5 ～ 10 次元）
29/38
IKA-SST の GPU タスク並列実装
処理フロー
CPU 通信 GPU
① データ
② 行列 T� を計
③ 行列 T� 算
④ 変化度スコア z を計算 行列サイズが
小さすぎる (5~10) ので GPU 化は難しい
GPU では 複数の行列を タスク並列で計算 時系列データ
フィードバック GPU メモリ
GPU 内に
過去のデータを保持 入力センサー
データのみ GPU に送信
対称３重対角行列 なので対角要素と 上対角要素のみ転送
30/38
100 150 200 250 300 350 400 450 500 0
5 10 15 20 25
30 GeForce GTX460
Tesla C1060
GeForce 8800 GTS 512
Window Size(N x N)
Speed up
IKA-SST の GPU タスク並列実装
性能評価： CPU 実装と比較
SIMD を使って最適化した CPU 実装 1 コアと比較
=100~500, 単精度 , データ転送時間を含む
CPU:AMD Phenom X4 9850(4 コア ) MEM:4GB OS:Cent OS 5.4
接続 :PCI Express 2.0 x16 Software:CUDA 3.2, GCC 4.1.2 31/38
10 20 30 40 50 60 70 80 90 100 0
10 20 30 40 50 60 70 80
90 GeForce GTX460
Tesla C1060
GeForce 8800 GTS 512
Matrix Size(N x N)
Speed up
IKA-SST の GPU タスク並列実装
性能評価： CPU 実装と比較
SIMD を使って最適化した CPU 実装 1 コアと比較
=10~100, 単精度 , データ転送時間を含む
CPU:AMD Phenom X4 9850(4 コア )
CPU 実装 (ATLAS) が
謎の速度低下
最大
76
倍CPU:AMD Phenom X4 9850(4 コア ) MEM:4GB OS:Cent OS 5.4
接続 :PCI Express 2.0 x16 Software:CUDA 3.2, GCC 4.1.2 32/38
IKA-SST の GPU タスク並列実装
System S 上に実装して評価
SST の実験と同じマシン構成
最大４台マシン (16GPU,16CPU コア ) で分散処理
CPU:AMD Phenom X4 9850(4 コア ) GPU:GeForce 8800 GTS 512
0 4 8 12 16
0 100,000 200,000 300,000 400,000 500,000 600,000 700,000 800,000
900,000 GPU
CPU
CPU(SIMD)
Number of GPUs/CPUs
Throughput(# of scores/sec)
最適化された CPU 実装と
約 12 倍
比較してに高速化ウィンドウサイズ：
250
ウィンドウサイズ：250
33/38
GPU タスク並列の実装と評価まとめ
SST と IKA-SST を GPU タスク並列で実装
データストリーム処理系 System S に両者を実装
16CPU+16GPU 環境における性能
スループット ( 処理データ数 / 秒 )
IKA-SST は 8 万センサーを 0.1 秒間隔で処理可能な性
能
CPU のみ CPU+GPU Speed up
SST 40 305 7.6 倍
IKA-SST 7 万 89 万 12.3 倍
34/38
目次
研究背景
既存手法とその問題点
GPU タスク並列
実装と評価
関連研究
まとめと今後の展望
35/38
関連研究
SVD の CUDA への移植
Sheetal Lahabar, P. J. Narayanan. Singular value decomposition
on GPU using CUDA. IPDPS. 2009.
小さい行列には最適化されていない
CUDA におけるタスク並列
Marisabel Guevara, et al. Enabling Task Parallelism in the CUDA Schedule
r. PMEA, pp.69-76, September 2009.
コンカレントカーネルと同様の機能を実現可能
CUDA を使ったリアルタイム処理
Sangjin Han, et al. PacketShader: a GPU-Accelerated Software R
outer, SIGCOMM 2010, New Delhi, India, Aug 30 - Sep 3, 2010.
GPU を使ったソフトウェアルータのフレームワーク
36/38
目次
研究背景
既存手法とその問題点
GPU タスク並列
実装と評価
関連研究
まとめと今後の展望
37/38
まとめと今後の展望
まとめ
変化点検知アルゴリズム SST と IKA-SST を GPU に実装
GPU タスク並列を使うことで高速化に成功
データストリーム処理系 System S に両者を実装
各実装を評価
今後の展望
GPU タスク並列化における実装コスト軽減方法の検討
38/38
39/38
カーネル関数内のタスク並列数
（コンカレントカーネルとの比較）
40/38
横軸 : 1= 完全なコンカレントカーネル
ウィンドウサイズは 30-300 から乱数で 100 個選択
1.0 2.0 5.0 10.0 20.0 50.0 100.0200.0500.0 0.0 0.1
0.2 0.3 0.4 0.5 0.6 0.7 0.8 0.9 1.0
GeForce GTX460 Tesla C1060 GeForce 8800 GTX
カーネル関数内のタスク並列数
相対速度
GPU タスク並列の実現方法
GPU 部分の並列化方法
依存関係のある複数のカーネルを連続で呼び出す
タスクごとに計算量やループ回数が異なる場合
GPUGPU
GPUGPU GPUGPU
GPUGPU
GPUGPU GPUGPU GPUGPU Task1 Task2 Task3
理想的な並列実行
Task1 Task2 Task3
GPU タスク並列による実行
（ A,B 両者とも） GPUGPU
GPUGPU GPUGPU
GPUGPU GPUGPU GPUGPU
GPUGPU
効率が悪い 効率が悪い
41/38
GPU タスク並列の実現方法
GPU 部分の並列化方法
2 つ方法を比較
性能や実装コストはどちらもあまり変わらない
本研究では異なるカーネルの同時実行の必要がない
ので利用できる環境の制約が少ない A を選択
方法 (A) ブロックごとに タスクを割り当て
る
(B) コンカレント カーネルを利用す
る 最大同時実行数 Tesla C2050: 112
Tesla C1060: 240 FermiFermiより前 : 1以降 : 16 異なるカーネルの
同時実行 不可 可
42/38
SST の GPU タスク並列実装
SVD の性能（実行時間内訳）
43/38
0% 10% 20% 30% 40% 50% 60% 70% 80% 90% 100%
data transfers diagonalization (GPU) diagonalization (CPU) bidiag. (GPU)
Matrix Size (N x N)
Rato of executon tme
GPU とは
主にリアルタイム 3D 画像生成に使われる
GPGPU
３Ｄグラフィックス以外の計算目的に GPU を利用すること
(Ubisoft’s AVARTAR, from http://ubi.com)
44/38
GPU と CPU の比較
GPU: コア数が多い
それぞれのコアは CPU と比較して遅い
トータルの性能は CPU より高い ( 演算性能 , メモリ速度）
数は少ないが高速なコアCPU コア数が多いGPU
45/38
