659 Views
October 09, 24
スライド概要
2024/10/7のセミナー資料です
かぴばら
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) シミュレーション CG(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より低速だった ● ● ● ● 繰り返し ● ● ● ● 各スレッドが1行分のデータをロード GiST-Index探索用のキー値を抽出する。 キー値を用いてGiST-Indexを探索する GiST-Indexにヒットした? スレッドグループ(512スレッド)内に、 1個でも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スレッド)毎に1行分のデータをロード 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要素を無効
化するため、何度も何度も自明な条件句の評価を行う必要はない。
➔ 『1週間分のデータから直近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対応