diff --git a/bench/mo_exact_decimal.cpp b/bench/mo_exact_decimal.cpp index a9df58f6e6..b5672f721d 100644 --- a/bench/mo_exact_decimal.cpp +++ b/bench/mo_exact_decimal.cpp @@ -1,9 +1,11 @@ /* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ #include "expression/ast/node.hpp" #include "expression_evaluator/expression_evaluator.hpp" +#include "numeric/decimal_aggregate_layout.hpp" #include "numeric/exact_decimal_gpu.hpp" #include +#include #include #include @@ -79,6 +81,16 @@ int main() auto b = input(narrow, stream.view(), mr); auto x = input(wide, stream.view(), mr); auto y = input(wide, stream.view(), mr); + // Ramp idle-device clocks before the fixed per-case warm-ups. Keep this + // bounded and identical for every run; do not change device clock settings. + auto warm_until = std::chrono::steady_clock::now() + std::chrono::seconds(3); + do { + auto ordinary = cudf::binary_operation( + a->view(), b->view(), cudf::binary_operator::ADD, a->type(), stream.view(), mr); + auto checked = evaluate_decimal_columns( + decimal_op::divide, x->view(), wide, y->view(), wide, wide, nullptr, stream.view(), mr); + stream.synchronize(); + } while (std::chrono::steady_clock::now() < warm_until); std::cout << "case,rows,repetition,wall_ms,rows_per_second\n"; for (auto operation : {decimal_op::add, decimal_op::multiply}) { auto out = operation == decimal_op::add ? narrow : decimal_type{128, 38, 0}; @@ -138,4 +150,32 @@ int main() return evaluator.evaluate(cudf::table_view{{a->view(), b->view()}}); }); } + // At this fixed row bound, even the entire signed physical input domain + // fits the declared 128-bit result. Widen before the ordinary reduction so + // both paths include safe accumulation, allocation and stream completion. + measure("cudf_sum64_to128", stream.view(), [&] { + auto widened = + cudf::cast(a->view(), cudf::data_type{cudf::type_id::DECIMAL128, 0}, stream.view(), mr); + auto result = cudf::reduce(widened->view(), + *cudf::make_sum_aggregation(), + widened->type(), + stream.view(), + mr); + stream.synchronize(); + return result; + }); + duckdb::vector> aggregates; + std::vector> args; + args.push_back(std::make_unique( + sirius::ast::reference{0, sirius::logical_type::make_mo_decimal({64, 15, 0}, false)})); + aggregates.push_back(std::make_unique( + sirius::ast::aggregate{sirius::aggregate_id::mo_decimal_sum, + std::move(args), + sirius::logical_type::make_mo_decimal({128, 37, 0}, true), + false})); + auto layout = make_aggregate_layout({}, aggregates); + measure("exact_sum64_to128", stream.view(), [&] { + auto state = local_aggregate_table(cudf::table_view{{a->view()}}, *layout, stream.view(), mr); + return merge_aggregate_tables({state->view()}, *layout, stream.view(), mr); + }); } diff --git a/cmake/sirius-embedding-tests.cmake b/cmake/sirius-embedding-tests.cmake index 7cb4f6105e..bf57d75721 100644 --- a/cmake/sirius-embedding-tests.cmake +++ b/cmake/sirius-embedding-tests.cmake @@ -1,8 +1,10 @@ # Embedding control tests need neither a GPU nor a DuckDB runtime. Keep them # independently runnable from the linked ABI and GPU integration tests. add_executable( - sirius_exact_decimal_unittest test/cpp/exec/execution_evidence_unittest.cpp - test/cpp/embedding/test_exact_decimal.cpp) + sirius_exact_decimal_unittest + test/cpp/exec/execution_evidence_unittest.cpp + test/cpp/embedding/test_exact_decimal.cpp + test/cpp/embedding/test_decimal_aggregate.cpp) target_compile_features(sirius_exact_decimal_unittest PRIVATE cxx_std_20) target_include_directories( sirius_exact_decimal_unittest PRIVATE "${CMAKE_SOURCE_DIR}/third_party/catch" @@ -45,6 +47,17 @@ target_include_directories( "${SIRIUS_SUBSTRAIT_DIR}/third_party/substrait") target_link_libraries(sirius_native_result_integration PRIVATE Sirius::embed Threads::Threads) +add_executable( + sirius_native_numeric_unittest test/cpp/exec/execution_evidence_unittest.cpp + test/cpp/embedding/test_native_numeric.cpp) +target_compile_features(sirius_native_numeric_unittest PRIVATE cxx_std_20) +target_include_directories( + sirius_native_numeric_unittest + PRIVATE "${CMAKE_SOURCE_DIR}/third_party/catch" + "${SIRIUS_SUBSTRAIT_DIR}/third_party" + "${SIRIUS_SUBSTRAIT_DIR}/third_party/substrait") +target_link_libraries(sirius_native_numeric_unittest PRIVATE Sirius::embed + Threads::Threads) add_executable( sirius_native_binding_unittest @@ -76,6 +89,7 @@ add_executable( test/cpp/embedding/test_native_result_codec.cpp test/cpp/embedding/test_exact_decimal_gpu.cpp test/cpp/embedding/test_decimal_expression_gpu.cpp + test/cpp/embedding/test_decimal_aggregate_gpu.cpp test/cpp/embedding/test_tae_gpu.cpp src/embedding/c_api.cpp) target_compile_features(sirius_native_gpu_unittest PRIVATE cxx_std_20) diff --git a/docs/mo-exact-decimal-import.md b/docs/mo-exact-decimal-import.md index c739a59b55..b4d5148660 100644 --- a/docs/mo-exact-decimal-import.md +++ b/docs/mo-exact-decimal-import.md @@ -6,6 +6,10 @@ This is the first of the two native integration increments approved for merged Sirius #25 and importer #4. The semantic contract remains the approved MO #29449 document blob `42a89f09a1d168d02b9583cb3ea7b4de6dbb5634`. +This document records the Native A boundary and its historical measurements. +[Native B](mo-exact-decimal-integration.md) supersedes the exposure boundary +below with complete embedded preparation, aggregates, keys and capability 16u. + ## Contract and ownership Physical coefficient width, declared precision, scale, and nullability survive diff --git a/docs/mo-exact-decimal-integration-kernels.csv b/docs/mo-exact-decimal-integration-kernels.csv new file mode 100644 index 0000000000..52b1d5b4f0 --- /dev/null +++ b/docs/mo-exact-decimal-integration-kernels.csv @@ -0,0 +1,211 @@ +run,case,rows,repetition,wall_ms,rows_per_second +1,cudf_add64,262144,0,0.20497,1.27894e+09 +1,cudf_add64,262144,1,0.188959,1.38731e+09 +1,cudf_add64,262144,2,0.193369,1.35567e+09 +1,cudf_add64,262144,3,0.1946,1.34709e+09 +1,cudf_add64,262144,4,0.194839,1.34544e+09 +1,cudf_add64,262144,5,0.199599,1.31335e+09 +1,cudf_add64,262144,6,0.197809,1.32524e+09 +1,checked_add64,262144,0,0.02917,8.98677e+09 +1,checked_add64,262144,1,0.03414,7.6785e+09 +1,checked_add64,262144,2,0.02703,9.69826e+09 +1,checked_add64,262144,3,0.03307,7.92694e+09 +1,checked_add64,262144,4,0.03386,7.742e+09 +1,checked_add64,262144,5,0.0338,7.75574e+09 +1,checked_add64,262144,6,0.02713,9.66251e+09 +1,cudf_widen_multiply64_to128,262144,0,0.501388,5.22837e+08 +1,cudf_widen_multiply64_to128,262144,1,0.494758,5.29843e+08 +1,cudf_widen_multiply64_to128,262144,2,0.496258,5.28241e+08 +1,cudf_widen_multiply64_to128,262144,3,0.496198,5.28305e+08 +1,cudf_widen_multiply64_to128,262144,4,0.495969,5.28549e+08 +1,cudf_widen_multiply64_to128,262144,5,0.505178,5.18914e+08 +1,cudf_widen_multiply64_to128,262144,6,0.495688,5.28849e+08 +1,checked_multiply64_to128,262144,0,0.502158,5.22035e+08 +1,checked_multiply64_to128,262144,1,0.498948,5.25393e+08 +1,checked_multiply64_to128,262144,2,0.504449,5.19664e+08 +1,checked_multiply64_to128,262144,3,0.494338,5.30293e+08 +1,checked_multiply64_to128,262144,4,0.492638,5.32123e+08 +1,checked_multiply64_to128,262144,5,0.494428,5.30197e+08 +1,checked_multiply64_to128,262144,6,0.493968,5.3069e+08 +1,checked_add256,262144,0,0.267809,9.78847e+08 +1,checked_add256,262144,1,0.261119,1.00393e+09 +1,checked_add256,262144,2,0.260409,1.00666e+09 +1,checked_add256,262144,3,0.258879,1.01261e+09 +1,checked_add256,262144,4,0.259119,1.01167e+09 +1,checked_add256,262144,5,0.266299,9.84397e+08 +1,checked_add256,262144,6,0.265619,9.86917e+08 +1,checked_divide256,262144,0,4.72385,5.54937e+07 +1,checked_divide256,262144,1,4.72027,5.55358e+07 +1,checked_divide256,262144,2,4.71809,5.55614e+07 +1,checked_divide256,262144,3,4.72171,5.55188e+07 +1,checked_divide256,262144,4,4.72278,5.55063e+07 +1,checked_divide256,262144,5,4.71955,5.55442e+07 +1,checked_divide256,262144,6,4.71993,5.55398e+07 +1,ordinary_expression_add64,262144,0,0.199669,1.31289e+09 +1,ordinary_expression_add64,262144,1,0.196089,1.33686e+09 +1,ordinary_expression_add64,262144,2,0.20103,1.304e+09 +1,ordinary_expression_add64,262144,3,0.196079,1.33693e+09 +1,ordinary_expression_add64,262144,4,0.195619,1.34007e+09 +1,ordinary_expression_add64,262144,5,0.196169,1.33632e+09 +1,ordinary_expression_add64,262144,6,0.19765,1.3263e+09 +1,exact_expression_add64,262144,0,0.16399,1.59854e+09 +1,exact_expression_add64,262144,1,0.161689,1.62129e+09 +1,exact_expression_add64,262144,2,0.160069,1.63769e+09 +1,exact_expression_add64,262144,3,0.16064,1.63187e+09 +1,exact_expression_add64,262144,4,0.163799,1.6004e+09 +1,exact_expression_add64,262144,5,0.15962,1.6423e+09 +1,exact_expression_add64,262144,6,0.159089,1.64778e+09 +1,cudf_sum64_to128,262144,0,0.285259,9.18968e+08 +1,cudf_sum64_to128,262144,1,0.270139,9.70404e+08 +1,cudf_sum64_to128,262144,2,0.276389,9.4846e+08 +1,cudf_sum64_to128,262144,3,0.271249,9.66433e+08 +1,cudf_sum64_to128,262144,4,0.275289,9.5225e+08 +1,cudf_sum64_to128,262144,5,0.272939,9.60449e+08 +1,cudf_sum64_to128,262144,6,0.270939,9.67539e+08 +1,exact_sum64_to128,262144,0,0.221459,1.18371e+09 +1,exact_sum64_to128,262144,1,0.222239,1.17956e+09 +1,exact_sum64_to128,262144,2,0.21954,1.19406e+09 +1,exact_sum64_to128,262144,3,0.218509,1.19969e+09 +1,exact_sum64_to128,262144,4,0.218459,1.19997e+09 +1,exact_sum64_to128,262144,5,0.218719,1.19854e+09 +1,exact_sum64_to128,262144,6,0.216399,1.21139e+09 +2,cudf_add64,262144,0,0.204999,1.27876e+09 +2,cudf_add64,262144,1,0.186999,1.40185e+09 +2,cudf_add64,262144,2,0.194959,1.34461e+09 +2,cudf_add64,262144,3,0.19545,1.34123e+09 +2,cudf_add64,262144,4,0.194719,1.34627e+09 +2,cudf_add64,262144,5,0.198849,1.31831e+09 +2,cudf_add64,262144,6,0.19526,1.34254e+09 +2,checked_add64,262144,0,0.02703,9.69826e+09 +2,checked_add64,262144,1,0.02729,9.60586e+09 +2,checked_add64,262144,2,0.03485,7.52207e+09 +2,checked_add64,262144,3,0.02695,9.72705e+09 +2,checked_add64,262144,4,0.02753,9.52212e+09 +2,checked_add64,262144,5,0.034279,7.64736e+09 +2,checked_add64,262144,6,0.03381,7.75345e+09 +2,cudf_widen_multiply64_to128,262144,0,0.509678,5.14333e+08 +2,cudf_widen_multiply64_to128,262144,1,0.496268,5.28231e+08 +2,cudf_widen_multiply64_to128,262144,2,0.500418,5.2385e+08 +2,cudf_widen_multiply64_to128,262144,3,0.501069,5.23169e+08 +2,cudf_widen_multiply64_to128,262144,4,0.504228,5.19892e+08 +2,cudf_widen_multiply64_to128,262144,5,0.499158,5.25172e+08 +2,cudf_widen_multiply64_to128,262144,6,0.503338,5.20811e+08 +2,checked_multiply64_to128,262144,0,0.499178,5.25151e+08 +2,checked_multiply64_to128,262144,1,0.511978,5.12022e+08 +2,checked_multiply64_to128,262144,2,0.497248,5.2719e+08 +2,checked_multiply64_to128,262144,3,0.497919,5.26479e+08 +2,checked_multiply64_to128,262144,4,0.498268,5.2611e+08 +2,checked_multiply64_to128,262144,5,0.497968,5.26427e+08 +2,checked_multiply64_to128,262144,6,0.494468,5.30154e+08 +2,checked_add256,262144,0,0.258559,1.01387e+09 +2,checked_add256,262144,1,0.264199,9.92222e+08 +2,checked_add256,262144,2,0.264149,9.9241e+08 +2,checked_add256,262144,3,0.268029,9.78043e+08 +2,checked_add256,262144,4,0.266119,9.85063e+08 +2,checked_add256,262144,5,0.258619,1.01363e+09 +2,checked_add256,262144,6,0.264479,9.91171e+08 +2,checked_divide256,262144,0,4.7162,5.55837e+07 +2,checked_divide256,262144,1,4.7215,5.55213e+07 +2,checked_divide256,262144,2,4.71662,5.55787e+07 +2,checked_divide256,262144,3,4.71515,5.55961e+07 +2,checked_divide256,262144,4,4.71569,5.55897e+07 +2,checked_divide256,262144,5,4.71772,5.55658e+07 +2,checked_divide256,262144,6,4.71491,5.55989e+07 +2,ordinary_expression_add64,262144,0,0.197489,1.32739e+09 +2,ordinary_expression_add64,262144,1,0.195739,1.33925e+09 +2,ordinary_expression_add64,262144,2,0.197019,1.33055e+09 +2,ordinary_expression_add64,262144,3,0.20012,1.30993e+09 +2,ordinary_expression_add64,262144,4,0.200429,1.30791e+09 +2,ordinary_expression_add64,262144,5,0.198379,1.32143e+09 +2,ordinary_expression_add64,262144,6,0.197039,1.33042e+09 +2,exact_expression_add64,262144,0,0.1585,1.65391e+09 +2,exact_expression_add64,262144,1,0.159569,1.64283e+09 +2,exact_expression_add64,262144,2,0.15803,1.65882e+09 +2,exact_expression_add64,262144,3,0.168509,1.55567e+09 +2,exact_expression_add64,262144,4,0.159799,1.64046e+09 +2,exact_expression_add64,262144,5,0.15778,1.66145e+09 +2,exact_expression_add64,262144,6,0.163029,1.60796e+09 +2,cudf_sum64_to128,262144,0,0.279139,9.39116e+08 +2,cudf_sum64_to128,262144,1,0.272269,9.62813e+08 +2,cudf_sum64_to128,262144,2,0.284509,9.21391e+08 +2,cudf_sum64_to128,262144,3,0.274519,9.54921e+08 +2,cudf_sum64_to128,262144,4,0.273739,9.57642e+08 +2,cudf_sum64_to128,262144,5,0.273149,9.59711e+08 +2,cudf_sum64_to128,262144,6,0.276819,9.46987e+08 +2,exact_sum64_to128,262144,0,0.217059,1.20771e+09 +2,exact_sum64_to128,262144,1,0.224169,1.1694e+09 +2,exact_sum64_to128,262144,2,0.218149,1.20167e+09 +2,exact_sum64_to128,262144,3,0.218899,1.19756e+09 +2,exact_sum64_to128,262144,4,0.21724,1.2067e+09 +2,exact_sum64_to128,262144,5,0.217039,1.20782e+09 +2,exact_sum64_to128,262144,6,0.217799,1.20361e+09 +3,cudf_add64,262144,0,0.20639,1.27014e+09 +3,cudf_add64,262144,1,0.191469,1.36912e+09 +3,cudf_add64,262144,2,0.198879,1.31811e+09 +3,cudf_add64,262144,3,0.201479,1.3011e+09 +3,cudf_add64,262144,4,0.19953,1.31381e+09 +3,cudf_add64,262144,5,0.204569,1.28145e+09 +3,cudf_add64,262144,6,0.199409,1.3146e+09 +3,checked_add64,262144,0,0.02819,9.29918e+09 +3,checked_add64,262144,1,0.034299,7.64291e+09 +3,checked_add64,262144,2,0.02779,9.43303e+09 +3,checked_add64,262144,3,0.03435,7.63156e+09 +3,checked_add64,262144,4,0.03699,7.08689e+09 +3,checked_add64,262144,5,0.02679,9.78514e+09 +3,checked_add64,262144,6,0.03394,7.72375e+09 +3,cudf_widen_multiply64_to128,262144,0,0.509878,5.14131e+08 +3,cudf_widen_multiply64_to128,262144,1,0.509238,5.14777e+08 +3,cudf_widen_multiply64_to128,262144,2,0.512789,5.11212e+08 +3,cudf_widen_multiply64_to128,262144,3,0.508449,5.15576e+08 +3,cudf_widen_multiply64_to128,262144,4,0.508378,5.15648e+08 +3,cudf_widen_multiply64_to128,262144,5,0.508818,5.15202e+08 +3,cudf_widen_multiply64_to128,262144,6,0.511298,5.12703e+08 +3,checked_multiply64_to128,262144,0,0.507528,5.16511e+08 +3,checked_multiply64_to128,262144,1,0.506148,5.1792e+08 +3,checked_multiply64_to128,262144,2,0.514158,5.09851e+08 +3,checked_multiply64_to128,262144,3,0.506589,5.17469e+08 +3,checked_multiply64_to128,262144,4,0.512928,5.11074e+08 +3,checked_multiply64_to128,262144,5,0.505498,5.18586e+08 +3,checked_multiply64_to128,262144,6,0.507998,5.16034e+08 +3,checked_add256,262144,0,0.264629,9.90609e+08 +3,checked_add256,262144,1,0.273349,9.59008e+08 +3,checked_add256,262144,2,0.267969,9.78262e+08 +3,checked_add256,262144,3,0.261749,1.00151e+09 +3,checked_add256,262144,4,0.266089,9.85174e+08 +3,checked_add256,262144,5,0.264099,9.92597e+08 +3,checked_add256,262144,6,0.262539,9.98495e+08 +3,checked_divide256,262144,0,4.72434,5.54879e+07 +3,checked_divide256,262144,1,4.72097,5.55275e+07 +3,checked_divide256,262144,2,4.72129,5.55238e+07 +3,checked_divide256,262144,3,4.72736,5.54525e+07 +3,checked_divide256,262144,4,4.72303,5.55033e+07 +3,checked_divide256,262144,5,4.7266,5.54614e+07 +3,checked_divide256,262144,6,4.72237,5.55111e+07 +3,ordinary_expression_add64,262144,0,0.202799,1.29263e+09 +3,ordinary_expression_add64,262144,1,0.20388,1.28578e+09 +3,ordinary_expression_add64,262144,2,0.201169,1.3031e+09 +3,ordinary_expression_add64,262144,3,0.200309,1.3087e+09 +3,ordinary_expression_add64,262144,4,0.200629,1.30661e+09 +3,ordinary_expression_add64,262144,5,0.2,1.31072e+09 +3,ordinary_expression_add64,262144,6,0.203589,1.28761e+09 +3,exact_expression_add64,262144,0,0.165029,1.58847e+09 +3,exact_expression_add64,262144,1,0.16826,1.55797e+09 +3,exact_expression_add64,262144,2,0.172959,1.51564e+09 +3,exact_expression_add64,262144,3,0.172239,1.52198e+09 +3,exact_expression_add64,262144,4,0.1656,1.583e+09 +3,exact_expression_add64,262144,5,0.165249,1.58636e+09 +3,exact_expression_add64,262144,6,0.164699,1.59166e+09 +3,cudf_sum64_to128,262144,0,0.282999,9.26307e+08 +3,cudf_sum64_to128,262144,1,0.279929,9.36466e+08 +3,cudf_sum64_to128,262144,2,0.284129,9.22623e+08 +3,cudf_sum64_to128,262144,3,0.282809,9.26929e+08 +3,cudf_sum64_to128,262144,4,0.283009,9.26274e+08 +3,cudf_sum64_to128,262144,5,0.279079,9.39318e+08 +3,cudf_sum64_to128,262144,6,0.283039,9.26176e+08 +3,exact_sum64_to128,262144,0,0.223419,1.17333e+09 +3,exact_sum64_to128,262144,1,0.232619,1.12692e+09 +3,exact_sum64_to128,262144,2,0.224789,1.16618e+09 +3,exact_sum64_to128,262144,3,0.222589,1.1777e+09 +3,exact_sum64_to128,262144,4,0.22315,1.17474e+09 +3,exact_sum64_to128,262144,5,0.226679,1.15645e+09 +3,exact_sum64_to128,262144,6,0.231379,1.13296e+09 diff --git a/docs/mo-exact-decimal-integration.md b/docs/mo-exact-decimal-integration.md new file mode 100644 index 0000000000..2092659094 --- /dev/null +++ b/docs/mo-exact-decimal-integration.md @@ -0,0 +1,206 @@ +# Embedded MO exact-decimal native completion + +Native B follows merged Sirius #26 (`c440013c89f1a54a6e5d934aabd0f9d3ef15ed3c`) +and implements the owner-approved next increment for MO #28968 / #28966. +The semantic authority remains MO #29449 document blob +`42a89f09a1d168d02b9583cb3ea7b4de6dbb5634`; MO #29690 records the B-F +delivery sequence. The detailed implementation plan was approved before code. + +## Contracts and owners + +The validated import descriptor owns width, precision, scale and nullability. +Private aggregate bindings preserve that immutable contract through cloning +and reconstruction. Unsupported aggregate modifiers fail preparation. + +Local and merge aggregation share one typed layout. SUM/AVG partial states +use the full selected signed 128/256-bit physical range, not a scalar decimal +precision limit. AVG retains a non-NULL count and rounds only at finalization. +Empty/all-NULL groups produce NULL. Every public result is finalized, including +a single partial batch. MIN/MAX preserve their operand type and ignore NULLs. +Ordinary and exact aggregate slots share group ordinals. + +A single local partial goes directly to finalization. Multiple partials are +concatenated, regrouped and reduced before the same finalizer. Initial +coefficients can skip aggregate error polling when the accumulator is at least +32 bits wider than the full signed input domain: at most INT32_MAX input rows +fit every intermediate subset. MIN/MAX have the same infallible initial +reduction property. Same-width wide states and all merges retain checked error +polling. Final declared precision checks are never skipped. This proof does not +depend on sampled coefficient values. + +Grouped tile bounds are computed directly from sorted row offsets, with at +most one identity tile added per segment. Prefix indices use 64 bits. This +avoids a CUB scan, duplicate CUDA template-kernel registrations and a host +tile-count transfer; nonzero-origin subranges and tile-boundary transitions +have independent value/count controls. + +Partial states remain ordinary column-backed cuCascade batches. Existing task +reservations own coefficients, counts, grouping/permutation/reduction scratch, +and output buffers; existing partition/spill/restore owns retained batches. +There is no separate aggregate registry, cache or spill service. + +Equality partitioning and hash probing use the same exact key encoding. +Mixed scales normalize trailing decimal zeros and zero itself; ordering uses +the signed coefficient at the column's fixed scale. Original columns retain +their declared schema. Canonical Decimal256 validity is restored only after +operations whose canonical inputs prove child masks redundant. + +The embedded preparation path installs scoped import only for exact plans, +validates the complete reachable graph and schema, then permits readers to +start. Ordinary import and standalone execution retain their existing path. +Read filters are checked against the base schema before flat read projections +select their output schema. CASE/COALESCE producer descriptors and outer-join +nullability survive binding through private checked identity casts. Alias +inspection constructs a fresh carrier; DuckDB's shared type metadata must not +be mutated while inspecting an immutable function binding. + +The exact embedded path keeps MO's optimized logical plan and skips DuckDB's +ordinary optimizer, whose coefficient-carrier folding cannot interpret MO +scales. It reconstructs typed equality join conditions for the native planner. +Schema propagation is charged against the existing query metadata admission; +original and normalized serialized plans remain bounded at 16 MiB. + +Runtime numeric errors retain the failing operation's owner through the task +and session boundary, including producer cancellation during cleanup. Fatal +quiescence failures retain their existing precedence. ABI-v1 layouts stay +unchanged; statuses 12/13 are appended and capability mask 16u is enabled. The +complete current capability mask is 31u. + +Unsupported exact-plan shapes fail preparation: raw numeric casts, ordinary +arithmetic/comparison over exact operands, exact IN/nested tuples, non-equality +decimal join keys, grouping sets, aggregate filters/options/order/intermediate +phases, best-effort read filters and zero-column intermediate schemas. Exact aggregates accept ALL; ordinary +grouped COUNT DISTINCT retains the existing set-state representation. Public +Decimal256 results remain limited to precision 65. These restrictions must be +respected by the MO lowering increment. + +## Change and validation map + +| Closure | Risk | Required proof | +| --- | --- | --- | +| Raw accumulators and aggregate finalization | R2 | Independent CPU vectors, physical versus declared bounds, cancellation, signed rounding, NULL and owner-specific errors | +| Aggregate import/binding/clone and preparation | R2 | Real importer and planner, malformed identities/modifiers, full graph rejection before reader work, unchanged ordinary IDs and C layouts | +| GPU local/merge aggregation and spill states | R3 | Actual values across widths, mixed slots, batch/merge permutations, single/empty batches, spill/restore and bounded cleanup | +| Keys, predicates and Decimal256 validity | R3 | Cross-width/scale partition/hash equality, NULL joins, signed ordering and following exact operations after gather | +| Runtime exposure and errors | R3 | Real C ABI queries, statuses 12/13 by operation, cancellation/async failure, quiescence and healthy reuse/fatal quarantine | +| SDK/build/docs | R0/R1 | Frozen native build, C consumer/schema fingerprint, pinned hooks and exact artifact provenance | + +New GPU owners are constructed before submission, retain operands and scratch, +and prove task-stream quiescence before destruction. Unprovable owners transfer +to the existing fatal quarantine. Allocation, cancellation and partial failure +must leave one cleanup owner. All scratch and partial-state growth is charged +before allocation and released or spilled through existing resource ownership. +Host buffers used by asynchronous copies belong to that same owner and survive +exception unwinding until synchronization or fatal quarantine. + +The ownership and wait edges are: + +| Resource | Transfer and terminal owner | Failure/progress bound | +| --- | --- | --- | +| Typed binding/layout | Query import → immutable function data → physical local/merge operators | Query metadata admission and plan lifetime; no shared alias mutation | +| Coefficients/counts/tile scratch | Task resource → reducer owner → returned columns → cuCascade batch | Same task stream is synchronized on success/error; fatal quarantine retains unprovable owners | +| Borrowed operands/partial tables | Read-only task batches → reducer/concat/merge views | Caller retains inputs until task completion; concatenated input survives all queued consumers | +| Retained aggregate batches | Local output → partition queue → existing host/disk spill → restore → merge | Existing reservation/backpressure and batch retirement; no separate retained-state registry | +| Exact join keys | Borrowed input → normalized key columns → partition/probe owners | Identical encoding on both sides; task-stream ordering retains keys through consumers | +| Public output | Finalized columns → admitted codec scratch/output → native result lease → batch release | Existing bounded result credit and query close; numeric failure survives producer cleanup | + +The new helpers introduce no locks, independent worker queues or retry loops. +Cancellation continues through the existing task/control owners. Stream waits +must complete or make the runtime fatal; a failed synchronization never +authorizes buffer destruction or healthy reuse. Growth is bounded by admitted +schemas, cuDF row counts and the caller's resource. The allocation-cutpoint +matrix checks both proven/checked reductions and single/multiple segments, +with zero live charged bytes after each success or injected failure. + +Decimal256 publication verifies child validity against the parent using the +already admitted codec scratch slab and bounded stack bitmap chunks. A NULL +child under a valid parent is rejected; redundant child masks beneath NULL +parents preserve the parent NULL. No additional GPU allocation is introduced +by result validation. + +## Delivery boundary + +MO lowering remains opt-in work in C after this increment merges and is pinned. +D owns the real native-MO/embedded-MO all-22 SF1/SF10 public, resource and timing +campaign. E owns strict legacy Flight configuration/recovery cutover; F waits +for E and its verified release artifact. Neither native tests nor capability +advertisement establish migration completion. + +## Local terminal evidence + +Validated on 2026-10-08 in the frozen `mo` Pixi environment: release `-O3 +-DNDEBUG`, CUDA 13.3.73, GCC 14.4.0, cuDF/RMM 26.08, RTX 3070 and driver +615.71.09. No Go production code, durable user format or C ABI-v1 struct layout +changes in this increment. The importer pin remains +`95d9ce8d78490db3991ab6145653716aa3ec42c9`. + +| Check | Terminal result | +| --- | --- | +| CPU exact arithmetic and raw aggregate boundaries | 52,645 assertions / 8 cases passed | +| Binding, complete-plan preparation and ordinary type/function controls | 831 assertions / 80 cases passed | +| Native cancellation, ownership and numeric error controls | 606 assertions / 78 cases passed | +| Complete affected GPU group, including native input/admission/result controls | 32,349 assertions / 30 cases passed | +| Production C ABI numeric queries | 1,468 assertions / 8 cases passed under memcheck; zero sanitizer errors | +| Numeric kernels/expressions/aggregates and result codecs under memcheck | 21,294 assertions / 26 cases passed; zero sanitizer errors | +| SDK exporter | 8 tests passed | +| Independent C consumer, ordinary GPU result/release/shutdown/reuse | Passed | + +The selected native targets built successfully. Pinned hooks cover every +changed file. The build-tree SDK exports the unchanged literal schema and +updated header, with source revision and dependency/artifact fingerprints; +the delivery export is regenerated from the clean committed source. + +The existing nonnumeric `[native_gpu]` control selection also passed 11,034 +assertions / 3 cases, but its full leak-check invocation is **not clean**: +it reports a 41,386,248-byte cuDF default pinned-pool allocation and a generic +CUB `EmptyKernel` registration diagnostic in the unchanged input-converter +path. The new aggregate CUDA object contains no CUB scan/empty-kernel symbols. +The two scan-registration diagnostics initially introduced by the aggregate +scan were removed by direct tile-prefix construction. This is not full leak +or resource acceptance; D must reconcile process baselines and complete the +required lifecycle/resource campaign. + +Reproduce from the Sirius source directory. `MO_SIRIUS_TEST_CONFIG` must name a +Sirius YAML configuration with GPU 0, bounded GPU/host/disk pools and at least +four scan-manager threads; the production tests fail when it is absent. + +```bash +pixi run --frozen -e mo build/mo/extension/sirius/sirius_exact_decimal_unittest +pixi run --frozen -e mo build/mo/extension/sirius/sirius_native_binding_unittest +pixi run --frozen -e mo build/mo/extension/sirius/sirius_native_control_unittest +pixi run --frozen -e mo build/mo/extension/sirius/sirius_native_gpu_unittest \ + '[decimal_import_gpu],[decimal_aggregate_gpu],[exact_decimal_gpu],[native_gpu],[native_result_gpu],[native_admission]' +pixi run --frozen -e mo compute-sanitizer --tool memcheck --error-exitcode 99 \ + build/mo/extension/sirius/sirius_native_gpu_unittest \ + '[decimal_import_gpu],[decimal_aggregate_gpu],[exact_decimal_gpu],[native_result_gpu]' +pixi run --frozen -e mo compute-sanitizer --tool memcheck --error-exitcode 99 \ + build/mo/extension/sirius/sirius_native_numeric_unittest '[native_numeric]' +``` + +## Retained kernel measurements + +The [raw repetitions](mo-exact-decimal-integration-kernels.csv) retain all three +complete benchmark invocations. Each uses a three-second mixed-operation +device warmup, two excluded warmups per case, then seven measured repetitions +at 262,144 rows. Inputs are uploaded before timing. Both routes use the same +caller-owned 64 MiB initial / 128 MiB maximum pool. Allocation, masks, error +storage and required stream completion are timed; output destruction is +excluded symmetrically. No GPU clock settings were changed. GPU measurements +ran sequentially after builds/tests completed. + +| Matched case | Ordinary control median ms | Exact median ms | Exact/control | +| --- | --- | --- | --- | +| Evaluator add64, identical values | 0.199669 | 0.163029 | 0.8165 | +| SUM64 → 128, safe full-domain accumulation | 0.276819 | 0.219540 | 0.7931 | + +The SUM control widens before cuDF reduction. The exact route includes local +state, counts and mandatory final precision/NULL publication. These are local +kernel/evaluator controls, not full SQL or resource acceptance. + +Historical primitive comparison remains separate. The unchanged checked-add64 +kernel's three run medians were 0.033070, 0.027530 and 0.033940 ms; the pooled +median is 18.15% above Native A's noncontemporaneous sample. All runs are +retained, and this difference is not declared an accepted performance tradeoff +or a passed rollout gate. D still requires fresh matched baseline/query +measurements and the approved 10% policy before rollout, together with the +required native-MO/embedded-MO SF1/SF10 campaign. diff --git a/src/embedding/c_api.cpp b/src/embedding/c_api.cpp index 93ebcebd26..e96aad8ce2 100644 --- a/src/embedding/c_api.cpp +++ b/src/embedding/c_api.cpp @@ -70,7 +70,7 @@ extern "C" uint32_t sirius_abi_version(void) { return SIRIUS_ABI_VERSION; } extern "C" uint64_t sirius_capabilities(void) { return SIRIUS_CAP_ENGINE_CONTROL | SIRIUS_CAP_MO_INPUT | SIRIUS_CAP_TAE_INPUT | - SIRIUS_CAP_NATIVE_RESULTS; + SIRIUS_CAP_NATIVE_RESULTS | SIRIUS_CAP_MO_EXACT_DECIMAL_V1; } extern "C" sirius_status sirius_engine_create(const sirius_engine_options* options, diff --git a/src/embedding/control.cpp b/src/embedding/control.cpp index 32721bb770..852092d5b1 100644 --- a/src/embedding/control.cpp +++ b/src/embedding/control.cpp @@ -7,6 +7,7 @@ #include "embedding/input.hpp" #include "embedding/result.hpp" #include "embedding/tae_demand.hpp" +#include "numeric/decimal_error.hpp" #include "pipeline/gpu_stream_quiescence_error.hpp" #include @@ -45,6 +46,10 @@ sirius_error current_error() noexcept return error(e.deadline_expired ? SIRIUS_TIMEOUT : SIRIUS_CANCELLED, e.what()); } catch (pipeline::gpu_stream_quiescence_error const& e) { return error(SIRIUS_GPU_UNAVAILABLE, e.what()); + } catch (mo_decimal::numeric_error const& e) { + return error(e.code() == mo_decimal::decimal_error::out_of_range ? SIRIUS_NUMERIC_OUT_OF_RANGE + : SIRIUS_NUMERIC_INVALID_INPUT, + e.what()); } catch (std::bad_alloc const&) { return error(SIRIUS_RESOURCE_EXHAUSTED, "native allocation failed"); } catch (std::exception const& e) { @@ -532,7 +537,9 @@ void engine_control::process(engine_backend& backend, } else { driver.reset(); auto producer_error = q->inputs->outcome(); - if (producer_error.code) result = producer_error; + if (producer_error.code && result.code != SIRIUS_NUMERIC_OUT_OF_RANGE && + result.code != SIRIUS_NUMERIC_INVALID_INPUT) + result = producer_error; q->inputs->stop(); q->inputs->discard(); } diff --git a/src/embedding/plan.cpp b/src/embedding/plan.cpp index 3dc4de8b47..e0ef615410 100644 --- a/src/embedding/plan.cpp +++ b/src/embedding/plan.cpp @@ -1,5 +1,8 @@ /* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ #include "embedding/control.hpp" +#include "numeric/decimal_aggregate_bind.hpp" +#include "numeric/decimal_functions.hpp" +#include "numeric/decimal_import.hpp" #include "substrait/algebra.pb.h" #include "substrait/plan.pb.h" @@ -13,6 +16,11 @@ namespace sirius::embedding { namespace { constexpr std::string_view prefix = "__sirius_embedded_v1"; +struct numeric_profile { + bool exact{}; + std::vector types; + std::unordered_set functions; +}; uint64_t binding_id(std::string const& text) { @@ -80,8 +88,19 @@ bool supported_aggregate_function(std::string const& name) return supported.contains(normalized_function(name)); } -bool read_type_matches(substrait::Type const& type, owned_column const& column) +bool read_type_matches(substrait::Type const& type, + owned_column const& column, + numeric_profile const& profile) { + if (profile.exact && (column.oid == 32 || column.oid == 33 || column.oid == 34)) { + if (!type.has_user_defined()) return false; + auto logical = mo_decimal::import_result_type(type, profile.types); + auto decimal = logical.mo_decimal_type(); + return decimal.bits == (column.oid == 32 ? 64 + : column.oid == 33 ? 128 + : 256) && + decimal.precision == column.width && decimal.scale == column.scale; + } using Kind = substrait::Type::KindCase; switch (column.oid) { case 10: return type.kind_case() == Kind::kBool; @@ -107,8 +126,11 @@ bool read_type_matches(substrait::Type const& type, owned_column const& column) } } -bool supported_embedded_type(substrait::Type const& type) +bool supported_embedded_type(substrait::Type const& type, numeric_profile const& profile) { + if (type.has_user_defined()) { + return profile.exact && mo_decimal::import_result_type(type, profile.types).is_mo_decimal(); + } using Kind = substrait::Type::KindCase; switch (type.kind_case()) { case Kind::kBool: @@ -126,14 +148,18 @@ bool supported_embedded_type(substrait::Type const& type) return type.kind_case() != Kind::kPrecisionTimestamp || type.precision_timestamp().precision() == 6; case Kind::kDecimal: - return type.decimal().precision() > 0 && type.decimal().precision() <= 38 && + return !profile.exact && type.decimal().precision() > 0 && type.decimal().precision() <= 38 && type.decimal().scale() >= 0 && type.decimal().scale() <= type.decimal().precision(); default: return false; } } -bool supported_literal(substrait::Expression_Literal const& literal) +bool supported_literal(substrait::Expression_Literal const& literal, numeric_profile const& profile) { + if (literal.has_user_defined()) + return profile.exact && + std::binary_search( + profile.types.begin(), profile.types.end(), literal.user_defined().type_reference()); using Literal = substrait::Expression_Literal; switch (literal.literal_type_case()) { case Literal::kNull: @@ -146,8 +172,9 @@ bool supported_literal(substrait::Expression_Literal const& literal) case Literal::kDate: case Literal::kVarChar: return true; case Literal::kDecimal: - return literal.decimal().value().size() == 16 && literal.decimal().precision() > 0 && - literal.decimal().precision() <= 38 && literal.decimal().scale() >= 0 && + return !profile.exact && literal.decimal().value().size() == 16 && + literal.decimal().precision() > 0 && literal.decimal().precision() <= 38 && + literal.decimal().scale() >= 0 && literal.decimal().scale() <= literal.decimal().precision(); default: return false; } @@ -172,17 +199,23 @@ void inspect_message(duckdb::google::protobuf::Message const& message, int relation_ordinal, query_state const& query, std::set& reads, - std::unordered_map const& functions) + std::unordered_map const& functions, + numeric_profile const& profile) { auto const name = message.GetDescriptor()->full_name(); if (name == "substrait.Type") { - if (!supported_embedded_type(static_cast(message))) + if (!supported_embedded_type(static_cast(message), profile)) throw failure(SIRIUS_UNSUPPORTED, "unsupported Substrait type in embedded plan"); } else if (name == "substrait.Expression.Literal") { - if (!supported_literal(static_cast(message))) + if (!supported_literal(static_cast(message), profile)) throw failure(SIRIUS_UNSUPPORTED, "unsupported Substrait literal in embedded plan"); } else if (name == "substrait.Rel") { auto const& rel = static_cast(message); + if (rel.has_aggregate()) + for (auto const& measure : rel.aggregate().measures()) + if (measure.has_filter() && + profile.functions.contains(measure.measure().function_reference())) + throw failure(SIRIUS_UNSUPPORTED, "MO aggregate filters are unsupported"); switch (rel.rel_type_case()) { case substrait::Rel::kRead: case substrait::Rel::kFilter: @@ -245,7 +278,11 @@ void inspect_message(duckdb::google::protobuf::Message const& message, } else if (name == "substrait.Expression.ScalarFunction") { auto const& function = static_cast(message); auto found = functions.find(function.function_reference()); - if (found == functions.end() || !supported_scalar_function(found->second)) + auto exact = found == functions.end() ? std::optional{} + : from_duckdb_function_name("__sirius_" + found->second); + bool numeric = exact && mo_decimal::is_decimal_function(*exact) && + profile.functions.contains(function.function_reference()); + if (found == functions.end() || (!numeric && !supported_scalar_function(found->second))) throw failure(SIRIUS_UNSUPPORTED, "unsupported scalar function in embedded plan"); auto const normalized = normalized_function(found->second); static const std::unordered_set extract_fields{"year", @@ -268,7 +305,11 @@ void inspect_message(duckdb::google::protobuf::Message const& message, } else if (name == "substrait.AggregateFunction") { auto const& function = static_cast(message); auto found = functions.find(function.function_reference()); - if (found == functions.end() || !supported_aggregate_function(found->second)) + auto exact = found == functions.end() ? std::optional{} + : from_duckdb_aggregate_name("__sirius_" + found->second); + bool numeric = exact && mo_decimal::is_decimal_aggregate(*exact) && + profile.functions.contains(function.function_reference()); + if (found == functions.end() || (!numeric && !supported_aggregate_function(found->second))) throw failure(SIRIUS_UNSUPPORTED, "unsupported aggregate function in embedded plan"); for (auto const& argument : function.arguments()) if (!argument.has_value()) @@ -301,7 +342,8 @@ void inspect_message(duckdb::google::protobuf::Message const& message, auto actual_nullability = read_nullability(actual); auto expected_nullability = expected.nullable ? substrait::Type::NULLABILITY_NULLABLE : substrait::Type::NULLABILITY_REQUIRED; - if (!read_type_matches(actual, expected) || actual_nullability != expected_nullability) + if (!read_type_matches(actual, expected, profile) || + actual_nullability != expected_nullability) throw failure(SIRIUS_INVALID_ARGUMENT, "plan read column type does not match its binding"); } } else if (name == "substrait.ReferenceRel") { @@ -326,10 +368,11 @@ void inspect_message(duckdb::google::protobuf::Message const& message, relation_ordinal, query, reads, - functions); + functions, + profile); } else { inspect_message( - reflection->GetMessage(message, field), relation_ordinal, query, reads, functions); + reflection->GetMessage(message, field), relation_ordinal, query, reads, functions, profile); } } } @@ -345,9 +388,34 @@ void validate_embedded_plan(std::string_view bytes, query_state const& query) plan.version().minor_number() != 78) throw failure(SIRIUS_UNSUPPORTED, "embedded plans require Substrait version 0.78"); std::unordered_map functions; + numeric_profile profile; + std::unordered_set numeric_uris; + for (auto const& uri : plan.extension_urns()) + if (uri.urn() == mo_decimal::extension_uri) { + profile.exact = true; + numeric_uris.insert(uri.extension_urn_anchor()); + } + for (auto const& extension : plan.extensions()) + if (extension.has_extension_type()) { + auto const& type = extension.extension_type(); + if (numeric_uris.contains(type.extension_urn_reference()) && + type.name() == "mo_exact_decimal") + profile.types.push_back(type.type_anchor()); + } + std::sort(profile.types.begin(), profile.types.end()); + if (profile.exact) + for (auto const& binding : query.bindings) + if (binding.source_kind != SIRIUS_READ_MO) + throw failure(SIRIUS_UNSUPPORTED, "exact decimal v1 requires MO readers"); + if (profile.exact) + for (auto const& output : query.contract->outputs) + if (output.oid == 34 && output.width > 65) + throw failure(SIRIUS_UNSUPPORTED, "public Decimal256 precision exceeds 65"); for (auto const& extension : plan.extensions()) { if (!extension.has_extension_function()) continue; auto const& mapping = extension.extension_function(); + if (numeric_uris.contains(mapping.extension_urn_reference())) + profile.functions.insert(mapping.function_anchor()); if (!functions.emplace(mapping.function_anchor(), mapping.name()).second) throw failure(SIRIUS_INVALID_ARGUMENT, "duplicate Substrait function anchor"); } @@ -366,7 +434,7 @@ void validate_embedded_plan(std::string_view bytes, query_state const& query) throw failure(SIRIUS_INVALID_ARGUMENT, "plan output name does not match query contract"); std::set reads; for (int i = 0; i < plan.relations_size(); ++i) - inspect_message(plan.relations(i), i, query, reads, functions); + inspect_message(plan.relations(i), i, query, reads, functions, profile); if (reads.size() != query.bindings.size()) throw failure(SIRIUS_INVALID_ARGUMENT, "registered and planned read binding sets differ"); } diff --git a/src/embedding/result_codec.cpp b/src/embedding/result_codec.cpp index 549584982f..387414d6c0 100644 --- a/src/embedding/result_codec.cpp +++ b/src/embedding/result_codec.cpp @@ -80,6 +80,63 @@ cudf::data_type expected(owned_column const& c) } } std::size_t align8(std::size_t n) { return (n + 7) & ~std::size_t(7); } +bool wide_shape_matches(cudf::column_view const& column, owned_column const& type) +{ + std::vector children; + for (int i = 0; i < column.num_children(); ++i) { + auto child = column.child(i); + children.emplace_back( + child.type(), child.size(), child.head(), nullptr, 0, child.offset()); + } + return mo_decimal::decimal_column_matches( + cudf::column_view(column.type(), + column.size(), + nullptr, + column.null_mask(), + column.null_count(), + column.offset(), + children), + {256, static_cast(type.width), static_cast(type.scale)}); +} +void validate_wide_validity(cudf::column_view const& column, + uint32_t begin, + uint32_t rows, + result_batch& scratch, + rmm::cuda_stream_view stream) +{ + bool masks = false; + for (int i = 0; i < column.num_children(); ++i) + masks |= column.child(i).nullable(); + if (!masks) return; + // Use the already admitted codec slab, not an unreserved GPU allocation. + // Child NULLs are redundant only beneath a NULL parent; verify before this + // output slice can become visible. The temporary bitmaps stay on the stack. + for (uint32_t first = 0; first < rows; first += 1024) { + auto count = std::min(1024, rows - first); + auto base = static_cast(column.offset()) + begin + first; + std::array parent, child; + parent.fill(UINT32_MAX); + if (column.nullable()) { + auto bytes = (base % 32 + count + 31) / 32 * 4; + download(*scratch.storage, 0, column.null_mask() + base / 32, bytes, stream); + scratch.storage->read(0, {reinterpret_cast(parent.data()), bytes}); + } + for (int i = 0; i < column.num_children(); ++i) { + auto limb = column.child(i); + if (limb.nullable()) { + auto origin = base + limb.offset(); + auto bytes = (origin % 32 + count + 31) / 32 * 4; + download(*scratch.storage, 0, limb.null_mask() + origin / 32, bytes, stream); + scratch.storage->read(0, {reinterpret_cast(child.data()), bytes}); + for (uint32_t row = 0; row < count; ++row) { + auto p = base % 32 + row, c = origin % 32 + row; + if ((parent[p / 32] & (1u << (p % 32))) && !(child[c / 32] & (1u << (c % 32)))) + throw failure(SIRIUS_EXECUTION_FAILED, "Decimal256 result has independent child NULLs"); + } + } + } + } +} sirius_input_vector column_layout(cudf::column_view const& col, owned_column const& type, uint32_t begin, @@ -118,10 +175,7 @@ result_slice_layout size_native_result(cudf::table_view table, throw failure(SIRIUS_EXECUTION_FAILED, "native result column count mismatch"); for (std::size_t c = 0; c < schema.size(); ++c) if (table.column(c).type() != expected(schema[c]) || - (schema[c].oid == 34 && - !mo_decimal::decimal_column_matches( - table.column(c), - {256, static_cast(schema[c].width), static_cast(schema[c].scale)}))) + (schema[c].oid == 34 && !wide_shape_matches(table.column(c), schema[c]))) throw failure(SIRIUS_UNSUPPORTED, "native result physical type does not match MO contract"); auto size = [&](uint32_t rows) { std::size_t position{}; @@ -158,6 +212,7 @@ void encode_native_result(cudf::table_view table, vector = column_layout(col, schema[c], begin, layout.rows, position, scratch, stream); auto width = input_element_size(schema[c].oid); if (schema[c].oid == 34) { + validate_wide_validity(col, begin, layout.rows, scratch, stream); // Reuse the admitted scratch slab; never allocate a second complete // result while interleaving the four fixed-width device children. if (scratch.storage->size() < 32) diff --git a/src/expression/aggregate_id.cpp b/src/expression/aggregate_id.cpp index 7e6e15e729..345948b8a5 100644 --- a/src/expression/aggregate_id.cpp +++ b/src/expression/aggregate_id.cpp @@ -30,7 +30,7 @@ namespace { // Forward table: DuckDB aggregate name -> Sirius aggregate id. // Linear scan; called once per BoundAggregateExpression at translator entry. -constexpr std::array, 8> kForwardTable = {{ +constexpr std::array, 12> kForwardTable = {{ {"sum", aggregate_id::sum}, {"sum_no_overflow", aggregate_id::sum_no_overflow}, {"count", aggregate_id::count}, @@ -39,11 +39,15 @@ constexpr std::array, 8> kForwardTable {"max", aggregate_id::max}, {"avg", aggregate_id::avg}, {"first", aggregate_id::first}, + {"__sirius_mo_decimal_sum", aggregate_id::mo_decimal_sum}, + {"__sirius_mo_decimal_avg", aggregate_id::mo_decimal_avg}, + {"__sirius_mo_decimal_min", aggregate_id::mo_decimal_min}, + {"__sirius_mo_decimal_max", aggregate_id::mo_decimal_max}, }}; // Reverse table: Sirius aggregate id -> canonical DuckDB aggregate name. // Indexed directly by enum value; never searched. -constexpr std::array kReverseTable = { +constexpr std::array kReverseTable = { "sum", "sum_no_overflow", "count", @@ -52,13 +56,17 @@ constexpr std::array kReverseTable = { "max", "avg", "first", + "__sirius_mo_decimal_sum", + "__sirius_mo_decimal_avg", + "__sirius_mo_decimal_min", + "__sirius_mo_decimal_max", }; static_assert(static_cast(aggregate_id::first) + 1 == 8, "aggregate_id::first must be the last entry; cardinality locked at 8."); -static_assert(kReverseTable.size() == 8, +static_assert(kReverseTable.size() == 12, "kReverseTable must have one slot per aggregate_id value."); -static_assert(kForwardTable.size() == 8, +static_assert(kForwardTable.size() == 12, "kForwardTable must have one entry per aggregate_id value."); // Walks both tables to ensure every enum value has exactly one canonical diff --git a/src/expression/aggregate_id.hpp b/src/expression/aggregate_id.hpp index d65a0ac402..515a06dfb8 100644 --- a/src/expression/aggregate_id.hpp +++ b/src/expression/aggregate_id.hpp @@ -43,6 +43,10 @@ enum class aggregate_id : uint16_t { max, avg, first, + mo_decimal_sum, + mo_decimal_avg, + mo_decimal_min, + mo_decimal_max, }; /** diff --git a/src/expression/from_duckdb.cpp b/src/expression/from_duckdb.cpp index 1ee8bc53e5..2355549e89 100644 --- a/src/expression/from_duckdb.cpp +++ b/src/expression/from_duckdb.cpp @@ -16,6 +16,7 @@ #include "expression/ast/from_duckdb.hpp" +#include "numeric/decimal_aggregate_bind.hpp" #include "numeric/decimal_functions.hpp" // sirius @@ -115,6 +116,34 @@ std::unique_ptr translate_comparison(duckdb::BoundComparisonExpression con auto left = from_duckdb(*expr.left); auto right = from_duckdb(*expr.right); if (!left || !right) { return nullptr; } + if (left->return_type().is_mo_decimal() || right->return_type().is_mo_decimal()) { + std::optional op; + using T = duckdb::ExpressionType; + switch (expr.GetExpressionType()) { + case T::COMPARE_EQUAL: op = mo_decimal::decimal_op::equal; break; + case T::COMPARE_NOTEQUAL: op = mo_decimal::decimal_op::not_equal; break; + case T::COMPARE_LESSTHAN: op = mo_decimal::decimal_op::less; break; + case T::COMPARE_LESSTHANOREQUALTO: op = mo_decimal::decimal_op::less_equal; break; + case T::COMPARE_GREATERTHAN: op = mo_decimal::decimal_op::greater; break; + case T::COMPARE_GREATERTHANOREQUALTO: op = mo_decimal::decimal_op::greater_equal; break; + default: break; + } + if (op) { + auto result = sirius::from_duckdb(expr.return_type); + // Join construction/cloning introduces ordinary comparison nodes, whose + // Boolean type has no nullability carrier. Their strictness follows the + // already validated exact key descriptors. Public scalar markers retain + // their original immutable binding and declared output type instead. + if (!result.nullability()) + result = result.with_nullability(left->return_type().nullability() == 2 || + right->return_type().nullability() == 2); + std::vector> arguments; + arguments.push_back(std::move(left)); + arguments.push_back(std::move(right)); + return std::make_unique( + function_call{mo_decimal::function(*op), std::move(arguments), result}); + } + } return std::make_unique(comparison{ /*op=*/sirius::from_duckdb(expr.GetExpressionType()), /*left=*/std::move(left), @@ -203,6 +232,7 @@ std::unique_ptr translate_aggregate(duckdb::BoundAggregateExpression const { auto agg_id_opt = sirius::from_duckdb_aggregate_name(aggr.function.name); if (!agg_id_opt.has_value()) { return nullptr; } + if (mo_decimal::is_decimal_aggregate(*agg_id_opt)) mo_decimal::validate_bound_aggregate(aggr); auto arguments = translate_children(aggr.children); if (!arguments) { return nullptr; } auto return_type = sirius::from_duckdb(aggr.return_type); diff --git a/src/expression/to_duckdb.cpp b/src/expression/to_duckdb.cpp index f25ebb7bb2..dc15f2be6c 100644 --- a/src/expression/to_duckdb.cpp +++ b/src/expression/to_duckdb.cpp @@ -16,6 +16,7 @@ #include "expression/ast/to_duckdb.hpp" +#include "numeric/decimal_aggregate_bind.hpp" #include "numeric/decimal_functions.hpp" // sirius @@ -261,8 +262,16 @@ std::unique_ptr to_duckdb(function_call const& alt) /*bind_info=*/nullptr)); } -std::unique_ptr to_duckdb(aggregate const& /*alt*/) +std::unique_ptr to_duckdb(aggregate const& alt) { + if (mo_decimal::is_decimal_aggregate(alt.function())) { + if (alt.distinct()) throw std::invalid_argument("MO aggregate DISTINCT is unsupported"); + duckdb::vector> children; + for (auto const& child : alt.arguments()) + children.push_back(to_duck_ptr(to_duckdb(*child))); + return from_duck_ptr(mo_decimal::bound_aggregate( + mo_decimal::aggregate_operation(alt.function()), std::move(children), alt.return_type())); + } throw not_implemented_exception( "[sirius::ast::to_duckdb] aggregate reconstruction is not implemented; consumers read the " "aggregate node natively (#863)."); diff --git a/src/expression_evaluator/expression_evaluator.cpp b/src/expression_evaluator/expression_evaluator.cpp index 9dc1183d40..def3c82019 100644 --- a/src/expression_evaluator/expression_evaluator.cpp +++ b/src/expression_evaluator/expression_evaluator.cpp @@ -15,6 +15,7 @@ */ // sirius +#include "numeric/decimal_aggregate_gpu.hpp" #include "pipeline/gpu_stream_quiescence_error.hpp" #include @@ -319,6 +320,9 @@ std::unique_ptr expression_evaluator::evaluate(cudf::table_view inp } else { result_column = result.release_column(); } + if (expr.return_type().is_mo_decimal()) + result_column = mo_decimal::restore_decimal_validity( + std::move(result_column), expr.return_type().mo_decimal_type(), _stream, _mr); if (expr.return_type().is_mo_decimal() && !mo_decimal::decimal_column_matches(result_column->view(), expr.return_type().mo_decimal_type())) @@ -339,6 +343,9 @@ std::unique_ptr expression_evaluator::evaluate(cudf::table_view inp } _output_columns.push_back(std::move(result_column)); } + if (expr.return_type().is_mo_decimal()) + _output_columns.back() = mo_decimal::restore_decimal_validity( + std::move(_output_columns.back()), expr.return_type().mo_decimal_type(), _stream, _mr); }; // Iterate _ast_expressions and route through the std::visit dispatcher, then diff --git a/src/expression_evaluator/specializations/mo_decimal.cpp b/src/expression_evaluator/specializations/mo_decimal.cpp index 9944eb8d8f..166c4bd424 100644 --- a/src/expression_evaluator/specializations/mo_decimal.cpp +++ b/src/expression_evaluator/specializations/mo_decimal.cpp @@ -22,6 +22,11 @@ expression_evaluator::evaluate_result expression_evaluator::evaluate_decimal_fun types.push_back(argument->return_type()); } validate_signature(op, types, expression.return_type()); + if (op == decimal_op::cast && types[0].is_mo_decimal()) { + auto a = types[0].mo_decimal_type(), b = expression.return_type().mo_decimal_type(); + if (a.bits == b.bits && a.precision == b.precision && a.scale == b.scale) + return evaluate(*arguments[0], mode); + } struct owners { std::vector arguments; std::vector> lifted; diff --git a/src/expression_evaluator/specializations/reference.cpp b/src/expression_evaluator/specializations/reference.cpp index 4b91e122c6..0d231fd673 100644 --- a/src/expression_evaluator/specializations/reference.cpp +++ b/src/expression_evaluator/specializations/reference.cpp @@ -15,6 +15,8 @@ */ // sirius +#include "numeric/decimal_aggregate_gpu.hpp" + #include #include @@ -76,7 +78,9 @@ evaluate_result expression_evaluator::evaluate(sirius::ast::reference const& alt if (logical.is_mo_decimal() && alt.column_index >= static_cast(_input_table.num_columns())) throw std::invalid_argument("MO exact-decimal reference is outside its input schema"); - auto const source = _input_table.column(alt.column_index); + auto source = _input_table.column(alt.column_index); + if (logical.is_mo_decimal()) + source = mo_decimal::canonical_decimal_view(source, logical.mo_decimal_type(), _stream, _mr); if (logical.is_mo_decimal() && !mo_decimal::decimal_column_matches(source, logical.mo_decimal_type())) throw std::invalid_argument("MO exact-decimal reference carrier does not match its descriptor"); diff --git a/src/numeric/CMakeLists.txt b/src/numeric/CMakeLists.txt index fdfd831856..a73cfada44 100644 --- a/src/numeric/CMakeLists.txt +++ b/src/numeric/CMakeLists.txt @@ -1,2 +1,10 @@ -target_sources(sirius_objects PRIVATE exact_decimal_gpu.cu decimal_types.cpp - decimal_functions.cpp decimal_import.cpp) +target_sources( + sirius_objects + PRIVATE exact_decimal_gpu.cu + decimal_types.cpp + decimal_functions.cpp + decimal_import.cpp + decimal_aggregate_bind.cpp + decimal_aggregate_gpu.cu + decimal_aggregate_layout.cpp + decimal_plan.cpp) diff --git a/src/numeric/decimal_aggregate.hpp b/src/numeric/decimal_aggregate.hpp new file mode 100644 index 0000000000..eb62ccffc8 --- /dev/null +++ b/src/numeric/decimal_aggregate.hpp @@ -0,0 +1,138 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include "numeric/exact_decimal.hpp" + +#ifdef __CUDACC__ +#define SIRIUS_AGG_HD __host__ __device__ +#else +#define SIRIUS_AGG_HD +#endif +namespace sirius::mo_decimal { +enum class aggregate_op : uint8_t { sum, avg, min, max }; + +// Raw physical state: its coefficient deliberately has no declared precision. +// All partials retain the input scale, and count zero is the NULL identity. +struct aggregate_state { + coefficient value{}; + int64_t count{}; + decimal_error error{decimal_error::none}; +}; +SIRIUS_AGG_HD inline bool fits_physical(coefficient const& value, uint16_t bits) +{ + if (bits != 128 && bits != 256) return false; + auto words = bits / 32; + auto extension = value.words[words - 1] >> 31 ? UINT32_MAX : 0; + for (int i = words; i < 8; ++i) + if (value.words[i] != extension) return false; + return true; +} +SIRIUS_AGG_HD inline int compare_coefficient(coefficient const& a, coefficient const& b) +{ + if (a.negative() != b.negative()) return a.negative() ? -1 : 1; + for (int i = 7; i >= 0; --i) { + if (a.words[i] < b.words[i]) return -1; + if (a.words[i] > b.words[i]) return 1; + } + return 0; +} +SIRIUS_AGG_HD inline aggregate_state merge_aggregate(aggregate_op op, + aggregate_state left, + aggregate_state const& right, + uint16_t bits) +{ + if (left.error != decimal_error::none) return left; + if (right.error != decimal_error::none) return right; + if (op > aggregate_op::max || (bits != 128 && bits != 256) || left.count < 0 || right.count < 0 || + right.count > INT64_MAX - left.count) + return {{}, 0, decimal_error::invalid_input}; + if (!right.count) return left; + if (!left.count) { + if (!fits_physical(right.value, bits)) return {{}, 0, decimal_error::invalid_input}; + return right; + } + if (op == aggregate_op::min || op == aggregate_op::max) { + auto cmp = compare_coefficient(left.value, right.value); + if ((op == aggregate_op::min && cmp > 0) || (op == aggregate_op::max && cmp < 0)) + left.value = right.value; + } else { + auto negative = left.value.negative(); + uint64_t carry = 0; + for (int i = 0; i < 8; ++i) { + auto sum = uint64_t(left.value.words[i]) + right.value.words[i] + carry; + left.value.words[i] = sum; + carry = sum >> 32; + } + if ((negative == right.value.negative() && left.value.negative() != negative) || + !fits_physical(left.value, bits)) + return {{}, 0, decimal_error::invalid_input}; + } + left.count += right.count; + return left; +} +SIRIUS_AGG_HD inline aggregate_state accumulate_aggregate( + aggregate_op op, aggregate_state state, coefficient value, bool valid, uint16_t bits) +{ + return valid ? merge_aggregate(op, state, {value, 1, decimal_error::none}, bits) : state; +} +SIRIUS_AGG_HD inline decimal_result finalize_aggregate(aggregate_op op, + aggregate_state const& state, + decimal_type input, + decimal_type output) +{ + if (state.error != decimal_error::none) return {{}, state.error, false}; + if (!input.valid() || !output.valid() || op > aggregate_op::max || state.count < 0) + return {{}, decimal_error::invalid_input, false}; + if (!state.count) return {}; + if (op == aggregate_op::min || op == aggregate_op::max) + return {state.value, decimal_error::none, true}; + if (op == aggregate_op::sum) { + if (input.scale != output.scale) return {{}, decimal_error::invalid_input, false}; + return decimal_detail::finish(decimal_detail::absolute(state.value), + state.value.negative(), + output, + decimal_error::invalid_input); + } + coefficient count; + count.words[0] = static_cast(state.count); + count.words[1] = static_cast(state.count) >> 32; + // The scalar division core never validates the input coefficient against + // precision. This descriptor expresses only the raw physical state/scale. + auto result = evaluate_decimal(decimal_op::divide, + state.value, + {256, 76, input.scale}, + true, + count, + {64, 18, 0}, + true, + output); + if (result.error != decimal_error::none) result.error = decimal_error::invalid_input; + return result; +} + +struct normalized_decimal_key { + coefficient value{}; + uint8_t scale{}; +}; +SIRIUS_AGG_HD inline normalized_decimal_key normalize_key(coefficient value, uint8_t scale) +{ + auto magnitude = decimal_detail::absolute(value); + if (magnitude.zero()) return {}; + while (scale) { + auto quotient = magnitude; + uint64_t remainder = 0; + for (int i = 7; i >= 0; --i) { + auto word = (remainder << 32) | magnitude.words[i]; + quotient.words[i] = word / 10; + remainder = word % 10; + } + if (remainder) break; + magnitude = quotient; + --scale; + } + return {decimal_detail::finish( + magnitude, value.negative(), {256, 76, 0}, decimal_error::invalid_input, false) + .value, + scale}; +} +} // namespace sirius::mo_decimal +#undef SIRIUS_AGG_HD diff --git a/src/numeric/decimal_aggregate_bind.cpp b/src/numeric/decimal_aggregate_bind.cpp new file mode 100644 index 0000000000..a78c011a52 --- /dev/null +++ b/src/numeric/decimal_aggregate_bind.cpp @@ -0,0 +1,172 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "numeric/decimal_aggregate_bind.hpp" + +#include "helper/type_conversions.hpp" +#include "numeric/decimal_plan.hpp" + +#include +#include +#include +#include +#include + +namespace sirius::mo_decimal { +bool is_decimal_aggregate(aggregate_id id) noexcept +{ + return id >= aggregate_id::mo_decimal_sum && id <= aggregate_id::mo_decimal_max; +} +aggregate_op aggregate_operation(aggregate_id id) +{ + if (!is_decimal_aggregate(id)) throw std::invalid_argument("not an MO decimal aggregate"); + return static_cast(static_cast(id) - + static_cast(aggregate_id::mo_decimal_sum)); +} +aggregate_id aggregate_function(aggregate_op op) +{ + if (op > aggregate_op::max) throw std::invalid_argument("unknown MO decimal aggregate"); + return static_cast(static_cast(aggregate_id::mo_decimal_sum) + + static_cast(op)); +} +decimal_type aggregate_input_type(logical_type const& input) +{ + if (input.is_mo_decimal()) return input.mo_decimal_type(); + if (input.id() == type_id::BIGINT) return {64, 18, 0}; + throw std::invalid_argument("unsupported MO decimal aggregate operand"); +} +void validate_aggregate_signature(aggregate_op op, + logical_type const& input, + logical_type const& output) +{ + if (op > aggregate_op::max || !output.is_mo_decimal()) + throw std::invalid_argument("MO aggregate requires an exact result descriptor"); + auto a = aggregate_input_type(input), r = output.mo_decimal_type(); + if (r.precision > 65) throw std::invalid_argument("MO aggregate result exceeds public precision"); + if (op == aggregate_op::min || op == aggregate_op::max) { + if (!input.is_mo_decimal() || a.bits != r.bits || a.precision != r.precision || + a.scale != r.scale) + throw std::invalid_argument("MO MIN/MAX must preserve the argument descriptor"); + } else if (r.bits < 128 || r.bits < a.bits || (op == aggregate_op::sum && r.scale != a.scale) || + (op == aggregate_op::avg && r.scale < a.scale)) + throw std::invalid_argument("invalid MO SUM/AVG result descriptor"); +} +namespace { +struct aggregate_data final : duckdb::FunctionData { + aggregate_op op; + duckdb::LogicalType input, output; + aggregate_data(aggregate_op o, duckdb::LogicalType a, duckdb::LogicalType r) + : op(o), input(std::move(a)), output(std::move(r)) + { + } + duckdb::unique_ptr Copy() const override + { + return duckdb::make_uniq(op, input, output); + } + bool Equals(duckdb::FunctionData const& other) const override + { + auto data = dynamic_cast(&other); + return data && op == data->op && input == data->input && output == data->output; + } +}; +[[noreturn]] void never_execute() +{ + throw duckdb::NotImplementedException("MO exact-decimal aggregate cannot execute in DuckDB"); +} +duckdb::idx_t state_size(duckdb::AggregateFunction const&) { return 1; } +void initialize(duckdb::AggregateFunction const&, duckdb::data_ptr_t) { never_execute(); } +void update( + duckdb::Vector[], duckdb::AggregateInputData&, duckdb::idx_t, duckdb::Vector&, duckdb::idx_t) +{ + never_execute(); +} +void simple_update( + duckdb::Vector[], duckdb::AggregateInputData&, duckdb::idx_t, duckdb::data_ptr_t, duckdb::idx_t) +{ + never_execute(); +} +void combine(duckdb::Vector&, duckdb::Vector&, duckdb::AggregateInputData&, duckdb::idx_t) +{ + never_execute(); +} +void finalize( + duckdb::Vector&, duckdb::AggregateInputData&, duckdb::Vector&, duckdb::idx_t, duckdb::idx_t) +{ + never_execute(); +} +duckdb::unique_ptr bind( + duckdb::ClientContext&, + duckdb::AggregateFunction& fn, + duckdb::vector>& arguments) +{ + auto id = from_duckdb_aggregate_name(fn.name); + if (!id || !is_decimal_aggregate(*id) || arguments.size() != 2 || + arguments.back()->GetExpressionClass() != duckdb::ExpressionClass::BOUND_CONSTANT) + throw duckdb::InvalidInputException("invalid MO aggregate binding marker"); + auto const& marker = arguments.back()->Cast().value; + auto output = from_duckdb_type(marker.type()); + if (!marker.IsNull() || !output) + throw duckdb::InvalidInputException("MO aggregate lost its output marker"); + restore_exact_bound_types(arguments[0]); + validate_aggregate_signature( + aggregate_operation(*id), sirius::from_duckdb(arguments[0]->return_type), *output); + fn.return_type = marker.type(); + arguments.pop_back(); + fn.arguments = {arguments[0]->return_type}; + return duckdb::make_uniq( + aggregate_operation(*id), fn.arguments[0], fn.return_type); +} +duckdb::AggregateFunction make_function(aggregate_op op) +{ + duckdb::AggregateFunction fn(std::string(to_duckdb_aggregate_name(aggregate_function(op))), + {duckdb::LogicalType::ANY, duckdb::LogicalType::ANY}, + duckdb::LogicalType::ANY, + state_size, + initialize, + update, + combine, + finalize, + duckdb::FunctionNullHandling::SPECIAL_HANDLING, + simple_update, + bind); + fn.stability = duckdb::FunctionStability::VOLATILE; + fn.order_dependent = duckdb::AggregateOrderDependent::NOT_ORDER_DEPENDENT; + return fn; +} +} // namespace +void register_aggregate_functions(duckdb::DatabaseInstance& instance) +{ + auto transaction = duckdb::CatalogTransaction::GetSystemTransaction(instance); + auto& catalog = duckdb::Catalog::GetSystemCatalog(instance); + for (unsigned i = 0; i <= static_cast(aggregate_op::max); ++i) { + duckdb::CreateAggregateFunctionInfo info(make_function(static_cast(i))); + catalog.CreateFunction(transaction, info); + } +} +void validate_bound_aggregate(duckdb::BoundAggregateExpression const& expression) +{ + auto data = dynamic_cast(expression.bind_info.get()); + auto id = from_duckdb_aggregate_name(expression.function.name); + if (!data || !id || !is_decimal_aggregate(*id) || data->op != aggregate_operation(*id) || + expression.children.size() != 1 || data->input != expression.children[0]->return_type || + data->output != expression.return_type || expression.function.bind != bind || + expression.IsDistinct() || expression.filter || expression.order_bys) + throw std::invalid_argument( + "MO aggregate lost its checked signature or has unsupported modifiers"); + validate_aggregate_signature( + data->op, sirius::from_duckdb(data->input), sirius::from_duckdb(data->output)); +} +duckdb::unique_ptr bound_aggregate( + aggregate_op op, + duckdb::vector> children, + logical_type const& output) +{ + if (children.size() != 1) throw std::invalid_argument("MO aggregate needs one operand"); + validate_aggregate_signature(op, sirius::from_duckdb(children[0]->return_type), output); + auto fn = make_function(op); + fn.arguments = {children[0]->return_type}; + fn.return_type = duckdb_type(output); + auto result = duckdb::make_uniq( + fn, std::move(children), nullptr, nullptr, duckdb::AggregateType::NON_DISTINCT); + result->bind_info = duckdb::make_uniq(op, fn.arguments[0], fn.return_type); + return result; +} +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_aggregate_bind.hpp b/src/numeric/decimal_aggregate_bind.hpp new file mode 100644 index 0000000000..983cbeb8fb --- /dev/null +++ b/src/numeric/decimal_aggregate_bind.hpp @@ -0,0 +1,26 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include "expression/aggregate_id.hpp" +#include "numeric/decimal_aggregate.hpp" +#include "numeric/decimal_types.hpp" + +namespace duckdb { +class DatabaseInstance; +class BoundAggregateExpression; +class Expression; +} // namespace duckdb +namespace sirius::mo_decimal { +bool is_decimal_aggregate(aggregate_id id) noexcept; +aggregate_op aggregate_operation(aggregate_id id); +aggregate_id aggregate_function(aggregate_op op); +decimal_type aggregate_input_type(logical_type const& input); +void validate_aggregate_signature(aggregate_op op, + logical_type const& input, + logical_type const& output); +void register_aggregate_functions(duckdb::DatabaseInstance& instance); +void validate_bound_aggregate(duckdb::BoundAggregateExpression const& expression); +duckdb::unique_ptr bound_aggregate( + aggregate_op op, + duckdb::vector> children, + logical_type const& output); +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_aggregate_gpu.cu b/src/numeric/decimal_aggregate_gpu.cu new file mode 100644 index 0000000000..e7fc7e3023 --- /dev/null +++ b/src/numeric/decimal_aggregate_gpu.cu @@ -0,0 +1,569 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "numeric/decimal_aggregate_gpu.hpp" +#include "numeric/decimal_error.hpp" +#include "pipeline/gpu_stream_quiescence_error.hpp" + +#include +#include +#include + +#include +#include + +namespace sirius::mo_decimal { +namespace { +struct operand { + uint8_t const* data{}; + uint64_t const* limbs[4]{}; + cudf::bitmask_type const* mask{}; + cudf::size_type offset{}; + uint32_t bytes{}; + __device__ bool valid(cudf::size_type row) const + { + auto bit = row + offset; + return !mask || ((mask[bit / 32] >> (bit % 32)) & 1); + } + __device__ coefficient value(cudf::size_type row) const + { + if (bytes != 32) return load_coefficient(data + uint64_t(row + offset) * bytes, bytes); + coefficient value; + for (int i = 0; i < 4; ++i) { + auto word = limbs[3 - i][row + offset]; + value.words[2 * i] = word; + value.words[2 * i + 1] = word >> 32; + } + return value; + } +}; +operand view(cudf::column_view column, decimal_type type) +{ + operand result; + result.bytes = type.bytes(); + result.offset = column.offset(); + result.mask = column.null_mask(); + if (type.bits != 256) + result.data = column.head(); + else + for (int i = 0; i < 4; ++i) + result.limbs[i] = column.child(i).head() + column.child(i).offset(); + return result; +} +struct destination { + uint8_t* data{}; + uint64_t* limbs[4]{}; + uint32_t bytes{}; + __device__ void store(coefficient const& value, cudf::size_type row) const + { + if (bytes != 32) + store_coefficient(value, data + uint64_t(row) * bytes, bytes); + else + for (int i = 0; i < 4; ++i) + limbs[3 - i][row] = uint64_t(value.words[2 * i]) | (uint64_t(value.words[2 * i + 1]) << 32); + } +}; +destination output_view(cudf::mutable_column_view column, uint32_t bytes) +{ + destination result; + result.bytes = bytes; + if (bytes != 32) + result.data = column.data(); + else + for (int i = 0; i < 4; ++i) + result.limbs[i] = column.child(i).data(); + return result; +} +constexpr int block_size = 256, tile_rows = 4096; +__global__ void tile_prefix(cudf::size_type const* offsets, + cudf::size_type groups, + int64_t* tile_offsets) +{ + for (int64_t group = blockIdx.x * blockDim.x + threadIdx.x; group <= groups; + group += blockDim.x * gridDim.x) + // For every segment, floor(end/T)-floor(start/T)+1 >= ceil(length/T). + // Extra tiles are identities. This bounded prefix needs neither a device + // scan nor a host rendezvous; subtract the origin to support subranges. + tile_offsets[group] = (int64_t(offsets[group]) - offsets[0]) / tile_rows + group; +} +__global__ void reduce_tiles(aggregate_op op, + operand values, + uint16_t bits, + int64_t const* counts, + cudf::size_type const* order, + cudf::size_type const* offsets, + int64_t const* tile_offsets, + cudf::size_type groups, + int64_t tiles, + aggregate_state* partials) +{ + __shared__ aggregate_state states[block_size]; + for (int64_t tile = blockIdx.x; tile < tiles; tile += gridDim.x) { + cudf::size_type lo = 0, hi = groups; + while (lo + 1 < hi) { + auto mid = lo + (hi - lo) / 2; + if (tile_offsets[mid] <= tile) + lo = mid; + else + hi = mid; + } + auto start = int64_t(offsets[lo]) + int64_t(tile - tile_offsets[lo]) * tile_rows; + auto end = start + tile_rows; + if (end > offsets[lo + 1]) end = offsets[lo + 1]; + aggregate_state state; + for (auto position = start + threadIdx.x; position < end; position += block_size) { + auto row = order ? order[position] : position; + if (values.valid(row)) + state = merge_aggregate( + op, state, {values.value(row), counts ? counts[row] : 1, decimal_error::none}, bits); + } + states[threadIdx.x] = state; + __syncthreads(); + for (int stride = block_size / 2; stride; stride /= 2) { + if (threadIdx.x < stride) + states[threadIdx.x] = + merge_aggregate(op, states[threadIdx.x], states[threadIdx.x + stride], bits); + __syncthreads(); + } + if (!threadIdx.x) partials[tile] = states[0]; + __syncthreads(); + } +} +__global__ void reduce_groups(aggregate_op op, + uint16_t bits, + aggregate_state const* partials, + int64_t const* offsets, + cudf::size_type groups, + destination output, + int64_t* counts, + uint8_t* errors) +{ + __shared__ aggregate_state states[block_size]; + for (int64_t group = blockIdx.x; group < groups; group += gridDim.x) { + aggregate_state state; + for (int64_t tile = int64_t(offsets[group]) + threadIdx.x; tile < offsets[group + 1]; + tile += block_size) + state = merge_aggregate(op, state, partials[tile], bits); + states[threadIdx.x] = state; + __syncthreads(); + for (int stride = block_size / 2; stride; stride /= 2) { + if (threadIdx.x < stride) + states[threadIdx.x] = + merge_aggregate(op, states[threadIdx.x], states[threadIdx.x + stride], bits); + __syncthreads(); + } + if (!threadIdx.x) { + output.store(states[0].value, group); + counts[group] = states[0].count; + if (errors) errors[group] = static_cast(states[0].error); + } + __syncthreads(); + } +} +__global__ void finalize_rows(aggregate_op op, + operand values, + int64_t const* counts, + decimal_type input, + decimal_type output, + cudf::size_type rows, + destination destination, + cudf::bitmask_type* mask, + uint32_t* valid, + uint8_t* errors) +{ + for (int64_t row = blockIdx.x * blockDim.x + threadIdx.x; row < rows; + row += blockDim.x * gridDim.x) { + auto result = + finalize_aggregate(op, {values.value(row), counts[row], decimal_error::none}, input, output); + destination.store(result.value, row); + errors[row] = static_cast(result.error); + if (result.valid) { + atomicOr(mask + row / 32, uint32_t(1) << (row % 32)); + atomicAdd(valid, 1u); + } + } +} +__global__ void key_rows( + operand values, uint8_t scale, cudf::size_type rows, destination output, uint8_t* scales) +{ + for (int64_t row = blockIdx.x * blockDim.x + threadIdx.x; row < rows; + row += blockDim.x * gridDim.x) { + auto key = + values.valid(row) ? normalize_key(values.value(row), scale) : normalized_decimal_key{}; + output.store(key.value, row); + scales[row] = key.scale; + } +} +struct child_validity { + cudf::bitmask_type const* masks[4]{}; + cudf::size_type offsets[4]{}; +}; +__global__ void check_children(operand values, + child_validity children, + cudf::size_type rows, + uint32_t* bad) +{ + for (int64_t row = blockIdx.x * blockDim.x + threadIdx.x; row < rows; + row += blockDim.x * gridDim.x) + if (values.valid(row)) { + for (int i = 0; i < 4; ++i) { + auto bit = row + values.offset + children.offsets[i]; + auto mask = children.masks[i]; + if (mask && !(mask[bit / 32] & (uint32_t(1) << (bit % 32)))) atomicExch(bad, 1u); + } + } +} +__global__ void group_offsets(cudf::size_type const* labels, + cudf::size_type const* order, + cudf::size_type rows, + cudf::size_type* offsets, + cudf::size_type groups) +{ + for (int64_t row = blockIdx.x * blockDim.x + threadIdx.x; row < rows; + row += blockDim.x * gridDim.x) { + auto label = labels[order[row]]; + if (!row || labels[order[row - 1]] != label) offsets[label] = row; + } + if (!blockIdx.x && !threadIdx.x) offsets[groups] = rows; +} +template +[[noreturn]] void rethrow_quiescent(std::unique_ptr& owner, rmm::cuda_stream_view stream) +{ + auto failure = std::current_exception(); + if (cudaStreamSynchronize(stream.value()) != cudaSuccess) { + (void)owner.release(); + throw pipeline::gpu_stream_quiescence_error( + "MO aggregate/key buffers could not prove quiescence"); + } + std::rethrow_exception(failure); +} +} // namespace +aggregate_columns reduce_aggregate_columns(aggregate_op op, + cudf::column_view values, + decimal_type input, + uint16_t state_bits, + cudf::column_view counts, + cudf::column_view order, + cudf::size_type const* offsets, + cudf::size_type groups, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + values = canonical_decimal_view(values, input, stream, mr); + if (!decimal_column_matches(values, input) || (state_bits != 128 && state_bits != 256) || + state_bits < input.bits || op > aggregate_op::max || groups < 0 || !offsets || + (counts.type().id() != cudf::type_id::EMPTY && + (counts.type().id() != cudf::type_id::INT64 || counts.size() != values.size() || + counts.has_nulls())) || + (order.type().id() != cudf::type_id::EMPTY && + (order.type().id() != cudf::type_id::INT32 || order.size() != values.size() || + order.has_nulls()))) + throw std::invalid_argument("invalid MO aggregate reduction inputs"); + struct owners { + aggregate_columns result; + std::unique_ptr> tile_offsets; + std::unique_ptr> partials; + std::optional errors; + int64_t single_offsets[2]{}; + }; + auto owner = std::make_unique(); + try { + // At most INT32_MAX initial rows: an accumulator at least 32 bits + // wider than the signed physical input contains every subset. MIN/MAX + // are also infallible. This skips error polling only, never publication + // precision validation, and depends on domains rather than sampled data. + bool proven = + counts.type().id() == cudf::type_id::EMPTY && + (state_bits >= input.bits + 32 || op == aggregate_op::min || op == aggregate_op::max); + decimal_type state_type{ + state_bits, static_cast(state_bits == 128 ? 38 : 76), input.scale}; + owner->result.values = + make_decimal_column(state_type, groups, cudf::mask_state::UNALLOCATED, stream, mr); + owner->result.counts = cudf::make_numeric_column( + cudf::data_type{cudf::type_id::INT64}, groups, cudf::mask_state::UNALLOCATED, stream, mr); + owner->errors.emplace(decimal_column_result{ + nullptr, rmm::device_uvector(proven ? 0 : groups, stream, mr)}); + if (groups) { + owner->tile_offsets = + std::make_unique>(int64_t(groups) + 1, stream, mr); + int64_t tiles{}; + if (groups == 1) { + // Empty extra tiles are identities even for a subrange. No device + // prefix scan or host tile-count rendezvous is needed for one group. + tiles = (int64_t(values.size()) + tile_rows - 1) / tile_rows; + owner->single_offsets[1] = tiles; + CUDF_CUDA_TRY(cudaMemcpyAsync(owner->tile_offsets->data(), + owner->single_offsets, + sizeof(owner->single_offsets), + cudaMemcpyHostToDevice, + stream.value())); + } else { + tiles = int64_t(values.size()) / tile_rows + groups; + tile_prefix<<<128, block_size, 0, stream.value()>>>( + offsets, groups, owner->tile_offsets->data()); + CUDF_CUDA_TRY(cudaGetLastError()); + } + owner->partials = std::make_unique>(tiles, stream, mr); + if (tiles) + reduce_tiles<<(tiles, 4096), block_size, 0, stream.value()>>>( + op, + view(values, input), + state_bits, + counts.type().id() == cudf::type_id::EMPTY ? nullptr : counts.data(), + order.type().id() == cudf::type_id::EMPTY ? nullptr : order.data(), + offsets, + owner->tile_offsets->data(), + groups, + tiles, + owner->partials->data()); + reduce_groups<<(groups, 4096), block_size, 0, stream.value()>>>( + op, + state_bits, + owner->partials->data(), + owner->tile_offsets->data(), + groups, + output_view(owner->result.values->mutable_view(), state_type.bytes()), + owner->result.counts->mutable_view().data(), + proven ? nullptr : owner->errors->errors.data()); + CUDF_CUDA_TRY(cudaGetLastError()); + if (!proven) { + auto error = column_error(*owner->errors, stream, mr); + if (error != decimal_error::none) throw numeric_error(error); + } + } + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + return std::move(owner->result); + } catch (...) { + rethrow_quiescent(owner, stream); + } +} +std::unique_ptr finalize_aggregate_columns(aggregate_op op, + cudf::column_view values, + cudf::column_view counts, + decimal_type input, + decimal_type output, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + if (values.size() != counts.size() || counts.type().id() != cudf::type_id::INT64 || + counts.has_nulls() || !input.valid() || !output.valid()) + throw std::invalid_argument("invalid MO aggregate finalization inputs"); + struct owners { + std::optional result; + std::unique_ptr> valid; + uint32_t host_valid{}; + }; + auto owner = std::make_unique(); + try { + auto rows = values.size(); + owner->result.emplace( + decimal_column_result{nullptr, rmm::device_uvector(0, stream, mr)}); + owner->result->values = + make_decimal_column(output, rows, cudf::mask_state::ALL_NULL, stream, mr); + owner->result->errors = rmm::device_uvector(rows, stream, mr); + owner->valid = std::make_unique>(1, stream, mr); + CUDF_CUDA_TRY(cudaMemsetAsync(owner->valid->data(), 0, sizeof(uint32_t), stream.value())); + auto bits = values.type().id() == cudf::type_id::STRUCT ? 256 : 128; + decimal_type state_type{ + static_cast(bits), static_cast(bits == 256 ? 76 : 38), input.scale}; + if (!decimal_column_matches(values, state_type)) + throw std::invalid_argument("invalid MO aggregate state carrier"); + if (rows) + finalize_rows<<<128, block_size, 0, stream.value()>>>( + op, + view(values, state_type), + counts.data(), + input, + output, + rows, + output_view(owner->result->values->mutable_view(), output.bytes()), + owner->result->values->mutable_view().null_mask(), + owner->valid->data(), + owner->result->errors.data()); + CUDF_CUDA_TRY(cudaGetLastError()); + auto error = column_error(*owner->result, stream, mr); + if (error != decimal_error::none) throw numeric_error(error); + CUDF_CUDA_TRY(cudaMemcpyAsync(&owner->host_valid, + owner->valid->data(), + sizeof(owner->host_valid), + cudaMemcpyDeviceToHost, + stream.value())); + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + owner->result->values->set_null_count(rows - owner->host_valid); + return std::move(owner->result->values); + } catch (...) { + rethrow_quiescent(owner, stream); + } +} +std::unique_ptr make_equality_key(cudf::column_view values, + decimal_type type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + values = canonical_decimal_view(values, type, stream, mr); + if (!decimal_column_matches(values, type)) + throw std::invalid_argument("invalid MO decimal equality key"); + struct owners { + std::unique_ptr limbs, scales, result; + }; + auto owner = std::make_unique(); + try { + owner->limbs = + make_decimal_column({256, 76, 0}, values.size(), cudf::mask_state::UNALLOCATED, stream, mr); + owner->scales = cudf::make_numeric_column(cudf::data_type{cudf::type_id::UINT8}, + values.size(), + cudf::mask_state::UNALLOCATED, + stream, + mr); + if (values.size()) + key_rows<<<128, block_size, 0, stream.value()>>>( + view(values, type), + type.scale, + values.size(), + output_view(owner->limbs->mutable_view(), 32), + owner->scales->mutable_view().data()); + CUDF_CUDA_TRY(cudaGetLastError()); + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + auto contents = owner->limbs->release(); + contents.children.push_back(std::move(owner->scales)); + owner->result = std::make_unique(cudf::data_type{cudf::type_id::STRUCT}, + values.size(), + rmm::device_buffer{}, + cudf::copy_bitmask(values, stream, mr), + values.null_count(), + std::move(contents.children)); + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + return std::move(owner->result); + } catch (...) { + rethrow_quiescent(owner, stream); + } +} +std::unique_ptr restore_decimal_validity(std::unique_ptr column, + decimal_type type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + if (type.bits != 256) { + if (!decimal_column_matches(column->view(), type)) + throw std::invalid_argument("MO decimal carrier mismatch"); + return column; + } + struct owners { + std::unique_ptr column; + std::unique_ptr> bad; + uint32_t host_bad{}; + }; + auto owner = std::make_unique(); + owner->column = std::move(column); + try { + auto col = owner->column->view(); + if (col.type().id() != cudf::type_id::STRUCT || col.num_children() != 4 || col.offset()) + throw std::invalid_argument("MO Decimal256 carrier shape mismatch"); + bool masks = false; + for (int i = 0; i < 4; ++i) { + auto child = col.child(i); + if (child.type().id() != (i ? cudf::type_id::UINT64 : cudf::type_id::INT64) || + child.offset() || child.size() != col.size()) + throw std::invalid_argument("MO Decimal256 child shape mismatch"); + masks |= child.nullable(); + } + if (!masks) return std::move(owner->column); + owner->bad = std::make_unique>(1, stream, mr); + CUDF_CUDA_TRY(cudaMemsetAsync(owner->bad->data(), 0, sizeof(uint32_t), stream.value())); + child_validity validity; + for (int i = 0; i < 4; ++i) { + validity.masks[i] = col.child(i).null_mask(); + validity.offsets[i] = col.child(i).offset(); + } + if (col.size()) + check_children<<<128, block_size, 0, stream.value()>>>( + view(col, type), validity, col.size(), owner->bad->data()); + CUDF_CUDA_TRY(cudaGetLastError()); + CUDF_CUDA_TRY(cudaMemcpyAsync(&owner->host_bad, + owner->bad->data(), + sizeof(owner->host_bad), + cudaMemcpyDeviceToHost, + stream.value())); + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + if (owner->host_bad) throw std::invalid_argument("MO Decimal256 has independent child NULLs"); + auto rows = col.size(), nulls = col.null_count(); + auto contents = owner->column->release(); + for (auto& child : contents.children) { + auto child_type = child->type(); + auto data = child->release(); + child = std::make_unique( + child_type, rows, std::move(*data.data), rmm::device_buffer{}, 0); + } + owner->column = std::make_unique( + cudf::data_type{cudf::type_id::STRUCT}, + rows, + rmm::device_buffer{}, + contents.null_mask ? std::move(*contents.null_mask) : rmm::device_buffer{}, + nulls, + std::move(contents.children)); + return std::move(owner->column); + } catch (...) { + rethrow_quiescent(owner, stream); + } +} +void make_group_offsets(cudf::column_view labels, + cudf::column_view order, + cudf::size_type* offsets, + cudf::size_type groups, + rmm::cuda_stream_view stream) +{ + group_offsets<<<128, block_size, 0, stream.value()>>>( + labels.data(), order.data(), labels.size(), offsets, groups); + CUDF_CUDA_TRY(cudaGetLastError()); +} +cudf::column_view canonical_decimal_view(cudf::column_view column, + decimal_type type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + if (decimal_column_matches(column, type)) return column; + if (!type.valid() || type.bits != 256 || column.type().id() != cudf::type_id::STRUCT || + column.num_children() != 4) + throw std::invalid_argument("MO exact decimal carrier mismatch"); + child_validity validity; + std::vector children; + for (int i = 0; i < 4; ++i) { + auto child = column.child(i); + if (child.type().id() != (i ? cudf::type_id::UINT64 : cudf::type_id::INT64) || + child.size() < column.offset() + column.size()) + throw std::invalid_argument("MO exact decimal child shape mismatch"); + validity.masks[i] = child.null_mask(); + validity.offsets[i] = child.offset(); + children.emplace_back( + child.type(), child.size(), child.head(), nullptr, 0, child.offset()); + } + struct owners { + std::unique_ptr> bad; + uint32_t host_bad{}; + }; + auto owner = std::make_unique(); + try { + owner->bad = std::make_unique>(1, stream, mr); + CUDF_CUDA_TRY(cudaMemsetAsync(owner->bad->data(), 0, sizeof(uint32_t), stream.value())); + if (column.size()) + check_children<<<128, block_size, 0, stream.value()>>>( + view(column, type), validity, column.size(), owner->bad->data()); + CUDF_CUDA_TRY(cudaGetLastError()); + CUDF_CUDA_TRY(cudaMemcpyAsync(&owner->host_bad, + owner->bad->data(), + sizeof(owner->host_bad), + cudaMemcpyDeviceToHost, + stream.value())); + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + if (owner->host_bad) + throw std::invalid_argument("MO exact decimal contains independent child NULLs"); + return cudf::column_view(column.type(), + column.size(), + nullptr, + column.null_mask(), + column.null_count(), + column.offset(), + std::move(children)); + } catch (...) { + rethrow_quiescent(owner, stream); + } +} +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_aggregate_gpu.hpp b/src/numeric/decimal_aggregate_gpu.hpp new file mode 100644 index 0000000000..f9670dedca --- /dev/null +++ b/src/numeric/decimal_aggregate_gpu.hpp @@ -0,0 +1,51 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include "numeric/decimal_aggregate.hpp" +#include "numeric/exact_decimal_gpu.hpp" + +namespace sirius::mo_decimal { +struct aggregate_columns { + std::unique_ptr values; + std::unique_ptr counts; +}; + +// Offsets cover dense groups in the permutation (or original rows when order +// is empty). Empty counts means initial rows; otherwise these are raw states. +aggregate_columns reduce_aggregate_columns(aggregate_op op, + cudf::column_view values, + decimal_type input, + uint16_t state_bits, + cudf::column_view counts, + cudf::column_view order, + cudf::size_type const* offsets, + cudf::size_type groups, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); +std::unique_ptr finalize_aggregate_columns(aggregate_op op, + cudf::column_view values, + cudf::column_view counts, + decimal_type input, + decimal_type output, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); +// Build the same scale-independent key for equality partitioning and probing. +std::unique_ptr make_equality_key(cudf::column_view values, + decimal_type type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); +// cuDF gather/selection can superimpose parent validity on STRUCT children. +// Verify there are no independent child NULLs before removing those masks. +std::unique_ptr restore_decimal_validity(std::unique_ptr column, + decimal_type type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); +cudf::column_view canonical_decimal_view(cudf::column_view column, + decimal_type type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); +void make_group_offsets(cudf::column_view labels, + cudf::column_view order, + cudf::size_type* offsets, + cudf::size_type groups, + rmm::cuda_stream_view stream); +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_aggregate_layout.cpp b/src/numeric/decimal_aggregate_layout.cpp new file mode 100644 index 0000000000..f087165bae --- /dev/null +++ b/src/numeric/decimal_aggregate_layout.cpp @@ -0,0 +1,440 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "numeric/decimal_aggregate_layout.hpp" + +#include "cudf/cudf_utils.hpp" +#include "numeric/decimal_aggregate_gpu.hpp" +#include "op/aggregate/group_key_labels.hpp" +#include "pipeline/gpu_stream_quiescence_error.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +namespace sirius::mo_decimal { +namespace { +logical_type ordinary_sum_type(logical_type const& input) +{ + if (input.id() == type_id::FLOAT) return logical_type::make(type_id::DOUBLE); + if (input.is_integer() && input.fixed_width_byte_size() < 8) + return logical_type::make(input.id() == type_id::UTINYINT || input.id() == type_id::USMALLINT || + input.id() == type_id::UINTEGER + ? type_id::UBIGINT + : type_id::BIGINT); + return input; +} +} // namespace +std::shared_ptr make_aggregate_layout( + duckdb::vector> const& groups, + duckdb::vector> const& aggregates) +{ + bool needed = false; + for (auto const& group : groups) + needed |= group && group->return_type().is_mo_decimal(); + for (auto const& value : aggregates) { + auto const& aggr = ast::require_aggregate(value.get(), "MO aggregate layout"); + needed |= is_decimal_aggregate(aggr.function()); + for (auto const& child : aggr.arguments()) + needed |= child && child->return_type().is_mo_decimal(); + } + if (!needed) return {}; + auto layout = std::make_shared(); + for (auto const& group : groups) { + auto const& ref = ast::require_reference(group.get(), "MO aggregate group"); + layout->group_indices.push_back(ref.column_index); + layout->group_types.push_back(ref.return_type()); + } + size_t next = groups.size(); + for (auto const& value : aggregates) { + auto const& aggr = ast::require_aggregate(value.get(), "MO aggregate layout"); + aggregate_slot slot{aggr.function(), -1, {}, aggr.return_type(), {}, aggr.distinct(), next}; + auto const& args = aggr.arguments(); + if (args.empty()) { + if (slot.id != aggregate_id::count_star) + throw std::invalid_argument("aggregate operand missing"); + } else if (args.size() != 1 || !args[0]) + throw std::invalid_argument("invalid aggregate arity"); + else if (args[0]->is_reference()) { + slot.input = args[0]->as_reference().column_index; + slot.input_type = args[0]->return_type(); + } else if (slot.distinct && slot.id == aggregate_id::count && args[0]->is_function_call() && + args[0]->as_function_call().function() == function_id::struct_pack) { + for (auto const& child : args[0]->as_function_call().arguments()) + slot.struct_inputs.push_back( + ast::require_reference(child.get(), "COUNT DISTINCT tuple").column_index); + slot.input_type = logical_type::make(type_id::STRUCT); + } else + throw std::invalid_argument("aggregate operand was not materialized"); + if (slot.exact()) { + if (slot.distinct) throw std::invalid_argument("MO aggregate DISTINCT is unsupported"); + validate_aggregate_signature(aggregate_operation(slot.id), slot.input_type, slot.output_type); + } else { + if (slot.input_type.is_mo_decimal() && slot.id != aggregate_id::count) + throw std::invalid_argument("ordinary aggregate cannot infer an MO decimal signature"); + switch (slot.id) { + case aggregate_id::sum: + case aggregate_id::sum_no_overflow: + case aggregate_id::avg: + case aggregate_id::min: + case aggregate_id::max: + case aggregate_id::count: + case aggregate_id::count_star: break; + default: throw std::invalid_argument("unsupported aggregate in MO exact plan"); + } + if (slot.distinct && slot.id != aggregate_id::count) + throw std::invalid_argument("unsupported DISTINCT aggregate"); + if (slot.distinct && groups.empty()) + throw std::invalid_argument("ungrouped DISTINCT aggregate is unsupported"); + } + next += slot.exact() || slot.avg() ? 2 : 1; + layout->slots.push_back(std::move(slot)); + } + return layout; +} +duckdb::vector aggregate_layout::local_types() const +{ + auto result = group_types; + for (auto const& slot : slots) { + if (slot.exact()) { + auto input = aggregate_input_type(slot.input_type); + result.push_back(logical_type::make_mo_decimal( + {slot.state_bits(), static_cast(slot.state_bits() == 256 ? 76 : 38), input.scale}, + false)); + result.push_back(logical_type::make(type_id::BIGINT)); + } else if (slot.distinct) + result.push_back(logical_type::make(type_id::LIST)); + else { + result.push_back(slot.avg() ? ordinary_sum_type(slot.input_type) : slot.output_type); + if (slot.avg()) result.push_back(logical_type::make(type_id::BIGINT)); + } + } + return result; +} +namespace { +struct owners { + std::unique_ptr combined; + std::optional keys; + std::unique_ptr order, sorted_labels; + std::unique_ptr> offsets; + std::unique_ptr grouping; + std::vector> columns, temporaries, finalized; + std::vector> scalars; + std::vector> intermediates; + std::optional pending; + cudf::size_type ungrouped_offsets[2]{}; +}; +template +std::unique_ptr aggregation(cudf::aggregation::Kind kind) +{ + switch (kind) { + case cudf::aggregation::Kind::SUM: return cudf::make_sum_aggregation(); + case cudf::aggregation::Kind::MIN: return cudf::make_min_aggregation(); + case cudf::aggregation::Kind::MAX: return cudf::make_max_aggregation(); + case cudf::aggregation::Kind::COUNT_ALL: + return cudf::make_count_aggregation(cudf::null_policy::INCLUDE); + case cudf::aggregation::Kind::COUNT_VALID: + return cudf::make_count_aggregation(cudf::null_policy::EXCLUDE); + case cudf::aggregation::Kind::COLLECT_SET: + return cudf::make_collect_set_aggregation(cudf::null_policy::EXCLUDE); + case cudf::aggregation::Kind::MERGE_SETS: return cudf::make_merge_sets_aggregation(); + default: throw std::invalid_argument("unsupported ordinary aggregate kind"); + } +} +std::unique_ptr ordinary(owners& owner, + cudf::column_view input, + cudf::aggregation::Kind kind, + cudf::data_type result_type, + bool grouped, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + std::unique_ptr result; + if (kind == cudf::aggregation::Kind::SUM && input.type() != result_type && + cudf::is_fixed_width(input.type()) && cudf::is_fixed_width(result_type)) { + owner.temporaries.push_back(cudf::cast(input, result_type, stream, mr)); + input = owner.temporaries.back()->view(); + } + if (grouped) { + auto gathered = cudf::gather(cudf::table_view{{input}}, + owner.order->view(), + cudf::out_of_bounds_policy::DONT_CHECK, + stream, + mr) + ->release(); + owner.temporaries.push_back(std::move(gathered[0])); + std::vector requests(1); + requests[0].values = owner.temporaries.back()->view(); + requests[0].aggregations.push_back(aggregation(kind)); + auto out = owner.grouping->aggregate(requests, stream, mr); + owner.intermediates.push_back(std::move(out.first)); + result = std::move(out.second[0].results[0]); + } else { + owner.scalars.push_back( + cudf::reduce(input, *aggregation(kind), result_type, stream, mr)); + result = cudf::make_column_from_scalar(*owner.scalars.back(), 1, stream, mr); + } + if (result->type() != result_type && kind != cudf::aggregation::Kind::COLLECT_SET && + kind != cudf::aggregation::Kind::MERGE_SETS) { + owner.temporaries.push_back(std::move(result)); + result = cudf::cast(owner.temporaries.back()->view(), result_type, stream, mr); + } + return result; +} +std::unique_ptr finish(cudf::table_view input, + aggregate_layout const& layout, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr, + std::unique_ptr owner = std::make_unique()) +{ + try { + owner->finalized.reserve(layout.group_types.size() + layout.slots.size()); + owner->temporaries.reserve(2 * layout.slots.size()); + for (size_t i = 0; i < layout.group_types.size(); ++i) { + owner->finalized.push_back(std::make_unique(input.column(i), stream, mr)); + if (layout.group_types[i].is_mo_decimal()) + owner->finalized.back() = restore_decimal_validity( + std::move(owner->finalized.back()), layout.group_types[i].mo_decimal_type(), stream, mr); + } + for (auto const& slot : layout.slots) { + auto source = input.column(slot.state_index); + if (slot.exact()) + owner->finalized.push_back(finalize_aggregate_columns(aggregate_operation(slot.id), + source, + input.column(slot.state_index + 1), + aggregate_input_type(slot.input_type), + slot.output_type.mo_decimal_type(), + stream, + mr)); + else if (slot.avg()) + owner->finalized.push_back(cudf::binary_operation(source, + input.column(slot.state_index + 1), + cudf::binary_operator::DIV, + get_cudf_type(slot.output_type), + stream, + mr)); + else if (slot.distinct) { + owner->temporaries.push_back( + cudf::lists::count_elements(cudf::lists_column_view(source), stream, mr)); + owner->finalized.push_back(cudf::cast( + owner->temporaries.back()->view(), get_cudf_type(slot.output_type), stream, mr)); + } else + owner->finalized.push_back(std::make_unique(source, stream, mr)); + } + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + return std::make_unique(std::move(owner->finalized)); + } catch (...) { + auto error = std::current_exception(); + if (cudaStreamSynchronize(stream.value()) != cudaSuccess) { + (void)owner.release(); + throw pipeline::gpu_stream_quiescence_error( + "MO aggregate finalization could not prove quiescence"); + } + std::rethrow_exception(error); + } +} +std::unique_ptr compute(cudf::table_view input, + aggregate_layout const& layout, + bool merging, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr, + std::unique_ptr owner = std::make_unique()) +{ + try { + owner->columns.reserve(layout.group_types.size() + 2 * layout.slots.size()); + owner->temporaries.reserve(4 * layout.slots.size()); + owner->scalars.reserve(2 * layout.slots.size()); + owner->intermediates.reserve(2 * layout.slots.size()); + bool grouped = !layout.group_indices.empty(); + cudf::size_type groups = 1; + std::vector group_views; + for (size_t i = 0; i < layout.group_indices.size(); ++i) { + auto key = input.column(merging ? i : layout.group_indices[i]); + if (layout.group_types[i].is_mo_decimal()) + key = canonical_decimal_view(key, layout.group_types[i].mo_decimal_type(), stream, mr); + group_views.push_back(key); + } + if (grouped) { + owner->keys.emplace( + op::detail::make_group_key_labels(cudf::table_view(group_views), stream, mr)); + groups = owner->keys->sorted_unique_keys->num_rows(); + owner->order = + cudf::sorted_order(cudf::table_view{{owner->keys->labels->view()}}, {}, {}, stream, mr); + auto sorted = cudf::gather(cudf::table_view{{owner->keys->labels->view()}}, + owner->order->view(), + cudf::out_of_bounds_policy::DONT_CHECK, + stream, + mr) + ->release(); + owner->sorted_labels = std::move(sorted[0]); + owner->grouping = + std::make_unique(cudf::table_view{{owner->sorted_labels->view()}}, + cudf::null_policy::INCLUDE, + cudf::sorted::YES); + owner->offsets = + std::make_unique>(groups + 1, stream, mr); + make_group_offsets( + owner->keys->labels->view(), owner->order->view(), owner->offsets->data(), groups, stream); + auto keys = owner->keys->sorted_unique_keys->release(); + for (size_t i = 0; i < keys.size(); ++i) { + if (layout.group_types[i].is_mo_decimal()) + keys[i] = restore_decimal_validity( + std::move(keys[i]), layout.group_types[i].mo_decimal_type(), stream, mr); + owner->columns.push_back(std::move(keys[i])); + } + } else { + owner->offsets = std::make_unique>(2, stream, mr); + owner->ungrouped_offsets[1] = input.num_rows(); + CUDF_CUDA_TRY(cudaMemcpyAsync(owner->offsets->data(), + owner->ungrouped_offsets, + sizeof(owner->ungrouped_offsets), + cudaMemcpyHostToDevice, + stream.value())); + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + } + for (auto const& slot : layout.slots) { + auto source_index = merging ? static_cast(slot.state_index) : slot.input; + cudf::column_view source; + if (!merging && slot.id == aggregate_id::count_star) + source = input.column(0); + else if (source_index >= 0) + source = input.column(source_index); + else { + std::vector> children; + for (int index : slot.struct_inputs) + children.push_back(std::make_unique(input.column(index), stream, mr)); + owner->temporaries.push_back(cudf::make_structs_column( + input.num_rows(), std::move(children), 0, rmm::device_buffer{}, stream, mr)); + source = owner->temporaries.back()->view(); + } + if (slot.exact()) { + auto original = aggregate_input_type(slot.input_type); + auto descriptor = merging + ? decimal_type{slot.state_bits(), + static_cast(slot.state_bits() == 256 ? 76 : 38), + original.scale} + : original; + if (!merging && !slot.input_type.is_mo_decimal()) + source = cudf::column_view(cudf::data_type{cudf::type_id::DECIMAL64, 0}, + source.size(), + source.head(), + source.null_mask(), + source.null_count(), + source.offset()); + owner->pending.emplace(reduce_aggregate_columns( + aggregate_operation(slot.id), + source, + descriptor, + slot.state_bits(), + merging ? input.column(slot.state_index + 1) : cudf::column_view{}, + owner->order ? owner->order->view() : cudf::column_view{}, + owner->offsets->data(), + groups, + stream, + mr)); + owner->columns.push_back(std::move(owner->pending->values)); + owner->columns.push_back(std::move(owner->pending->counts)); + owner->pending.reset(); + continue; + } + auto type = get_cudf_type(slot.output_type); + if (slot.distinct) { + owner->columns.push_back(ordinary( + *owner, + source, + merging ? cudf::aggregation::Kind::MERGE_SETS : cudf::aggregation::Kind::COLLECT_SET, + cudf::data_type{cudf::type_id::LIST}, + grouped, + stream, + mr)); + } else if (slot.avg()) { + auto sum_type = get_cudf_type(ordinary_sum_type(slot.input_type)); + owner->columns.push_back( + ordinary(*owner, source, cudf::aggregation::Kind::SUM, sum_type, grouped, stream, mr)); + owner->columns.push_back( + ordinary(*owner, + merging ? input.column(slot.state_index + 1) : source, + merging ? cudf::aggregation::Kind::SUM : cudf::aggregation::Kind::COUNT_VALID, + cudf::data_type{cudf::type_id::INT64}, + grouped, + stream, + mr)); + } else { + auto kind = + slot.id == aggregate_id::min ? cudf::aggregation::Kind::MIN + : slot.id == aggregate_id::max ? cudf::aggregation::Kind::MAX + : slot.id == aggregate_id::count_star && !merging ? cudf::aggregation::Kind::COUNT_ALL + : slot.id == aggregate_id::count && !merging ? cudf::aggregation::Kind::COUNT_VALID + : cudf::aggregation::Kind::SUM; + owner->columns.push_back(ordinary(*owner, source, kind, type, grouped, stream, mr)); + } + } + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + if (merging) { + owner->combined = std::make_unique(std::move(owner->columns)); + auto states = owner->combined->view(); + return finish(states, layout, stream, mr, std::move(owner)); + } + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + return std::make_unique(std::move(owner->columns)); + } catch (...) { + auto error = std::current_exception(); + if (cudaStreamSynchronize(stream.value()) != cudaSuccess) { + (void)owner.release(); + throw pipeline::gpu_stream_quiescence_error("MO aggregate layout could not prove quiescence"); + } + std::rethrow_exception(error); + } +} +} // namespace +std::unique_ptr local_aggregate_table(cudf::table_view input, + aggregate_layout const& layout, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + return compute(input, layout, false, stream, mr); +} +std::unique_ptr merge_aggregate_tables(std::vector const& inputs, + aggregate_layout const& layout, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + if (inputs.empty()) + throw std::invalid_argument("MO aggregate merge requires a schema-bearing partial batch"); + // A single local partial already has one row per group. It still needs + // final precision/NULL publication, but no second grouping or reduction. + if (inputs.size() == 1) return finish(inputs[0], layout, stream, mr); + auto owner = std::make_unique(); + try { + owner->combined = cudf::concatenate(inputs, stream, mr); + // Concatenation can superimpose validity on Decimal256 group keys. The + // accumulator coefficients themselves are always nonnullable raw states. + owner->columns = owner->combined->release(); + for (size_t i = 0; i < layout.group_types.size(); ++i) + if (layout.group_types[i].is_mo_decimal()) + owner->columns[i] = restore_decimal_validity( + std::move(owner->columns[i]), layout.group_types[i].mo_decimal_type(), stream, mr); + owner->combined = std::make_unique(std::move(owner->columns)); + auto view = owner->combined->view(); + return compute(view, layout, true, stream, mr, std::move(owner)); + } catch (...) { + auto error = std::current_exception(); + if (cudaStreamSynchronize(stream.value()) != cudaSuccess) { + (void)owner.release(); + throw pipeline::gpu_stream_quiescence_error( + "MO aggregate concatenation could not prove quiescence"); + } + std::rethrow_exception(error); + } +} +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_aggregate_layout.hpp b/src/numeric/decimal_aggregate_layout.hpp new file mode 100644 index 0000000000..757a7dca55 --- /dev/null +++ b/src/numeric/decimal_aggregate_layout.hpp @@ -0,0 +1,46 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include "expression/ast/node.hpp" +#include "numeric/decimal_aggregate_bind.hpp" + +#include + +#include +#include + +namespace sirius::mo_decimal { +struct aggregate_slot { + aggregate_id id; + int input; + logical_type input_type, output_type; + std::vector struct_inputs; + bool distinct{}; + std::size_t state_index{}; + bool exact() const { return is_decimal_aggregate(id); } + bool avg() const { return id == aggregate_id::avg || id == aggregate_id::mo_decimal_avg; } + uint16_t state_bits() const + { + return exact() + ? static_cast(std::max(128, output_type.mo_decimal_type().bits)) + : 0; + } +}; +struct aggregate_layout { + std::vector group_indices; + duckdb::vector group_types; + std::vector slots; + duckdb::vector local_types() const; +}; +// Returns null when the unchanged ordinary aggregate implementation applies. +std::shared_ptr make_aggregate_layout( + duckdb::vector> const& groups, + duckdb::vector> const& aggregates); +std::unique_ptr local_aggregate_table(cudf::table_view input, + aggregate_layout const& layout, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); +std::unique_ptr merge_aggregate_tables(std::vector const& inputs, + aggregate_layout const& layout, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_error.hpp b/src/numeric/decimal_error.hpp index 1f96771c5e..647bda54dc 100644 --- a/src/numeric/decimal_error.hpp +++ b/src/numeric/decimal_error.hpp @@ -8,7 +8,7 @@ class numeric_error final : public std::runtime_error { public: explicit numeric_error(decimal_error code) : std::runtime_error(code == decimal_error::invalid_input - ? "MO exact-decimal checked cast failed" + ? "MO exact-decimal checked numeric operation failed" : "MO exact-decimal arithmetic is out of range"), code_(code) { diff --git a/src/numeric/decimal_functions.cpp b/src/numeric/decimal_functions.cpp index 8eb2425c9c..f9d8573d08 100644 --- a/src/numeric/decimal_functions.cpp +++ b/src/numeric/decimal_functions.cpp @@ -3,6 +3,7 @@ #include "expression/ast/node.hpp" #include "helper/type_conversions.hpp" +#include "numeric/decimal_plan.hpp" #include "numeric/exact_decimal.hpp" #include @@ -142,6 +143,7 @@ duckdb::unique_ptr bind_scalar( std::vector inputs; duckdb::vector carriers; for (std::size_t i = 0; i + 1 < arguments.size(); ++i) { + restore_exact_bound_types(arguments[i]); inputs.push_back(sirius::from_duckdb(arguments[i]->return_type)); carriers.push_back(arguments[i]->return_type); } diff --git a/src/numeric/decimal_import.cpp b/src/numeric/decimal_import.cpp index 88a1604017..20028c4d49 100644 --- a/src/numeric/decimal_import.cpp +++ b/src/numeric/decimal_import.cpp @@ -2,6 +2,7 @@ #include "numeric/decimal_import.hpp" #include "from_substrait.hpp" +#include "numeric/decimal_aggregate_bind.hpp" #include "numeric/decimal_functions.hpp" #include "numeric/exact_decimal.hpp" @@ -85,7 +86,7 @@ class importer final : public duckdb::SubstraitExtensionHandler { auto type = descriptor(user); auto const& any = user.value(); auto const& payload = any.value(); - if (any.type_url() != literal_type_url) + if (any.type_url() != literal_type_url || payload.size() > (16u << 20)) throw std::invalid_argument("invalid MO exact-decimal literal encoding"); // Parse the approved protobuf using the pinned reader, including ordinary // unknown-field and singular-field semantics. Only the final coefficient @@ -135,8 +136,42 @@ class importer final : public duckdb::SubstraitExtensionHandler { return duckdb::make_uniq(std::string(to_duckdb_function_name(*id)), std::move(children)); } + duckdb::unique_ptr Aggregate( + duckdb::ClientContext&, + duckdb::SubstraitExtensionIdentity const& identity, + substrait::AggregateFunction const& aggregate, + duckdb::vector> children) const override + { + if (!Handles(identity) || !aggregate.has_output_type() || children.size() != 1 || + aggregate.arguments_size() != 1 || aggregate.options_size() || aggregate.sorts_size() || + aggregate.invocation() != substrait::AggregateFunction::AGGREGATION_INVOCATION_ALL || + aggregate.phase() != substrait::AGGREGATION_PHASE_INITIAL_TO_RESULT) + throw std::invalid_argument("unsupported MO exact-decimal aggregate signature or modifiers"); + auto id = from_duckdb_aggregate_name("__sirius_" + identity.name); + if (!id || !is_decimal_aggregate(*id)) + throw std::invalid_argument("unknown MO aggregate identity"); + auto output = result_type(aggregate.output_type(), anchors_); + children.push_back( + duckdb::make_uniq(duckdb::Value(duckdb_type(output)))); + return duckdb::make_uniq(std::string(to_duckdb_aggregate_name(*id)), + std::move(children)); + } }; } // namespace +logical_type import_result_type(substrait::Type const& type, std::vector const& anchors) +{ + return result_type(type, anchors); +} +bool uses_exact_decimal(std::string_view bytes) +{ + if (bytes.size() > (16u << 20)) throw std::invalid_argument("MO numeric plan exceeds 16 MiB"); + substrait::Plan plan; + if (!plan.ParseFromArray(bytes.data(), static_cast(bytes.size()))) + throw std::invalid_argument("invalid MO numeric plan"); + for (auto const& uri : plan.extension_urns()) + if (uri.urn() == extension_uri) return true; + return false; +} duckdb::shared_ptr make_import_handler(std::string_view bytes) { if (bytes.size() > (16u << 20)) throw std::invalid_argument("MO numeric plan exceeds 16 MiB"); diff --git a/src/numeric/decimal_import.hpp b/src/numeric/decimal_import.hpp index 32da20f14c..6797efb61d 100644 --- a/src/numeric/decimal_import.hpp +++ b/src/numeric/decimal_import.hpp @@ -1,14 +1,21 @@ /* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ #pragma once +#include "numeric/decimal_types.hpp" + #include #include namespace duckdb { class SubstraitExtensionHandler; } +namespace substrait { +class Type; +} namespace sirius::mo_decimal { inline constexpr char extension_uri[] = "urn:matrixone:sirius:exact-decimal:v1"; inline constexpr char literal_type_url[] = "type.googleapis.com/matrixone.sirius.numeric.v1.ExactDecimalLiteral"; duckdb::shared_ptr make_import_handler(std::string_view plan); +bool uses_exact_decimal(std::string_view plan); +logical_type import_result_type(substrait::Type const& type, std::vector const& anchors); } // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_plan.cpp b/src/numeric/decimal_plan.cpp new file mode 100644 index 0000000000..91e9dd6580 --- /dev/null +++ b/src/numeric/decimal_plan.cpp @@ -0,0 +1,625 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "numeric/decimal_plan.hpp" + +#include "helper/type_conversions.hpp" +#include "numeric/decimal_aggregate_bind.hpp" +#include "numeric/decimal_functions.hpp" +#include "numeric/decimal_import.hpp" +#include "substrait/plan.pb.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include + +namespace sirius::mo_decimal { +namespace { +bool same_decimal(logical_type const& a, logical_type const& b) +{ + if (!a.is_mo_decimal() || !b.is_mo_decimal()) return false; + auto x = a.mo_decimal_type(), y = b.mo_decimal_type(); + return x.bits == y.bits && x.precision == y.precision && x.scale == y.scale; +} +struct wire_validator { + substrait::Plan plan; + std::vector anchors; + std::unordered_map functions; + std::unordered_set exact_functions; + uint32_t cast_anchor{}, uri{}; + std::vector> relations; + std::size_t schema_limit{}, schema_bytes{}; + void charge(std::size_t cells) + { + if (cells > (schema_limit - schema_bytes) / sizeof(logical_type)) throw std::bad_alloc(); + schema_bytes += cells * sizeof(logical_type); + } + logical_type type(substrait::Type const& t) + { + if (t.has_user_defined()) return import_result_type(t, anchors); + using K = substrait::Type::KindCase; + type_id id; + bool nullable = false; + switch (t.kind_case()) { + case K::kBool: + id = type_id::BOOLEAN; + nullable = t.bool_().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + case K::kI8: + id = type_id::TINYINT; + nullable = t.i8().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + case K::kI16: + id = type_id::SMALLINT; + nullable = t.i16().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + case K::kI32: + id = type_id::INTEGER; + nullable = t.i32().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + case K::kI64: + id = type_id::BIGINT; + nullable = t.i64().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + case K::kFp32: + id = type_id::FLOAT; + nullable = t.fp32().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + case K::kFp64: + id = type_id::DOUBLE; + nullable = t.fp64().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + case K::kString: + id = type_id::VARCHAR; + nullable = t.string().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + case K::kVarchar: + id = type_id::VARCHAR; + nullable = t.varchar().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + case K::kBinary: + id = type_id::VARCHAR; + nullable = t.binary().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + case K::kDate: + id = type_id::DATE; + nullable = t.date().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + case K::kPrecisionTimestamp: + id = type_id::TIMESTAMP; + nullable = t.precision_timestamp().nullability() == substrait::Type::NULLABILITY_NULLABLE; + break; + default: throw std::invalid_argument("unsupported type in exact-decimal closure"); + } + return logical_type::make(id).with_nullability(nullable); + } + substrait::Type wire_type(logical_type const& logical) + { + substrait::Type t; + auto user = t.mutable_user_defined(); + user->set_type_reference(anchors.front()); + auto d = logical.mo_decimal_type(); + for (auto value : {int(d.bits), int(d.precision), int(d.scale)}) + user->add_type_parameters()->set_integer(value); + user->set_nullability(logical.nullability() == 2 ? substrait::Type::NULLABILITY_NULLABLE + : substrait::Type::NULLABILITY_REQUIRED); + return t; + } + void annotate(substrait::Expression& expression, logical_type const& logical) + { + substrait::Expression original; + original.Swap(&expression); + auto function = expression.mutable_scalar_function(); + function->set_function_reference(cast_anchor); + *function->mutable_output_type() = wire_type(logical); + function->add_arguments()->mutable_value()->Swap(&original); + } + logical_type expression(substrait::Expression& e, std::vector const& input) + { + if (e.has_selection()) { + auto const& s = e.selection(); + if (!s.has_direct_reference() || !s.direct_reference().has_struct_field() || + s.direct_reference().struct_field().has_child() || !s.has_root_reference()) + throw std::invalid_argument("unsupported exact-plan field reference"); + auto index = s.direct_reference().struct_field().field(); + if (index < 0 || static_cast(index) >= input.size()) + throw std::invalid_argument("exact-plan field outside schema"); + auto result = input[index]; + if (result.is_mo_decimal()) annotate(e, result); + return result; + } + if (e.has_literal()) { + auto const& l = e.literal(); + if (l.has_null()) return type(l.null()); + if (l.has_user_defined()) { + substrait::Type t; + auto u = t.mutable_user_defined(); + u->set_type_reference(l.user_defined().type_reference()); + *u->mutable_type_parameters() = l.user_defined().type_parameters(); + u->set_nullability(l.nullable() ? substrait::Type::NULLABILITY_NULLABLE + : substrait::Type::NULLABILITY_REQUIRED); + return type(t); + } + type_id id = l.has_boolean() ? type_id::BOOLEAN + : l.has_i8() ? type_id::TINYINT + : l.has_i16() ? type_id::SMALLINT + : l.has_i32() ? type_id::INTEGER + : l.has_i64() ? type_id::BIGINT + : l.has_fp32() ? type_id::FLOAT + : l.has_fp64() ? type_id::DOUBLE + : l.has_date() ? type_id::DATE + : type_id::VARCHAR; + return logical_type::make(id).with_nullability(l.nullable()); + } + if (e.has_scalar_function()) { + auto function = e.mutable_scalar_function(); + std::vector args; + charge(function->arguments_size()); + for (auto& argument : *function->mutable_arguments()) + if (argument.has_value()) args.push_back(expression(*argument.mutable_value(), input)); + auto result = type(function->output_type()); + auto found = functions.find(function->function_reference()); + if (found == functions.end()) + throw std::invalid_argument("unknown exact-plan function anchor"); + if (exact_functions.contains(function->function_reference())) { + auto id = from_duckdb_function_name("__sirius_" + found->second); + if (!id || !is_decimal_function(*id)) + throw std::invalid_argument("invalid MO scalar identity"); + validate_signature(operation(*id), args, result); + } else { + bool exact = result.is_mo_decimal(); + for (auto const& a : args) + exact |= a.is_mo_decimal(); + auto name = found->second.substr(0, found->second.find(':')); + if (exact && name != "coalesce" && name != "is_null" && name != "is_not_null") + throw std::invalid_argument("ordinary function cannot coerce an MO exact decimal"); + if (name == "coalesce" && result.is_mo_decimal()) + for (auto const& a : args) + if (!same_decimal(a, result)) + throw std::invalid_argument("MO COALESCE descriptors differ"); + if (name == "coalesce" && result.is_mo_decimal()) annotate(e, result); + } + return result; + } + if (e.has_if_then()) { + auto branch = e.mutable_if_then(); + std::optional result; + bool nullable = false; + for (auto& pair : *branch->mutable_ifs()) { + if (expression(*pair.mutable_if_(), input).id() != type_id::BOOLEAN) + throw std::invalid_argument("CASE condition is not boolean"); + auto value = expression(*pair.mutable_then(), input); + nullable |= value.nullability() == 2; + if (result && (result->is_mo_decimal() || value.is_mo_decimal()) && + !same_decimal(*result, value)) + throw std::invalid_argument("MO CASE branches require matching exact descriptors"); + result = value; + } + if (!result || !branch->has_else_()) + throw std::invalid_argument("CASE requires typed branches and ELSE"); + auto other = expression(*branch->mutable_else_(), input); + nullable |= other.nullability() == 2; + if ((result->is_mo_decimal() || other.is_mo_decimal()) && !same_decimal(*result, other)) + throw std::invalid_argument("MO CASE ELSE descriptor mismatch"); + auto output = result->with_nullability(nullable); + if (output.is_mo_decimal()) annotate(e, output); + return output; + } + if (e.has_cast()) { + auto child = expression(*e.mutable_cast()->mutable_input(), input), + target = type(e.cast().type()); + if (child.is_mo_decimal() || target.is_mo_decimal()) + throw std::invalid_argument("MO casts require the checked versioned function"); + return target; + } + if (e.has_singular_or_list()) { + auto list = e.mutable_singular_or_list(); + auto probe = expression(*list->mutable_value(), input); + if (probe.is_mo_decimal()) + throw std::invalid_argument("exact decimal IN lists are unsupported"); + for (auto& value : *list->mutable_options()) + if (expression(value, input).is_mo_decimal()) + throw std::invalid_argument("mixed exact IN list"); + return logical_type::make(type_id::BOOLEAN).with_nullability(true); + } + if (e.has_nested()) { + if (!e.nested().has_struct_()) + throw std::invalid_argument("unsupported nested exact-plan value"); + for (auto& field : *e.mutable_nested()->mutable_struct_()->mutable_fields()) + if (expression(field, input).is_mo_decimal()) + throw std::invalid_argument("nested exact-decimal tuple is unsupported"); + return logical_type::make(type_id::STRUCT); + } + throw std::invalid_argument("unsupported expression in exact-plan closure"); + } + std::vector relation(substrait::Rel& rel) + { + std::vector result; + substrait::RelCommon const* common = nullptr; + if (rel.has_read()) { + auto read = rel.mutable_read(); + common = &read->common(); + charge(read->base_schema().struct_().types_size()); + for (auto const& t : read->base_schema().struct_().types()) + result.push_back(type(t)); + if (read->has_filter() && + expression(*read->mutable_filter(), result).id() != type_id::BOOLEAN) + throw std::invalid_argument("read filter is not boolean"); + if (read->has_best_effort_filter()) + throw std::invalid_argument("exact-plan best-effort read filters are unsupported"); + if (read->has_projection()) { + if (!read->projection().has_select() || !read->projection().select().struct_items_size()) + throw std::invalid_argument("exact-plan read projection requires a flat selection"); + charge(read->projection().select().struct_items_size()); + std::vector projected; + for (auto const& item : read->projection().select().struct_items()) { + auto index = item.field(); + if (item.has_child() || index < 0 || static_cast(index) >= result.size()) + throw std::invalid_argument("invalid exact-plan read projection field"); + projected.push_back(result[index]); + } + result = std::move(projected); + } + } else if (rel.has_project()) { + auto p = rel.mutable_project(); + common = &p->common(); + result = relation(*p->mutable_input()); + charge(result.size() + p->expressions_size()); + auto input = result; + for (auto& e : *p->mutable_expressions()) + result.push_back(expression(e, input)); + } else if (rel.has_filter()) { + auto p = rel.mutable_filter(); + common = &p->common(); + result = relation(*p->mutable_input()); + if (expression(*p->mutable_condition(), result).id() != type_id::BOOLEAN) + throw std::invalid_argument("filter is not boolean"); + } else if (rel.has_fetch()) { + auto p = rel.mutable_fetch(); + common = &p->common(); + result = relation(*p->mutable_input()); + } else if (rel.has_sort()) { + auto p = rel.mutable_sort(); + common = &p->common(); + result = relation(*p->mutable_input()); + for (auto& field : *p->mutable_sorts()) + expression(*field.mutable_expr(), result); + } else if (rel.has_aggregate()) { + auto p = rel.mutable_aggregate(); + common = &p->common(); + auto input = relation(*p->mutable_input()); + charge(p->grouping_expressions_size() + p->measures_size()); + if (p->groupings_size() > 1) + throw std::invalid_argument("exact-plan grouping sets are unsupported"); + if (p->grouping_expressions_size()) { + if (p->groupings_size() != 1 || + p->groupings(0).expression_references_size() != p->grouping_expressions_size()) + throw std::invalid_argument("exact-plan grouping set must cover every key"); + std::unordered_set keys; + for (auto index : p->groupings(0).expression_references()) + if (index >= static_cast(p->grouping_expressions_size()) || + !keys.insert(index).second) + throw std::invalid_argument("invalid exact-plan grouping key reference"); + } + for (auto& e : *p->mutable_grouping_expressions()) + result.push_back(expression(e, input)); + for (auto& measure : *p->mutable_measures()) { + auto f = measure.mutable_measure(); + if (measure.has_filter() || f->sorts_size() || f->options_size() || + f->phase() != substrait::AGGREGATION_PHASE_INITIAL_TO_RESULT || + (f->invocation() != substrait::AggregateFunction::AGGREGATION_INVOCATION_ALL && + f->invocation() != substrait::AggregateFunction::AGGREGATION_INVOCATION_DISTINCT)) + throw std::invalid_argument("unsupported exact-plan aggregate modifiers"); + std::vector args; + charge(f->arguments_size()); + for (auto& arg : *f->mutable_arguments()) { + if (!arg.has_value()) throw std::invalid_argument("invalid aggregate argument"); + args.push_back(expression(*arg.mutable_value(), input)); + } + auto output = type(f->output_type()); + auto found = functions.find(f->function_reference()); + if (found == functions.end()) throw std::invalid_argument("missing aggregate anchor"); + if (exact_functions.contains(f->function_reference())) { + auto id = from_duckdb_aggregate_name("__sirius_" + found->second); + if (!id || !is_decimal_aggregate(*id) || args.size() != 1 || measure.has_filter()) + throw std::invalid_argument("invalid exact aggregate"); + validate_aggregate_signature(aggregate_operation(*id), args[0], output); + } else { + auto name = found->second.substr(0, found->second.find(':')); + bool exact = output.is_mo_decimal(); + for (auto const& a : args) + exact |= a.is_mo_decimal(); + if (exact && name != "count") + throw std::invalid_argument("ordinary aggregate cannot infer an exact descriptor"); + } + result.push_back(output); + } + } else if (rel.has_join()) { + auto p = rel.mutable_join(); + common = &p->common(); + auto left = relation(*p->mutable_left()), right = relation(*p->mutable_right()); + charge(2 * (left.size() + right.size()) + 1); + auto both = left; + both.insert(both.end(), right.begin(), right.end()); + if (expression(*p->mutable_expression(), both).id() != type_id::BOOLEAN) + throw std::invalid_argument("join condition is not boolean"); + auto kind = p->type(); + if (kind == substrait::JoinRel::JOIN_TYPE_LEFT || + kind == substrait::JoinRel::JOIN_TYPE_LEFT_SINGLE || + kind == substrait::JoinRel::JOIN_TYPE_OUTER) + for (auto& t : right) + t = t.with_nullability(true); + if (kind == substrait::JoinRel::JOIN_TYPE_RIGHT || + kind == substrait::JoinRel::JOIN_TYPE_OUTER) + for (auto& t : left) + t = t.with_nullability(true); + if (kind == substrait::JoinRel::JOIN_TYPE_RIGHT_SEMI || + kind == substrait::JoinRel::JOIN_TYPE_RIGHT_ANTI) + result = right; + else { + result = left; + if (kind == substrait::JoinRel::JOIN_TYPE_LEFT_MARK) + result.push_back(logical_type::make(type_id::BOOLEAN).with_nullability(true)); + else if (kind != substrait::JoinRel::JOIN_TYPE_LEFT_SEMI && + kind != substrait::JoinRel::JOIN_TYPE_LEFT_ANTI) + result.insert(result.end(), right.begin(), right.end()); + } + } else if (rel.has_cross()) { + auto p = rel.mutable_cross(); + common = &p->common(); + result = relation(*p->mutable_left()); + auto right = relation(*p->mutable_right()); + charge(right.size()); + result.insert(result.end(), right.begin(), right.end()); + } else if (rel.has_reference()) { + auto index = rel.reference().subtree_ordinal(); + if (index < 0 || static_cast(index) >= relations.size()) + throw std::invalid_argument("invalid exact-plan relation reference"); + charge(relations[index].size()); + result = relations[index]; + } else if (rel.has_set()) { + auto p = rel.mutable_set(); + common = &p->common(); + for (auto& input : *p->mutable_inputs()) { + auto schema = relation(input); + if (result.empty()) + result = std::move(schema); + else { + if (result.size() != schema.size()) + throw std::invalid_argument("set schema width differs"); + for (size_t i = 0; i < result.size(); ++i) { + if ((result[i].is_mo_decimal() || schema[i].is_mo_decimal()) && result[i] != schema[i]) + throw std::invalid_argument("set exact descriptors differ"); + result[i] = result[i].with_nullability(result[i].nullability() == 2 || + schema[i].nullability() == 2); + } + } + } + } else + throw std::invalid_argument("unsupported exact-plan relation"); + if (result.empty()) + throw std::invalid_argument("exact-plan relations require a nonempty row schema"); + if (common && common->has_emit()) { + charge(common->emit().output_mapping_size()); + std::vector emitted; + for (int index : common->emit().output_mapping()) { + if (index < 0 || static_cast(index) >= result.size()) + throw std::invalid_argument("exact-plan emit outside schema"); + emitted.push_back(result[index]); + } + if (emitted.empty()) + throw std::invalid_argument("exact-plan emit requires a nonempty row schema"); + return emitted; + } + return result; + } +}; +} // namespace +std::string normalize_exact_substrait(std::string_view bytes, std::size_t schema_bytes_limit) +{ + if (bytes.size() > (16u << 20)) throw std::invalid_argument("exact plan exceeds 16 MiB"); + wire_validator v; + v.schema_limit = schema_bytes_limit; + if (!v.plan.ParseFromArray(bytes.data(), static_cast(bytes.size()))) + throw std::invalid_argument("invalid exact plan"); + std::unordered_set uris; + uint32_t maximum = 0; + for (auto const& urn : v.plan.extension_urns()) + if (urn.urn() == extension_uri) { + uris.insert(urn.extension_urn_anchor()); + v.uri = urn.extension_urn_anchor(); + } + for (auto const& ext : v.plan.extensions()) { + if (ext.has_extension_type() && uris.contains(ext.extension_type().extension_urn_reference()) && + ext.extension_type().name() == "mo_exact_decimal") + v.anchors.push_back(ext.extension_type().type_anchor()); + if (ext.has_extension_function()) { + auto const& f = ext.extension_function(); + maximum = std::max(maximum, f.function_anchor()); + if (!v.functions.emplace(f.function_anchor(), f.name()).second) + throw std::invalid_argument("duplicate function anchor"); + if (uris.contains(f.extension_urn_reference())) { + v.exact_functions.insert(f.function_anchor()); + if (f.name() == "mo_decimal_cast") v.cast_anchor = f.function_anchor(); + } + } + } + std::sort(v.anchors.begin(), v.anchors.end()); + if (v.anchors.empty()) + throw std::invalid_argument("exact plan has no canonical type declaration"); + if (!v.cast_anchor) { + if (maximum == UINT32_MAX) throw std::invalid_argument("exact plan function anchor overflow"); + v.cast_anchor = maximum + 1; + auto f = v.plan.add_extensions()->mutable_extension_function(); + f->set_function_anchor(v.cast_anchor); + f->set_extension_urn_reference(v.uri); + f->set_name("mo_decimal_cast"); + } + for (auto& relation : *v.plan.mutable_relations()) + v.relations.push_back(v.relation(relation.has_root() ? *relation.mutable_root()->mutable_input() + : *relation.mutable_rel())); + if (v.plan.ByteSizeLong() > (16u << 20)) + throw std::invalid_argument("normalized exact plan exceeds 16 MiB"); + return v.plan.SerializeAsString(); +} +void restore_exact_bound_types(duckdb::unique_ptr& expression) +{ + duckdb::ExpressionIterator::EnumerateChildren( + *expression, + [](duckdb::unique_ptr& child) { restore_exact_bound_types(child); }); + if (expression->GetExpressionClass() == duckdb::ExpressionClass::BOUND_CAST) { + auto& cast = expression->Cast(); + auto source = from_duckdb_type(cast.child->return_type); + if (source) { + auto plain = duckdb_type(*source); + plain.SetAlias(""); + auto target = from_duckdb_type(cast.return_type); + // All wire-level MO casts were checked versioned functions. These + // BoundCast nodes are DuckDB's conditional alias/nullability coercions; + // restore the original declared branch rather than a generic STRUCT cast. + if (plain == cast.return_type || cast.child->return_type == cast.return_type || + (target && same_decimal(*source, *target))) + expression = std::move(cast.child); + } + } + std::vector branches; + bool coalesce = false; + if (expression->GetExpressionClass() == duckdb::ExpressionClass::BOUND_CASE) { + auto& c = expression->Cast(); + for (auto& branch : c.case_checks) + branches.push_back(branch.then_expr.get()); + branches.push_back(c.else_expr.get()); + } else if (expression->GetExpressionClass() == duckdb::ExpressionClass::BOUND_OPERATOR && + expression->GetExpressionType() == duckdb::ExpressionType::OPERATOR_COALESCE) { + coalesce = true; + for (auto& child : expression->Cast().children) + branches.push_back(child.get()); + } + std::optional result; + bool nullable = coalesce; + for (auto branch : branches) { + auto t = from_duckdb_type(branch->return_type); + if (!t || !t->is_mo_decimal()) { + result.reset(); + break; + } + if (result && !same_decimal(*result, *t)) + throw std::invalid_argument("bound conditional changed an exact descriptor"); + result = t; + if (coalesce) + nullable &= t->nullability() == 2; + else + nullable |= t->nullability() == 2; + } + if (result) expression->return_type = duckdb_type(result->with_nullability(nullable)); +} +namespace { +duckdb::ExpressionType comparison_type(decimal_op op) +{ + using T = duckdb::ExpressionType; + switch (op) { + case decimal_op::equal: return T::COMPARE_EQUAL; + case decimal_op::not_equal: return T::COMPARE_NOTEQUAL; + case decimal_op::less: return T::COMPARE_LESSTHAN; + case decimal_op::less_equal: return T::COMPARE_LESSTHANOREQUALTO; + case decimal_op::greater: return T::COMPARE_GREATERTHAN; + case decimal_op::greater_equal: return T::COMPARE_GREATERTHANOREQUALTO; + default: throw std::invalid_argument("not an exact comparison"); + } +} +void rewrite_expression(duckdb::unique_ptr& expression, bool comparisons) +{ + duckdb::ExpressionIterator::EnumerateChildren( + *expression, + [&](duckdb::unique_ptr& child) { rewrite_expression(child, comparisons); }); + if (expression->GetExpressionClass() == duckdb::ExpressionClass::BOUND_FUNCTION) { + auto& fn = expression->Cast(); + auto id = from_duckdb_function_name(fn.function.name); + if (id && is_decimal_function(*id)) { + validate_bound_scalar(fn); + if (operation(*id) == decimal_op::cast && + sirius::from_duckdb(fn.children[0]->return_type) == sirius::from_duckdb(fn.return_type)) { + expression = std::move(fn.children[0]); + } else if (comparisons && operation(*id) >= decimal_op::equal) { + auto output = fn.return_type; + expression = duckdb::make_uniq( + comparison_type(operation(*id)), std::move(fn.children[0]), std::move(fn.children[1])); + expression->return_type = std::move(output); + } + } + } else if (expression->GetExpressionClass() == duckdb::ExpressionClass::BOUND_AGGREGATE) { + auto& agg = expression->Cast(); + auto id = from_duckdb_aggregate_name(agg.function.name); + if (id && is_decimal_aggregate(*id)) validate_bound_aggregate(agg); + } +} +} // namespace +void rewrite_exact_comparisons(duckdb::unique_ptr& plan, + duckdb::ClientContext& context) +{ + for (auto& child : plan->children) + rewrite_exact_comparisons(child, context); + duckdb::column_binding_map_t bindings; + for (auto& child : plan->children) { + auto columns = child->GetColumnBindings(); + if (columns.size() != child->types.size()) + throw std::invalid_argument("exact-plan child schema lost bindings"); + for (size_t i = 0; i < columns.size(); ++i) + bindings.emplace(columns[i], child->types[i]); + } + std::function&)> restore_references; + restore_references = [&](duckdb::unique_ptr& expression) { + if (expression->GetExpressionClass() == duckdb::ExpressionClass::BOUND_COLUMN_REF && + !from_duckdb_type(expression->return_type)) { + auto& ref = expression->Cast(); + auto found = bindings.find(ref.binding); + if (found != bindings.end() && from_duckdb_type(found->second)) { + auto plain = duckdb_type(*from_duckdb_type(found->second)); + plain.SetAlias(""); + if (plain == ref.return_type) ref.return_type = found->second; + } + } + duckdb::ExpressionIterator::EnumerateChildren(*expression, restore_references); + }; + duckdb::LogicalOperatorVisitor::EnumerateExpressions( + *plan, + [&](duckdb::unique_ptr* expression) { restore_references(*expression); }); + bool comparisons = plan->type == duckdb::LogicalOperatorType::LOGICAL_ANY_JOIN; + duckdb::LogicalOperatorVisitor::EnumerateExpressions( + *plan, [&](duckdb::unique_ptr* expression) { + restore_exact_bound_types(*expression); + rewrite_expression(*expression, comparisons); + }); + if (plan->type == duckdb::LogicalOperatorType::LOGICAL_ANY_JOIN) { + auto& join = plan->Cast(); + plan = duckdb::LogicalComparisonJoin::CreateJoin(context, + join.join_type, + duckdb::JoinRefType::REGULAR, + std::move(join.children[0]), + std::move(join.children[1]), + std::move(join.condition)); + } + if (plan->type == duckdb::LogicalOperatorType::LOGICAL_COMPARISON_JOIN) { + auto& join = plan->Cast(); + for (auto const& condition : join.conditions) { + auto left = from_duckdb_type(condition.left->return_type), + right = from_duckdb_type(condition.right->return_type); + if (((left && left->is_mo_decimal()) || (right && right->is_mo_decimal())) && + condition.comparison != duckdb::ExpressionType::COMPARE_EQUAL) + throw std::invalid_argument("MO decimal joins require exact equality keys"); + } + } + plan->ResolveOperatorTypes(); +} +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_plan.hpp b/src/numeric/decimal_plan.hpp new file mode 100644 index 0000000000..6454143a86 --- /dev/null +++ b/src/numeric/decimal_plan.hpp @@ -0,0 +1,18 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include "numeric/decimal_types.hpp" + +#include + +#include +namespace duckdb { +class LogicalOperator; +class ClientContext; +} // namespace duckdb +namespace sirius::mo_decimal { +std::string normalize_exact_substrait(std::string_view bytes, + std::size_t schema_bytes_limit = 16u << 20); +void restore_exact_bound_types(duckdb::unique_ptr& expression); +void rewrite_exact_comparisons(duckdb::unique_ptr& plan, + duckdb::ClientContext& context); +} // namespace sirius::mo_decimal diff --git a/src/op/aggregate/aggregate_op_util.cpp b/src/op/aggregate/aggregate_op_util.cpp index 7e4e782c86..cc51a8fb4e 100644 --- a/src/op/aggregate/aggregate_op_util.cpp +++ b/src/op/aggregate/aggregate_op_util.cpp @@ -20,6 +20,7 @@ #include "duckdb/common/assert.hpp" #include "expression/aggregate_id.hpp" #include "expression/ast/node.hpp" +#include "numeric/decimal_aggregate_layout.hpp" #include #include @@ -64,6 +65,7 @@ CudfAggregateDefinitions convert_duckdb_aggregates_to_cudf( const duckdb::vector>& expressions) { CudfAggregateDefinitions result; + result.exact_layout = mo_decimal::make_aggregate_layout(groups_p, expressions); // 1. Extract group_idx from groups_p for (const auto& group : groups_p) { @@ -72,6 +74,8 @@ CudfAggregateDefinitions convert_duckdb_aggregates_to_cudf( result.group_idx.push_back(static_cast(ref.column_index)); } + if (result.exact_layout) return result; + // 2. Extract aggregates (cudf::aggregation::Kind) from expressions for (const auto& aggregate : expressions) { auto const& aggr = sirius::ast::require_aggregate( diff --git a/src/op/aggregate/aggregate_op_util.hpp b/src/op/aggregate/aggregate_op_util.hpp index b25832b8b5..816593eb71 100644 --- a/src/op/aggregate/aggregate_op_util.hpp +++ b/src/op/aggregate/aggregate_op_util.hpp @@ -27,6 +27,9 @@ #include namespace sirius { +namespace mo_decimal { +struct aggregate_layout; +} namespace op { /** @@ -60,6 +63,7 @@ struct AggregateSlot { * @brief Result of converting DuckDB aggregate expressions to cuDF compute definitions. */ struct CudfAggregateDefinitions { + std::shared_ptr exact_layout; std::vector group_idx; ///< Column indices for GROUP BY keys std::vector cudf_aggregates; ///< cuDF aggregation types (expanded: 2 ///< entries per AVG) diff --git a/src/op/partition/gpu_partition_impl.cpp b/src/op/partition/gpu_partition_impl.cpp index 1ef7be66fe..1b27fa9e03 100644 --- a/src/op/partition/gpu_partition_impl.cpp +++ b/src/op/partition/gpu_partition_impl.cpp @@ -18,6 +18,7 @@ #include "data/data_batch_utils.hpp" #include "helper/numeric_narrowing.hpp" +#include "numeric/decimal_aggregate_gpu.hpp" #include @@ -33,7 +34,8 @@ std::vector> gpu_partition_impl::hash_par int num_partitions, ::cuda::stream_ref stream, cucascade::memory::memory_space& memory_space, - const telemetry::batch_telemetry_info& telemetry_info) + const telemetry::batch_telemetry_info& telemetry_info, + const std::vector& exact_types) { // Sanity check. if (num_partitions < 2) { @@ -54,8 +56,22 @@ std::vector> gpu_partition_impl::hash_par all_col_views.push_back(input_table.column(i)); } std::vector effective_key_idx = partition_key_idx; + if (!exact_types.empty() && exact_types.size() != partition_key_idx.size()) + throw std::invalid_argument("MO partition exact-key descriptors do not match keys"); + owned_cast_cols.reserve(partition_key_idx.size() + partition_key_cast_types.size()); + for (size_t i = 0; i < exact_types.size(); ++i) + if (exact_types[i].bits) { + auto key = mo_decimal::make_equality_key(input_table.column(partition_key_idx[i]), + exact_types[i], + stream, + memory_space.get_default_allocator()); + effective_key_idx[i] = static_cast(all_col_views.size()); + all_col_views.push_back(key->view()); + owned_cast_cols.push_back(std::move(key)); + } for (size_t i = 0; i < partition_key_cast_types.size(); i++) { - if (partition_key_cast_types[i].id() != cudf::type_id::EMPTY) { + if ((exact_types.empty() || !exact_types[i].bits) && + partition_key_cast_types[i].id() != cudf::type_id::EMPTY) { auto cast_col = sirius::cast_through_rep( input_table.column(partition_key_idx[i]), partition_key_cast_types[i], stream); effective_key_idx[i] = static_cast(all_col_views.size()); diff --git a/src/op/partition/gpu_partition_impl.hpp b/src/op/partition/gpu_partition_impl.hpp index a3ede38c7b..0198f653cc 100644 --- a/src/op/partition/gpu_partition_impl.hpp +++ b/src/op/partition/gpu_partition_impl.hpp @@ -65,7 +65,8 @@ class gpu_partition_impl { int num_partitions, ::cuda::stream_ref stream, cucascade::memory::memory_space& memory_space, - const telemetry::batch_telemetry_info& telemetry_info = {}); + const telemetry::batch_telemetry_info& telemetry_info = {}, + const std::vector& exact_types = {}); /// Overload without cast types (all keys hashed as-is). Kept for backward compatibility. static std::vector> hash_partition( diff --git a/src/op/sirius_physical_grouped_aggregate.cpp b/src/op/sirius_physical_grouped_aggregate.cpp index 9bea2d8dd7..52f8d9d2ec 100644 --- a/src/op/sirius_physical_grouped_aggregate.cpp +++ b/src/op/sirius_physical_grouped_aggregate.cpp @@ -18,6 +18,7 @@ #include "config.hpp" #include "data/data_batch_utils.hpp" +#include "numeric/decimal_aggregate_layout.hpp" #include "op/aggregate/aggregate_op_util.hpp" #include "op/aggregate/gpu_aggregate_impl.hpp" #include "telemetry/nvtx.hpp" @@ -63,6 +64,7 @@ sirius_physical_grouped_aggregate::sirius_physical_grouped_aggregate( { auto cudf_defs = convert_duckdb_aggregates_to_cudf(groups_p, expressions); group_idx = std::move(cudf_defs.group_idx); + exact_layout = std::move(cudf_defs.exact_layout); cudf_aggregates = std::move(cudf_defs.cudf_aggregates); cudf_aggregate_idx = std::move(cudf_defs.cudf_aggregate_idx); cudf_aggregate_struct_col_indices = std::move(cudf_defs.cudf_aggregate_struct_col_indices); @@ -88,6 +90,11 @@ sirius_physical_grouped_aggregate::get_count_distinct_local_output_types() const } return local_types; } +duckdb::vector sirius_physical_grouped_aggregate::get_local_output_types() + const +{ + return exact_layout ? exact_layout->local_types() : types; +} std::unique_ptr sirius_physical_grouped_aggregate::execute( const operator_data& input_data, ::cuda::stream_ref stream) @@ -99,6 +106,15 @@ std::unique_ptr sirius_physical_grouped_aggregate::execute( for (auto const& input_batch : input_batches) { auto* space = input_batch.get_memory_space(); if (!space) { continue; } + if (exact_layout) { + results.push_back(make_data_batch( + mo_decimal::local_aggregate_table( + get_cudf_table_view(input_batch), *exact_layout, stream, space->get_default_allocator()), + *space, + stream, + batch_telemetry())); + continue; + } auto result = gpu_aggregate_impl::local_grouped_aggregate(input_batch, group_idx, cudf_aggregates, diff --git a/src/op/sirius_physical_grouped_aggregate.hpp b/src/op/sirius_physical_grouped_aggregate.hpp index f167e7bfc1..0b3e1b2797 100644 --- a/src/op/sirius_physical_grouped_aggregate.hpp +++ b/src/op/sirius_physical_grouped_aggregate.hpp @@ -77,6 +77,7 @@ class sirius_physical_grouped_aggregate : public sirius_physical_operator { // Grouped aggregatge definitions for cudf compute std::vector group_idx; + std::shared_ptr exact_layout; std::vector cudf_aggregates; std::vector cudf_aggregate_idx; std::vector> cudf_aggregate_struct_col_indices; @@ -97,6 +98,7 @@ class sirius_physical_grouped_aggregate : public sirius_physical_operator { //! Runtime schema of the local COUNT(DISTINCT) accumulator. The local aggregate and PARTITION //! carry LIST sets; MERGE_GROUP_BY later converts those sets to the declared BIGINT count. [[nodiscard]] duckdb::vector get_count_distinct_local_output_types() const; + [[nodiscard]] duckdb::vector get_local_output_types() const; // Source interface bool is_source() const override { return true; } diff --git a/src/op/sirius_physical_grouped_aggregate_merge.cpp b/src/op/sirius_physical_grouped_aggregate_merge.cpp index 081b0eab4e..ebc3398e3e 100644 --- a/src/op/sirius_physical_grouped_aggregate_merge.cpp +++ b/src/op/sirius_physical_grouped_aggregate_merge.cpp @@ -18,6 +18,7 @@ #include "cudf/cudf_utils.hpp" #include "data/data_batch_utils.hpp" #include "duckdb/planner/expression/bound_reference_expression.hpp" +#include "numeric/decimal_aggregate_layout.hpp" #include "op/aggregate/aggregate_op_util.hpp" #include "op/merge/gpu_merge_impl.hpp" #include "pipeline/sirius_meta_pipeline.hpp" @@ -75,7 +76,8 @@ sirius_physical_grouped_aggregate_merge::sirius_physical_grouped_aggregate_merge grouped_aggregate->aggregate_slots, grouped_aggregate->has_avg, grouped_aggregate->has_count_distinct, - grouped_aggregate->estimated_cardinality) + grouped_aggregate->estimated_cardinality, + grouped_aggregate->exact_layout) { child_op = grouped_aggregate; _hash_partition_bytes = hash_partition_bytes; @@ -90,10 +92,12 @@ sirius_physical_grouped_aggregate_merge::sirius_physical_grouped_aggregate_merge std::vector aggregate_slots, bool has_avg, bool has_count_distinct, - std::size_t estimated_cardinality) + std::size_t estimated_cardinality, + std::shared_ptr exact_layout) : sirius_physical_partition_consumer_operator( SiriusPhysicalOperatorType::MERGE_GROUP_BY, std::move(types), estimated_cardinality), group_idx(std::move(group_idx)), + exact_layout(std::move(exact_layout)), cudf_aggregates(std::move(cudf_aggregates)), cudf_aggregate_idx(std::move(cudf_aggregate_idx)), cudf_aggregate_struct_col_indices(std::move(cudf_aggregate_struct_col_indices)), @@ -142,6 +146,7 @@ sirius_physical_grouped_aggregate_merge::sirius_physical_grouped_aggregate_merge // Convert input parameters to cudf compute definitions BEFORE moving them auto cudf_defs = convert_duckdb_aggregates_to_cudf(groups_p, expressions); group_idx = std::move(cudf_defs.group_idx); + exact_layout = std::move(cudf_defs.exact_layout); cudf_aggregates = std::move(cudf_defs.cudf_aggregates); cudf_aggregate_idx = std::move(cudf_defs.cudf_aggregate_idx); cudf_aggregate_struct_col_indices = std::move(cudf_defs.cudf_aggregate_struct_col_indices); @@ -208,6 +213,19 @@ std::unique_ptr sirius_physical_grouped_aggregate_merge::execute( throw std::runtime_error( "We expect at least one input batch for grouped aggregate merge operator"); } + if (exact_layout) { + auto* space = input_batches[0].get_memory_space(); + std::vector tables; + for (auto const& batch : input_batches) + tables.push_back(get_cudf_table_view(batch)); + auto result = make_data_batch(mo_decimal::merge_aggregate_tables( + tables, *exact_layout, stream, space->get_default_allocator()), + *space, + stream, + batch_telemetry()); + return std::make_unique( + std::vector>{std::move(result)}); + } // Fast path: single batch with no post-processing needed if (input_batches.size() == 1 && !has_avg && !has_count_distinct) { diff --git a/src/op/sirius_physical_grouped_aggregate_merge.hpp b/src/op/sirius_physical_grouped_aggregate_merge.hpp index fe387feec6..19863e9cec 100644 --- a/src/op/sirius_physical_grouped_aggregate_merge.hpp +++ b/src/op/sirius_physical_grouped_aggregate_merge.hpp @@ -59,7 +59,8 @@ class sirius_physical_grouped_aggregate_merge : public sirius_physical_partition std::vector aggregate_slots, bool has_avg, bool has_count_distinct, - std::size_t estimated_cardinality); + std::size_t estimated_cardinality, + std::shared_ptr exact_layout = {}); sirius_physical_grouped_aggregate_merge( duckdb::vector types, @@ -96,6 +97,7 @@ class sirius_physical_grouped_aggregate_merge : public sirius_physical_partition // Grouped aggregatge definitions for cudf compute std::vector group_idx; + std::shared_ptr exact_layout; std::vector cudf_aggregates; std::vector cudf_aggregate_idx; std::vector> cudf_aggregate_struct_col_indices; diff --git a/src/op/sirius_physical_hash_join.cpp b/src/op/sirius_physical_hash_join.cpp index 663a2827d8..605e22a876 100644 --- a/src/op/sirius_physical_hash_join.cpp +++ b/src/op/sirius_physical_hash_join.cpp @@ -44,6 +44,7 @@ #include "helper/numeric_narrowing.hpp" #include "helper/type_conversions.hpp" #include "log/logging.hpp" +#include "numeric/decimal_aggregate_gpu.hpp" #include "op/dynamic_filter/dynamic_filter_publisher.hpp" #include "op/dynamic_filter/sirius_dynamic_filter.hpp" #include "op/sirius_physical_concat.hpp" @@ -476,6 +477,15 @@ sirius_physical_hash_join::sirius_physical_hash_join( // Extract left key index (may be BOUND_REF or BOUND_CAST wrapping a BOUND_REF) key_cast_info cast_info; + if (condition.left->return_type().is_mo_decimal() || + condition.right->return_type().is_mo_decimal()) { + if (!condition.left->return_type().is_mo_decimal() || + !condition.right->return_type().is_mo_decimal()) + throw std::invalid_argument("MO decimal join key requires two exact operands"); + cast_info.left_exact = condition.left->return_type().mo_decimal_type(); + cast_info.right_exact = condition.right->return_type().mo_decimal_type(); + cast_necessary = true; + } auto left_class = left_expr->GetExpressionClass(); auto right_class = right_expr->GetExpressionClass(); @@ -1460,6 +1470,8 @@ static join_side_keys_result prepare_join_keys( ::cuda::stream_ref stream) { join_side_keys_result result; + result.owned_cast_columns.reserve(key_col_indices.size()); + result.key_views.reserve(key_col_indices.size()); cudf::table_view table = get_cudf_table_view(input_batch); @@ -1497,7 +1509,13 @@ static join_side_keys_result prepare_join_keys( cudf::data_type target_type = is_left_side ? cast_info.left_target_type : cast_info.right_target_type; - if (needs_cast) { + auto exact = is_left_side ? cast_info.left_exact : cast_info.right_exact; + if (exact.bits) { + auto key = mo_decimal::make_equality_key( + col, exact, stream, input_batch.get_memory_space()->get_default_allocator()); + result.key_views.push_back(key->view()); + result.owned_cast_columns.push_back(std::move(key)); + } else if (needs_cast) { auto cast_col = sirius::cast_through_rep(col, target_type, stream); result.key_views.push_back(cast_col->view()); result.owned_cast_columns.push_back(std::move(cast_col)); diff --git a/src/op/sirius_physical_hash_join.hpp b/src/op/sirius_physical_hash_join.hpp index ea607ebbcc..c80598c8f9 100644 --- a/src/op/sirius_physical_hash_join.hpp +++ b/src/op/sirius_physical_hash_join.hpp @@ -519,6 +519,7 @@ class sirius_physical_hash_join : public sirius_physical_partition_consumer_oper bool cast_right = false; cudf::data_type left_target_type{cudf::type_id::EMPTY}; cudf::data_type right_target_type{cudf::type_id::EMPTY}; + mo_decimal::decimal_type left_exact{}, right_exact{}; }; protected: diff --git a/src/op/sirius_physical_partition.cpp b/src/op/sirius_physical_partition.cpp index 0d4d1508ed..667d483de7 100644 --- a/src/op/sirius_physical_partition.cpp +++ b/src/op/sirius_physical_partition.cpp @@ -142,6 +142,9 @@ void sirius_physical_partition::get_partition_keys_and_type(sirius_physical_oper // keys would land in different partitions. We apply the same cast used by the join // condition so both sides hash identically. const auto& key_expr = is_build ? *right_owned : *left_owned; + auto logical = is_build ? condition.right->return_type() : condition.left->return_type(); + _partition_exact_types.push_back(logical.is_mo_decimal() ? logical.mo_decimal_type() + : mo_decimal::decimal_type{}); if (is_build) { _partition_keys.push_back(right_index.value()); } else { @@ -298,7 +301,8 @@ std::unique_ptr sirius_physical_partition::execute(const operator _num_partitions.value(), stream, *space, - batch_telemetry()); + batch_telemetry(), + _partition_exact_types); break; case PartitionType::RANGE: throw std::runtime_error("Range partitioning is not implemented yet"); diff --git a/src/op/sirius_physical_partition.hpp b/src/op/sirius_physical_partition.hpp index c7b37765da..7a7f55e8c8 100644 --- a/src/op/sirius_physical_partition.hpp +++ b/src/op/sirius_physical_partition.hpp @@ -179,6 +179,7 @@ class sirius_physical_partition : public sirius_physical_operator { /// cast the key column to this type before hashing. Used to align hash values when the /// two join sides have different physical column types for the same logical key. std::vector _partition_key_cast_types; + std::vector _partition_exact_types; std::optional _num_partitions; bool _is_build; bool _drives_partition_count{false}; diff --git a/src/op/sirius_physical_ungrouped_aggregate.cpp b/src/op/sirius_physical_ungrouped_aggregate.cpp index 2c2fbb7199..5b40440119 100644 --- a/src/op/sirius_physical_ungrouped_aggregate.cpp +++ b/src/op/sirius_physical_ungrouped_aggregate.cpp @@ -24,6 +24,7 @@ #include "expression/ast/reference.hpp" #include "expression/ast/utils.hpp" #include "helper/type_conversions.hpp" +#include "numeric/decimal_aggregate_layout.hpp" #include "op/merge/gpu_merge_impl.hpp" #include "op/sirius_physical_ungrouped_aggregate_merge.hpp" #include "sirius/exception.hpp" @@ -276,6 +277,7 @@ std::unique_ptr make_avg_column(const cudf::column_view& sum_view, duckdb::vector sirius_physical_ungrouped_aggregate::get_local_output_types() const { + if (auto exact = mo_decimal::make_aggregate_layout({}, aggregates)) return exact->local_types(); aggregate_layout layout; try { layout = build_aggregate_layout(aggregates); @@ -304,6 +306,21 @@ std::unique_ptr sirius_physical_ungrouped_aggregate::execute( std::vector>{}); } + auto exact = mo_decimal::make_aggregate_layout({}, aggregates); + if (exact) { + std::vector> outputs; + for (auto const& batch : input_batches) { + auto* space = batch.get_memory_space(); + if (!space) continue; + outputs.push_back(make_data_batch( + mo_decimal::local_aggregate_table( + get_cudf_table_view(batch), *exact, stream, space->get_default_allocator()), + *space, + stream, + batch_telemetry())); + } + return std::make_unique(std::move(outputs)); + } auto layout = build_aggregate_layout(aggregates); std::vector> outputs; outputs.reserve(input_batches.size()); @@ -497,6 +514,18 @@ std::unique_ptr sirius_physical_ungrouped_aggregate_merge::execut std::vector>{}); } + if (auto exact = mo_decimal::make_aggregate_layout({}, aggregates)) { + std::vector tables; + for (auto const& batch : input_batches) + tables.push_back(get_cudf_table_view(batch)); + auto result = make_data_batch( + mo_decimal::merge_aggregate_tables(tables, *exact, stream, space->get_default_allocator()), + *space, + stream, + batch_telemetry()); + return std::make_unique( + std::vector>{std::move(result)}); + } auto layout = build_aggregate_layout(aggregates); std::shared_ptr merged_batch; if (input_batches.size() == 1) { diff --git a/src/planner/sirius_physical_plan_generator.cpp b/src/planner/sirius_physical_plan_generator.cpp index d848cbaf9c..160a535aec 100644 --- a/src/planner/sirius_physical_plan_generator.cpp +++ b/src/planner/sirius_physical_plan_generator.cpp @@ -39,6 +39,7 @@ #include "helper/type_conversions.hpp" #include "io/uri_parser.hpp" #include "log/logging.hpp" +#include "numeric/decimal_types.hpp" #include "op/dynamic_filter/sirius_dynamic_filter.hpp" #include "op/scan/duckdb_native_gpu_ingestible.hpp" #include "op/scan/iceberg_gpu_ingestible.hpp" @@ -117,6 +118,8 @@ std::vector resolve_parquet_scan_file_paths( namespace { bool is_nested_logical_type(duckdb::LogicalType const& type) { + if (auto exact = mo_decimal::from_duckdb_type(type); exact && exact->is_mo_decimal()) + return false; auto const id = type.id(); return id == duckdb::LogicalTypeId::STRUCT || id == duckdb::LogicalTypeId::LIST || id == duckdb::LogicalTypeId::MAP; @@ -616,7 +619,9 @@ void wrap_hash_group_by(duckdb::unique_ptr bool const has_supported_count_distinct_layout = grouped.has_count_distinct && !grouped.has_avg && !hgb_ptr->has_physical_overrides() && hgb_ptr->types.size() == grouped.group_idx.size() + grouped.aggregate_slots.size(); - if (has_supported_count_distinct_layout) { + if (grouped.exact_layout) { + hgb_ptr->types = grouped.get_local_output_types(); + } else if (has_supported_count_distinct_layout) { hgb_ptr->types = grouped.get_count_distinct_local_output_types(); } diff --git a/src/sirius_c.h b/src/sirius_c.h index 2d24aa629d..0d3b1bf22a 100644 --- a/src/sirius_c.h +++ b/src/sirius_c.h @@ -116,18 +116,20 @@ typedef struct sirius_result_stats { typedef uint32_t sirius_status; enum { - SIRIUS_OK = 0, - SIRIUS_UNSUPPORTED = 1, - SIRIUS_INVALID_ARGUMENT = 2, - SIRIUS_INVALID_STATE = 3, - SIRIUS_RESOURCE_EXHAUSTED = 4, - SIRIUS_CANCELLED = 5, - SIRIUS_TIMEOUT = 6, - SIRIUS_GPU_UNAVAILABLE = 7, - SIRIUS_EOF = 8, - SIRIUS_NOT_NEEDED = 9, - SIRIUS_EXECUTION_FAILED = 10, - SIRIUS_BUSY = 11 + SIRIUS_OK = 0, + SIRIUS_UNSUPPORTED = 1, + SIRIUS_INVALID_ARGUMENT = 2, + SIRIUS_INVALID_STATE = 3, + SIRIUS_RESOURCE_EXHAUSTED = 4, + SIRIUS_CANCELLED = 5, + SIRIUS_TIMEOUT = 6, + SIRIUS_GPU_UNAVAILABLE = 7, + SIRIUS_EOF = 8, + SIRIUS_NOT_NEEDED = 9, + SIRIUS_EXECUTION_FAILED = 10, + SIRIUS_BUSY = 11, + SIRIUS_NUMERIC_OUT_OF_RANGE = 12, + SIRIUS_NUMERIC_INVALID_INPUT = 13 }; enum { SIRIUS_QUERY_SOURCE_MO = 1u, SIRIUS_QUERY_SOURCE_TAE = 2u }; @@ -158,10 +160,11 @@ typedef struct sirius_query_execution_stats { } sirius_query_execution_stats; enum { - SIRIUS_CAP_ENGINE_CONTROL = 1u, - SIRIUS_CAP_MO_INPUT = 2u, - SIRIUS_CAP_TAE_INPUT = 4u, - SIRIUS_CAP_NATIVE_RESULTS = 8u + SIRIUS_CAP_ENGINE_CONTROL = 1u, + SIRIUS_CAP_MO_INPUT = 2u, + SIRIUS_CAP_TAE_INPUT = 4u, + SIRIUS_CAP_NATIVE_RESULTS = 8u, + SIRIUS_CAP_MO_EXACT_DECIMAL_V1 = 16u }; typedef struct sirius_error { diff --git a/src/sirius_ffi.cpp b/src/sirius_ffi.cpp index 05c4d59b70..bbc2b52a37 100644 --- a/src/sirius_ffi.cpp +++ b/src/sirius_ffi.cpp @@ -47,7 +47,11 @@ #include "from_substrait.hpp" // duckdb::SubstraitToDuckDB (compiled into libsirius) #include "helper/type_conversions.hpp" // sirius::from_duckdb #include "log/logging.hpp" // SIRIUS_LOG_INFO -#include "parquet_extension.hpp" // duckdb::ParquetExtension +#include "numeric/decimal_aggregate_bind.hpp" +#include "numeric/decimal_functions.hpp" +#include "numeric/decimal_import.hpp" +#include "numeric/decimal_plan.hpp" +#include "parquet_extension.hpp" // duckdb::ParquetExtension #include "pipeline/gpu_stream_quiescence_error.hpp" #include "planner/sirius_physical_plan_generator.hpp" // sirius::planner::sirius_physical_plan_generator #include "sirius/ffi.hpp" @@ -118,11 +122,14 @@ struct lowered_plan { duckdb::unique_ptr plan; }; -lowered_plan lower_substrait(duckdb::Connection& conn, const std::string& substrait_plan) +lowered_plan lower_substrait(duckdb::Connection& conn, + const std::string& substrait_plan, + duckdb::shared_ptr extensions = {}) { auto& client = *conn.context; - duckdb::SubstraitToDuckDB transformer(conn.context, substrait_plan, /*json=*/false); + duckdb::SubstraitToDuckDB transformer( + conn.context, substrait_plan, /*json=*/false, /*acquire_lock=*/true, extensions); auto relation = transformer.TransformPlan(); duckdb::Planner planner(client); @@ -135,11 +142,15 @@ lowered_plan lower_substrait(duckdb::Connection& conn, const std::string& substr prepared->value_map = std::move(planner.value_map); auto logical_plan = std::move(planner.plan); - if (client.config.enable_optimizer) { + if (extensions) mo_decimal::rewrite_exact_comparisons(logical_plan, client); + // MO already optimized this logical plan. DuckDB's ordinary comparison + // folding/coercion cannot interpret private coefficient carriers or scales. + if (client.config.enable_optimizer && !extensions) { duckdb::Optimizer optimizer(*planner.binder, client); logical_plan = optimizer.Optimize(std::move(logical_plan)); } logical_plan->ResolveOperatorTypes(); + if (extensions) prepared->types = logical_plan->types; duckdb::ColumnBindingResolver resolver; duckdb::ColumnBindingResolver::Verify(*logical_plan); resolver.VisitOperator(*logical_plan); @@ -180,6 +191,8 @@ struct Context::Impl { db = duckdb::make_uniq(nullptr, &duckdb_config); db->LoadStaticExtension(); db->LoadStaticExtension(); + mo_decimal::register_scalar_functions(*db->instance); + mo_decimal::register_aggregate_functions(*db->instance); // Cache parquet footers across binds. The Substrait consumer builds the plan through the // Relation API, which re-binds the whole subtree at every level, so a read of one file is // bound once per operator above it; with the cache off each bind re-parses the footer and @@ -334,8 +347,16 @@ void EmbeddedPrepared::run(std::stop_token stop, std::chrono::steady_clock::time namespace { std::atomic embedded_generation{1}; -duckdb::LogicalType embedded_type(sirius::embedding::owned_column const& column) +duckdb::LogicalType embedded_type(sirius::embedding::owned_column const& column, bool exact = false) { + if (exact && (column.oid == 32 || column.oid == 33 || column.oid == 34)) + return mo_decimal::duckdb_type( + logical_type::make_mo_decimal({static_cast(column.oid == 32 ? 64 + : column.oid == 33 ? 128 + : 256), + static_cast(column.width), + static_cast(column.scale)}, + column.nullable)); tae::MOType type{}; type.oid = static_cast(column.oid); type.width = column.width; @@ -387,6 +408,9 @@ std::unique_ptr Context::prepare_embedded(const std::string& b { if (!embedded_runtime_available()) throw embedding::failure(SIRIUS_GPU_UNAVAILABLE, "native runtime is unavailable"); + bool exact = mo_decimal::uses_exact_decimal(bytes); + auto native_bytes = + exact ? mo_decimal::normalize_exact_substrait(bytes, query.metadata_bytes) : bytes; impl_->embedded_catalog->clear(); auto const generation = embedded_generation.fetch_add(1, std::memory_order_relaxed); auto const has_tae = std::any_of(query.bindings.begin(), query.bindings.end(), [](auto const& b) { @@ -410,7 +434,7 @@ std::unique_ptr Context::prepare_embedded(const std::string& b embedding::embedded_binding binding; for (auto const& column : read.columns) { binding.names.push_back(column.logical.name); - binding.types.push_back(embedded_type(column.logical)); + binding.types.push_back(embedded_type(column.logical, exact)); } const char* function = nullptr; if (read.source_kind == SIRIUS_READ_MO) { @@ -453,14 +477,22 @@ std::unique_ptr Context::prepare_embedded(const std::string& b relation->CreateView(view_name, true, true); views.push_back(std::move(view_name)); } - auto lowered = lower_substrait(*impl_->conn, rewrite_embedded_reads(bytes, generation)); + auto rewritten = rewrite_embedded_reads(native_bytes, generation); + auto lowered = lower_substrait( + *impl_->conn, rewritten, exact ? mo_decimal::make_import_handler(rewritten) : nullptr); if (lowered.prepared->names.size() != query.contract->outputs.size() || lowered.prepared->types.size() != query.contract->outputs.size()) throw embedding::failure(SIRIUS_INVALID_ARGUMENT, "prepared output schema does not match query contract"); for (std::size_t i = 0; i < query.contract->outputs.size(); ++i) { + auto expected = embedded_type(query.contract->outputs[i], exact); + if (exact) + if (auto tagged = mo_decimal::from_duckdb_type(lowered.prepared->types[i]); + tagged && !tagged->is_mo_decimal()) + expected = mo_decimal::duckdb_type( + sirius::from_duckdb(expected).with_nullability(query.contract->outputs[i].nullable)); if (lowered.prepared->names[i] != query.contract->outputs[i].name || - lowered.prepared->types[i] != embedded_type(query.contract->outputs[i])) + lowered.prepared->types[i] != expected) throw embedding::failure(SIRIUS_INVALID_ARGUMENT, "prepared output schema does not match query contract"); } diff --git a/test/cpp/embedding/decimal_import_fixture.hpp b/test/cpp/embedding/decimal_import_fixture.hpp index b1ab492304..f33d5dccad 100644 --- a/test/cpp/embedding/decimal_import_fixture.hpp +++ b/test/cpp/embedding/decimal_import_fixture.hpp @@ -2,13 +2,16 @@ #pragma once #include "expression/ast/from_duckdb.hpp" #include "from_substrait.hpp" +#include "numeric/decimal_aggregate_bind.hpp" #include "numeric/decimal_functions.hpp" #include "numeric/decimal_import.hpp" #include "numeric/exact_decimal.hpp" #include #include +#include #include +#include #include #include #include @@ -103,6 +106,7 @@ class importer { { db.LoadStaticExtension(); sirius::mo_decimal::register_scalar_functions(*db.instance); + sirius::mo_decimal::register_aggregate_functions(*db.instance); } std::unique_ptr bind(substrait::Plan const& plan) { @@ -114,10 +118,21 @@ class importer { auto relation = converter.TransformPlan(); duckdb::Planner planner(*connection.context); planner.CreatePlan(duckdb::make_uniq(relation)); + planner.plan->ResolveOperatorTypes(); + duckdb::ColumnBindingResolver resolver; + resolver.VisitOperator(*planner.plan); std::unique_ptr found; std::function expression; expression = [&](duckdb::Expression const& value) { if (found) return; + if (value.GetExpressionClass() == duckdb::ExpressionClass::BOUND_AGGREGATE) { + auto const& aggregate = value.Cast(); + auto id = sirius::from_duckdb_aggregate_name(aggregate.function.name); + if (id && sirius::mo_decimal::is_decimal_aggregate(*id)) { + found = sirius::ast::from_duckdb(value); + return; + } + } if (value.GetExpressionClass() == duckdb::ExpressionClass::BOUND_FUNCTION) { auto const& function = value.Cast(); auto id = sirius::from_duckdb_function_name(function.function.name); diff --git a/test/cpp/embedding/test_decimal_aggregate.cpp b/test/cpp/embedding/test_decimal_aggregate.cpp new file mode 100644 index 0000000000..986253c570 --- /dev/null +++ b/test/cpp/embedding/test_decimal_aggregate.cpp @@ -0,0 +1,97 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "numeric/decimal_aggregate.hpp" + +#include + +#include +#include + +using namespace sirius::mo_decimal; +namespace { +coefficient integer(int64_t value) +{ + return load_coefficient(reinterpret_cast(&value), 8); +} +void equal(coefficient const& observed, coefficient const& expected) +{ + for (int i = 0; i < 8; ++i) + CHECK(observed.words[i] == expected.words[i]); +} +} // namespace +TEST_CASE("MO aggregate states retain physical range until final publication", + "[decimal_aggregate]") +{ + auto large = decimal_detail::power10(65); + decimal_detail::magnitude one; + one.words[0] = 1; + decimal_detail::subtract(large, one); + auto positive = + decimal_detail::finish(large, false, {256, 65, 0}, decimal_error::invalid_input).value; + auto negative = + decimal_detail::finish(large, true, {256, 65, 0}, decimal_error::invalid_input).value; + std::array order{0, 1, 2, 3}; + std::array values{positive, positive, negative, negative}; + do { + aggregate_state state; + for (auto index : order) + state = accumulate_aggregate(aggregate_op::sum, state, values[index], true, 256); + REQUIRE(state.error == decimal_error::none); + auto result = finalize_aggregate(aggregate_op::sum, state, {256, 65, 0}, {256, 65, 0}); + REQUIRE(result.valid); + equal(result.value, {}); + } while (std::next_permutation(order.begin(), order.end())); + auto state = accumulate_aggregate(aggregate_op::sum, {}, positive, true, 256); + state = accumulate_aggregate(aggregate_op::sum, state, integer(1), true, 256); + CHECK(state.error == decimal_error::none); + CHECK(finalize_aggregate(aggregate_op::sum, state, {256, 65, 0}, {256, 65, 0}).error == + decimal_error::invalid_input); + coefficient physical_max; + for (int i = 0; i < 8; ++i) + physical_max.words[i] = UINT32_MAX; + physical_max.words[7] >>= 1; + state = accumulate_aggregate(aggregate_op::sum, {}, physical_max, true, 256); + state = accumulate_aggregate(aggregate_op::sum, state, integer(1), true, 256); + CHECK(state.error == decimal_error::invalid_input); + coefficient physical128; + for (int i = 0; i < 4; ++i) + physical128.words[i] = UINT32_MAX; + physical128.words[3] >>= 1; + state = accumulate_aggregate(aggregate_op::sum, {}, physical128, true, 128); + state = accumulate_aggregate(aggregate_op::sum, state, integer(1), true, 128); + CHECK(state.error == decimal_error::invalid_input); +} +TEST_CASE("MO aggregate empty identities and signed AVG round only once", "[decimal_aggregate]") +{ + for (auto op : {aggregate_op::sum, aggregate_op::avg, aggregate_op::min, aggregate_op::max}) { + auto empty = accumulate_aggregate(op, {}, integer(999), false, 128); + CHECK(empty.count == 0); + CHECK_FALSE(finalize_aggregate(op, empty, {64, 9, 2}, {128, 19, 6}).valid); + } + for (auto sign : {-1, 1}) { + auto left = accumulate_aggregate(aggregate_op::avg, {}, integer(sign * 124), true, 128); + auto right = accumulate_aggregate(aggregate_op::avg, {}, integer(sign * 126), true, 128); + auto state = merge_aggregate(aggregate_op::avg, left, right, 128); + auto result = finalize_aggregate(aggregate_op::avg, state, {64, 9, 2}, {128, 19, 1}); + REQUIRE(result.valid); + equal(result.value, integer(sign * 13)); + } + auto min = accumulate_aggregate(aggregate_op::min, {}, integer(7), true, 128); + min = accumulate_aggregate(aggregate_op::min, min, integer(-9), true, 128); + equal(finalize_aggregate(aggregate_op::min, min, {64, 9, 0}, {64, 9, 0}).value, integer(-9)); + auto max = accumulate_aggregate(aggregate_op::max, {}, integer(-9), true, 128); + max = accumulate_aggregate(aggregate_op::max, max, integer(7), true, 128); + equal(finalize_aggregate(aggregate_op::max, max, {64, 9, 0}, {64, 9, 0}).value, integer(7)); +} +TEST_CASE("MO equality keys normalize scales signs and zero without loss", "[decimal_aggregate]") +{ + for (auto sign : {-1, 1}) { + auto a = normalize_key(integer(sign * 1200), 3), b = normalize_key(integer(sign * 12), 1); + equal(a.value, b.value); + CHECK(a.scale == b.scale); + auto distinct = normalize_key(integer(sign * 1201), 3); + CHECK(compare_coefficient(a.value, distinct.value) != 0); + } + auto zero = normalize_key({}, 76); + equal(zero.value, {}); + CHECK(zero.scale == 0); +} diff --git a/test/cpp/embedding/test_decimal_aggregate_gpu.cpp b/test/cpp/embedding/test_decimal_aggregate_gpu.cpp new file mode 100644 index 0000000000..8fa3d64d8e --- /dev/null +++ b/test/cpp/embedding/test_decimal_aggregate_gpu.cpp @@ -0,0 +1,470 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "data/convertible_data_batch.hpp" +#include "numeric/decimal_aggregate_gpu.hpp" +#include "numeric/decimal_aggregate_layout.hpp" +#include "numeric/decimal_error.hpp" +#include "operator/operator_test_utils.hpp" + +#include +#include +#include + +#include +#include + +#include + +#include + +using namespace sirius; +using namespace sirius::mo_decimal; +namespace { +coefficient small(int64_t value) +{ + return load_coefficient(reinterpret_cast(&value), 8); +} +coefficient read(cudf::column_view column, uint32_t bytes, rmm::cuda_stream_view stream) +{ + coefficient result; + if (bytes == 32) { + for (int i = 0; i < 4; ++i) { + uint64_t word = 0; + REQUIRE(cudaMemcpyAsync(&word, + column.child(3 - i).head() + column.child(3 - i).offset() + + column.offset(), + 8, + cudaMemcpyDeviceToHost, + stream.value()) == cudaSuccess); + stream.synchronize(); + result.words[2 * i] = word; + result.words[2 * i + 1] = word >> 32; + } + } else { + uint8_t data[16]{}; + REQUIRE(cudaMemcpyAsync(data, + column.head() + uint64_t(column.offset()) * bytes, + bytes, + cudaMemcpyDeviceToHost, + stream.value()) == cudaSuccess); + stream.synchronize(); + result = load_coefficient(data, bytes); + } + return result; +} +void expect(cudf::column_view column, + decimal_type type, + int64_t value, + rmm::cuda_stream_view stream) +{ + REQUIRE(column.size() == 1); + REQUIRE(column.null_count() == 0); + CHECK(decimal_column_matches(column, type)); + auto observed = read(column, type.bytes(), stream), expected = small(value); + for (int i = 0; i < 8; ++i) + CHECK(observed.words[i] == expected.words[i]); +} +std::shared_ptr layout(aggregate_op op, + logical_type input, + logical_type output, + bool grouped = false, + bool count = false) +{ + duckdb::vector> groups, aggregates; + if (grouped) + groups.push_back( + std::make_unique(ast::reference{1, logical_type::make(type_id::INTEGER)})); + std::vector> args; + args.push_back(std::make_unique(ast::reference{0, input})); + aggregates.push_back(std::make_unique( + ast::aggregate{aggregate_function(op), std::move(args), output, false})); + if (count) + aggregates.push_back(std::make_unique( + ast::aggregate{aggregate_id::count_star, {}, logical_type::make(type_id::BIGINT), false})); + return make_aggregate_layout(groups, aggregates); +} +} // namespace +TEST_CASE("MO GPU aggregate local and merge retain all widths and ordinary slots", + "[decimal_aggregate_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + constexpr int rows = 10003; + for (uint16_t bits : {64, 128, 256}) + for (auto op : {aggregate_op::sum, aggregate_op::avg, aggregate_op::min, aggregate_op::max}) { + INFO("bits=" << bits << " op=" << static_cast(op)); + decimal_type input{bits, 15, 2}, output = input; + if (op == aggregate_op::sum) output = {static_cast(bits == 64 ? 128 : 256), 37, 2}; + if (op == aggregate_op::avg) output = {static_cast(bits == 64 ? 128 : 256), 19, 6}; + auto definition = layout(op, + logical_type::make_mo_decimal(input, false), + logical_type::make_mo_decimal(output, true), + false, + true); + auto values = make_decimal_literal(input, small(125), true, rows, stream, mr); + auto partial = + local_aggregate_table(cudf::table_view{{values->view()}}, *definition, stream, mr); + REQUIRE(partial->num_rows() == 1); + REQUIRE(partial->num_columns() == 3); + auto result = + merge_aggregate_tables({partial->view(), partial->view()}, *definition, stream, mr); + expect(result->view().column(0), + output, + op == aggregate_op::sum ? 125 * rows * 2 + : op == aggregate_op::avg ? 1250000 + : 125, + stream); + auto count = read(result->view().column(1), 8, stream); + CHECK(count.words[0] == rows * 2); + auto single = merge_aggregate_tables({partial->view()}, *definition, stream, mr); + expect(single->view().column(0), + output, + op == aggregate_op::sum ? 125 * rows + : op == aggregate_op::avg ? 1250000 + : 125, + stream); + } +} +TEST_CASE("MO GPU grouped states preserve common ordinals and NULL identities", + "[decimal_aggregate_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + decimal_type type{256, 15, 2}, output{256, 37, 2}; + for (int rows : {0, 3}) + for (bool valid : {false, true}) { + auto values = make_decimal_literal(type, small(125), valid, rows, stream, mr); + auto groups = cudf::make_numeric_column( + cudf::data_type{cudf::type_id::INT32}, rows, cudf::mask_state::UNALLOCATED, stream, mr); + int32_t labels[]{1, 0, 1}; + if (rows) + REQUIRE(cudaMemcpyAsync(groups->mutable_view().data(), + labels, + sizeof(labels), + cudaMemcpyHostToDevice, + stream.value()) == cudaSuccess); + stream.synchronize(); + auto definition = layout(aggregate_op::sum, + logical_type::make_mo_decimal(type, true), + logical_type::make_mo_decimal(output, true), + true, + true); + auto partial = local_aggregate_table( + cudf::table_view{{values->view(), groups->view()}}, *definition, stream, mr); + auto result = + merge_aggregate_tables({partial->view(), partial->view()}, *definition, stream, mr); + CHECK(result->num_rows() == (rows ? 2 : 0)); + if (rows) { + CHECK(result->view().column(1).null_count() == (valid ? 0 : 2)); + auto count = read(result->view().column(2), 8, stream); + CHECK(count.words[0] == 2); + if (valid) { + auto first = read(result->view().column(1), 32, stream); + CHECK(first.words[0] == 250); + } + } + auto ungrouped = layout(aggregate_op::avg, + logical_type::make_mo_decimal(type, true), + logical_type::make_mo_decimal({256, 19, 6}, true)); + auto empty_partial = + local_aggregate_table(cudf::table_view{{values->view()}}, *ungrouped, stream, mr); + auto empty_result = merge_aggregate_tables({empty_partial->view()}, *ungrouped, stream, mr); + CHECK(empty_result->num_rows() == 1); + CHECK(empty_result->view().column(0).null_count() == (!rows || !valid ? 1 : 0)); + } +} +TEST_CASE("MO GPU SUM checks final precision after partial cancellation", "[decimal_aggregate_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + decimal_type type{256, 65, 0}; + auto definition = layout(aggregate_op::sum, + logical_type::make_mo_decimal(type, false), + logical_type::make_mo_decimal(type, true)); + auto magnitude = decimal_detail::power10(65); + decimal_detail::magnitude one; + one.words[0] = 1; + decimal_detail::subtract(magnitude, one); + std::vector> partials; + for (bool negative : {false, false, true, true}) { + auto coefficient = + decimal_detail::finish(magnitude, negative, type, decimal_error::invalid_input).value; + auto values = make_decimal_literal(type, coefficient, true, 1, stream, mr); + partials.push_back( + local_aggregate_table(cudf::table_view{{values->view()}}, *definition, stream, mr)); + } + auto result = merge_aggregate_tables( + {partials[0]->view(), partials[1]->view(), partials[2]->view(), partials[3]->view()}, + *definition, + stream, + mr); + expect(result->view().column(0), type, 0, stream); + auto one_value = make_decimal_literal(type, small(1), true, 1, stream, mr); + auto one_partial = + local_aggregate_table(cudf::table_view{{one_value->view()}}, *definition, stream, mr); + try { + merge_aggregate_tables({partials[0]->view(), one_partial->view()}, *definition, stream, mr); + FAIL("final SUM precision overflow was ignored"); + } catch (numeric_error const& error) { + CHECK(error.code() == decimal_error::invalid_input); + } + auto recovered = + merge_aggregate_tables({partials[0]->view(), partials[2]->view()}, *definition, stream, mr); + expect(recovered->view().column(0), type, 0, stream); +} +TEST_CASE("MO segmented reductions retain boundaries and nonzero subranges", + "[decimal_aggregate_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + auto values = make_decimal_literal({64, 15, 0}, small(1), true, 8202, stream, mr); + for (auto host : {std::array{0, 4095, 4096, 8197}, + std::array{4093, 4095, 4096, 8197}}) { + rmm::device_uvector offsets(host.size(), stream, mr); + REQUIRE(cudaMemcpyAsync( + offsets.data(), host.data(), sizeof(host), cudaMemcpyHostToDevice, stream.value()) == + cudaSuccess); + stream.synchronize(); + auto state = reduce_aggregate_columns( + aggregate_op::sum, values->view(), {64, 15, 0}, 128, {}, {}, offsets.data(), 3, stream, mr); + auto result = finalize_aggregate_columns(aggregate_op::sum, + state.values->view(), + state.counts->view(), + {64, 15, 0}, + {128, 37, 0}, + stream, + mr); + for (int group = 0; group < 3; ++group) { + auto row = cudf::slice(result->view(), {group, group + 1}, stream); + expect(row[0], {128, 37, 0}, host[group + 1] - host[group], stream); + auto count = cudf::slice(state.counts->view(), {group, group + 1}, stream); + CHECK(read(count[0], 8, stream).words[0] == host[group + 1] - host[group]); + } + } +} +TEST_CASE("MO GPU equality keys match across widths and scales", "[decimal_aggregate_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + for (int sign : {-1, 0, 1}) { + auto narrow = make_decimal_literal({64, 18, 2}, small(sign * 120), true, 2, stream, mr); + auto wide = make_decimal_literal({256, 65, 3}, small(sign * 1200), true, 1, stream, mr); + auto a = make_equality_key(narrow->view(), {64, 18, 2}, stream, mr); + auto b = make_equality_key(wide->view(), {256, 65, 3}, stream, mr); + cudf::hash_join join(cudf::table_view{{b->view()}}, cudf::null_equality::UNEQUAL, stream); + auto indices = join.inner_join(cudf::table_view{{a->view()}}, std::nullopt, stream, mr); + stream.synchronize(); + CHECK(indices.first->size() == 2); + CHECK(indices.second->size() == 2); + } +} +TEST_CASE("MO Decimal256 validity restoration retains outer gather NULLs", + "[decimal_aggregate_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + decimal_type type{256, 15, 2}; + auto value = make_decimal_literal(type, small(125), true, 1, stream, mr); + auto indices = cudf::make_numeric_column( + cudf::data_type{cudf::type_id::INT32}, 2, cudf::mask_state::UNALLOCATED, stream, mr); + int32_t host[]{0, 1}; + REQUIRE(cudaMemcpyAsync(indices->mutable_view().data(), + host, + sizeof(host), + cudaMemcpyHostToDevice, + stream.value()) == cudaSuccess); + stream.synchronize(); + auto gathered = cudf::gather(cudf::table_view{{value->view()}}, + indices->view(), + cudf::out_of_bounds_policy::NULLIFY, + stream, + mr) + ->release(); + auto restored = restore_decimal_validity(std::move(gathered[0]), type, stream, mr); + CHECK(decimal_column_matches(restored->view(), type)); + CHECK(restored->null_count() == 1); + auto zero = make_decimal_literal(type, {}, true, 2, stream, mr); + auto result = evaluate_decimal_columns( + decimal_op::add, restored->view(), type, zero->view(), type, type, nullptr, stream, mr); + CHECK(column_error(result, stream, mr) == decimal_error::none); + CHECK(result.values->null_count() == 1); +} +TEST_CASE("MO aggregate allocation cutpoints retire submitted buffers", "[decimal_aggregate_gpu]") +{ + rmm::cuda_stream stream; + auto upstream = cudf::get_current_device_resource_ref(); + for (bool wide : {false, true}) + for (int groups : {1, 2}) { + decimal_type type = wide ? decimal_type{256, 65, 2} : decimal_type{64, 15, 2}; + auto values = make_decimal_literal(type, small(125), true, 3, stream, upstream); + rmm::device_uvector offsets(groups + 1, stream, upstream); + cudf::size_type host[]{0, groups == 1 ? 3 : 1, 3}; + REQUIRE(cudaMemcpyAsync(offsets.data(), + host, + (groups + 1) * sizeof(*host), + cudaMemcpyHostToDevice, + stream.value()) == cudaSuccess); + stream.synchronize(); + int failures = 0, successes = 0; + for (int cutpoint = 0; cutpoint < 24; ++cutpoint) { + size_t live = 0; + int remaining = cutpoint; + bool bad_free = false; + rmm::mr::callback_memory_resource resource( + [&](size_t bytes, rmm::cuda_stream_view stream, void*) { + if (!remaining--) throw std::bad_alloc(); + auto p = upstream.allocate(stream, bytes, 256); + live += bytes; + return p; + }, + [&](void* p, size_t bytes, rmm::cuda_stream_view stream, void*) { + if (bytes > live) + bad_free = true; + else + live -= bytes; + upstream.deallocate(stream, p, bytes, 256); + }); + auto mr = rmm::to_device_async_resource_ref_checked(&resource); + try { + auto result = reduce_aggregate_columns(aggregate_op::sum, + values->view(), + type, + wide ? 256 : 128, + {}, + {}, + offsets.data(), + groups, + stream, + mr); + REQUIRE(result.counts->size() == groups); + ++successes; + } catch (std::bad_alloc const&) { + ++failures; + } + stream.synchronize(); + CHECK(live == 0); + CHECK_FALSE(bad_free); + } + CHECK(failures > 0); + CHECK(successes > 0); + } +} +TEST_CASE("MO proven local reduction contains the full signed 64-bit domain", + "[decimal_aggregate_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + auto values = cudf::make_numeric_column( + cudf::data_type{cudf::type_id::INT64}, 3, cudf::mask_state::UNALLOCATED, stream, mr); + int64_t endpoints[]{INT64_MAX, INT64_MAX, INT64_MIN}; + REQUIRE(cudaMemcpyAsync(values->mutable_view().data(), + endpoints, + sizeof(endpoints), + cudaMemcpyHostToDevice, + stream.value()) == cudaSuccess); + stream.synchronize(); + auto decimal = cudf::column_view( + cudf::data_type{cudf::type_id::DECIMAL64, 0}, 3, values->view().head(), nullptr, 0); + rmm::device_uvector offsets(2, stream, mr); + cudf::size_type host[]{0, 3}; + REQUIRE( + cudaMemcpyAsync(offsets.data(), host, sizeof(host), cudaMemcpyHostToDevice, stream.value()) == + cudaSuccess); + stream.synchronize(); + for (uint16_t bits : {128, 256}) { + auto state = reduce_aggregate_columns( + aggregate_op::sum, decimal, {64, 18, 0}, bits, {}, {}, offsets.data(), 1, stream, mr); + decimal_type output{bits, static_cast(bits == 128 ? 38 : 65), 0}; + auto result = finalize_aggregate_columns(aggregate_op::sum, + state.values->view(), + state.counts->view(), + {64, 18, 0}, + output, + stream, + mr); + expect(result->view(), output, INT64_MAX - 1, stream); + auto nulls = make_decimal_literal({64, 15, 2}, {}, false, 3, stream, mr); + auto empty = reduce_aggregate_columns( + aggregate_op::sum, nulls->view(), {64, 15, 2}, bits, {}, {}, offsets.data(), 1, stream, mr); + output.scale = 2; + auto absent = finalize_aggregate_columns(aggregate_op::sum, + empty.values->view(), + empty.counts->view(), + {64, 15, 2}, + output, + stream, + mr); + CHECK(absent->null_count() == 1); + } +} +TEST_CASE("MO local widening contains the full signed 128-bit domain", "[decimal_aggregate_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + coefficient maximum, minimum; + for (int i = 0; i < 4; ++i) + maximum.words[i] = UINT32_MAX; + maximum.words[3] >>= 1; + minimum.words[3] = uint32_t(1) << 31; + for (int i = 4; i < 8; ++i) + minimum.words[i] = UINT32_MAX; + std::array bytes{}; + store_coefficient(maximum, bytes.data(), 16); + store_coefficient(maximum, bytes.data() + 16, 16); + store_coefficient(minimum, bytes.data() + 32, 16); + auto values = make_decimal_column({128, 38, 0}, 3, cudf::mask_state::UNALLOCATED, stream, mr); + REQUIRE(cudaMemcpyAsync(values->mutable_view().data(), + bytes.data(), + bytes.size(), + cudaMemcpyHostToDevice, + stream.value()) == cudaSuccess); + rmm::device_uvector offsets(2, stream, mr); + cudf::size_type host[]{0, 3}; + REQUIRE( + cudaMemcpyAsync(offsets.data(), host, sizeof(host), cudaMemcpyHostToDevice, stream.value()) == + cudaSuccess); + stream.synchronize(); + auto state = reduce_aggregate_columns( + aggregate_op::sum, values->view(), {128, 38, 0}, 256, {}, {}, offsets.data(), 1, stream, mr); + auto result = finalize_aggregate_columns(aggregate_op::sum, + state.values->view(), + state.counts->view(), + {128, 38, 0}, + {256, 65, 0}, + stream, + mr); + auto observed = read(result->view(), 32, stream); + maximum.words[0] -= 1; + for (int i = 0; i < 8; ++i) + CHECK(observed.words[i] == maximum.words[i]); +} +TEST_CASE("MO full-width aggregate states survive existing host spill and restore", + "[decimal_aggregate_gpu]") +{ + auto manager = test::operator_utils::initialize_memory_manager(); + auto* gpu = manager->get_memory_space(cucascade::memory::Tier::GPU, 0); + auto* host = manager->get_memory_space(cucascade::memory::Tier::HOST, 0); + REQUIRE(gpu); + REQUIRE(host); + rmm::cuda_stream stream; + auto mr = gpu->get_default_allocator(); + decimal_type type{256, 65, 0}; + auto definition = layout(aggregate_op::sum, + logical_type::make_mo_decimal(type, false), + logical_type::make_mo_decimal(type, true)); + auto value = make_decimal_literal(type, small(125), true, 3, stream, mr); + auto partial = local_aggregate_table(cudf::table_view{{value->view()}}, *definition, stream, mr); + auto batch = make_data_batch(std::move(partial), *gpu, stream, {}); + convertible_data_batch wrapper(batch); + REQUIRE(wrapper.convert({host}, stream, *manager, true).has_value()); + stream.synchronize(); + { + auto ro = batch->to_read_only(); + CHECK(ro.get_memory_space()->get_tier() == cucascade::memory::Tier::HOST); + } + REQUIRE(wrapper.convert({gpu}, stream, *manager, true).has_value()); + stream.synchronize(); + auto ro = batch->to_read_only(); + auto restored = get_cudf_table_view(ro); + auto result = merge_aggregate_tables({restored}, *definition, stream, mr); + expect(result->view().column(0), type, 375, stream); +} diff --git a/test/cpp/embedding/test_decimal_expression_gpu.cpp b/test/cpp/embedding/test_decimal_expression_gpu.cpp index 72a917a2ec..4e5aaa9590 100644 --- a/test/cpp/embedding/test_decimal_expression_gpu.cpp +++ b/test/cpp/embedding/test_decimal_expression_gpu.cpp @@ -173,7 +173,7 @@ TEST_CASE("MO imported scalar signatures execute on the GPU without coercion", CHECK(decimal_column_matches(result->view().column(0), output)); } } - CHECK((sirius_capabilities() & 16u) == 0); + CHECK((sirius_capabilities() & SIRIUS_CAP_MO_EXACT_DECIMAL_V1) != 0); } TEST_CASE("MO imported literals preserve distinct high limbs and signs", "[decimal_import_gpu]") diff --git a/test/cpp/embedding/test_decimal_import.cpp b/test/cpp/embedding/test_decimal_import.cpp index 9d20d0b843..dee0756989 100644 --- a/test/cpp/embedding/test_decimal_import.cpp +++ b/test/cpp/embedding/test_decimal_import.cpp @@ -5,10 +5,13 @@ #include "expression/ast/utils.hpp" #include "expression/value.hpp" #include "helper/type_conversions.hpp" +#include "numeric/decimal_plan.hpp" #include "numeric/decimal_types.hpp" #include "numeric/exact_decimal.hpp" #include +#include +#include using namespace sirius; @@ -163,3 +166,216 @@ TEST_CASE("MO scalar import rejects unrelated identities and malformed literals" {logical_type::make(type_id::BIGINT)}, logical_type::make_mo_decimal({64, 18, 0}, false))); } + +TEST_CASE("MO aggregate import preserves checked signatures and rejects modifiers", + "[decimal_import]") +{ + decimal_fixture::importer importer; + for (auto op : {mo_decimal::aggregate_op::sum, + mo_decimal::aggregate_op::avg, + mo_decimal::aggregate_op::min, + mo_decimal::aggregate_op::max}) { + for (uint16_t bits : {64, 128, 256}) { + mo_decimal::decimal_type input{bits, 15, 2}; + auto output = input; + if (op == mo_decimal::aggregate_op::sum) + output = {static_cast(bits == 64 ? 128 : 256), 37, 2}; + if (op == mo_decimal::aggregate_op::avg) + output = {static_cast(bits == 64 ? 128 : 256), 19, 6}; + auto original = decimal_fixture::plan( + "mo_decimal_cast", decimal_fixture::type(input), {decimal_fixture::literal(input, 125)}); + auto root = original.mutable_relations(0)->mutable_root(); + auto seed = root->input().project().input(); + auto aggregation = root->mutable_input()->mutable_aggregate(); + *aggregation->mutable_input() = std::move(seed); + auto measure = aggregation->add_measures()->mutable_measure(); + measure->set_function_reference(21); + measure->set_phase(substrait::AGGREGATION_PHASE_INITIAL_TO_RESULT); + measure->set_invocation(substrait::AggregateFunction::AGGREGATION_INVOCATION_ALL); + *measure->mutable_output_type() = decimal_fixture::type(output, true); + *measure->add_arguments()->mutable_value() = decimal_fixture::literal(input, 125); + auto name = + std::string(to_duckdb_aggregate_name(mo_decimal::aggregate_function(op))).substr(9); + original.mutable_extensions(1)->mutable_extension_function()->set_name(name); + auto expression = importer.bind(original); + REQUIRE(expression->is_aggregate()); + CHECK(expression->as_aggregate().function() == mo_decimal::aggregate_function(op)); + CHECK(expression->return_type() == logical_type::make_mo_decimal(output, true)); + auto cloned = ast::clone(*expression); + auto reverse = ast::to_duckdb(*cloned); + auto reconstructed = ast::from_duckdb(*reverse); + REQUIRE(reconstructed->is_aggregate()); + CHECK(reconstructed->return_type() == expression->return_type()); + auto bad = original; + bad.mutable_relations(0) + ->mutable_root() + ->mutable_input() + ->mutable_aggregate() + ->mutable_measures(0) + ->mutable_measure() + ->set_invocation(substrait::AggregateFunction::AGGREGATION_INVOCATION_DISTINCT); + REQUIRE_THROWS(importer.bind(bad)); + bad = original; + bad.mutable_relations(0) + ->mutable_root() + ->mutable_input() + ->mutable_aggregate() + ->mutable_measures(0) + ->mutable_measure() + ->set_phase(substrait::AGGREGATION_PHASE_INTERMEDIATE_TO_RESULT); + REQUIRE_THROWS(importer.bind(bad)); + } + } +} +TEST_CASE("MO alias-only inspection preserves shared immutable comparison bindings", + "[decimal_import]") +{ + auto input = logical_type::make_mo_decimal({64, 9, 2}, false); + auto output = logical_type::make(type_id::BOOLEAN).with_nullability(false); + duckdb::vector> args; + for (int64_t number : {125, 200}) + args.push_back(duckdb::make_uniq( + mo_decimal::duckdb_value(decimal_fixture::small(number), input))); + auto expression = mo_decimal::bound_scalar(mo_decimal::decimal_op::less, std::move(args), output); + auto original = expression->return_type; + auto alias = original.GetAlias(); + expression = duckdb::make_uniq( + std::move(expression), duckdb::LogicalType::BOOLEAN, duckdb::BoundCastInfo(nullptr)); + mo_decimal::restore_exact_bound_types(expression); + REQUIRE(expression->GetExpressionClass() == duckdb::ExpressionClass::BOUND_FUNCTION); + CHECK(original.GetAlias() == alias); + CHECK(from_duckdb(original) == output); + REQUIRE_NOTHROW( + mo_decimal::validate_bound_scalar(expression->Cast())); +} +TEST_CASE("exact wire normalization is bounded and rejects ordinary numeric coercion", + "[decimal_import]") +{ + mo_decimal::decimal_type input{64, 9, 2}, output{128, 15, 2}; + auto plan = decimal_fixture::plan( + "mo_decimal_add", + decimal_fixture::type(output), + {decimal_fixture::literal(input, 125), decimal_fixture::literal(input, 25)}); + auto bytes = plan.SerializeAsString(); + REQUIRE_NOTHROW(mo_decimal::normalize_exact_substrait(bytes, 4096)); + REQUIRE_THROWS_AS(mo_decimal::normalize_exact_substrait(bytes, 0), std::bad_alloc); + auto uri = plan.add_extension_urns(); + uri->set_extension_urn_anchor(2); + uri->set_urn("extension:io.substrait:functions_arithmetic"); + auto function = plan.mutable_extensions(1)->mutable_extension_function(); + function->set_extension_urn_reference(2); + function->set_name("add"); + REQUIRE_THROWS(mo_decimal::normalize_exact_substrait(plan.SerializeAsString(), 4096)); +} +TEST_CASE("exact preparation rejects grouping loss and ignored aggregate filters", + "[decimal_import]") +{ + mo_decimal::decimal_type input{64, 9, 2}; + auto plan = decimal_fixture::plan("mo_decimal_sum", decimal_fixture::type(input), {}); + auto root = plan.mutable_relations(0)->mutable_root(); + auto seed = root->input().project().input(); + auto aggregate = root->mutable_input()->mutable_aggregate(); + *aggregate->mutable_input() = std::move(seed); + for (int64_t value : {125, 250}) + *aggregate->add_grouping_expressions() = decimal_fixture::literal(input, value); + auto grouping = aggregate->add_groupings(); + grouping->add_expression_references(0); + grouping->add_expression_references(1); + auto measure = aggregate->add_measures()->mutable_measure(); + measure->set_function_reference(21); + measure->set_phase(substrait::AGGREGATION_PHASE_INITIAL_TO_RESULT); + measure->set_invocation(substrait::AggregateFunction::AGGREGATION_INVOCATION_ALL); + *measure->mutable_output_type() = decimal_fixture::type({128, 15, 2}, true); + *measure->add_arguments()->mutable_value() = decimal_fixture::literal(input, 125); + REQUIRE_NOTHROW(mo_decimal::normalize_exact_substrait(plan.SerializeAsString(), 4096)); + auto bad = plan; + auto get = [](substrait::Plan& p) { + return p.mutable_relations(0)->mutable_root()->mutable_input()->mutable_aggregate(); + }; + get(bad)->mutable_groupings(0)->mutable_expression_references()->RemoveLast(); + REQUIRE_THROWS(mo_decimal::normalize_exact_substrait(bad.SerializeAsString(), 4096)); + bad = plan; + get(bad)->mutable_groupings(0)->set_expression_references(1, 0); + REQUIRE_THROWS(mo_decimal::normalize_exact_substrait(bad.SerializeAsString(), 4096)); + bad = plan; + get(bad)->mutable_groupings(0)->set_expression_references(1, 2); + REQUIRE_THROWS(mo_decimal::normalize_exact_substrait(bad.SerializeAsString(), 4096)); + bad = plan; + get(bad)->mutable_measures(0)->mutable_filter()->mutable_literal()->set_boolean(false); + REQUIRE_THROWS(mo_decimal::normalize_exact_substrait(bad.SerializeAsString(), 4096)); + // Ordinary COUNT shares the typed layout, so its filter must also be + // rejected rather than silently counting rows the caller excluded. + auto urn = bad.add_extension_urns(); + urn->set_extension_urn_anchor(2); + urn->set_urn("extension:io.substrait:functions_aggregate_generic"); + auto function = bad.mutable_extensions(1)->mutable_extension_function(); + function->set_extension_urn_reference(2); + function->set_name("count"); + get(bad) + ->mutable_measures(0) + ->mutable_measure() + ->mutable_output_type() + ->mutable_i64() + ->set_nullability(substrait::Type::NULLABILITY_REQUIRED); + REQUIRE_THROWS(mo_decimal::normalize_exact_substrait(bad.SerializeAsString(), 4096)); + get(bad)->mutable_measures(0)->clear_filter(); + REQUIRE_NOTHROW(mo_decimal::normalize_exact_substrait(bad.SerializeAsString(), 4096)); +} +TEST_CASE("exact read projection carries the selected descriptor into checked binding", + "[decimal_import]") +{ + mo_decimal::decimal_type first{64, 9, 2}, second{256, 15, 3}; + substrait::Expression reference; + reference.mutable_selection()->mutable_root_reference(); + reference.mutable_selection()->mutable_direct_reference()->mutable_struct_field()->set_field(0); + auto plan = decimal_fixture::plan("mo_decimal_add", + decimal_fixture::type(second), + {reference, decimal_fixture::literal(second, 25)}); + auto read = plan.mutable_relations(0) + ->mutable_root() + ->mutable_input() + ->mutable_project() + ->mutable_input() + ->mutable_read(); + read->mutable_base_schema()->clear_names(); + read->mutable_base_schema()->add_names("first"); + read->mutable_base_schema()->add_names("second"); + auto schema = read->mutable_base_schema()->mutable_struct_(); + schema->clear_types(); + *schema->add_types() = decimal_fixture::type(first); + *schema->add_types() = decimal_fixture::type(second); + auto row = read->mutable_virtual_table()->mutable_expressions(0); + row->clear_fields(); + *row->add_fields() = decimal_fixture::literal(first, 125); + *row->add_fields() = decimal_fixture::literal(second, 2500); + read->mutable_projection()->mutable_select()->add_struct_items()->set_field(1); + auto normalized = mo_decimal::normalize_exact_substrait(plan.SerializeAsString(), 4096); + substrait::Plan parsed; + REQUIRE(parsed.ParseFromString(normalized)); + decimal_fixture::importer importer; + auto expression = importer.bind(parsed); + REQUIRE(expression->is_function_call()); + CHECK(expression->return_type() == logical_type::make_mo_decimal(second, false)); + CHECK(expression->as_function_call().arguments()[0]->return_type() == + logical_type::make_mo_decimal(second, false)); + read->mutable_projection()->mutable_select()->mutable_struct_items(0)->set_field(2); + REQUIRE_THROWS(mo_decimal::normalize_exact_substrait(plan.SerializeAsString(), 4096)); + // A zero-column cuDF table cannot retain input row cardinality for a later + // COUNT(*). Decline this shape before it could silently turn N rows into 0. + read->mutable_projection()->mutable_select()->clear_struct_items(); + auto seed = plan.relations(0).root().input().project().input(); + auto aggregate = plan.mutable_relations(0)->mutable_root()->mutable_input()->mutable_aggregate(); + *aggregate->mutable_input() = std::move(seed); + auto measure = aggregate->add_measures()->mutable_measure(); + measure->set_function_reference(21); + measure->set_phase(substrait::AGGREGATION_PHASE_INITIAL_TO_RESULT); + measure->set_invocation(substrait::AggregateFunction::AGGREGATION_INVOCATION_ALL); + measure->mutable_output_type()->mutable_i64()->set_nullability( + substrait::Type::NULLABILITY_REQUIRED); + auto urn = plan.add_extension_urns(); + urn->set_extension_urn_anchor(2); + urn->set_urn("extension:io.substrait:functions_aggregate_generic"); + plan.mutable_extensions(1)->mutable_extension_function()->set_extension_urn_reference(2); + plan.mutable_extensions(1)->mutable_extension_function()->set_name("count"); + REQUIRE_THROWS(mo_decimal::normalize_exact_substrait(plan.SerializeAsString(), 4096)); +} diff --git a/test/cpp/embedding/test_native_control.cpp b/test/cpp/embedding/test_native_control.cpp index d2770b525d..0bc5837347 100644 --- a/test/cpp/embedding/test_native_control.cpp +++ b/test/cpp/embedding/test_native_control.cpp @@ -4,6 +4,7 @@ #include "embedding/buffer_budget.hpp" #include "embedding/control.hpp" #include "embedding/input.hpp" +#include "numeric/decimal_error.hpp" #include "pipeline/completion_handler.hpp" #include @@ -46,6 +47,8 @@ struct recording { bool fail_finish{false}; bool startable{true}; bool runtime_available{true}; + sirius::mo_decimal::decimal_error numeric{sirius::mo_decimal::decimal_error::none}; + std::function cleanup; }; class test_driver final : public query_driver { public: @@ -67,11 +70,14 @@ class test_driver final : public query_driver { changed.wait_until(lock, stop, deadline, [] { return false; }); } if (r_->fail_run) throw std::runtime_error("injected execution error"); + if (r_->numeric != sirius::mo_decimal::decimal_error::none) + throw sirius::mo_decimal::numeric_error(r_->numeric); } void finish() override { check_thread(); if (r_->fail_finish) throw std::runtime_error("injected unprovable cleanup"); + if (r_->cleanup) r_->cleanup(); ++r_->finished; } bool startable() const noexcept override { return r_->startable; } @@ -145,6 +151,30 @@ TEST_CASE("native coordinator keeps preparation execution and destruction on one CHECK_FALSE(f.record->wrong_thread); CHECK(f.control.start(q).code == SIRIUS_INVALID_STATE); } +TEST_CASE("native numeric owner survives producer cancellation during cleanup", "[native_control]") +{ + for (auto error : {sirius::mo_decimal::decimal_error::out_of_range, + sirius::mo_decimal::decimal_error::invalid_input}) { + fixture f; + auto q = f.create(); + auto input = std::make_shared( + 1, std::vector{{23, 0, 0, 0}}, q->stop.get_token(), q->deadline); + q->inputs->reads.push_back(input); + f.record->numeric = error; + f.record->cleanup = [input] { + input->stop(SIRIUS_CANCELLED, "producer cancellation after numeric failure"); + }; + REQUIRE(f.control.prepare(q, 10s).code == SIRIUS_OK); + REQUIRE(f.control.start(q).code == SIRIUS_OK); + auto expected = error == sirius::mo_decimal::decimal_error::out_of_range + ? SIRIUS_NUMERIC_OUT_OF_RANGE + : SIRIUS_NUMERIC_INVALID_INPUT; + CHECK(f.control.wait(q, 10s).code == expected); + CHECK(input->outcome().code == SIRIUS_CANCELLED); + CHECK(f.record->destroyed == 1); + CHECK(f.control.close_query(q, 10s).code == SIRIUS_OK); + } +} TEST_CASE("prepared admission can remain explicitly non-startable", "[native_control]") { diff --git a/test/cpp/embedding/test_native_numeric.cpp b/test/cpp/embedding/test_native_numeric.cpp new file mode 100644 index 0000000000..91880e832f --- /dev/null +++ b/test/cpp/embedding/test_native_numeric.cpp @@ -0,0 +1,723 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +// This target links the production C ABI, not native_gpu_unittest's test backend. +#include "numeric/decimal_aggregate.hpp" +#include "numeric/exact_decimal.hpp" +#include "sirius_c.h" +#include "substrait/plan.pb.h" + +#include + +#include +#include +#include +#include +#include +#include + +using namespace sirius::mo_decimal; +namespace { +void ok(sirius_status status, sirius_error const& error) +{ + INFO(error.message); + REQUIRE(status == SIRIUS_OK); +} +coefficient small(int64_t value) +{ + return load_coefficient(reinterpret_cast(&value), 8); +} +substrait::Type type(decimal_type value, bool nullable) +{ + substrait::Type result; + auto t = result.mutable_user_defined(); + t->set_type_reference(11); + t->set_nullability(nullable ? substrait::Type::NULLABILITY_NULLABLE + : substrait::Type::NULLABILITY_REQUIRED); + for (int parameter : {int(value.bits), int(value.precision), int(value.scale)}) + t->add_type_parameters()->set_integer(parameter); + return result; +} +substrait::Expression field(int index) +{ + substrait::Expression result; + auto s = result.mutable_selection(); + s->mutable_direct_reference()->mutable_struct_field()->set_field(index); + s->mutable_root_reference(); + return result; +} +substrait::Expression literal(decimal_type type, int64_t value) +{ + substrait::Expression result; + auto u = result.mutable_literal()->mutable_user_defined(); + u->set_type_reference(11); + for (int parameter : {int(type.bits), int(type.precision), int(type.scale)}) + u->add_type_parameters()->set_integer(parameter); + auto any = u->mutable_value(); + any->set_type_url("type.googleapis.com/matrixone.sirius.numeric.v1.ExactDecimalLiteral"); + std::string payload(type.bytes() + 2, '\0'); + payload[0] = 10; + payload[1] = type.bytes(); + store_coefficient(small(value), reinterpret_cast(payload.data() + 2), type.bytes()); + any->set_value(payload); + return result; +} +substrait::Expression scalar(uint32_t anchor, + substrait::Type output, + std::vector args) +{ + substrait::Expression result; + auto f = result.mutable_scalar_function(); + f->set_function_reference(anchor); + *f->mutable_output_type() = std::move(output); + for (auto& arg : args) + *f->add_arguments()->mutable_value() = std::move(arg); + return result; +} +struct plan_case { + substrait::Plan plan; + decimal_type input, output; + bool input_nullable{}, output_nullable{true}; + uint32_t output_oid{}; + plan_case(decimal_type input, + decimal_type output, + std::string name, + bool aggregate, + bool nullable = false) + : input(input), + output(output), + input_nullable(nullable), + output_oid(output.bits == 64 ? 32 + : output.bits == 128 ? 33 + : 34) + { + plan.mutable_version()->set_minor_number(78); + auto urn = plan.add_extension_urns(); + urn->set_extension_urn_anchor(1); + urn->set_urn("urn:matrixone:sirius:exact-decimal:v1"); + auto t = plan.add_extensions()->mutable_extension_type(); + t->set_type_anchor(11); + t->set_extension_urn_reference(1); + t->set_name("mo_exact_decimal"); + function(21, name); + auto root = plan.add_relations()->mutable_root(); + root->add_names("answer"); + substrait::Rel* source; + if (aggregate) { + auto a = root->mutable_input()->mutable_aggregate(); + source = a->mutable_input(); + auto m = a->add_measures()->mutable_measure(); + m->set_function_reference(21); + m->set_phase(substrait::AGGREGATION_PHASE_INITIAL_TO_RESULT); + m->set_invocation(substrait::AggregateFunction::AGGREGATION_INVOCATION_ALL); + *m->mutable_output_type() = type(output, true); + *m->add_arguments()->mutable_value() = field(0); + } else { + auto p = root->mutable_input()->mutable_project(); + source = p->mutable_input(); + p->mutable_common()->mutable_emit()->add_output_mapping(1); + *p->add_expressions() = scalar(21, type(output, true), {field(0), literal(input, 1)}); + } + auto read = source->mutable_read(); + read->mutable_named_table()->add_names("__sirius_embedded_v1"); + read->mutable_named_table()->add_names("1"); + read->mutable_base_schema()->add_names("v"); + *read->mutable_base_schema()->mutable_struct_()->add_types() = type(input, nullable); + } + void function(uint32_t anchor, std::string const& name) + { + auto f = plan.add_extensions()->mutable_extension_function(); + f->set_function_anchor(anchor); + f->set_extension_urn_reference(1); + f->set_name(name); + } +}; +struct engine { + sirius_engine_handle* handle{}; + engine() + { + auto config = std::getenv("MO_SIRIUS_TEST_CONFIG"); + REQUIRE(config); + sirius_engine_options options{sizeof(options), + SIRIUS_ABI_VERSION, + config, + static_cast(std::strlen(config)), + 16, + 2, + 0}; + sirius_error error{}; + ok(sirius_engine_create(&options, &handle, &error), error); + } + ~engine() + { + sirius_error error{}; + if (handle) { + auto status = sirius_engine_close(&handle, 30000, &error); + if (status != SIRIUS_OK) + std::cerr << "numeric engine cleanup: " << status << " " << error.message << '\n'; + } + } +}; +struct query { + sirius_query_handle* handle{}; + sirius_input_handle* input{}; + sirius_batch_handle* batch{}; + std::vector extra_inputs; + std::string prepare_error; + explicit query(engine& engine, plan_case const& plan) + { + extra_inputs.reserve(1); + auto wire = plan.plan.SerializeAsString(); + sirius_error error{}; + sirius_query_options options{sizeof(options), SIRIUS_ABI_VERSION, 30000, 0}; + ok(sirius_query_create(engine.handle, &options, wire.data(), wire.size(), &handle, &error), + error); + try { + sirius_column result{plan.output_oid, + plan.output_oid == 10 ? 0 : plan.output.precision, + plan.output_oid == 10 ? 0 : plan.output.scale, + plan.output_nullable, + "answer", + 6, + 0}; + sirius_query_contract contract{ + sizeof(contract), SIRIUS_ABI_VERSION, 0, 0, "numeric", 7, {}, &result, 1}; + ok(sirius_query_bind(handle, &contract, &error), error); + sirius_column column{plan.input.bits == 64 ? 32u + : plan.input.bits == 128 ? 33u + : 34u, + plan.input.precision, + plan.input.scale, + plan.input_nullable, + "v", + 1, + 0}; + sirius_read_column physical{column, 7, 0, 0}; + sirius_read_binding binding{sizeof(binding), + SIRIUS_ABI_VERSION, + 1, + SIRIUS_READ_MO, + 0, + "db", + 2, + "t", + 1, + "s", + 1, + &physical, + 1, + nullptr, + 0, + nullptr, + 0}; + ok(sirius_read_register(handle, &binding, &error), error); + sirius_input_column schema{column.oid, column.width, column.scale, column.nullable}; + ok(sirius_input_register(handle, 1, &schema, 1, &input, &error), error); + } catch (...) { + if (input) sirius_input_close(&input, &error); + sirius_query_close(&handle, 30000, &error); + throw; + } + } + ~query() + { + sirius_error error{}; + if (handle) sirius_query_cancel(handle, &error); + if (batch) sirius_batch_release(&batch, &error); + if (input) sirius_input_close(&input, &error); + for (auto& extra : extra_inputs) + if (extra) sirius_input_close(&extra, &error); + if (handle) { + auto status = sirius_query_close(&handle, 30000, &error); + if (status != SIRIUS_OK) + std::cerr << "numeric query cleanup: " << status << " " << error.message << '\n'; + } + } + sirius_status prepare() + { + sirius_error error{}; + auto status = sirius_query_prepare(handle, 30000, &error); + prepare_error = error.message; + return status; + } + void publish(plan_case const& plan, + std::vector const& values, + bool last_null = false) + { + sirius_error error{}; + ok(sirius_query_start(handle, &error), error); + write(input, plan.input, values, last_null); + } + void register_extra(decimal_type type, bool nullable) + { + sirius_error error{}; + sirius_column column{type.bits == 64 ? 32u + : type.bits == 128 ? 33u + : 34u, + type.precision, + type.scale, + nullable, + "v", + 1, + 0}; + sirius_read_column physical{column, 8, 0, 0}; + sirius_read_binding binding{sizeof(binding), + SIRIUS_ABI_VERSION, + 2, + SIRIUS_READ_MO, + 0, + "db", + 2, + "u", + 1, + "s", + 1, + &physical, + 1, + nullptr, + 0, + nullptr, + 0}; + ok(sirius_read_register(handle, &binding, &error), error); + sirius_input_column schema{column.oid, column.width, column.scale, column.nullable}; + extra_inputs.push_back(nullptr); + ok(sirius_input_register(handle, 2, &schema, 1, &extra_inputs.back(), &error), error); + } + void write(sirius_input_handle* target, + decimal_type type, + std::vector const& values, + bool last_null = false) + { + sirius_error error{}; + if (!values.empty()) { + std::vector data(values.size() * type.bytes()); + for (size_t i = 0; i < values.size(); ++i) + store_coefficient(values[i], data.data() + i * type.bytes(), type.bytes()); + auto null_bytes = last_null ? 8u : 0u; + uint64_t mask = last_null ? uint64_t(1) << (values.size() - 1) : 0; + ok(sirius_input_acquire(target, data.size() + null_bytes, 30000, &batch, &error), error); + ok(sirius_input_write(batch, 0, data.data(), data.size(), &error), error); + if (last_null) ok(sirius_input_write(batch, data.size(), &mask, 8, &error), error); + sirius_input_vector vector{}; + vector.data_bytes = data.size(); + vector.null_offset = data.size(); + vector.null_bytes = null_bytes; + ok(sirius_input_publish(target, &batch, values.size(), &vector, 1, &error), error); + } + ok(sirius_input_finish(target, &error), error); + } + std::vector results(uint32_t bytes, std::vector* nulls = nullptr) + { + std::vector values; + sirius_error error{}; + for (;;) { + auto code = sirius_query_next_result(handle, 30000, &batch, &error); + if (code == SIRIUS_EOF) break; + ok(code, error); + sirius_result_batch_info info{sizeof(info), SIRIUS_ABI_VERSION}; + ok(sirius_result_describe(batch, &info, &error), error); + REQUIRE(info.column_count == 1); + std::vector data(info.columns[0].data_bytes), validity(info.columns[0].null_bytes); + ok(sirius_result_read(batch, info.columns[0].data_offset, data.data(), data.size(), &error), + error); + if (!validity.empty()) + ok(sirius_result_read( + batch, info.columns[0].null_offset, validity.data(), validity.size(), &error), + error); + for (uint32_t row = 0; row < info.rows; ++row) { + if (bytes == 1) { + coefficient value; + value.words[0] = data[row]; + values.push_back(value); + } else + values.push_back(load_coefficient(data.data() + row * bytes, bytes)); + if (nulls) nulls->push_back(!validity.empty() && (validity[row / 8] & (1u << (row % 8)))); + } + ok(sirius_batch_release(&batch, &error), error); + } + ok(sirius_query_wait(handle, 30000, &error), error); + return values; + } +}; +void equal(coefficient const& value, int64_t expected) +{ + auto reference = small(expected); + for (int i = 0; i < 8; ++i) + CHECK(value.words[i] == reference.words[i]); +} +} // namespace +TEST_CASE("production ABI executes exact aggregate widths and empty NULL metadata", + "[native_numeric]") +{ + CHECK(sirius_abi_version() == 1); + CHECK((sirius_capabilities() & SIRIUS_CAP_MO_EXACT_DECIMAL_V1) != 0); + engine engine; + for (uint16_t bits : {64, 128, 256}) + for (auto name : {"mo_decimal_sum", "mo_decimal_avg", "mo_decimal_min", "mo_decimal_max"}) { + INFO("bits=" << bits << " name=" << name); + decimal_type input{bits, 15, 2}, output = input; + if (std::string_view(name) == "mo_decimal_sum") + output = {static_cast(bits == 64 ? 128 : 256), 37, 2}; + if (std::string_view(name) == "mo_decimal_avg") + output = {static_cast(bits == 64 ? 128 : 256), 19, 6}; + for (bool empty : {false, true}) { + plan_case plan(input, output, name, true); + query run(engine, plan); + auto status = run.prepare(); + INFO(run.prepare_error); + REQUIRE(status == SIRIUS_OK); + run.publish( + plan, + empty ? std::vector{} : std::vector{small(125), small(-25)}); + std::vector nulls; + auto values = run.results(output.bytes(), &nulls); + REQUIRE(values.size() == 1); + REQUIRE(nulls.size() == 1); + CHECK(nulls[0] == empty); + if (!empty) + equal(values[0], + std::string_view(name) == "mo_decimal_sum" ? 100 + : std::string_view(name) == "mo_decimal_avg" ? 500000 + : std::string_view(name) == "mo_decimal_min" ? -25 + : 125); + sirius_query_execution_stats stats{sizeof(stats), SIRIUS_ABI_VERSION}; + sirius_error error{}; + ok(sirius_query_get_execution_stats(run.handle, &stats, &error), error); + CHECK(stats.gpu_tasks_started > 0); + CHECK(stats.terminal_status == SIRIUS_OK); + CHECK(stats.result_retained_charged_bytes == 0); + } + } +} +TEST_CASE("production exact ABI rejects raw numeric casts before GPU work", "[native_numeric]") +{ + engine engine; + plan_case plan({64, 9, 2}, {64, 9, 2}, "mo_decimal_add", false); + auto project = plan.plan.mutable_relations(0)->mutable_root()->mutable_input()->mutable_project(); + project->mutable_expressions(0)->Clear(); + auto cast = project->mutable_expressions(0)->mutable_cast(); + *cast->mutable_input() = field(0); + cast->mutable_type()->mutable_i64()->set_nullability(substrait::Type::NULLABILITY_REQUIRED); + query run(engine, plan); + CHECK(run.prepare() == SIRIUS_UNSUPPORTED); + sirius_query_execution_stats stats{sizeof(stats), SIRIUS_ABI_VERSION}; + sirius_error error{}; + ok(sirius_query_get_execution_stats(run.handle, &stats, &error), error); + CHECK(stats.gpu_tasks_started == 0); + CHECK(stats.mo_input_units == 0); +} +TEST_CASE("production numeric statuses retain scalar aggregate and cast owners", "[native_numeric]") +{ + engine engine; + auto magnitude = decimal_detail::power10(65); + decimal_detail::magnitude one; + one.words[0] = 1; + decimal_detail::subtract(magnitude, one); + auto maximum = + decimal_detail::finish(magnitude, false, {256, 65, 0}, decimal_error::invalid_input).value; + for (int mode : {0, 1, 2, 3}) { + plan_case plan(mode == 3 ? decimal_type{64, 3, 0} : decimal_type{256, 65, 0}, + mode == 3 ? decimal_type{64, 2, 0} : decimal_type{256, 65, 0}, + mode == 3 ? "mo_decimal_cast" + : mode == 1 ? "mo_decimal_add" + : "mo_decimal_sum", + mode == 0 || mode == 2); + if (mode == 2) { + plan.function(22, "mo_decimal_add"); + auto measure = plan.plan.mutable_relations(0) + ->mutable_root() + ->mutable_input() + ->mutable_aggregate() + ->mutable_measures(0) + ->mutable_measure(); + *measure->mutable_arguments(0)->mutable_value() = + scalar(22, type({256, 65, 0}, false), {field(0), literal({256, 65, 0}, 1)}); + } + if (mode == 3) + plan.plan.mutable_relations(0) + ->mutable_root() + ->mutable_input() + ->mutable_project() + ->mutable_expressions(0) + ->mutable_scalar_function() + ->mutable_arguments() + ->RemoveLast(); + { + query run(engine, plan); + auto status = run.prepare(); + INFO(run.prepare_error); + REQUIRE(status == SIRIUS_OK); + run.publish(plan, + mode == 3 ? std::vector{small(999)} + : std::vector{maximum, small(1)}); + sirius_error error{}; + auto result = sirius_query_next_result(run.handle, 30000, &run.batch, &error); + INFO(error.message); + CHECK(result == + (mode == 0 || mode == 3 ? SIRIUS_NUMERIC_INVALID_INPUT : SIRIUS_NUMERIC_OUT_OF_RANGE)); + CHECK(run.batch == nullptr); + } + plan_case control({64, 15, 2}, {128, 37, 2}, "mo_decimal_sum", true); + query recovered(engine, control); + auto status = recovered.prepare(); + INFO(recovered.prepare_error); + REQUIRE(status == SIRIUS_OK); + recovered.publish(control, {small(125)}); + auto result = recovered.results(16); + REQUIRE(result.size() == 1); + equal(result[0], 125); + } +} +TEST_CASE("production exact read filters use checked types before projection", "[native_numeric]") +{ + engine engine; + for (uint16_t bits : {64, 128, 256}) { + decimal_type domain{bits, 15, 2}; + plan_case plan(domain, domain, "mo_decimal_add", false); + plan.function(22, "mo_decimal_less"); + auto read = plan.plan.mutable_relations(0) + ->mutable_root() + ->mutable_input() + ->mutable_project() + ->mutable_input() + ->mutable_read(); + substrait::Type boolean; + boolean.mutable_bool_()->set_nullability(substrait::Type::NULLABILITY_REQUIRED); + *read->mutable_filter() = scalar(22, boolean, {field(0), literal(domain, 200)}); + read->mutable_projection()->mutable_select()->add_struct_items()->set_field(0); + { + query run(engine, plan); + auto status = run.prepare(); + INFO(run.prepare_error); + REQUIRE(status == SIRIUS_OK); + run.publish(plan, {small(125), small(999)}); + auto values = run.results(domain.bytes()); + REQUIRE(values.size() == 1); + equal(values[0], 126); + } + auto urn = plan.plan.add_extension_urns(); + urn->set_extension_urn_anchor(2); + urn->set_urn("extension:io.substrait:functions_comparison"); + auto ordinary = plan.plan.mutable_extensions(2)->mutable_extension_function(); + ordinary->set_extension_urn_reference(2); + ordinary->set_name("lt"); + query rejected(engine, plan); + CHECK(rejected.prepare() == SIRIUS_UNSUPPORTED); + sirius_query_execution_stats stats{sizeof(stats), SIRIUS_ABI_VERSION}; + sirius_error error{}; + ok(sirius_query_get_execution_stats(rejected.handle, &stats, &error), error); + CHECK(stats.gpu_tasks_started == 0); + CHECK(stats.mo_input_units == 0); + } +} +TEST_CASE("production exact CASE suppresses inactive failures and COALESCE retains types", + "[native_numeric]") +{ + engine engine; + for (uint16_t bits : {64, 128, 256}) { + decimal_type domain{bits, 3, 0}; + plan_case plan(domain, domain, "mo_decimal_add", false); + plan.output_nullable = false; + plan.function(22, "mo_decimal_less"); + auto expression = plan.plan.mutable_relations(0) + ->mutable_root() + ->mutable_input() + ->mutable_project() + ->mutable_expressions(0); + expression->Clear(); + auto c = expression->mutable_if_then(); + auto branch = c->add_ifs(); + substrait::Type boolean; + boolean.mutable_bool_()->set_nullability(substrait::Type::NULLABILITY_REQUIRED); + *branch->mutable_if_() = scalar(22, boolean, {field(0), literal(domain, 200)}); + *branch->mutable_then() = scalar(21, type(domain, false), {field(0), literal(domain, 1)}); + *c->mutable_else_() = literal(domain, 9); + { + query run(engine, plan); + auto status = run.prepare(); + INFO(run.prepare_error); + REQUIRE(status == SIRIUS_OK); + run.publish(plan, {small(999), small(125)}); + auto values = run.results(domain.bytes()); + REQUIRE(values.size() == 2); + equal(values[0], 9); + equal(values[1], 126); + } + plan_case coalesce(domain, domain, "coalesce", false, true); + coalesce.output_nullable = false; + auto urn = coalesce.plan.add_extension_urns(); + urn->set_extension_urn_anchor(2); + urn->set_urn("extension:io.substrait:functions_boolean"); + coalesce.plan.mutable_extensions(1)->mutable_extension_function()->set_extension_urn_reference( + 2); + auto f = coalesce.plan.mutable_relations(0) + ->mutable_root() + ->mutable_input() + ->mutable_project() + ->mutable_expressions(0) + ->mutable_scalar_function(); + *f->mutable_output_type() = type(domain, false); + *f->mutable_arguments(1)->mutable_value() = literal(domain, 9); + { + query run(engine, coalesce); + auto status = run.prepare(); + INFO(run.prepare_error); + REQUIRE(status == SIRIUS_OK); + run.publish(coalesce, {small(999), small(125)}, true); + auto values = run.results(domain.bytes()); + REQUIRE(values.size() == 2); + equal(values[0], 999); + equal(values[1], 9); + } + auto lower = coalesce.plan.relations(0).root().input(); + auto upper = + coalesce.plan.mutable_relations(0)->mutable_root()->mutable_input()->mutable_project(); + upper->Clear(); + *upper->mutable_input() = std::move(lower); + upper->mutable_common()->mutable_emit()->add_output_mapping(1); + coalesce.output = {bits, 4, 0}; + coalesce.function(23, "mo_decimal_add"); + *upper->add_expressions() = + scalar(23, type(coalesce.output, false), {field(0), literal(domain, 1)}); + query projected(engine, coalesce); + auto status = projected.prepare(); + INFO(projected.prepare_error); + REQUIRE(status == SIRIUS_OK); + projected.publish(coalesce, {small(999), small(125)}, true); + auto values = projected.results(domain.bytes()); + REQUIRE(values.size() == 2); + equal(values[0], 1000); + equal(values[1], 10); + } +} +TEST_CASE("production exact join keys and outer NULLs preserve values across scales", + "[native_numeric]") +{ + engine engine; + for (bool outer : {false, true}) { + decimal_type left{64, 15, 2}, right{256, 15, 3}; + plan_case plan(left, outer ? right : left, "mo_decimal_equal", false); + plan.output_nullable = outer; + auto project = + plan.plan.mutable_relations(0)->mutable_root()->mutable_input()->mutable_project(); + project->clear_expressions(); + *project->add_expressions() = field(outer ? 1 : 0); + project->mutable_common()->mutable_emit()->set_output_mapping(0, 2); + auto read = project->input(); + auto join = project->mutable_input()->mutable_join(); + *join->mutable_left() = read; + *join->mutable_right() = read; + join->mutable_right()->mutable_read()->mutable_named_table()->set_names(1, "2"); + *join->mutable_right()->mutable_read()->mutable_base_schema()->mutable_struct_()->mutable_types( + 0) = type(right, false); + join->set_type(outer ? substrait::JoinRel::JOIN_TYPE_LEFT + : substrait::JoinRel::JOIN_TYPE_INNER); + substrait::Type boolean; + boolean.mutable_bool_()->set_nullability(substrait::Type::NULLABILITY_REQUIRED); + *join->mutable_expression() = scalar(21, boolean, {field(0), field(1)}); + query run(engine, plan); + run.register_extra(right, false); + auto status = run.prepare(); + INFO(run.prepare_error); + REQUIRE(status == SIRIUS_OK); + run.publish(plan, {small(125), small(500)}); + run.write(run.extra_inputs[0], right, {small(1250)}); + std::vector nulls; + auto values = run.results((outer ? right : left).bytes(), &nulls); + REQUIRE(values.size() == (outer ? 2 : 1)); + int valid = 0, missing = 0; + for (size_t i = 0; i < values.size(); ++i) + if (nulls[i]) + ++missing; + else { + equal(values[i], outer ? 1250 : 125); + ++valid; + } + CHECK(valid == 1); + CHECK(missing == (outer ? 1 : 0)); + } +} +TEST_CASE("production Decimal256 ordering retains signed coefficients and NULL position", + "[native_numeric]") +{ + engine engine; + decimal_type domain{256, 15, 2}; + plan_case plan(domain, domain, "mo_decimal_add", false, true); + auto root = plan.plan.mutable_relations(0)->mutable_root(); + auto project = root->mutable_input()->mutable_project(); + *project->mutable_expressions(0) = field(0); + auto input = root->input(); + auto sort = root->mutable_input()->mutable_sort(); + *sort->mutable_input() = std::move(input); + auto key = sort->add_sorts(); + *key->mutable_expr() = field(0); + key->set_direction(substrait::SortField::SORT_DIRECTION_DESC_NULLS_LAST); + query run(engine, plan); + auto status = run.prepare(); + INFO(run.prepare_error); + REQUIRE(status == SIRIUS_OK); + run.publish(plan, {small(-125), small(500), small(99)}, true); + std::vector nulls; + auto values = run.results(32, &nulls); + REQUIRE(values.size() == 3); + equal(values[0], 500); + equal(values[1], -125); + CHECK(nulls[2]); +} +TEST_CASE("production grouped exact and ordinary slots finalize through the pipeline", + "[native_numeric]") +{ + engine engine; + for (bool count : {false, true}) + for (bool ordered : {false, true}) { + decimal_type input{256, 15, 2}, output{256, 37, 2}; + plan_case plan(input, output, "mo_decimal_sum", true); + auto root = plan.plan.mutable_relations(0)->mutable_root(); + auto aggregate = root->mutable_input()->mutable_aggregate(); + *aggregate->add_grouping_expressions() = field(0); + aggregate->add_groupings()->add_expression_references(0); + plan.function(22, "count"); + plan.plan.mutable_extensions(2)->mutable_extension_function()->set_extension_urn_reference(2); + auto urn = plan.plan.add_extension_urns(); + urn->set_extension_urn_anchor(2); + urn->set_urn("extension:io.substrait:functions_aggregate_generic"); + auto ordinary = aggregate->add_measures()->mutable_measure(); + ordinary->set_function_reference(22); + ordinary->set_phase(substrait::AGGREGATION_PHASE_INITIAL_TO_RESULT); + ordinary->set_invocation(substrait::AggregateFunction::AGGREGATION_INVOCATION_ALL); + ordinary->mutable_output_type()->mutable_i64()->set_nullability( + substrait::Type::NULLABILITY_REQUIRED); + auto grouped = root->input(); + auto projection = root->mutable_input()->mutable_project(); + projection->mutable_common()->mutable_emit()->add_output_mapping(3); + *projection->add_expressions() = field(count ? 2 : 1); + if (ordered) { + auto sort = projection->mutable_input()->mutable_sort(); + *sort->mutable_input() = std::move(grouped); + auto key = sort->add_sorts(); + *key->mutable_expr() = field(0); + key->set_direction(substrait::SortField::SORT_DIRECTION_DESC_NULLS_LAST); + } else + *projection->mutable_input() = std::move(grouped); + if (count) { + plan.output_oid = 23; + plan.output = {}; + plan.output_nullable = false; + } + query run(engine, plan); + auto status = run.prepare(); + INFO(run.prepare_error); + REQUIRE(status == SIRIUS_OK); + run.publish(plan, {small(125), small(-25), small(125)}); + auto values = run.results(count ? 8 : 32); + REQUIRE(values.size() == 2); + if (ordered) { + equal(values[0], count ? 2 : 250); + equal(values[1], count ? 1 : -25); + } else { + std::sort(values.begin(), values.end(), [](auto const& a, auto const& b) { + return compare_coefficient(a, b) < 0; + }); + equal(values[0], count ? 1 : -25); + equal(values[1], count ? 2 : 250); + } + } +} diff --git a/test/cpp/embedding/test_native_result_codec.cpp b/test/cpp/embedding/test_native_result_codec.cpp index 71254294df..5daf8a3749 100644 --- a/test/cpp/embedding/test_native_result_codec.cpp +++ b/test/cpp/embedding/test_native_result_codec.cpp @@ -1,5 +1,8 @@ /* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ #include "embedding/result_codec.hpp" +#include "numeric/exact_decimal_gpu.hpp" + +#include #include #include @@ -141,3 +144,39 @@ TEST_CASE("native result codec keeps decimal integer bits and unsigned high bits CHECK(read(*batch, batch->columns[1].data_offset + 8) == unsigned_values[1]); result->cancel(); } +TEST_CASE("wide result codec validates child NULLs against parent before publication", + "[native_result_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + for (bool parent_null : {false, true}) { + auto column = + sirius::mo_decimal::make_decimal_literal({256, 65, 2}, {}, !parent_null, 2, stream, mr); + auto contents = column->release(); + contents.children[0]->set_null_mask( + cudf::create_null_mask(2, cudf::mask_state::ALL_NULL, stream, mr), 2); + column = std::make_unique( + cudf::data_type{cudf::type_id::STRUCT}, + 2, + rmm::device_buffer{}, + contents.null_mask ? std::move(*contents.null_mask) : rmm::device_buffer{}, + parent_null ? 2 : 0, + std::move(contents.children)); + stream.synchronize(); + cudf::table_view table{{column->view()}}; + std::vector schema{{34, 65, 2, true, "wide"}}; + auto result = std::make_shared(); + result->activate(std::make_shared()); + std::shared_ptr scratch, batch; + REQUIRE(result->try_allocate(64u << 10, 0, scratch) == SIRIUS_OK); + auto layout = size_native_result(table, 0, schema, *scratch, stream, 4096); + REQUIRE(result->try_allocate(layout.bytes, 2, batch) == SIRIUS_OK); + if (parent_null) { + REQUIRE_NOTHROW(encode_native_result(table, 0, schema, layout, *batch, *scratch, stream)); + CHECK(read(*batch, batch->columns[0].null_offset) == 3); + } else + REQUIRE_THROWS_AS(encode_native_result(table, 0, schema, layout, *batch, *scratch, stream), + failure); + result->cancel(); + } +}