search

heterodb / pg-strom

PG-Strom - Master development repository
http://heterodb.github.io/pg-strom/
Other
1.29k stars 161 forks source link

Over-usage of stack frame in recursive calls #818

Open kaigai opened 2 months ago

kaigai commented 2 months ago

812の問題を調査中、再帰呼出しのやりすぎでstack frameの使い果たしが発生。

再帰呼出しの場合にスタック使用量をチェックする仕組みが必要。 LIKE句でも同様。

CUDA Exception: Warp Out-of-range Address
The exception was triggered at PC 0x7f62137d8190 (xpu_postgis.cu:73)
[Current focus set to CUDA kernel 0, grid 1207, block (21,0,0), thread (1,0,0), device 0, sm 42, warp 0, lane 1]
#0  0x00007f62137d8200 in _INTERNAL_61659473_14_xpu_postgis_cu_f2ffba3a::setup_geometry_rawsize ()
    at /home/kaigai/pg-strom/src/xpu_postgis.cu:75
75              switch (geom->type)
(cuda-gdb) bt
#0  0x00007f62137d8200 in _INTERNAL_61659473_14_xpu_postgis_cu_f2ffba3a::setup_geometry_rawsize ()
    at /home/kaigai/pg-strom/src/xpu_postgis.cu:75
#1  0x00007f6213b162e0 in _INTERNAL_61659473_14_xpu_postgis_cu_f2ffba3a::__geom_relate_seg_polygon ()
    at /home/kaigai/pg-strom/src/xpu_postgis.cu:165 in _ZN45_INTERNAL_61659473_14_xpu_postgis_cu_f2ffba3a21geometry_load_subitemEP14xpu_geometry_tPKS0_PKciP12kern_context inlined from xpu_postgis.cu:4937
#2  0x00007f6213b18530 in _INTERNAL_61659473_14_xpu_postgis_cu_f2ffba3a::__geom_relate_seg_polygon ()
    at /home/kaigai/pg-strom/src/xpu_postgis.cu:5249
#3  0x00007f6213b18530 in _INTERNAL_61659473_14_xpu_postgis_cu_f2ffba3a::__geom_relate_seg_polygon ()
    at /home/kaigai/pg-strom/src/xpu_postgis.cu:5249
#4  0x00007f6213b18530 in _INTERNAL_61659473_14_xpu_postgis_cu_f2ffba3a::__geom_relate_seg_polygon ()
    at /home/kaigai/pg-strom/src/xpu_postgis.cu:5249
#5  0x00007f6213cf4dc0 in pgfn_st_crosses ()
    at /home/kaigai/pg-strom/src/xpu_postgis.cu:5249 in _ZN45_INTERNAL_61659473_14_xpu_postgis_cu_f2ffba3a25__geom_relate_seg_polygonEP12kern_contextRK7POINT2DbS4_bPK14xpu_geometry_tib inlined from xpu_postgis.cu:5249
#6  0x00007f6213b73c00 in _INTERNAL_ffb4507b_13_xpu_common_cu___kcxt::pgfn_JoinQuals ()
    at /home/kaigai/pg-strom/src/xpu_common.cu:1184
#7  0x00007f6213bb99a0 in ExecGiSTIndexPostQuals () at /home/kaigai/pg-strom/src/xpu_common.cu:1939
#8  0x00007f6213b940a0 in kern_gpujoin_main<<<(108,1,1),(512,1,1)>>> ()
    at /home/kaigai/pg-strom/src/cuda_gpujoin.cu:473 in _ZN46_INTERNAL_20381b8c_15_cuda_gpujoin_cu_5fc7014219execGpuJoinGiSTJoinEP12kern_contextP17kern_warp_contextP14kern_multirelsiPcS6_PK15kern_expressionS6_RmRb inlined from cuda_gpujoin.cu:915
kaigai commented 2 months ago

https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#stack-manipulation-instructions

stacksave 命令で現在のスタックポインタを取得する事ができる。 初期値 0x0100 0000から段々と減っていく感じだが、GPU間での差異は？

kaigai commented 2 months ago

fd4c07d7c36dc190d9ab99b3d56d5d3b2bcd9405 でCUDA Stack Frameの利用チェックを導入。しかし、PostGISのロジックの中にエラーとbool値判定で区別していない部分が少なからず残っているので、必ずしも fallback に落ちてくれるわけではない。