[VTJ-JP]FULL JOINが失敗する
Closed this issue · 14 comments
ytooyama commented
testdb=# SET pg_strom.enabled=off;
testdb=# SELECT * FROM uriage1 FULL JOIN uriage2 ON uriage1.number = uriage2.number LIMIT 10;
number | username | okashi | amount_sold | number | username | okashi | amount_sold
--------+----------+-----------+-------------+--------+----------+-----------+-------------
10049 | Bob | Candy | 830 | 10049 | Alice | Chocolate | 200
10050 | Jane | Cookie | 230 | 10050 | Alice | Candy | 10
10051 | Bob | Icecream | 170 | 10051 | Bob | Chocolate | 640
10052 | Bob | Cookie | 410 | 10052 | Ted | Icecream | 610
10053 | Alice | Chocolate | 290 | 10053 | Jane | Cookie | 790
10054 | Bob | Cookie | 190 | 10054 | Jane | Cookie | 730
10055 | Bob | Cookie | 410 | 10055 | Alice | Chocolate | 660
10056 | Ted | Candy | 950 | 10056 | Jane | Candy | 870
10057 | Alice | Candy | 660 | 10057 | Jane | Cookie | 320
10058 | Bob | Cookie | 100 | 10058 | Ted | Candy | 550
(10 rows)
testdb=# SET pg_strom.enabled=on;
testdb=# SELECT * FROM uriage1 FULL JOIN uriage2 ON uriage1.number = uriage2.number LIMIT 10;
ERROR: (codegen_expression_walker:2394) not a supported expression type: {VAR :varno 2 :varattno 1 :vartype 23 :vartypmod -1 :varcollid 0 :varnullingrels (b 3) :varlevelsup 0 :varnosyn 2 :varattnosyn 1 :location 7}
DETAIL: problematic expression: {VAR :varno 2 :varattno 4 :vartype 23 :vartypmod -1 :varcollid 0 :varnullingrels (b 3) :varlevelsup 0 :varnosyn 2 :varattnosyn 4 :location 7}
testdb=# SELECT * FROM uriage1 FULL JOIN uriage2 ON uriage1.username = uriage2.username LIMIT 10;
ERROR: (codegen_expression_walker:2394) not a supported expression type: {VAR :varno 2 :varattno 1 :vartype 23 :vartypmod -1 :varcollid 0 :varnullingrels (b 3) :varlevelsup 0 :varnosyn 2 :varattnosyn 1 :location 7}
DETAIL: problematic expression: {VAR :varno 2 :varattno 4 :vartype 23 :vartypmod -1 :varcollid 0 :varnullingrels (b 3) :varlevelsup 0 :varnosyn 2 :varattnosyn 4 :location 7}
testdb=# SELECT * FROM uriage1 FULL JOIN uriage2 ON uriage1.okashi = uriage2.okashi LIMIT 10;
number | username | okashi | amount_sold | number | username | okashi | amount_sold
--------+----------+--------+-------------+---------+----------+--------+-------------
10049 | Bob | Candy | 830 | 4999821 | Alice | Candy | 320
1 | Bob | Cookie | 80 | 4999978 | Alice | Cookie | 230
1 | Bob | Cookie | 80 | 4999817 | Ted | Cookie | 490
1 | Bob | Cookie | 80 | 4999816 | Alice | Cookie | 450
1 | Bob | Cookie | 80 | 4999814 | Alice | Cookie | 200
1 | Bob | Cookie | 80 | 4999964 | Bob | Cookie | 590
1 | Bob | Cookie | 80 | 4999810 | Jane | Cookie | 740
1 | Bob | Cookie | 80 | 4999963 | Bob | Cookie | 540
1 | Bob | Cookie | 80 | 4999961 | Jane | Cookie | 790
1 | Bob | Cookie | 80 | 4999807 | Ted | Cookie | 780
(10 rows)
testdb=# EXPLAIN VERBOSE SELECT * FROM uriage1 FULL JOIN uriage2 ON uriage1.number = uriage2.number LIMIT 10;
ERROR: (codegen_expression_walker:2394) not a supported expression type: {VAR :varno 2 :varattno 1 :vartype 23 :vartypmod -1 :varcollid 0 :varnullingrels (b 3) :varlevelsup 0 :varnosyn 2 :varattnosyn 1 :location 23}
DETAIL: problematic expression: {VAR :varno 2 :varattno 4 :vartype 23 :vartypmod -1 :varcollid 0 :varnullingrels (b 3) :varlevelsup 0 :varnosyn 2 :varattnosyn 4 :location 23}
テストケース作成
CREATE TABLE uriage1
(
number SERIAL,
username VARCHAR(128) NOT NULL,
okashi VARCHAR(128) NOT NULL,
amount_sold INTEGER NOT NULL
);
CREATE TABLE uriage2
(
number SERIAL,
username VARCHAR(128) NOT NULL,
okashi VARCHAR(128) NOT NULL,
amount_sold INTEGER NOT NULL
);
INSERT INTO
uriage1 (username,okashi,amount_sold)
SELECT
(array['Alice', 'Jane', 'Ted', 'Bob'])[ceil(random() * 4)] AS username,
(array['Candy', 'Cookie', 'Chocolate', 'Icecream'])[ceil(random() * 4)] AS okashi,
ceil(random() * 100) *10 amount_sold
FROM
GENERATE_SERIES(1, 10000000);
INSERT INTO
uriage2 (username,okashi,amount_sold)
SELECT
(array['Alice', 'Jane', 'Ted', 'Bob'])[ceil(random() * 4)] AS username,
(array['Candy', 'Cookie', 'Chocolate', 'Icecream'])[ceil(random() * 4)] AS okashi,
ceil(random() * 100) *10 amount_sold
FROM
GENERATE_SERIES(1, 5000000);
VACUUM FULL uriage1;
VACUUM FULL uriage2;
kaigai commented
これ、最新 3cd806bb56d91802ca7a94ce983c68ff7c3fa27c だと再現できませぬ。
最新版で再現確認をお願いできますか。（たしか割と最近、類似問題を対処した気がするので）
ytooyama commented
PG-Stromについては同じバージョンですが、相変わらず再現します。
testdb=# SELECT * FROM uriage1 FULL JOIN uriage2 ON uriage1.number = uriage2.number LIMIT 10;
ERROR: (codegen_expression_walker:2394) not a supported expression type: {VAR :varno 2 :varattno 1 :vartype 23 :vartypmod -1 :varcollid 0 :varnullingrels (b 3) :varlevelsup 0 :varnosyn 2 :varattnosyn 1 :location 7}
DETAIL: problematic expression: {VAR :varno 2 :varattno 4 :vartype 23 :vartypmod -1 :varcollid 0 :varnullingrels (b 3) :varlevelsup 0 :varnosyn 2 :varattnosyn 4 :location 7}
kaigai commented
make clean && make -j 12 した後でもどうでしょう？
プラナーの問題なので、あまり実行環境によって変わるという気もしないんですよね。
kaigai commented
なおこちらでは、全く発生機序の異なる以下のエラーが出ており、それを解析中です。
hoge=# SELECT * FROM uriage1 FULL JOIN uriage2 ON uriage1.number = uriage2.number LIMIT 10;
2023-12-27 10:13:02.251 JST [1935214] ERROR: gpu_service.c:1997 failed on cuEventSynchronize: CUDA_ERROR_ILLEGAL_INSTRUCTION
2023-12-27 10:13:02.251 JST [1935214] HINT: device at GPU-0, function at gpuservHandleGpuTaskExec
2023-12-27 10:13:02.251 JST [1935214] STATEMENT: SELECT * FROM uriage1 FULL JOIN uriage2 ON uriage1.number = uriage2.number LIMIT 10;
ytooyama commented
make clean && make -j 12 した後でもどうでしょう？
普段、make cleanは必ず実行するのですが、一応試してみました。
結果は変わらずです。ちなみにEXTENSIONのremove & createも変化ありませんでした。
kaigai commented
これ、ベースにしているPostgreSQLのバージョンっていくつでしょう？
オプティマイザなので、GPU側ともあまり思えないのです。
ytooyama commented
ここから取得したPostgreSQL 16の最新版です。
https://www.postgresql.org/download/linux/ubuntu/
postgres@kujira:~$ dpkg -l postgresql-16
Desired=Unknown/Install/Remove/Purge/Hold
| Status=Not/Inst/Conf-files/Unpacked/halF-conf/Half-inst/trig-aWait/Trig-pend
|/ Err?=(none)/Reinst-required (Status,Err: uppercase=bad)
||/ Name Version Architecture Description
+++-==============-==================-============-=========================================================
ii postgresql-16 16.1-1.pgdg22.04+1 amd64 The World's Most Advanced Open Source Relational Database
kaigai commented
なるほどPostgreSQL v16😃
ちょっと試してみます・・・が、先ず CUDA_ERROR_ILLEGAL_INSTRUCTION を見ています…。
kaigai commented
CUDA_ERROR_ILLEGAL_INSTRUCTIONに関しては32ce65ba042d3c17fec6e1942cc8f7bbed735cbaで修正。
kern_gpujoin_main() がサスペンド ⇒ レジュームした際、depthが未初期化となるケースがあり、
__syncthreads()が永遠に合流しないケースが発生。
一部のスレッドが既にexitしている状態、あるいは完全に合流の可能性がない場合に同期しようとすると、
（分かりにくいが）CUDA_ERROR_ILLEGAL_INSTRUCTION を起こすようだ。
大元の問題については、明日、PostgreSQL v16で試してみます。
kaigai commented
あと、GPU kernelのsuspend/resumeが発生するような状況で、実行結果が異なるので、それも。
hoge=# explain analyze SELECT * FROM uriage1 FULL OUTER JOIN uriage2 ON uriage1.number = uriage2.number;
QUERY PLAN
----------------------------------------------------------------------------------------------------------------------------------------------------------------
Custom Scan (GpuJoin) on uriage1 (cost=144447.82..311269.01 rows=10000115 width=40) (actual time=3166.901..4996.398 rows=8696185 loops=1)
GPU Projection: uriage1.number, uriage1.username, uriage1.okashi, uriage1.amount_sold, uriage2.number, uriage2.username, uriage2.okashi, uriage2.amount_sold
GPU Full Outer Join Quals [1]: (uriage1.number = uriage2.number) ... [plan: 10000120 -> 10000120, exec: 1402094 -> 1402094]
GPU Outer Hash [1]: uriage1.number
GPU Inner Hash [1]: uriage2.number
GPU-Direct SQL: enabled (GPU-0; direct=63695, ntuples=1402094)
-> Seq Scan on uriage2 (cost=0.00..81847.92 rows=4999992 width=20) (actual time=0.016..496.762 rows=5000000 loops=1)
Planning Time: 0.247 ms
Execution Time: 5564.046 ms
(9 rows)
kds_dstが満杯にならないようなケースだと、きちんとピッタリの行数を返す。
hoge=# explain analyze SELECT uriage1.number, uriage2.amount_sold FROM uriage1 FULL OUTER JOIN uriage2 ON uriage1.number = uriage2.number;
QUERY PLAN
--------------------------------------------------------------------------------------------------------------------------------------------
Custom Scan (GpuJoin) on uriage1 (cost=144447.82..311269.01 rows=10000115 width=8) (actual time=2452.021..3314.194 rows=10000000 loops=1)
GPU Projection: uriage1.number, uriage2.amount_sold
GPU Full Outer Join Quals [1]: (uriage1.number = uriage2.number) ... [plan: 10000120 -> 10000120, exec: 10000000 -> 10000000]
GPU Outer Hash [1]: uriage1.number
GPU Inner Hash [1]: uriage2.number
GPU-Direct SQL: enabled (GPU-0; direct=63695, ntuples=10000000)
-> Seq Scan on uriage2 (cost=0.00..81847.92 rows=4999992 width=8) (actual time=0.020..992.465 rows=5000000 loops=1)
Planning Time: 0.224 ms
Execution Time: 3660.088 ms
(9 rows)
kaigai commented
GPU kernelのsuspend/resumeの時に、再開時の depth が -1 のままだった。（→再開したGPU kernelは何もせず即終了になる）
suspendする時の depth は n_rels + 1 意外にあり得ないので、2caf58acf1016e9265695986899bd5a09851f3cc の修正でOK。
hoge=# explain analyze SELECT * FROM uriage1 FULL OUTER JOIN uriage2 ON uriage1.number = uriage2.number;
QUERY PLAN
----------------------------------------------------------------------------------------------------------------------------------------------------------------
Custom Scan (GpuJoin) on uriage1 (cost=144447.82..311269.01 rows=10000115 width=40) (actual time=3214.351..4483.935 rows=10000000 loops=1)
GPU Projection: uriage1.number, uriage1.username, uriage1.okashi, uriage1.amount_sold, uriage2.number, uriage2.username, uriage2.okashi, uriage2.amount_sold
GPU Full Outer Join Quals [1]: (uriage1.number = uriage2.number) ... [plan: 10000120 -> 10000120, exec: 10879399 -> 10879399]
GPU Outer Hash [1]: uriage1.number
GPU Inner Hash [1]: uriage2.number
GPU-Direct SQL: enabled (GPU-0; direct=63695, ntuples=10879399)
-> Seq Scan on uriage2 (cost=0.00..81847.92 rows=4999992 width=20) (actual time=0.010..497.704 rows=5000000 loops=1)
Planning Time: 1.598 ms
Execution Time: 5031.804 ms
(9 rows)
kaigai commented
これ、意外に根の深い問題で頭抱えてます。
PostgreSQL v16から、Varの定義が変更になってvarnullingrels というフィールドが追加されています。
このフィールドの追加によって、列（Var）の同一性をチェックするequal()関数の挙動が変わってしまい、
必要な列定義が内部のリストに追加されなくなってしまったと。
ちょっと他の問題を先に片づけます。。。
typedef struct Var
{
Expr xpr;
/*
* index of this var's relation in the range table, or
* INNER_VAR/OUTER_VAR/etc
*/
int varno;
/*
* attribute number of this var, or zero for all attrs ("whole-row Var")
*/
AttrNumber varattno;
/* pg_type OID for the type of this var */
Oid vartype pg_node_attr(query_jumble_ignore);
/* pg_attribute typmod value */
int32 vartypmod pg_node_attr(query_jumble_ignore);
/* OID of collation, or InvalidOid if none */
Oid varcollid pg_node_attr(query_jumble_ignore);
/*
* RT indexes of outer joins that can replace the Var's value with null.
* We can omit varnullingrels in the query jumble, because it's fully
* determined by varno/varlevelsup plus the Var's query location.
*/
Bitmapset *varnullingrels pg_node_attr(query_jumble_ignore);
/*
* for subquery variables referencing outer relations; 0 in a normal var,
* >0 means N levels up
*/
Index varlevelsup;
/*
* varnosyn/varattnosyn are ignored for equality, because Vars with
* different syntactic identifiers are semantically the same as long as
* their varno/varattno match.
*/
/* syntactic relation index (0 if unknown) */
Index varnosyn pg_node_attr(equal_ignore, query_jumble_ignore);
/* syntactic attribute number */
AttrNumber varattnosyn pg_node_attr(equal_ignore, query_jumble_ignore);
/* token location, or -1 if unknown */
int location;
} Var;
kaigai commented
da184e4232859cfc7d1bf3ca9cfcf44856a74fe5 で修正できたハズですが、ちょっと影響範囲が広いので要確認。
大西さんの方で調整してもらっている regression test の暫定完成部分だけでももう一度確認が必要ですね。