PG-StromにおけるGPU PostGIS対応

PG-StromにおけるGPU PostGIS対応 ～GPU版PostGISの実装～ HeteroDB,Inc Chief Architect & CEO 海外 浩平 <[email protected]>

自己紹介／HeteroDB社について ヘテロジニアスコンピューティング技術を データベース領域に適用し、 誰もが使いやすく、安価で高速なデータ解析基盤を提供する。 会社概要  商号  創業  拠点  事業内容 ヘテロDB株式会社 2017年7月4日 東京都品川区北品川 高速データベース製品の販売 GPU＆DB領域の技術コンサルティング 代表者プロフィール  海外 浩平（KaiGai Kohei）  OSS開発者コミュニティにおいて、PostgreSQLやLinux kernelの 開発に10年以上従事。主にセキュリティ・FDW等の分野でアッ プストリームへの貢献。  IPA未踏ソフト事業において“天才プログラマー”認定 (2006)  GPU Technology Conference Japan 2017でInception Awardを受賞 2 2024107セミナー PG-StromにおけるGPU PostGIS対応

PG-Stromとは？ PG-Strom: GPUとNVME/PMEMの能力を最大限に引き出し、 テラバイト級のデータを高速処理するPostgreSQL向け拡張モジュール off-loading for IoT/Big-Data App GPU for ML/Analytics ➢ GPU Direct SQL ➢ Columnar Store (Arrow_Fdw) ➢ Updatable GPU Cache ➢ Asymmetric Partition-wise JOIN/GROUP BY ➢ BRIN-Index Support ➢ Large Tables JOINs ➢ Inter-process Data Frame for Python scripts ➢ Procedural Language for GPU native code (w/ cuPy) ➢ PostGIS Support 【機能】  集計／解析ワークロードの透過的なGPU高速化  GPU Direct SQLによるPCIeバスレベルの最適化  Apache Arrow対応によるデータ交換、インポート時間をほぼゼロに  GPUメモリにデータを常駐。OLTPワークロードとの共存  PostGIS関数のサポート（一部）。位置情報分析を高速化 3 2024107セミナー PG-StromにおけるGPU PostGIS対応

GPUとはどんなプロセッサなのか？ 主にHPC分野で実績があり、機械学習用途で爆発的に普及 NVIDIA A100 スーパーコンピュータ （東京工業大学 TSUBAME3.0） シミュレーション ＣＧ（Computer Graphics） 数千コアの並列処理ユニット、TB/sに達する広帯域メモリを ワンチップに実装した半導体デバイス ➔ “同じ計算を大量のデータに並列実行” を最も得意とする 4 2024107セミナー PG-StromにおけるGPU PostGIS対応 機械学習

ターゲット：移動体デバイスのデータ検索 位置情報 エリア定義情報 （緯度、経度） （ポリゴン） エリア内に存在する 移動体デバイスの抽出 移動体デバイス ▌移動体デバイス  高頻度で位置（緯度、経度）の更新がかかる。  （デバイスID、タイムスタンプ、位置、その他の属性）というデータ構造が多い ▌エリア定義情報  比較的件数は少なく、ほぼ静的なデータだが、エリア定義情報（多角形）が 非常に複雑な構造を持ち、当たり判定が“重い” ▌用途  商圏分析、物流分析、広告配信、アラート通知、etc… 5 2024107セミナー PG-StromにおけるGPU PostGIS対応

GPU版PostGISについて 6 2024107セミナー PG-StromにおけるGPU PostGIS対応

PostGISって？（1/2） ▌概要  点（Point）、線（LineString）、多角形（Polygon）等のジオメトリ要素、 およびそれらの間の演算を PostgreSQL で実行するための拡張モジュール  演算の例 … 包含（Contains）、交差（Crosses）、接合（Touch）など  GiSTインデックスに対応して高速な検索を実現  2005年に PostGIS 1.0 がリリース。以降、20年近くにわたり継続的開発。 ▌PostGIS関数・演算子の例  st_distance(geometry g1, geometry g2) ✓ ジオメトリ間の距離を計算する  st_contains(geometry g1, geometry g2) ✓ ジオメトリ g2 が g1 に包含されていれば真  st_crosses(geometry g1, geometry g2) ✓ ジオメトリ g2 が g1 を横切っていれば真 …など © GAIA RESOURCES 7 2024107セミナー PG-StromにおけるGPU PostGIS対応

PostGISって？（2/2） St_Distance(a,b) ジオメトリ間の距離 St_Contains(a,b) ジオメトリの包含判定 St_Crosses(a,b) ジオメトリの交差判定 8 2024107セミナー PG-StromにおけるGPU PostGIS対応

高速化手法（1/2）- Bounding Box ▌Bounding Boxとは  複雑な形状のポリゴンを包摂する最も小さな矩形領域  (x1,y1) - (x2,y2) で表現できるのでデータ量が少ない  PostgreSQL / PostGISが geometry 型を保存する時に自動的に生成される。 ▌Bounding Boxの効果  空間関係演算を行う前に、『明らかに接していない』ものを識別できる。 ➔ 重なりを含むジオメトリのみ判定を行う事で、計算リソースを節約。 共通部分を持つ可能性 明らかに接していない 9 2024107セミナー PG-StromにおけるGPU PostGIS対応

高速化手法（2/2）- GiSTインデックス ▌GiSTインデックスとは  PostgreSQLのGiST（Generalized Search Tree）は、ツリー構造を持つ インデックスを一般化したフレームワーク。  PostGISのGeometry型は、GiST上にR木を実装している。 ▌GiSTのR木でできること  包含（@演算子）や重なり（&&演算子）判定で、インデックスを用いた絞込み  特に多数のポリゴン×ポイントの重なり判定で計算量が増大しやすい ➔ 事前の絞り込みで計算量を抑え込む # 国土地理院の市町村形状データをインポート $ shp2pgsql N03-20_200101.shp | psql gistest gistest=# ¥d+ List of relations Schema | Name | Type | Owner | Size | --------+-------------+----------+--------+------------+ public | geo_japan | table | kaigai | 243 MB | gistest=# ¥di+ List of relations Schema | Name | Type | Owner | Table | Size | --------+--------------------+-------+--------+-----------+---------+ public | geo_japan_pkey | index | kaigai | geo_japan | 2616 kB | public | geo_japan_geom_idx | index | kaigai | geo_japan | 14 MB | 10 2024107セミナー PG-StromにおけるGPU PostGIS対応

なぜGPUでSQLを高速化できるのか？ SQL処理の中でも、全件スキャンを伴うようなワークロードに強い ▌GPUの特徴  数千演算コアを搭載する計算能力  1.0TB/sに迫るメモリバンド ➔ 「同じ演算を多量のデータに対して実行する」ために作られたプロセッサ （e.g 行列演算） ▌SQLワークロードの特性  「同じ処理を大量の行に対して実行する」 （例：WHERE句の評価、JOINやGROUP BY処理） GPUに実装された数千コアの 実行ユニットが、個々の行を 並列に評価する。 11 2024107セミナー PG-StromにおけるGPU PostGIS対応

GPUでPostGIS関数を実行する（1/3） (100,100) (90,90) (10,90) (12,88) (90,88) (12,12) (90,12) この領域内に含まれる 点（Point）を抽出した (90,10) (10,10) (0,0) 12 テーブル ft および tt には (0,0)～(100,100)の範囲内に ランダムに 500万個の点を格納 2024107セミナー PG-StromにおけるGPU PostGIS対応

GPUでPostGIS関数を実行する（2/3）
postgres=# explain verbose
select count(*) from tt
where st_contains('polygon ((10 10,90 10,90 12,12 12,12 88,90 88,90 90,10 90,10 10))’,
st_makepoint(x,y));
QUERY PLAN
------------------------------------------------------------------------------------------------Aggregate (cost=3945406.28..3945406.29 rows=1 width=8)
Output: pgstrom.fcount((pgstrom.nrows()))
-> Custom Scan (GpuPreAgg) on public.tt (cost=3945406.27..3945406.28 rows=1 width=8)
Output: (pgstrom.nrows())
GPU Projection: pgstrom.nrows()
GPU Scan Quals: st_contains('0103.... 2440'::geometry,
st_makepoint(x, y)) [rows: 4999992 -> 500]
GPU Cache: NVIDIA A100-PCIE-40GB [phase: ready, max_num_rows: 10485760, main: 403.75MB, ...
KVars-Slot: <slot=0, type='float8', expr='x', kv_off=0x0000>, ¥
<slot=1, type='float8', expr='y', kv_off=0x2400>
KVecs-Buffer: nbytes: 18432, ndims: 2, items=[kvec0=<0x0000-23ff, type='float8', expr='x’>,
kvec1=<0x2400-47ff, type='float8', expr='y'>]
LoadVars OpCode: {Packed items[0]={LoadVars(depth=0): kvars=[<slot=0, type='float8' resno=2(x)>,
<slot=1, type='float8' resno=3(y)>]}}
Scan Quals OpCode: {Func(bool)::st_contains args=[{Const(geometry): value='0103.... 2440’},
{Func(geometry)::st_makepoint
args=[{Var(float8): slot=0, expr='x’},
{Var(float8): slot=1, expr='y'}]}]}
Partial Aggregation OpCode: {AggFuncs <nrows[*]>}
Fallback-desc: [<dest='1', expr='x', depth=0:0>, <dest='2', expr='y', depth=0:0>]
Partial Function BufSz: 8
CUDA Stack Size: 3904
(16 rows)
13

2024107セミナー PG-StromにおけるGPU PostGIS対応

GPUでPostGIS関数を実行する（3/3） --- GpuScan (GPU版PostGIS) + GPU-Cache # select count(*) from tt where st_contains('polygon ((10 10,90 10,90 12,12 12,12 88,90 88, 90 90,10 90,10 10))', st_makepoint(x,y)); count -------235985 500万個のPointから、 (1 row) 指定領域内の数をカウント Time: 11.365 ms --- 通常版PostGIS + PostgreSQLテーブル =# SET pg_strom.enabled = off; SET # select count(*) from tt where st_contains('polygon ((10 10,90 10,90 12,12 12,12 88,90 88, 90 90,10 90,10 10))', st_makepoint(x,y)); count -------235985 (1 row) Time: 1094.840 ms (00:01.095) 14 2024107セミナー PG-StromにおけるGPU PostGIS対応

ポリゴン × 座標の突合 位置情報 エリア定義情報 （緯度、経度） （ポリゴン） 移動体デバイス 数百～数万件 数百万～数千万件 ▌ナイーブなポリゴン×座標の突合を行おうとすると、 数万×数千万 = 数兆通りの組合せ、で計算量が爆発。 ➔ GPUが数千コアを搭載していても、 計算量が多すぎて焼け石に水。 焼け石（GPU）に ➔ 水冷モジュール 15 2024107セミナー PG-StromにおけるGPU PostGIS対応

GPUでGiSTインデックス 16 2024107セミナー PG-StromにおけるGPU PostGIS対応

GiSTインデックス（R木）の仕組み 検索キー（経度, 緯度） (xmin,ymin) 〇 〇 (xmax,ymax) ▌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並みに速くはできない。 17 2024107セミナー PG-StromにおけるGPU PostGIS対応

GiSTインデックス（R木）の仕組み 検索キー（経度, 緯度） (xmin,ymin) 〇 〇 × 〇 × (xmax,ymax) ▌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並みに速くはできない。 18 2024107セミナー PG-StromにおけるGPU PostGIS対応

GiSTインデックス（R木）の仕組み 検索キー（経度, 緯度） (xmin,ymin) 〇 〇 × 〇 × (xmax,ymax) × 〇 ▌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並みに速くはできない。 19 2024107セミナー PG-StromにおけるGPU PostGIS対応

GiSTインデックス（R木）の仕組み 検索キー（経度, 緯度） (xmin,ymin) 〇 〇 × 〇 × × × (xmax,ymax) × 〇 ▌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並みに速くはできない。 20 2024107セミナー PG-StromにおけるGPU PostGIS対応

GPU版GiSTインデックス 多角形 × 点の重なり判定などを、GpuJoinの一要素として実装 ポリゴン定義 位置データを含む テーブル GiST(R木)インデックス 数千スレッドが 並列に インデックスを 探索 ▌仕組み  多角形（エリア定義情報）を保持するテーブルと、位置情報（緯度経度）を 保持するテーブルとの JOIN にGiSTインデックスを使用できる。  JOIN処理時、テーブルとインデックスの双方をGPUにロード。 先ずGiSTインデックス上のBounding-Box（矩形領域）を使って荒く絞り込み、 次にテーブル上の多角形（ポリゴン）と「当たり判定」を行う。  GPUの数千コアをフル稼働してGiSTインデックスを探索する。 単純に並列度が高い分、検索速度も速くなるはず。 ➔ が、そうは問屋が卸さなかった。。。 21 2024107セミナー PG-StromにおけるGPU PostGIS対応

簡易テスト：ランダムな点とSt_Contains（初期実装） 国土地理院からDLした 全国市町村形状データ SELECT n03_001,n03_004,count(*) FROM geo_japan j, geopoint p WHERE st_contains(j.geom, st_makepoint(x,y)) AND j.n03_001 like '東京都’ GROUP BY n03_001,n03_004; (123.0, 20.0) n03_001 | n03_004 | count ---------+------------+------東京都 | あきる野市 | 105 東京都 | 三宅村 | 76 東京都 | 三鷹市 | 17 東京都 | 世田谷区 | 67 東京都 | 中央区 | 12 東京都 | 中野区 | 18 東京都 | 八丈町 | 105 ： ： ： 東京都 | 豊島区 | 14 東京都 | 足立区 | 55 東京都 | 青ヶ島村 | 7 東京都 | 青梅市 | 117 (63 rows) ランダムに 生成した座標 1000万個 (154.2, 46.2) 22 CPU版：24.892s GPU版：33.841s（遅い！） 2024107セミナー PG-StromにおけるGPU PostGIS対応

前提）GPUスレッドのスケジューリングについて ●●●…●●● ●●●…●●● ●●●…●●● ●●●…●●● Thread Group (1～1024 Threads) ●●●…●●● ●●●…●●● GPU Block Diagram (Tesla V100; 80SMs) Streaming Multiprocessors (64CUDA cores/SM) Warp (32 threads)  GPUはStreaming Multiprocessor(SM)と呼ばれる単位でコアがグループ化されている。 ✓ レジスタやL1キャッシュ（共有メモリ）はSM単位のリソース  同じSM上で実行されるスレッドの組を Thread Group と呼び、リソースが許せば 最大で 1024 スレッドまで並行で実行する事ができる。 ✓ スレッドの切り替えはH/W的に実現され、例えばDRAMアクセス待ちのタイミングで切り替わったり。  スレッドはWarp (32threads) 単位でスケジューリングされる。同一Warp内のス レッドは、同時には同じ命令しか実行できない。 ✓ 行列計算のように、各コアの負荷が均等になるタイプの処理であればGPUの使用効率は最大化。 ✓ スレッド毎に処理時間がバラバラだと、一部のコアが遊んでしまい、処理効率の低下を招く。 23 2024107セミナー PG-StromにおけるGPU PostGIS対応

前提）GpuJoinの並列処理デザイン Table-B SELECT * FROM A, B WHERE A.id = B.id; 512スレッドが 512行を一度に 取り出し ● ● ● Table-A 参照 GpuHashJoin / GpuNestLoop 処理 〇 × GpuHashJoin処理 • Hash値の計算 • Hash表の探索 • JOIN条件の評価 GpuNestLoop処理 • JOIN条件の評価 ➔スレッド間で大きな 処理時間の差が つきにくい。 24 ● ● ● 〇 × 〇 × N = __syncthreads_count(...) atomic演算を用いて、thread_id=0が 結果バッファにN個分の領域を確保 JOIN結果の書き出し 2024107セミナー PG-StromにおけるGPU PostGIS対応 次のフレームを 取り出して、 以降繰り返し

GPUでのインデックス探索の課題 スレッド間の処理時間の差が大きく、同期待ちを招く ● ● ● ● ● ● R木のRootノードを 見ただけで、即、 マッチする要素なし 参照 Index-A × R木のLeafノード まで探索したが ヒットせず 参照 Table-A × × × × R木を最後まで 探索し、かつ、 JOINの結合条件を評価 スレッドグループに属する 他のスレッドの完了待ちで GPUコアが遊んでしまう。 〇 N = __syncthreads_count(...) 25 2024107セミナー PG-StromにおけるGPU PostGIS対応

ナイーブなGPU版GiSTインデックス実装（初期実装） GPUの利用効率が上がらず、しばしばCPUより低速だった ● ● ● ● 繰り返し ● ● ● ● 各スレッドが１行分のデータをロード GiST-Index探索用のキー値を抽出する。 キー値を用いてGiST-Indexを探索する GiST-Indexにヒットした？ スレッドグループ（512スレッド）内に、 １個でもGiST-IndexのLeafまで降下して 探索したり、JOIN結合条件の評価まで 行うスレッドがいると、他の511スレッ ドはその完了を待たされてしまう。 結果として、GPUコアの使用率が全く 向上しないという事になる。 No JOIN結合条件を評価する JOIN結合条件は真？ No found=true nitems = __syncthreads_count(found); found=false nitems個分のJOIN結果を書き出し 26 2024107セミナー PG-StromにおけるGPU PostGIS対応

GPUのスレッドスケジューリングを意識した実装（現行版） できる限りコアが遊ばないよう、同期ポイントを最小化 ● ● ● ● 繰り返し ● ● ● ● 各Warp（32スレッド）毎に１行分のデータをロード GiST-Index探索用のキー値を抽出する。 キー値を用いてGiST-Indexを探索する No GiST-Indexにヒットした？ Atomic演算を用いて中間バッファの領域を確保 GiST-Indexにヒットした行のポインタを記録 No 中間バッファの使用量が 512個を越えた 中間バッファの使用量に余裕がある限り、 同期ポイントに達する前にどんどん Readポインタを進めて GiST-Index の 探索を進める。 ➔ 一部のWarp（スレッド）で探索に 時間を要しても、他のWarpは先に 次の行を処理できるため、 GPUコアの使用率を高くできる。 __syncthreads() 中間バッファを元に、JOIN結合条件を評価する JOIN結合条件は真？ found=false found=true nitems = __syncthreads_count(found) nitems個分のJOIN結果を書き出し 27 2024107セミナー PG-StromにおけるGPU PostGIS対応

簡易テスト：ランダムな点とSt_Contains（現行版） GPUコアのスケジューリングを 意識した実装に修正し再測定 国土地理院からDLした 全国市町村形状データ SELECT n03_001,n03_004,count(*) FROM geo_japan j, geopoint p WHERE st_contains(j.geom, st_makepoint(x,y)) AND j.n03_001 like '東京都’ GROUP BY n03_001,n03_004; (123.0, 20.0) n03_001 | n03_004 | count ---------+------------+------東京都 | あきる野市 | 105 東京都 | 三宅村 | 76 東京都 | 三鷹市 | 17 東京都 | 世田谷区 | 67 東京都 | 中央区 | 12 東京都 | 中野区 | 18 東京都 | 八丈町 | 105 ： ： ： 東京都 | 豊島区 | 14 東京都 | 足立区 | 55 東京都 | 青ヶ島村 | 7 東京都 | 青梅市 | 117 (63 rows) ランダムに 生成した座標 1000万個 (154.2, 46.2) 28 CPU版：24.892s GPU版： 1.154s 2024107セミナー PG-StromにおけるGPU PostGIS対応 21.5倍の 高速化

簡易テストの実行計画（CPU版）
postgres=# EXPLAIN (analyze, costs off)
SELECT n03_001,n03_004,count(*)
FROM geo_japan j, geopoint p
WHERE st_contains(j.geom, st_makepoint(x,y))
AND j.n03_001 like '東京都’
GROUP BY n03_001,n03_004;
QUERY PLAN
---------------------------------------------------------------------------------------Finalize GroupAggregate (actual time=24880.443..24880.682 rows=63 loops=1)
Group Key: j.n03_001, j.n03_004
-> Gather Merge (actual time=24880.430..24892.342 rows=299 loops=1)
Workers Planned: 4
Workers Launched: 4
-> Partial GroupAggregate (actual time=24855.860..24855.949 rows=60 loops=5)
Group Key: j.n03_001, j.n03_004
-> Sort (actual time=24855.846..24855.865 rows=529 loops=5)
Sort Key: j.n03_001, j.n03_004
Sort Method: quicksort Memory: 64kB
-> Nested Loop (actual time=50.124..24854.788 rows=529 loops=5)
-> Parallel Seq Scan on geopoint p (actual time=0.120..258.900
rows=2000000 loops=5)
-> Index Scan using geo_japan_geom_idx on geo_japan j
(actual time=0.012..0.012 rows=0 loops=10000000)
Index Cond: (geom ~ st_makepoint(p.x, p.y))
Filter: (((n03_001)::text ~~ '東京都'::text) AND
st_contains(geom, st_makepoint(p.x, p.y)))
Rows Removed by Filter: 0
Planning Time: 0.160 ms
Execution Time: 24892.586 ms

29

2024107セミナー PG-StromにおけるGPU PostGIS対応

簡易テストの実行計画（GPU版）
postgres=# EXPLAIN (analyze, costs off)
SELECT n03_001,n03_004,count(*)
FROM geo_japan j, fgeopoint p
WHERE st_contains(j.geom, st_makepoint(x,y))
AND j.n03_001 like '東京都’
GROUP BY n03_001,n03_004;
QUERY PLAN
----------------------------------------------------------------------------------------------GroupAggregate (actual time=1146.188..1146.212 rows=63 loops=1)
Group Key: j.n03_001, j.n03_004
-> Sort (actual time=1146.178..1146.181 rows=63 loops=1)
Sort Key: j.n03_001, j.n03_004
Sort Method: quicksort Memory: 29kB
-> Custom Scan (GpuPreAgg) (actual time=1146.004..1146.013 rows=63 loops=1)
Reduction: Local
Combined GpuJoin: enabled
-> Custom Scan (GpuJoin) on fgeopoint p (never executed)
Outer Scan: fgeopoint p (never executed)
Depth 1: GpuGiSTJoin
HeapSize: 7841.91KB (estimated: 3113.70KB), IndexSize: 13.28MB
IndexFilter: (j.geom ~ st_makepoint(p.x, p.y)) on geo_japan_geom_idx
Rows Fetched by Index: 5065
JoinQuals: st_contains(j.geom, st_makepoint(p.x, p.y))
-> Seq Scan on geo_japan j (actual time=0.136..16.067 rows=6173 loops=1)
Filter: ((n03_001)::text ~~ '東京都'::text)
Rows Removed by Filter: 112726
Planning Time: 0.307 ms
Execution Time: 1154.838 ms
30

2024107セミナー PG-StromにおけるGPU PostGIS対応

GPU版に固有の最適化

not in use

IndexTupleData

IndexTupleData

IndexTupleData

IndexTupleData

IndexTupleData

IndexTupleData

IndexTupleData

GiST Rootブロック

あらかじめ、明らかに条件に該当しないエントリ
（この場合は ‘東京都’ 以外）を無効化

EXPLAIN
SELECT n03_001,n03_004,count(*)
FROM geo_japan j, geopoint p
WHERE st_contains(j.geom, st_makepoint(x,y))
AND j.n03_001 like '東京都’
GROUP BY n03_001,n03_004;
QUERY PLAN
-----------------------------------------------------GroupAggregate
Group Key: j.n03_001, j.n03_004
-> Sort
Sort Key: j.n03_001, j.n03_004
-> Nested Loop
-> Seq Scan on geopoint p
-> Index Scan using geo_japan_geom_idx on geo_japan j
Index Cond: (geom ~ st_makepoint(p.x, p.y))
Filter: (((n03_001)::text ~~ '東京都') AND
st_contains(geom, st_makepoint(p.x, p.y)))
(9 rows)

▌ PostgreSQLの場合、Indexに付随する条件句の評価は、Indexを用いて候補となる行を取り出し
た後にしか行えない。つまり、結果的に無駄となる行フェッチが発生する。
▌ GPU版GiSTインデックスでは、予め「明らかに条件に該当しないエントリ」のLeaf要素を無効
化するため、何度も何度も自明な条件句の評価を行う必要はない。
➔ 『１週間分のデータから直近30分のイベントだけを取り出す』といった特性がある場合に、
通常のGiSTインデックス探索よりも有利に働く。
31

2024107セミナー PG-StromにおけるGPU PostGIS対応

小噺）ちょっと感動した話
postgres=# EXPLAIN

SELECT n03_001,n03_004,count(*)
FROM geo_japan j, geopoint p
WHERE st_dwithin(j.geom, st_makepoint(x,y), 0.004)
AND j.n03_001 like '東京都’
GROUP BY n03_001,n03_004;
QUERY PLAN
----------------------------------------------------------------------------------------------------HashAggregate (cost=1936845.15..1936893.73 rows=4858 width=29)
Group Key: j.n03_001, j.n03_004
-> Custom Scan (GpuJoin) on geopoint p (cost=159401.97..1480539.90 rows=60840700 width=21)
Outer Scan: geopoint p (cost=0.00..163696.15 rows=10000115 width=16)
Depth 1: GpuGiSTJoin(nrows 10000115...60840700)
HeapSize: 3113.70KB
IndexFilter: (j.geom && st_expand(st_makepoint(p.x, p.y),
'0.004'::double precision)) on geo_japan_geom_idx
JoinQuals: st_dwithin(j.geom, st_makepoint(p.x, p.y), '0.004'::double precision)
GPU Preference: GPU0 (Tesla V100-PCIE-16GB)
-> Seq Scan on geo_japan j (cost=0.00..8928.24 rows=6084 width=1868)
Filter: ((n03_001)::text ~~ '東京都'::text)
(11 rows)
st_expand()
✓ 引数1(点)を上下左右に引数2(実数)だけ拡張
した矩形領域を返す。
➔ これとGiSTインデックスのBounding-Boxが
共通領域を持てば、st_dwithin()が真となる
可能性がある。
32

0.004

st_makepoint(p.x, p.y)
0.004

2024107セミナー PG-StromにおけるGPU PostGIS対応

まとめ ▌GPU版PostGISとは  PG-StromがPostgreSQLのSQL関数を実行するのと同じ仕組みを利用して、 一部のPostGIS関数をGPUで実行するための実装。  非常に数多くのPostGIS関数が提供されているが、GPU版PostGISはそのうちの ごく一部を提供している。（リクエスト次第！） ▌GPU版GiSTインデックス  PostgreSQLでPostGISを使った検索に使われるGiSTインデックスをGPUに読出し、 GPU上の検索でもこれを利用する。 ✓ GISに対応したGPU-DBは他にもあるが、インデックス対応はPG-Stromのみ！  検索条件に合致しないGiSTインデックス上のエントリは、最初から無効化して GPUにロードするため、CPUのそれより効率が高い。 ▌GPUコアスケジューリングを意識した実装  GiSTインデックスによる絞込み後、PostGIS関数の評価の前にバッファリングを 挟む事で、32スレッドが同時に動作するようにする。 33 2024107セミナー PG-StromにおけるGPU PostGIS対応