2025-02-09 OSS基盤の上に実装するGPUデータベース

2025-02-09（筑波大学 情報科学類 集中講義）ソフトウェアサイエンス特別講義A
「OSS基盤の上に実装するGPUデータベース」の講義で使用した資料です。

※履修登録をした学生さん以外がレポート課題を提出しても単位は出ません。

Transcript

1. ソフトウェアサイエンス特別講義A OSS基盤の上に実装するGPUデータベース ヘテロDB株式会社 チーフアーキテクト 兼 CEO 海外 浩平 <kaigai@heterodb.com>
2. 自己紹介 講師略歴 ✓ 海外 浩平（KaiGai Kohei） ✓ ヘテロＤＢ株式会社 チーフアーキテクト兼CEO ✓ 経歴  1997年 私立洛南高等学校卒業  2001年 筑波大学 第三学群 情報学類卒業  2003年 筑波大学大学院 経営・政策科学研究科修了  2003年～ 日本電気株式会社（NEC） （2011年～2013年 NEC Europe） 主にLinux kernelやPostgreSQLのコア機能開発に従事。 SELinuxに関する一連の機能強化や、PostgreSQLのエグゼキュータ 内部APIの整備などで、本家の開発者コミュニティにも知られる。 2012年頃より、GPUを用いたPostgreSQL高速化モジュールである PG-Stromの開発を始める。以降、GPUとは10年以上の付き合い。 ✓ 2007年 IPA未踏ソフトウェア創造事業において、 天才プログラマ―／スーパークリエータ認定  2017年～ HeteroDB社を設立し、PG-Stromの事業化に取り組む。 GPU-Direct SQLやApache Arrow対応を中核機能とするPG-Stromを リリースし、複数のユーザで利用中。最新版はv5.2。 ✓ 2017年 NVIDIA主催 GPU Technology Conference 2017にて AI関連スタートアップ19社の中から Inception Award を受賞 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 2
3. 本日、お話しする事 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 3 ▌OSS（オープンソースソフトウェア）の開発  Linux kernelやPostgreSQLの開発を通じて経験した、OSSの開発について。  OSSを基盤として独自のソリューションを開発する事 ▌SQLワークロードへのGPUの適用  GPUとは何ぞや？  SQLワークロードへのGPUの適用  GPU-Direct SQL機構 シラバスより： 本講では大きく２つのテーマについてお伝えします。 一つは、ソフトウェアが高機能化・大規模化する現代において、OSS（オープンソースソフトウェア）という 巨人の肩に乗る事で、我々はどのようにユニークで価値あるソリューションを市場に届けてゆく事ができる のか。また、そのメリットの一方で、どういった課題があるのか。 もう一つは、CPUとGPUという異なった特性を持ったプロセッサを用いて高性能化や低コスト化を図ろうとい うヘテロジニアスコンピューティングの話題で、GPUを用いてSQLの検索・集計ワークロードを高速化する PG-Strom（これはOSSの基盤の上に実装されたユニーク機能を提供するソフトウェアでもあります）がどのよ うにGPUを利用しているのか、いくつかの機能についてその実装を掘り下げます。 受講に際しては、基本的な情報工学の全般の知識に加えて、並列処理やメモリ管理、データベースに関して の知識があると望ましいですが、必須という訳ではありません。
4. Linux kernel開発者コミュニティとの出会い（1/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 4 ▌内定先／配属先  『スパコン向けOS開発』のポジションで内定  HPC事業部・エンタープライズLinux-Gに配属 ➔ 当時、ベクトル型スパコン（SXシリーズ）の ファイルサーバとして、64bit CPUである Itanium 2搭載サーバを製品化 ▌当時の世相  x86_64はまだ世に出ていなかった。  64bitのエンプラ向けは Itanium (IA64) という 事で、Intel, HP, NEC, Fujitsuなどが共同で Linux kernelを開発（2001～; Atlas Project） ➔ 自社のハードを売るために、自社サーバで 動作する Linux kernel の開発が必要。 ▌新人のぼくがやってた事  Linux kernelのソースコードをひたすら読む  クラッシュダンプを眺めて、障害原因を探る  その他、部門内のローカルサーバのお守り ベクトル型スパコン SX-6 Technology: 150 nm CPU Freq: 500 MHz CPU Perf: 8.0 GFlops Memory Band: 32.0GB/sec スカラ型サーバ TX-7 CPU: Itanium 2 (1.0GHz) (max 32way) RAM: ～128GB OS: Windows, HP-UX, NEC IA-64 Linux Red Hat Enterprise Linux 4 計算ノード ファイルサーバ ここの開発・保守を する部隊 2003
5. Linux kernel開発者コミュニティとの出会い（2/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 5 Linux 2.6がやって来る！！➔ SELinuxの調査を任されると…。 https://www.kernel.org/pub/linux/kernel/v2.6/ChangeLog-2.6.11 ▌Red Hat Enterprise Linux 4対応  Linux kernel v2.6の新機能 ✓ O(1)スケジューラ、NUMA対応 ✓ bio(Block I/O)、aio(Asynchronous I/O) ✓ SELinx, Ext4, …等々  SELinuxにトンデモない性能劣化を 引き起こす実装を発見。 ▌開発者コミュニティへ この時の実装では、SELinuxのポリシー チェックは排他ロック（spinlock）下で 行われていた。 ➔ これをRCUを使って置き換え、CPU数が 多い環境でのスケーラビリティを確保 する提案。 ➔ 何度か実装し直したものの、 Stephen Smalley (NSA)や、James Morris (Red Hat) らの協力を得てメインライン化。 2004
6. [補足] RCU (Read Copy Update) について ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 6 リストを更新する際にロックを使う場合 RCUによるリスト更新 LOCK LOCK
7. [補足] RCU (Read Copy Update) について ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 7 リストを更新する際にロックを使う場合 RCUによるリスト更新 LOCK LOCK
8. [補足] RCU (Read Copy Update) について ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 8 リストを更新する際にロックを使う場合 RCUによるリスト更新
9. [補足] RCU (Read Copy Update) について ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 9 リストを更新する際にロックを使う場合 RCUによるリスト更新 LOCK LOCK Writerのみ止める next
10. [補足] RCU (Read Copy Update) について ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 10 リストを更新する際にロックを使う場合 RCUによるリスト更新 LOCK LOCK Writerのみ止める next next Readは止めない
11. [補足] RCU (Read Copy Update) について ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 11 リストを更新する際にロックを使う場合 RCUによるリスト更新 LOCK LOCK Writerのみ止める next next ポインタの付け替えは アトミックに行われる
12. [補足] RCU (Read Copy Update) について ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 12 リストを更新する際にロックを使う場合 RCUによるリスト更新 next ➔ Read-mostlyなデータ参照に向く排他方式として、Linux kernel v2.6で対応した 新機能で、SELinuxのポリシーチェックにはドンピシャ 後で解放する
13. Linux kernel開発者コミュニティとの出会い（3/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 13 ▌開発者コミュニティとは  その技術領域に強い関心を持って、 日々、改良・発展に取り組む人の 集まり。  技術のコンセンサスを作る ➔ そこに居ないと、合意形成に加わる ことができない。 一定の貢献（contribution）を続け る事が重要。  開発だけではなく、パッケージの 保守や、ドキュメントの整備、 イベントの運営などによって コミュニティに参加し、存在感を 発揮している人もいる。 当時、50通近くメールのやり取りがあり、Stephen Smalley（SELinuxの作者）、 James Morris（SELinuxのメンテナー）、Paul.E McKenney（RCUの作者）らの 協力を得て数回のリテイクと、10回程度の微修正を経て、Linus Torvaldsの 管理するメインラインへカーネルへ移行。 2004
14. ビジネスとの兼ね合い（1/2） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 14 ▌OSSの開発者は何をしてお給料をもらっているのか？  研究職（先進的なソフトウェアを作って論文を書く）  そのOSSから収益を得ている会社 ✓ クラウド事業者、H/Wベンダー、SIer、．．．など  そのOSSを改良して収益・投資を得ている会社  非営利法人・フェローのような立場 ▌一介のﾋﾗ社員はどう考えたか？  OSS（この場合はSELinuxプロジェクト）に貢献する事が、 会社の（部門の）事業に貢献するという“ストーリー”が必要である。 ※なお、部門長は部門長で別の思惑があったと思われる。 （Linux推進センターとしてのOSS貢献活動のKPI）  取り組んでみた事 ✓ 組み込みSELinux • JFFS2へのXATTR対応、BusyboxへのSELinux ✓ SE-PostgreSQL & Secure LAPPスタック • SE-PostgreSQL、mod_selinux 2006
15. ビジネスとの兼ね合い（2/2） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 15 当時、Linux、MySQL/PostgreSQL、Apache、PHPで作るWebシステムが流行 当時使っていた資料 IPA未踏ソフトウェア創造事業 2006年度：SELinuxによるPostgreSQLのアクセス制御強化 外部の権威（経産省）が 「コイツは凄いで」と お墨付きを与える事で、 社内でのOSS活動が とても進めやすくなった。 2007
16. SE-PostgreSQLと”Innovative Portion”（1/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 16 SQLパーサ オプティマイザ エグゼキュータ クエリ木 （内部形式） 実行計画 SQL構文 問い合わせ結果 buffer manager transaction control metadata cache IPC & Lock index access transaction logs sepgsql モジュール SELinux アクセス 可否？ allowed / denied object_access_hook()
17. SE-PostgreSQLと”Innovative Portion”（2/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 17 ▌SE-PostgreSQLの標準機能化に付随して 拡張された本体側機能  Object Access Hook機構  ClientAuthentication hook（認証フック）  Security Barrier ViewとLeakproof関数  行レベルアクセス制御  Large Objectへの権限付与  SECURITY LABELコマンド ➔ sepgsqlモジュールだけではなく、 他の拡張モジュールの共通インフラと なるものを共通化。 contrib/sepgsql 以下のスケールは5000行程度 PostgreSQL全体のスケールは180万行程度 （全体の僅か0.27%程度の拡張機能）
18. [補足] Security Barrier ViewとLeakproof関数（1/2） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 18 postgres=# SELECT * FROM customer; cid | cname | cmail | cpasswd -----+-------+-------------------+---------- 101 | alice | alice@example.com | abcdef 102 | bob | bob@example.com | xyz123 103 | eve | eve@example.com | deadbeaf (3 rows) postgres=# CREATE VIEW my_account AS SELECT * FROM customer WHERE cname = getpgusername(); CREATE VIEW postgres=# GRANT SELECT ON my_account TO public; GRANT ～ 一般ユーザの権限でログイン ～ postgres=> CREATE FUNCTION f_leak(text) RETURNS bool LANGUAGE plpgsql COST 0.00000001 AS 'BEGIN RAISE NOTICE ''f_leak => %'', $1; RETURN true; END'; CREATE FUNCTION postgres=> SELECT * FROM my_account WHERE f_leak(cmail); NOTICE: f_leak => alice@example.com NOTICE: f_leak => bob@example.com NOTICE: f_leak => eve@example.com cid | cname | cmail | cpasswd -----+-------+-------------------+--------- 101 | alice | alice@example.com | abcdef (1 row) ユーザー名が一致している 行しか見えないぞ！ おい！何か見えてるぞ！ （ぴえん）
19. [補足] Security Barrier ViewとLeakproof関数（2/2） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 19 postgres=> EXPLAIN SELECT * FROM my_account WHERE f_leak(cmail); QUERY PLAN ----------------------------------------------------------------- Seq Scan on customer (cost=0.00..20.85 rows=1 width=100) Filter: (f_leak(cmail) AND (cname = (getpgusername())::text)) (2 rows) postgres=# CREATE VIEW my_account_secure WITH (security_barrier) AS SELECT * FROM customer WHERE cname = getpgusername(); CREATE VIEW postgres=# GRANT SELECT ON my_credit_secure TO public; GRANT postgres=> SELECT * FROM my_account_secure WHERE f_leak(cmail); NOTICE: f_leak => alice@example.com cid | cname | cmail | cpasswd -----+-------+-------------------+--------- 101 | alice | alice@example.com | abcdef (1 row) postgres=> EXPLAIN SELECT * FROM my_account_secure WHERE f_leak(cmail); QUERY PLAN ------------------------------------------------------------------------- Subquery Scan on my_account_secure (cost=0.00..20.88 rows=1 width=100) Filter: f_leak(my_account_secure.cmail) -> Seq Scan on customer (cost=0.00..20.85 rows=3 width=100) Filter: (cname = (getpgusername())::text) (4 rows) f_leak()を先に実行している！ （マジか！？）
20. SE-PostgreSQLと”Innovative Portion”（3/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 20 ▌共通のソフトウェア部品を共有する（Shared Cost）とする事で、開発者は Unique Innovationに注力する事ができ、全体のコストを下げる事ができる。 ▌例）『セキュアなPostgreSQL』のために180万行フルスクラッチの開発と SW保守は許容できるか？ ➔ 普通は無理 How to Contribute to the Linux Kernel, and Why it Makes Economic Sense (James Bottomley, Novell; LinuxCon/Japan 2009)
21. OSSのエコシステムに乗っかる ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 21 ▌既に世界中で利用され、実績のあるOSSを基盤とする事で、、、  “ソフトウェアの開発コスト”だけではなく  既存のユーザや経験のあるエンジニアがチャネルとなる。  周辺ソフトウェアをそのまま利用する事が出来る。  中核機能以外のSW保守の恩恵を受ける事ができる。 Backup Replication JDBC/ODBC Admin tools DB monitor Applications Skilled Engineers Users base 周辺ソフト ウェアの蓄積 人的資本の 蓄積 Maintenance Stability
22. OSSを活用する上でのデメリットは？ ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 22 ▌合意形成には時間がかかる  特に問題意識を共有する人が少ない場合は、合意形成や再設計で年単位の 時間がかかる事も。  それを許容できる事業環境やスケジュールであるかどうか。  非互換の独自設計のままユーザ先に出荷したら、並行メンテナンスの悪夢…。 PostgreSQLコミュニティへ はじめてSE-PostgreSQLの アイデアを投げたのは 2007年4月 メインライン機能として マージされたのは 2011年1月
23. この頃、何をしていたかと言いますと・・・。 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 23 ドイツ駐在（Feb-2011～Dec-2013） 2011 SAP-AG社とのアライアンス担当として、 Heidelberg近郊のWalldorf市に赴任 （20%はOSS活動を継続するという条件）
24. PostgreSQLにGPUを試そうとしたきっかけ ▌PGconf 2011 (Ottawa) で聴講したあるセッション・・・  Parallel Image Searching Using PostgreSQL and PgOpenCL Running PostgreSQL Stored Procedures on a GPU ✓ https://www.pgcon.org/2011/schedule/events/352.en.html  画像処理用のストアドプロシジャを、GPU用のプログラミング環境 OpenCL で作成するための拡張モジュールについて。 A列 B列 C列 D列 E列 幅の小さいデータでも、 大量の行を並べれば 長大なBLOBと同じく GPU並列処理が効くのでは？ 2011 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 24
25. GPU(Graphics Processing Unit)とはどんなプロセッサなのか？ ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 25 元々は3Dゲームでポリゴンの座標を高速に計算するためのデバイス ➔ 汎用計算にも安く使えるという事で、HPCや機械学習の分野でも スーパーコンピュータ （東京工業大学 TSUBAME3.0） ＣＧ（Computer Graphics） 機械学習 数百～数千コアの並列処理ユニットと、TB/sを越える帯域のメモリを搭載。 大量の計算処理を得意とするが、専用のS/Wやアルゴリズムが必要。 CUDA Toolkitによって容易にSW開発が可能となった。 シミュレーション NVIDIA H100 RTX 4090Ti 3D Gaming
26. CPUとGPUの設計思想の違い（1/4） number of cores 128 core clocks 2.0GHz (boost: 3.9GHz) L1 cache (per core) data: 48kB / code: 64kB L2 cache (per socket) 256MB L3 cache (per socket) 504MB DRAM DDR5 6400 x 12 (614.4GB/s) 潤沢なキャッシュと高いコアクロックの他、 分岐予測なども含めてシングルスレッド性能 を追求する number of cores 7,296 (INT32) 14,592 (FP32) 7,296 (FP64) core clocks 1065MHz (boost: 1620MHz) L1 cache (per SM) 256 kB (*) combined use for shared memory L2 cache (per device) 50MB DRAM HBM2e 5120bit 80GB (2.0TB/s) 非常に多数の計算コアと、DRAMの高いスルー プットによって、大量のデータを計算する事に 特化した構造の計算機。 Intel Xeon 6980P (Granite Rapids; 2024-Q2) NVIDIA H100 (Hopper; 2022-Q3) ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 26
27. CPUとGPUの設計思想の違い（2/4） LZ77圧縮アルゴリズムの例 (F,8,’capybara’) (T,4,4) ’c’ ’a’ ’p’ ’y’ ’b’ ’a’ ’r’ ’a’ (T,4,3) (T,14,3) ’c’ ’a’ ’p’ ’y’ ’b’ ’a’ ’r’ ’a’ ’b’ ’a’ ’r’ ’a’ ’c’ ’a’ ’p’ ’y’ ’b’ ’a’ ’r’ ’a’ ’b’ ’a’ ’r’ ’a’ ’b’ ’a’ ’r’ ’c’ ’a’ ’p’ ’y’ ’b’ ’a’ ’r’ ’a’ ’b’ ’a’ ’r’ ’a’ ’b’ ’a’ ’r’ ’a’ ’p’ ’y’ 4文字戻って4文字を追記 4文字戻って3文字を追記 14文字戻って3文字を追記 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 27
28. CPUとGPUの設計思想の違い（3/4） 行列積の計算の例 x0 x1 x2 x3 x4 x5 x6 x7 xk xn × × × × × × × × × × y0 y1 y2 y3 y4 y5 y6 y7 yk yn ＋ ＋ ＋ ＋ ＋ ＋ ＋ ＋ 計算の依存性無し 計算の依存性は最小限 隣の要素の計算が 終わっていれば、 加算を実行可能 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 28
29. CPUとGPUの設計思想の違い（4/4） GPU DRAM DRAM Cache Cache ✓ データアクセスの局所性 ✓ キャッシュ＆高クロック ➔ 処理のレイテンシに優位性 ✓ 次から次へと、データを演算器に 流し込む事に特化した構造。 ✓ 広帯域メモリ＆大容量レジスタ ➔ 処理のスループットに優位性 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 29 小回りが利くが輸送力は 小さな乗用車のような プロセッサ 使える状況が限られるが、 大量輸送が可能な 高速鉄道のような プロセッサ
30. コレって検索ワークロードに似ていない？ テーブルスキャンの方向 x_val = 123 x_val = 234 x_val = 345 x_val = 456 x_val = 567 WHERE x_val % 2 = 1 〇 × 〇 × 〇 • • • • • GPUコア ✓ DBテーブルをある種のベクトルと見なせば、 一度の処理サイクルで数千～万行の処理が可能に 行列計算の場合 テーブルスキャンの場合 𝑣 × 𝐴 v0 v1 v2 v3 vn-4 vn-3 vn-2 vn-1 ak,0 ak,1 ak,2 ak,3 ak,n-4 ak,n-3 ak,n-2 ak,n-1 × × × × × × × × ベクトル v 行列A 異なるデータに 同じ演算を多数実行 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 30
31. 実際に作ってみた（1/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 31 実行イメージ：PostgreSQLは行単位、PG-Stromはバルクかつ非同期並列 実行イメージ：FDWのAPIを通じてオプティマイザ、エグゼキュータに介入 2012
32. 実際に作ってみた（2/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 32 void pgstrom_get_foreign_paths(PlannerInfo *root, RelOptInfo *baserel, Oid foreigntableid) { : /* check whether GPU/CPU executable qualifier, or not */ foreach (cell, baserel->baserestrictinfo) { RestrictInfo *rinfo = lfirst(cell); if (is_gpu_executable_qual(baserel, rinfo)) gpu_quals = lappend(gpu_quals, rinfo); : } /* Generate command series executed with GPU/CPU, if any */ if (gpu_quals) { cmds_bytea = make_gpu_commands(gpu_quals, &gpu_cols); defel = makeDefElem("gpu_cmds", (Node *) cmds_bytea); private = lappend(private, defel); } : f_path = create_foreignscan_path(root, baserel, ...); add_path(baserel, (Path *) f_path); } 2012 WHERE句がGPUで実行可能か どうかをチェックする。 GPU用のWHERE句処理関数を 自動生成する。 （当時は素のCUDA C++コードを 生成し、実行時ビルドしていた） 実行コストを計算して、 PostgreSQLのオプティマイザに GPUを使うパスを登録する。
33. 実際に作ってみた（3/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 33 ▌FDW: Foreign Data Wrapper  v8.4からある機能で、外部のデータソースをテーブルのように読み出す機能  『エグゼキュータを乗っ取る』ために利用できた PostgreSQLの実行計画作成、クエリ実行の流れ 2012 SQLパーサ オプティマイザ エグゼキュータ クエリ木 （内部形式） 実行計画 SQL構文 問い合わせ結果 buffer manager transaction control metadata cache IPC & Lock index access transaction logs GetForeignPaths() GetForeignPlan() FDW モジュール • 実行計画を作成 • 実行開始 • 次の１行を返す • 実行終了 BeginForeignScan() IterateForeignScan() EndForeignScan() External RDBMS CSV Files
34. 当時の課題 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 34 ▌FDWは読み込み専用だった  データのロードには専用関数を用いていた。  INSERT/UPDATE/DELETEが使用できない ▌FDWだとDDL構文が異なる  外部テーブルを作成するには CREATE FOREIGN TABLE name ( … ) SERVER … OPTIONS (…); と、やたらおまじないが多い。  SQL構文が互換でないと、アプリケーションが利用できない場合も。 ➔ 『利用できない』にはクエリの修正や再検証が必要も含む。 ▌GPU-Scanだけでは応用範囲が狭い  JOINやGroup-BYなどを高速化するにはどうするか？
35. Writable FDW ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 35 ▌Writable FDW機能  PostgreSQL v13 にて標準機能化  PostgreSQLのINSERT/UPDATE/DELETE処理に なぞらえる形で、更新処理用のAPIを追加 ▌概要  リモート（外部データソース）から行を 読み出す際に、FDWドライバは、ユニークな 識別子を含める必要がある。  UPDATE/DELETEの際は、その識別子を使って 更新/削除処理の対象行を識別する。  INSERT/UPDATEの際は、PostgreSQLが生成した 「行」をリモートの形式に合わせて成型する。 ▌課題  「DDL構文が非互換」という問題については 残ったまま 2013
36. Custom-Scan APIs（1/5） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 36 PostgreSQLのテーブルを Scan/Join する代替の方法を定義する 2014 SQLパーサ オプティマイザ エグゼキュータ クエリ木 （内部形式） 実行計画 SQL構文 問い合わせ結果 buffer manager transaction control metadata cache IPC & Lock index access transaction logs 拡張モジュール • 実行計画を作成・登録 • 実行開始 • 次の１行を返す • 実行終了 set_rel_pathlist_hook set_join_pathlist_hook create_upper_paths_hook それぞれPostgreSQLのテーブルを Scan/Join/Group-byするタイミングで 呼び出されるフックを追加 ExecCustomScan() どう実装するかは 拡張モジュールが 任意に決める事が できる。
37. Custom-Scan APIs（2/5） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 37 ssbm=# explain select sum(lo_revenue), d_year, p_brand1 from lineorder, date1, part, supplier where lo_orderdate = d_datekey and lo_partkey = p_partkey and lo_suppkey = s_suppkey and p_category = 'MFGR#12’ and s_region = 'AMERICA’ group by d_year, p_brand1; QUERY PLAN --------------------------------------------------------------------------------------------------- HashAggregate (cost=192049053.77..192049141.27 rows=7000 width=46) Group Key: date1.d_year, part.p_brand1 -> Hash Join (cost=373962.46..191695451.95 rows=47146910 width=20) Hash Cond: (lineorder.lo_orderdate = date1.d_datekey) -> Hash Join (cost=373857.95..191571381.58 rows=47146910 width=20) Hash Cond: (lineorder.lo_suppkey = supplier.s_suppkey) -> Hash Join (cost=55265.84..190631708.88 rows=236602030 width=26) Hash Cond: (lineorder.lo_partkey = part.p_partkey) -> Seq Scan on lineorder (cost=0.00..174826341.12 rows=6000026112 width=20) -> Hash (cost=54280.00..54280.00 rows=78867 width=14) -> Seq Scan on part (cost=0.00..54280.00 rows=78867 width=14) Filter: (p_category = 'MFGR#12'::bpchar) -> Hash (cost=293684.47..293684.47 rows=1992611 width=6) -> Seq Scan on supplier (cost=0.00..293684.47 rows=1992611 width=6) Filter: (s_region = 'AMERICA'::bpchar) -> Hash (cost=72.56..72.56 rows=2556 width=8) -> Seq Scan on date1 (cost=0.00..72.56 rows=2556 width=8) (17 rows) 2014
38. Custom-Scan APIs（3/5） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 38 ssbm=# explain select sum(lo_revenue), d_year, p_brand1 from lineorder, date1, part, supplier where lo_orderdate = d_datekey and lo_partkey = p_partkey and lo_suppkey = s_suppkey and p_category = 'MFGR#12’ and s_region = 'AMERICA’ group by d_year, p_brand1; QUERY PLAN ------------------------------------------------------------------------------------------------------------ HashAggregate (cost=30939287.03..30939374.53 rows=7000 width=46) Group Key: date1.d_year, part.p_brand1 -> Custom Scan (GpuPreAgg) on lineorder (cost=30939164.53..30939234.53 rows=7000 width=46) GPU Projection: pgstrom.psum(lineorder.lo_revenue), date1.d_year, part.p_brand1 GPU Join Quals [1]: (part.p_partkey = lineorder.lo_partkey) ... [nrows: 6000026000 -> 236602000] GPU Outer Hash [1]: lineorder.lo_partkey GPU Inner Hash [1]: part.p_partkey GPU Join Quals [2]: (supplier.s_suppkey = lineorder.lo_suppkey) ... [nrows: 236602000 -> 47146910] GPU Outer Hash [2]: lineorder.lo_suppkey GPU Inner Hash [2]: supplier.s_suppkey GPU Join Quals [3]: (date1.d_datekey = lineorder.lo_orderdate) ... [nrows: 47146910 -> 47146910] GPU Outer Hash [3]: lineorder.lo_orderdate GPU Inner Hash [3]: date1.d_datekey GPU Group Key: date1.d_year, part.p_brand1 GPU-Direct SQL: enabled (N=2,GPU0,1) -> Custom Scan (GpuScan) on part (cost=100.00..30481.17 rows=78867 width=14) GPU Projection: p_brand1, p_partkey GPU Scan Quals: (p_category = 'MFGR#12'::bpchar) [rows: 2000000 -> 78867] GPU-Direct SQL: enabled (N=2,GPU0,1) -> Custom Scan (GpuScan) on supplier (cost=100.00..190276.56 rows=1992611 width=6) GPU Projection: s_suppkey GPU Scan Quals: (s_region = 'AMERICA'::bpchar) [rows: 9999718 -> 1992611] GPU-Direct SQL: enabled (N=2,GPU0,1) -> Seq Scan on date1 (cost=0.00..72.56 rows=2556 width=8) (24 rows) 2014
39. Custom-Scan APIs（4/5） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 39 ▌FDWとCustomScanの違い  FDWはデータソース（What）を多様化する。  CustomScanはPostgreSQLテーブルを読み出す 方法（How）を多様化する。 ▌開発の副産物  JOINのパスを追加するためのフック  GROUP-BY等のパスを追加するためのフック ➔postgres_fdwのremote join pushdownなど  Background worker process ➔ CUDA C++のコードを実行時コンパイルするために 提案した機能 2014
40. Custom-Scan APIs（5/5） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 40 現在では、PG-Strom以外にもかなりのDBエンジンで利用されている Noriyuki Shinoda『列指向アクセスメソッド徹底比較』, PostgreSQL Conference Japan 2004 https://speakerdeck.com/nori_shinoda/postgresql-conference-japan-2024-a4-comparison-of-column-oriented-access-methods 2024
41. OLAPの性能評価に使われるTPC-DSの実行時間を分析 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 41  Scan(34.4%)、Join(36.9%)、Aggregate(23.3%)で総処理時間の95% ✓ 但し、Scanは予めバッファにロードした状態である事に留意。 ✓ この３つを倒せば、OLAP系のワークロードを大半をカバーできることに？  その他 (5.3%) の中で目立つのは、Sort、SetOp、Window関数 ✓ 但し、Sort/Window関数は最近（2024）依頼があり、対応に向けて作業中 2015
42. GPUとスレッド（1/4） 命令 SISD (Single Instruction Single Data) SIMD (Single Instruction Multiple Data) MISD (Multiple Instruction Single Data) MIMD (Multiple Instruction Multiple Data) コア データ 命令 コア データ コア データ コア データ 命令 コア コア データ コア 命令 命令 命令 コア コア コア 命令 命令 データ データ データ ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 42
43. GPUとスレッド（2/4） 命令 コア データ コア データ コア データ if (false) if (true) if (true) SIMT (Single Instruction, Multiple Threads) ✓ ほぼSIMDだが、条件分岐を含む事ができる。 ✓ データ形式の制約が少ない。 ※ ベクトル命令の場合、32bit x 16なら512bit境界への アライメントを要求するのが普通。 ✓ 32スレッドを束ねたWarpと呼ばれる単位で 命令を実行する。 NVIDIA H100 ブロックダイアグラム “Warp Scheduler”が、実行準備の できたWarp（32スレッド）を 次々と演算ユニットに投入する。 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 43
44. GPUとスレッド（3/4） core0 core1 core2 core3 INT Load FP32 INT Store 例：𝑣𝑖 = 𝑥𝑖 − 𝑦𝑖 𝑖 = 0 … 7 を計算するケース（模式図） INT Load FP32 INT Store INT Load FP32 INT Store INT Load FP32 INT Store INT Load FP32 INT Store INT Load FP32 INT Store INT Load FP32 INT Store INT Load FP32 INT Store core0 INT Load FP32 INT Store INT Load FP32 INT Store core1 INT Load FP32 INT Store INT Load FP32 INT Store core2 INT Load FP32 INT Store INT Load FP32 INT Store core3 INT Load FP32 INT Store INT Load FP32 INT Store ８要素の処理を４コア/４スレッドで実行 ８要素の処理を４コア/８スレッドで実行 Warp Schedulerが、 開いたリソースに 実行待ちのスレッドを 投入する。 ➔ 余剰リソースが減り、 実行時間を短くできる。 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 44
45. GPUとスレッド（4/4） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 45 NVIDIA H100 [PCI-E] (114SMs) SMあたり INT32: 64 FP32: 128 FP64: 64 LD/ST: 32 32bitレジスタ： 65536 L1 Cache/共有メモリ: 256kB ✓ 同一Warp内のスレッドは同時に異なる命令を実行する 事はできない。 例）あるスレッドがINT32コアに割り当てられてるのに、 別のスレッドがFP32命令を実行するのは不可 ✓ スレッド数が多いほど計算リソースを飽和させやすい。 ✓ スレッドあたりのレジスタ数や共有メモリが多いほど 性能は出しやすいので、その塩梅を調整する必要がある。 int index = threadIdx.x; z[index] = (float)index; if (index % 2 == 0) { z[index] += (x[index] + y[index]) } else { z[index] += (x[index] - y[index]) } Threads in Warp 32スレッド = 1Warp を単位としたスケジューリング
46. 多段GPU-Joinにおける工夫（1/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 46 ✓ 一度に512行を処理したとしても、GpuJoin(depth=2)や、結果セットを生成する GpuProjectionでは同時に処理できる行数が減ってしまう。 ➔ では、どのように処理したのか？ lineorder part supplier Join Join part supplier GpuJoin (depth=1) GpuJoin (depth=2) GpuProjection Projection GPU並列処理が 効いている 2/3のコアは 遊んでいる 大半のコアが 遊んでいる
47. 多段GPU-Joinにおける工夫（2/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 47 for (i=0; (outer_tuple = fetch_tuple_from_relation(outer_rel, i)) != NULL; i++) { // depth=1のJOIN hash = calculate_hash(outer_tuple); for (inner_tuple1 = fetch_tuple_from_hash(inner_rel_1, hash); inner_tuple1 != NULL; inner_tuple1 = fetch_next_tuple(innet_tuple1)) { // Hash値、またはJOIN-Keyが不一致ならスキップ if (hash != inner_tuple1->hash || hash_join_keycomp(outer_tuple, inner_tuple1) != 0) continue; // depth=2のJOIN hash = calculate_hash(outer_tuple, inner_tuple1); for (inner_tuple2 = fetch_tuple_from_hash(inner_rel_2, hash); inner_tuple2 != NULL; inner_tuple2 = fetch_next_tuple(innet_tuple2)) { // Hash値、またはJOIN-Keyが不一致ならスキップ if (hash != inner_tuple2->hash || hash_join_keycomp(outer_tuple, inner_tuple1, inner_tuple2) != 0) continue; // 結果セットを出力 write_out_projection(outer_tuple, inner_tuple1, inner_tuple2); } } } “普通に” JOIN処理を書くとなるとどうなるか？
48. 多段GPU-Joinにおける工夫（3/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 48 part supplier GpuProjection • • • • • • • • • • GpuJoin (depth=2) GpuJoin (depth=1) depth=1中間バッファ （環状バッファ） depth=2中間バッファ （環状バッファ） 全てのコアがActiveに 入力行を処理
49. 多段GPU-Joinにおける工夫（3/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 49 part supplier GpuProjection • • • • • • • • • • • • • • • • • • GpuJoin (depth=2) depth=1中間バッファ （環状バッファ） depth=2中間バッファ （環状バッファ） GpuJoin (depth=1) 全てのコアがActiveに 入力行を処理
50. 多段GPU-Joinにおける工夫（3/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 50 part supplier GpuProjection • • • • • • • • • • • • • • • • • • • • • • • • • depth=1中間バッファ （環状バッファ） depth=2中間バッファ （環状バッファ） GpuJoin (depth=1) 全てのコアがActiveに 入力行を処理 GpuJoin (depth=2)
51. 多段GPU-Joinにおける工夫（3/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 51 part supplier GpuProjection • • • • • • • • • • • • • • • • • • • • • • • • • • • depth=1中間バッファ （環状バッファ） depth=2中間バッファ （環状バッファ） GpuJoin (depth=2) 再び depth=1 に戻り、中間バッファを 埋める。 GpuJoin (depth=1)
52. GPUスケジューリングが顕著に性能影響した例（1/2） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 52 (123.0, 20.0) (154.2, 46.2) SELECT pref,city,count(*) FROM geo_japan j, geopoint p WHERE st_contains(j.geom, st_makepoint(x,y)) AND j.pref like '東京都’ GROUP BY pref,city; ランダムに 生成した座標 1000万個 pref | city | count ---------+------------+------- 東京都 | あきる野市 | 105 東京都 | 三宅村 | 76 東京都 | 三鷹市 | 17 東京都 | 世田谷区 | 67 東京都 | 中央区 | 12 東京都 | 中野区 | 18 ： ： ： 東京都 | 豊島区 | 14 東京都 | 足立区 | 55 東京都 | 青ヶ島村 | 7 東京都 | 青梅市 | 117 (63 rows) CPU版：24.892s GPU版：33.841s（遅い！） 国土地理院からDLした 全国市町村形状データ
53. GiSTインデックス（R木）の仕組み ▌GiSTインデックス（R木）の仕組み ✓ R1は(R3,R4,R5)を全て包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と、下位ノードへのポインタ ✓ R4は(R11,R12)を全て包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と、下位ノードへのポインタ ✓ R12は対象ジオメトリを包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と ItemPointer を含む。 ✓ 各ノード（BLCKSZ）内のエントリを順に評価。マッチしたものを次の階層へすすめる。 ✓ 階層ごとに繰り返し評価が発生するので、B-tree並みに速くはできない。 (xmin,ymin) (xmax,ymax) ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 53 検索キー（経度, 緯度） 〇 〇
54. GiSTインデックス（R木）の仕組み ▌GiSTインデックス（R木）の仕組み ✓ R1は(R3,R4,R5)を全て包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と、下位ノードへのポインタ ✓ R4は(R11,R12)を全て包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と、下位ノードへのポインタ ✓ R12は対象ジオメトリを包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と ItemPointer を含む。 ✓ 各ノード（BLCKSZ）内のエントリを順に評価。マッチしたものを次の階層へすすめる。 ✓ 階層ごとに繰り返し評価が発生するので、B-tree並みに速くはできない。 (xmin,ymin) (xmax,ymax) ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 54 検索キー（経度, 緯度） 〇 〇 × 〇 ×
55. GiSTインデックス（R木）の仕組み ▌GiSTインデックス（R木）の仕組み ✓ R1は(R3,R4,R5)を全て包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と、下位ノードへのポインタ ✓ R4は(R11,R12)を全て包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と、下位ノードへのポインタ ✓ R12は対象ジオメトリを包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と ItemPointer を含む。 ✓ 各ノード（BLCKSZ）内のエントリを順に評価。マッチしたものを次の階層へすすめる。 ✓ 階層ごとに繰り返し評価が発生するので、B-tree並みに速くはできない。 (xmin,ymin) (xmax,ymax) ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 55 検索キー（経度, 緯度） 〇 〇 × 〇 × 〇 ×
56. GiSTインデックス（R木）の仕組み ▌GiSTインデックス（R木）の仕組み ✓ R1は(R3,R4,R5)を全て包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と、下位ノードへのポインタ ✓ R4は(R11,R12)を全て包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と、下位ノードへのポインタ ✓ R12は対象ジオメトリを包含する矩形領域の(Xmin ,Ymin ) – (Xmax ,Ymax )と ItemPointer を含む。 ✓ 各ノード（BLCKSZ）内のエントリを順に評価。マッチしたものを次の階層へすすめる。 ✓ 階層ごとに繰り返し評価が発生するので、B-tree並みに速くはできない。 (xmin,ymin) (xmax,ymax) ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 56 検索キー（経度, 緯度） 〇 〇 × 〇 × 〇 × × ×
57. GPU版GiSTインデックス ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 57 ▌仕組み  多角形（エリア定義情報）を保持するテーブルと、位置情報（緯度経度）を 保持するテーブルとの JOIN にGiSTインデックスを使用できる。  先ずGiSTインデックス上のBounding-Box（矩形領域）によって荒く絞り込み、 次にテーブル上の多角形（ポリゴン）と「当たり判定」を行う。  GPUの数千コアをフル稼働してGiSTインデックスを探索する。 単純に並列度が高い分、検索速度も速くなるはず。 ➔ が、そうは問屋が卸さなかった。。。 多角形 × 点の重なり判定などを、GpuJoinの一要素として実装 GiST(R木)インデックス ポリゴン定義 位置データを含む テーブル 数千スレッドが 並列に インデックスを 探索
58. GPUスケジューリングが顕著に性能影響した例（2/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 58 スレッド間の処理時間の差が大きく、同期待ちを招く Table-A • • • • • • 参照 Index-A × × × × × 〇 N = __syncthreads_count(...) 参照 R木のLeafノード まで探索したが ヒットせず R木のRootノードを 見ただけで、即、 マッチする要素なし R木を最後まで 探索し、かつ、 JOINの結合条件を評価 スレッドグループに属する 他のスレッドの完了待ちで GPUコアが遊んでしまう。
59. GPUスケジューリングが顕著に性能影響した例（3/3） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 59 (123.0, 20.0) (154.2, 46.2) ランダムに 生成した座標 1000万個 pref | city | count ---------+------------+------- 東京都 | あきる野市 | 105 東京都 | 三宅村 | 76 東京都 | 三鷹市 | 17 東京都 | 世田谷区 | 67 東京都 | 中央区 | 12 東京都 | 中野区 | 18 ： ： ： 東京都 | 豊島区 | 14 東京都 | 足立区 | 55 東京都 | 青ヶ島村 | 7 東京都 | 青梅市 | 117 (63 rows) CPU版：24.892s GPU版： 1.154s 21.5倍の 高速化 GPUコアのスケジューリングを 意識した実装に修正し再測定 国土地理院からDLした 全国市町村形状データ SELECT pref,city,count(*) FROM geo_japan j, geopoint p WHERE st_contains(j.geom, st_makepoint(x,y)) AND j.pref like '東京都’ GROUP BY pref,city;
60. GPU-Direct SQLの着想と開発（1/3） GPUコア GPU Device Memory CPU Host Memory ストレージ （HDD/SSD） CPU 720GB/s 16GB/s 0.5GB/s 60GB/s NVME-SSD 一台あたり 3GB/s （当時の）PG-Stromを含む、 GPU-DBはデータがオンメモリ前提 ストレージに落ちたら「負け」 一方この頃、NVME-SSD製品が登場 安価な高速SSDがブレイクの兆し。 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 60 2015
61. GPU-Direct SQLの着想と開発（1/4） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 61 DMAコントローラは物理アドレスに対してデータを投げる 物理アドレス空間 論理アドレス空間 process-X process-Y 0x000000000000 PCIデバイス 0x037000000000 PCI-E Bar1 Linux kernel 0x038000000000 NVME-SSD上のブロック xxx 番から 20ブロック分を読み出して、 物理アドレス0x00... へと転送せよ NVME-SSD上のブロック xxx 番から 20ブロック分を読み出して、 物理アドレス0x37... へと転送せよ NVME READ BlkNo: xxxx Length: 20 Dest: 0x00.... NVME READ BlkNo: xxxx Length: xxx Dest: 0x037.... 2015
62. GPU-Direct SQLの着想と開発（2/4） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 62 static int ioctl_map_gpu_memory(StromCmd__MapGpuMemory __user *uarg) { mapped_gpu_memory *mgmem; : if (copy_from_user(&karg, uarg, sizeof(karg))) return -EFAULT; mgmem = kmalloc(sizeof(mapped_gpu_memory), GFP_KERNEL); if (!mgmem) return -ENOMEM; map_address = karg.vaddress & GPU_BOUND_MASK; map_offset = karg.vaddress & GPU_BOUND_OFFSET; handle = (unsigned long) mgmem; INIT_LIST_HEAD(&mgmem->chain); mgmem->hindex = strom_mapped_gpu_memory_index(handle); mgmem->refcnt = 0; mgmem->owner = current_euid(); mgmem->handle = handle; mgmem->map_address = map_address; mgmem->map_offset = map_offset; mgmem->map_length = map_offset + karg.length; mgmem->wait_task = NULL; : rc = __nvidia_p2p_get_pages(0, /* p2p_token; deprecated */ 0, /* va_space_token; deprecated */ mgmem->map_address, mgmem->map_length, &mgmem->page_table, callback_release_mapped_gpu_memory, mgmem); : } nvidiaドライバのkernel APIを用いて、 GPUデバイスメモリの論理⇒物理変換 それを利用して、NVME READリクエス トを偽装してドライブに送出する。 2015
63. GPU-Direct SQLの着想と開発（3/4） ✓ GPU-Direct SQL（当時は SSD-to-GPU Direct SQL）により、メモリサイズを越えた 大量データ（数TB～）を処理する事が可能に。 ➔ 他のGPU-DB製品と比べ、この特徴がPG-Stromを「ストレージとの密結合」という “異質な”進化をするきっかけに。 実験用に購入した Intel SSD 750 (400GB) の 理論帯域まで出ている (!) ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 63 2016
64. GPU-Direct SQLの着想と開発（4/4） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 64 同等の機能が、CUDA Toolkit 11.4で標準機能化し、 多くのストレージベンダの対応も。 ➔ これで独自ドライバの保守から解放される！ 2021 2017 GPU Technology Conference 2017において、 “An intelligent storage for PostgreSQL database” と題してポスター発表。 ➔ Top-5 Postersとして選出される。
65. GPU-Direct SQL機構（1/4） PCI-E Bus Buffer Copy Buffer Copy SCAN JOIN GROUP BY Storage Block Read 大量の ”ゴミデータ” も含む ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 65
66. GPU-Direct SQL機構（2/4） P2P-DMAを利用し、NVME-SSDとGPUを直結してデータ転送 PCI-E Bus SCAN JOIN GROUP BY Storage Block Read by NVIDIA GPUDirect Storage P2P-DMA P2P-DMA : Peer-to-Peer Direct Memory Access ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 66
67. GPU-Direct SQL機構（3/4） 測定環境 Supermicro AS-2014CS-TR CPU: AMD EPYC 7402P (24C, 2.8GHz) x1 RAM: 16GB DDR4-3200 ECC x8 GPU: NVIDIA A100 [PCI-E; 40GB] x1 SSD: Intel D7-P5510 [3.84TB] x4 DB: PostgreSQL 15.3 + PG-Strom v5.0dev SSBM (900GB; 60億行) select sum(lo_extendedprice*lo_discount) as revenue from lineorder,date1 where lo_orderdate = d_datekey and d_year = 1993 and lo_discount between 1 and 3 and lo_quantity < 25; ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 67
68. GPU-Direct SQL機構（4/4） ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 68 タスクを細分化し、I/Oと計算を並行させてリソースを使い尽くす GPU0 GPU1 ファイルの オープン （必要なら） ファイルの ダイレクト 読み出し GPUでの SQL処理の 実行 終了 ステータス 確認 ファイルの オープン （必要なら） ファイルの ダイレクト 読み出し GPUでの SQL処理の 実行 終了 ステータス 確認 ファイルの オープン （必要なら） ファイルの ダイレクト 読み出し GPUでの SQL処理の 実行 終了 ステータス 確認 ファイルの オープン （必要なら） ファイルの ダイレクト 読み出し GPUでの SQL処理の 実行 終了 ステータス 確認 ファイルの オープン （必要なら） ファイルの ダイレクト 読み出し GPUでの SQL処理の 実行 終了 ステータス 確認 ファイルの オープン （必要なら） ファイルの ダイレクト 読み出し GPUでの SQL処理の 実行 終了 ステータス 確認 PG-Strom GPU Service (background worker process) GPU0 Task Queue GPU1 Task Queue PostgreSQL Parallel worker process PostgreSQL Backend process PostgreSQL Backend process マルチスレッド処理
69. 本日、取り上げなかったトピック ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 69 ▌Apache Arrow対応  列指向の構造化データ形式  min/max統計情報  Virtual Column機構 ▌GPU Cache  GPUにテーブルのデータを常駐  ログを用いた更新/同期処理 ▌Large Tables Join  動的なパーティショニングによるマルチGPU対応のJOIN ▌GPU Sort/Window関数対応  v6.0に向けて開発中
70. まとめ ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 70 ▌OSSの活用について  きっかけは単なる偶然の気付き。  “Shared Cost”と”Unique Innovation”  “Unique Innovation”を追求する過程で、皆の利益になる副産物も ▌GPU対応PostgreSQL～PG-Strom  カンファレンスでの発見と、『とりあえず手を動かしてみる』  スループット重視のGPUと、SQL検索ワークロードとの類似性  ここでもOSSの利活用～CustomScan APIなど  GPUコア利用率を高めるための工夫 ▌GPU-Direct SQL  NVME-SSDを騙して、直接GPUへデータを転送する  以前は独自のkernel moduleを書いていたが、 現在ではCUDA Toolkitの標準機能で同等の機能を提供している。  非同期並行処理により、NVME-SSDの理論速度に近い処理速度
71. レポート課題 ソフトウェアサイエンス特別講義A～OSS基盤の上に実装するGPUデータベース～ 71 以下の①または②のテーマ、どちらか一つを選んで提出してください。 ① あなたにとって身近なOSS（オープンソースソフトウェア）を一つ選び、 その意思決定プロセスを調べて1～2枚程度で簡潔にまとめてください。 ✓ どういった方法で新機能の提案やバグ修正を行うのか（メール、プルリク） ✓ 意思決定を行っているのは誰か？（特定の個人か、チームか、多数決か） ✓ レポートの根拠となる情報源（開発者向けのガイダンス文書のURLなど） ② あなたの環境にPG-Stromをインストールして動作させ、 期待された動作と異なる、クラッシュするなどのバグを発見して、 GitHub上の issue tracker (※) に適切なバグレポートを登録してください。 ✓ どういった環境にインストールしたのか ✓ 問題を再現させるための手順 ✓ 期待される動作と、実際の動作の差異 ※ PG-Strom issue tracker: https://github.com/heterodb/pg-strom/issues ✓ とても良いバグレポートの例 https://github.com/heterodb/pg-strom/issues/647
72. オモシロ技術を、カタチにする。