---
title: PG-StromにおけるGPU PostGIS対応
tags: 
author: [KaiGai Kohei](https://convert.docswell.com/user/kaigai)
site: [Docswell](https://www.docswell.com/)
thumbnail: https://bcdn.docswell.com/page/VENGP2NVJ8.jpg?width=480
description: 2024/10/7のセミナー資料です
published: October 09, 24
canonical: https://convert.docswell.com/s/kaigai/KP2M3G-20241007-GPU-PostGIS
---
# Page. 1

![Page Image](https://bcdn.docswell.com/page/VENGP2NVJ8.jpg)

PG-StromにおけるGPU PostGIS対応
～GPU版PostGISの実装～
HeteroDB,Inc
Chief Architect &amp; CEO
海外 浩平 &lt;kaigai@heterodb.com&gt;


# Page. 2

![Page Image](https://bcdn.docswell.com/page/Y79W31RV73.jpg)

自己紹介／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対応


# Page. 3

![Page Image](https://bcdn.docswell.com/page/G78NM4W4JD.jpg)

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対応


# Page. 4

![Page Image](https://bcdn.docswell.com/page/L7LP36NVJR.jpg)

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


# Page. 5

![Page Image](https://bcdn.docswell.com/page/4EM4D5XNJW.jpg)

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


# Page. 6

![Page Image](https://bcdn.docswell.com/page/PERW43N4J9.jpg)

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


# Page. 7

![Page Image](https://bcdn.docswell.com/page/P7XP24NZJX.jpg)

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対応


# Page. 8

![Page Image](https://bcdn.docswell.com/page/37KDMZNVED.jpg)

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


# Page. 9

![Page Image](https://bcdn.docswell.com/page/LJ3RYDZQE5.jpg)

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


# Page. 10

![Page Image](https://bcdn.docswell.com/page/8JDY5PRWJG.jpg)

高速化手法（2/2）- GiSTインデックス
▌GiSTインデックスとは
 PostgreSQLのGiST（Generalized Search Tree）は、ツリー構造を持つ
インデックスを一般化したフレームワーク。
 PostGISのGeometry型は、GiST上にR木を実装している。
▌GiSTのR木でできること
 包含（@演算子）や重なり（&amp;&amp;演算子）判定で、インデックスを用いた絞込み
 特に多数のポリゴン×ポイントの重なり判定で計算量が増大しやすい
➔ 事前の絞り込みで計算量を抑え込む
# 国土地理院の市町村形状データをインポート
$ 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対応


# Page. 11

![Page Image](https://bcdn.docswell.com/page/VEP6GNWX78.jpg)

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


# Page. 12

![Page Image](https://bcdn.docswell.com/page/27V36K83JQ.jpg)

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対応


# Page. 13

![Page Image](https://bcdn.docswell.com/page/5JG3NQ5YEL.jpg)

GPUでPostGIS関数を実行する（2/3）
postgres=# explain verbose
select count(*) from tt
where st_contains(&#039;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()))
-&gt; 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(&#039;0103.... 2440&#039;::geometry,
st_makepoint(x, y)) [rows: 4999992 -&gt; 500]
GPU Cache: NVIDIA A100-PCIE-40GB [phase: ready, max_num_rows: 10485760, main: 403.75MB, ...
KVars-Slot: &lt;slot=0, type=&#039;float8&#039;, expr=&#039;x&#039;, kv_off=0x0000&gt;, ¥
&lt;slot=1, type=&#039;float8&#039;, expr=&#039;y&#039;, kv_off=0x2400&gt;
KVecs-Buffer: nbytes: 18432, ndims: 2, items=[kvec0=&lt;0x0000-23ff, type=&#039;float8&#039;, expr=&#039;x’&gt;,
kvec1=&lt;0x2400-47ff, type=&#039;float8&#039;, expr=&#039;y&#039;&gt;]
LoadVars OpCode: {Packed items[0]={LoadVars(depth=0): kvars=[&lt;slot=0, type=&#039;float8&#039; resno=2(x)&gt;,
&lt;slot=1, type=&#039;float8&#039; resno=3(y)&gt;]}}
Scan Quals OpCode: {Func(bool)::st_contains args=[{Const(geometry): value=&#039;0103.... 2440’},
{Func(geometry)::st_makepoint
args=[{Var(float8): slot=0, expr=&#039;x’},
{Var(float8): slot=1, expr=&#039;y&#039;}]}]}
Partial Aggregation OpCode: {AggFuncs &lt;nrows[*]&gt;}
Fallback-desc: [&lt;dest=&#039;1&#039;, expr=&#039;x&#039;, depth=0:0&gt;, &lt;dest=&#039;2&#039;, expr=&#039;y&#039;, depth=0:0&gt;]
Partial Function BufSz: 8
CUDA Stack Size: 3904
(16 rows)
13
2024107セミナー PG-StromにおけるGPU PostGIS対応


# Page. 14

![Page Image](https://bcdn.docswell.com/page/47Q5QKZ6EP.jpg)

GPUでPostGIS関数を実行する（3/3）
--- GpuScan (GPU版PostGIS) + GPU-Cache
# select count(*) from tt
where st_contains(&#039;polygon ((10 10,90 10,90 12,12 12,12 88,90 88,
90 90,10 90,10 10))&#039;, 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(&#039;polygon ((10 10,90 10,90 12,12 12,12 88,90 88,
90 90,10 90,10 10))&#039;, st_makepoint(x,y));
count
-------235985
(1 row)
Time: 1094.840 ms (00:01.095)
14
2024107セミナー PG-StromにおけるGPU PostGIS対応


# Page. 15

![Page Image](https://bcdn.docswell.com/page/KE49NP3MJ1.jpg)

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


# Page. 16

![Page Image](https://bcdn.docswell.com/page/L71K591YEG.jpg)

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


# Page. 17

![Page Image](https://bcdn.docswell.com/page/G7WD5Q81E2.jpg)

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対応


# Page. 18

![Page Image](https://bcdn.docswell.com/page/4JZGNW8X73.jpg)

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対応


# Page. 19

![Page Image](https://bcdn.docswell.com/page/YE6Y8GPPJV.jpg)

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対応


# Page. 20

![Page Image](https://bcdn.docswell.com/page/GE5G9VKPE4.jpg)

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対応


# Page. 21

![Page Image](https://bcdn.docswell.com/page/972Y23WZ7R.jpg)

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


# Page. 22

![Page Image](https://bcdn.docswell.com/page/DJYGY1LNJM.jpg)

簡易テスト：ランダムな点と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 &#039;東京都’
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対応


# Page. 23

![Page Image](https://bcdn.docswell.com/page/V7NGP24VE8.jpg)

前提）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対応


# Page. 24

![Page Image](https://bcdn.docswell.com/page/YJ9W31QVJ3.jpg)

前提）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対応
次のフレームを
取り出して、
以降繰り返し


# Page. 25

![Page Image](https://bcdn.docswell.com/page/GJ8NM4G4ED.jpg)

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


# Page. 26

![Page Image](https://bcdn.docswell.com/page/LJLP36GVER.jpg)

ナイーブな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対応


# Page. 27

![Page Image](https://bcdn.docswell.com/page/47M4D5QNEW.jpg)

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対応


# Page. 28

![Page Image](https://bcdn.docswell.com/page/P7RW4384E9.jpg)

簡易テスト：ランダムな点と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 &#039;東京都’
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倍の
高速化


# Page. 29

![Page Image](https://bcdn.docswell.com/page/PJXP248ZEX.jpg)

簡易テストの実行計画（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 &#039;東京都’
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
-&gt; Gather Merge (actual time=24880.430..24892.342 rows=299 loops=1)
Workers Planned: 4
Workers Launched: 4
-&gt; Partial GroupAggregate (actual time=24855.860..24855.949 rows=60 loops=5)
Group Key: j.n03_001, j.n03_004
-&gt; Sort (actual time=24855.846..24855.865 rows=529 loops=5)
Sort Key: j.n03_001, j.n03_004
Sort Method: quicksort Memory: 64kB
-&gt; Nested Loop (actual time=50.124..24854.788 rows=529 loops=5)
-&gt; Parallel Seq Scan on geopoint p (actual time=0.120..258.900
rows=2000000 loops=5)
-&gt; 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 ~~ &#039;東京都&#039;::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対応


# Page. 30

![Page Image](https://bcdn.docswell.com/page/3JKDMZKV7D.jpg)

簡易テストの実行計画（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 &#039;東京都’
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
-&gt; Sort (actual time=1146.178..1146.181 rows=63 loops=1)
Sort Key: j.n03_001, j.n03_004
Sort Method: quicksort Memory: 29kB
-&gt; Custom Scan (GpuPreAgg) (actual time=1146.004..1146.013 rows=63 loops=1)
Reduction: Local
Combined GpuJoin: enabled
-&gt; 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))
-&gt; Seq Scan on geo_japan j (actual time=0.136..16.067 rows=6173 loops=1)
Filter: ((n03_001)::text ~~ &#039;東京都&#039;::text)
Rows Removed by Filter: 112726
Planning Time: 0.307 ms
Execution Time: 1154.838 ms
30
2024107セミナー PG-StromにおけるGPU PostGIS対応


# Page. 31

![Page Image](https://bcdn.docswell.com/page/LE3RYD3Q75.jpg)

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 &#039;東京都’
GROUP BY n03_001,n03_004;
QUERY PLAN
-----------------------------------------------------GroupAggregate
Group Key: j.n03_001, j.n03_004
-&gt; Sort
Sort Key: j.n03_001, j.n03_004
-&gt; Nested Loop
-&gt; Seq Scan on geopoint p
-&gt; Index Scan using geo_japan_geom_idx on geo_japan j
Index Cond: (geom ~ st_makepoint(p.x, p.y))
Filter: (((n03_001)::text ~~ &#039;東京都&#039;) 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対応


# Page. 32

![Page Image](https://bcdn.docswell.com/page/8EDY5P4WEG.jpg)

小噺）ちょっと感動した話
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 &#039;東京都’
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
-&gt; 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 &amp;&amp; st_expand(st_makepoint(p.x, p.y),
&#039;0.004&#039;::double precision)) on geo_japan_geom_idx
JoinQuals: st_dwithin(j.geom, st_makepoint(p.x, p.y), &#039;0.004&#039;::double precision)
GPU Preference: GPU0 (Tesla V100-PCIE-16GB)
-&gt; Seq Scan on geo_japan j (cost=0.00..8928.24 rows=6084 width=1868)
Filter: ((n03_001)::text ~~ &#039;東京都&#039;::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対応


# Page. 33

![Page Image](https://bcdn.docswell.com/page/V7P6GNMXJ8.jpg)

まとめ
▌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対応


# Page. 34

![Page Image](https://bcdn.docswell.com/page/2JV36K93EQ.jpg)



