データストリーム処理における GPU タスク並列を
スケーラブルな異常検知機構の実現 用いた
上野 晃司（東京工業大学） , 鈴村 豊太郎（東京工業大学／ IBM 東京基礎究所）
目次
はじめに
変化点検知アルゴリズム SST の GPU 対応におけ
る
問題とその解決
SVD （特異値分解）の GPU タスク並列実装
性能評価
関連研究
まとめと今後の展望
データストリーム処理とは
データを蓄積せず、オンメモリ上でリアルタイムに
逐次処理を行う
従来は、データを蓄積してから計算するバッチ処理が主流
大規模なデータの処理やリアルタイムアプリケーションの実
現に有効
データストリーム処理系の例
IBM Research: System S
MIT: Borealis
意味のある必要なデー タ
時々刻々と流れてくる 大規模なデジタルデー タ
時々刻々と流れてくる 大規模なデジタルデー タ
データストリーム 管理システム
リアルタイムアプリ ケーションの実現 リアルタイムアプリ ケーションの実現 リアルタイムに
オンメモリ上で処理 リアルタイムに オンメモリ上で処理
異常・変化点検知
変化点検知とは
時系列データのデータ生成機構に、ある変化が生じた時点を
検知
応用例
トラフィックモニタリングによるネットワーク異常の検知
サーバのオンラインデータ観測によるサイバーテロ監視
株取引の自動処理
工場の生産ラインのセンサーによる監視
自動車のセンサーによる監視
ネットワークを使って自動車のセンサーデータを得る
SST （ Singular Spectrum Transformation, 特異スペクトル変
換）
変化点検知アルゴリズムの１つ
変化点検知アルゴリズム SST *1
*1 Tsuyoshi Ide, et al, Knowledge Discovery from Time-series Data using Nonlinear Transformations, The 4th Data Mining Workshop of JSSST 2004
SVD （特異値分解）を
計算し、時系列データ
の特徴を抽出
特徴ベクトルから変化
度スコア (z) を計算
時系列データから行列
（密行列）を作成
ウィンドウサイズが大きくなると SVD の計算量は O(n
3) で
増加 リアルタイム処理するには高速化が必要
モデルを使わない変化点検知手法
モデルを使う手法はモデルの想定していない場合に対応できな
い
より一般的な変化点検知に応用可能
GPGPU （ General Purpose GPU ）
GPU をグラフィック以外の計算目的に使用
GPU を使用するメリット
CPU の数倍の演算性能
高速なメモリ転送速度（ 100GB/s 超 cf. Core i7 970:25.6GB/s ）
GPGPU のデメリット
プログラミングの難易度が高い
高速化できないアプリケーションも多い
並列化が困難又は不可能な場合など
開発環境
NVIDIA CUDA 、 Open CL e.t.c.
本論文では最も一般的な NVIDIA の CUDA を使用する
NVIDIA GPU アーキテクチャ
数十個の Streaming Multiprocessor(SM) で構成される
各 SM に１つずつ SIMD 演算器がある
命令は warp （ 32 スレッド）ごとに発行される (SIMT)
１つの SM で複数の warp を処理可能
各 warp はそれぞれ命令ポインタを持っていて擬似的に並列実行
GPU 全体では数千スレッドを並列実行
スレッドブロックは warp の集まり
warp1 warp2 warp3
・・
・ Streaming Multiprocessor
SIMD 演算器
命令実行の準備ができた warpから処理
命令 GPU
SM
Global Memory SM
SM SM SM
SM SM SM
変化点検知アルゴリズム SST の GPU 対応
森田らによる研究
*1
SST 高速化に GPU を使用
CULA （ LAPACK の CUDA 実装）
を使って、 SVD （特異値分
解）を GPU で高速化
複数センサーに対応できない
監視対象が多い場合は？
数百個のセンサー
1 つ 2 秒だと、 100 個で 200 秒
小さい行列を高速化したい
行列サイズが大きいと GPU を
使っても遅すぎる
行列サイズ 400 以 下では、 CPU より
遅い GPU
CPU
*1 森田康介 , 鈴村豊太郎 . 「データストリーム処理を用いた変化点検知の実装と GPU による性能最適化」 . 電子情報通信学会 データ工学研究会 , 2010 年 6 月
行列サイズが大きくなると GPU の 速いが、時間がかかりすぎて方が
リアルタイム処理は不可能
性能評価グラフ
変化点検知アルゴリズム SST の GPU 対応における
問題
十分な並列性を得られない
密行列の SVD （特異値分解）
行列サイズが小さくて高速化できない
GPU の性能を引き出すには行列サイズ 1000 以上必要
SVD を並列化すれば解決
GPU スレッドの欠点
すべての GPU スレッドは１つの関数を実行 (SPMD)
複数の種類のプログラムを同時に実行できない
例外的に最新の Fermi コア GPU *1は 16 *2までの同時実行が可能
しかし、 SST では数十～数百レベルの並列化をするので、非効率
GPU スレッドは各スレッドが独立でない
SIMD で計算されるため
*1 GeForce GTX 480, Tesla C2050 などが該当
*2 GPU の SM 数により、同時実行可能数は異なる
GPU タスク並列で高速化
複数の時系列データ処理を並列化（＝タスク並列）
十分な並列性の確保が可能
スレッドブロックごとにタスクを割り当てる
各スレッドブロックは独立に動作
最新の Fermi コア GPU なら最大 128 タスクを同時実行可能
タスクスケジューリングは GPU ハードウェアで実行され
る
2
16(=65,536) 個までのタスク管理を CUDA フレームワーク
が処理
GPU タスク並列の実現方法
既存ライブラリはタスク並列非対応
カーネル関数を書き換える必要あり
CULA, CUBLAS などのライブラリはカーネル関数を隠蔽してし
まっているので、利用できない
タスク並列版 SVD を実装する必要がある
__global__ void kernel_function(float* A, int bda, … ){
/* ブロック ID を使ってデータを識別 */
float* bA = &A[bda*blockIdx.y];
/* 何かを実行する */
}
GPU タスク並列カーネル関数コード例
SVD アルゴリズム
密行列の SVD （特異値分解）
Sheetal Lahabar
*1の論文を参考にした
LAPACK で使われている Golub-Reinsch のアルゴリズム
前半： Householder 変換による二重対角化
後半： QR 反復による対角化
本発表では前半の Householder 変換よる二重対角
化のみを実装・評価
*1 Sheetal Lahabar, P. J. Narayanan. Singular value decomposition on GPU using CUDA. IEEE International Symposium on Parallel & Distributed Processing Symposium. 2009.
二重対角化タスク並列実装の詳細
計算をすべて GPU 化
Householder ベクトルを求める部分も含めてすべて GPU 化
CPU ・ GPU 間のデータ転送は遅い
ブロック化して実装
Householder 変換による二重対角化にはブロック化
*1が有効
ブロック化は実装コストが高いため、最初はブロック化なしで実装
しかし、ブロック化した方が速いことが判明
ブロックサイズは CUDA との相性から 32 とした
warp が 32 スレッドなので
*1 Jaeyoung Choi, Jack J. Dongrra, David W. Walker. The design of a parallel dense linear algebra software library: Reduction to Hessenberg, tridiagonal and bidiagonal form. Numerical
Algorithms, Volume 10, Number 2, 379-399. 1995.
性能評価
CPU AMD Phenom X4 9850 (4コア , 2.5GHz) メモリ 8 GB
GPU
Tesla C1060（ Geforce GTX 280 と同コ OS ア）CentOS 5.2
環境
評価方法
CPU
LAPACK(Version 3.2.1)+ATLAS(Version 3.8.3)
LAPACK の関数 sgebrd （密行列を二重対角化する関数）を使用
シングルスレッドで実行
GPU
データ転送時間も含めて計測
CPU ・ GPU ともに単精度
行列データは乱数で生成
0 64 128 192 256 320 384 448 512 576 0
2 4 6 8 10 12 14 16 18 20 22
ブロック化なし（タスク 数 64 ）
行列サイズ n
スピードアップ ( GPU / CPU 1コア )
性能評価
Partition Camping （メモ リアクセス競合）の影響
ブロック化によっ て２倍程度高速化 17.22
倍
行列サイズを変化させた時の高速化率
タスク数は 64, 256 で計測
CPU:AMD Phenom X4 9850 （ 1 core のみ使 用）GPU:Tesla C1060
性能評価
タスク数 32 以 上で十分高速
SM 数が 30 なので、タ スク数 16 以下では遅
い
タスク数を変化させた時の高速化率
行列サイズ 448 に固定
1 2 4 8 16 32 64 128 256 512
0 4 8 12 16 20
タスク数
スピードアップ ( GPU / CPU 1コア )
CPU:AMD Phenom X4 9850 （ 1 core のみ使 用）
関連研究
SST 高速化に関する研究
FELIX-SST
*1 SSTの発案者である井手氏による高速アルゴリズムの提案
ウィンドウサイズ 250 において、 130 倍の高速化に成功
GPU タスク並列なら一般的な行列計算を高速化可能
GPU タスク並列に関する研究
Marisabel Guevara らによる研究
*2 既存実装をコード修正なしで、タスク並列を実現
タスク数に比例して、条件分岐や引数の数が増えるという問題がある
GPU を使った行列計算の高速化に関する研究
ほとんどの研究は非常に大きな行列を対象としている
小さい行列の並列処理に関する研究はない
*1 T. Ide. Speeding up Change-Point Detection using Matrix Compression. IBIS, 2006.
*2 Marisabel Guevara, Chris Gregg, Kim Hazelwood, Kevin Skadron. Enabling Task Parallelism in the CUDA Scheduler. Proceedings of the Workshop on Programming Models for Emerging Architectures (PMEA). September 2009, pages 69-76.
まとめと今後の展望
まとめ
GPU タスク並列を提案した
SVD の前半の二重対角化を GPU タスク並列で実装
Householder 変換による二重対角化において、行列サ
イズ 320 ，タスク数 256 の場、ブロック化した実装で
17.22 倍の高速化
今後
SST 全体を GPU 実装する
データストリーム処理系に乗せる
ご清聴ありがとうございまし
た
Partition Camping
メモリアクセスの競合により、速度が低下する現象
発生メカニズム
GPU のグローバルメモリは８つの Partition に分かれている
*1
データは 64byte 幅で Partition に振り分けられる
ある Partition にアクセスが集中すると、その Partition がボトルネッ
クとなり速度が低下
行列サイズ 256,512 などの切りのいいサイズで発生しやすい
発生を抑制するには
行列の場合、アクセス順の対角化が有効
余計なインデックス計算が増える
二重対角化の実装において、行列転置を扱う場合などに発生
しやすいため、発生を抑える実装方法を取ったが、完全には
回避できなかった
*1 GeForce GTX 280, Tesla C1060 などの GPU の場合
