From 52af44b1e91267033f2cb7a66b4241469eed1404 Mon Sep 17 00:00:00 2001 From: bRong Njam Date: Thu, 8 Oct 2026 12:59:22 +0800 Subject: [PATCH] feat(numeric): bind MO exact-decimal plans and scalar expressions --- bench/mo_exact_decimal.cpp | 21 + cmake/sirius-embedding-tests.cmake | 17 +- cmake/sirius-embedding.cmake | 8 +- docs/mo-exact-decimal-import-kernels.csv | 57 +++ docs/mo-exact-decimal-import.md | 124 ++++++ .../sirius/numeric/v1/exact_decimal.proto | 10 + scripts/export_embed_link.py | 9 + scripts/tests/test_export_embed_link.py | 21 +- src/cudf/cudf_utils.hpp | 18 +- src/expression/from_duckdb.cpp | 3 + src/expression/function_id.cpp | 73 +++- src/expression/function_id.hpp | 18 +- src/expression/to_duckdb.cpp | 6 + src/expression/value.cpp | 16 +- src/expression/value.hpp | 3 +- src/expression_evaluator/CMakeLists.txt | 2 + .../expression_evaluator.cpp | 169 ++++---- .../expression_evaluator.hpp | 7 + .../specializations/case.cpp | 12 + .../specializations/constant.cpp | 10 + .../specializations/function.cpp | 3 + .../specializations/mo_conditional.cpp | 260 +++++++++++ .../specializations/mo_decimal.cpp | 93 ++++ .../specializations/operator.cpp | 10 + .../specializations/reference.cpp | 8 +- src/helper/logical_type.hpp | 46 +- src/helper/type_conversions.cpp | 5 + src/numeric/CMakeLists.txt | 3 +- src/numeric/decimal_error.hpp | 21 + src/numeric/decimal_functions.cpp | 244 +++++++++++ src/numeric/decimal_functions.hpp | 31 ++ src/numeric/decimal_import.cpp | 167 +++++++ src/numeric/decimal_import.hpp | 14 + src/numeric/decimal_type.hpp | 51 +++ src/numeric/decimal_types.cpp | 139 ++++++ src/numeric/decimal_types.hpp | 17 + src/numeric/exact_decimal.hpp | 42 +- src/numeric/exact_decimal_gpu.cu | 89 ++++ src/numeric/exact_decimal_gpu.hpp | 9 + substrait | 2 +- test/cpp/embedding/decimal_import_fixture.hpp | 147 +++++++ .../embedding/test_decimal_expression_gpu.cpp | 407 ++++++++++++++++++ test/cpp/embedding/test_decimal_import.cpp | 165 +++++++ 43 files changed, 2434 insertions(+), 143 deletions(-) create mode 100644 docs/mo-exact-decimal-import-kernels.csv create mode 100644 docs/mo-exact-decimal-import.md create mode 100644 proto/matrixone/sirius/numeric/v1/exact_decimal.proto create mode 100644 src/expression_evaluator/specializations/mo_conditional.cpp create mode 100644 src/expression_evaluator/specializations/mo_decimal.cpp create mode 100644 src/numeric/decimal_error.hpp create mode 100644 src/numeric/decimal_functions.cpp create mode 100644 src/numeric/decimal_functions.hpp create mode 100644 src/numeric/decimal_import.cpp create mode 100644 src/numeric/decimal_import.hpp create mode 100644 src/numeric/decimal_type.hpp create mode 100644 src/numeric/decimal_types.cpp create mode 100644 src/numeric/decimal_types.hpp create mode 100644 test/cpp/embedding/decimal_import_fixture.hpp create mode 100644 test/cpp/embedding/test_decimal_expression_gpu.cpp create mode 100644 test/cpp/embedding/test_decimal_import.cpp diff --git a/bench/mo_exact_decimal.cpp b/bench/mo_exact_decimal.cpp index 0554d3ccba..a9df58f6e6 100644 --- a/bench/mo_exact_decimal.cpp +++ b/bench/mo_exact_decimal.cpp @@ -1,4 +1,6 @@ /* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "expression/ast/node.hpp" +#include "expression_evaluator/expression_evaluator.hpp" #include "numeric/exact_decimal_gpu.hpp" #include @@ -117,4 +119,23 @@ int main() return evaluate_decimal_columns( operation, x->view(), wide, y->view(), wide, wide, nullptr, stream.view(), mr); }); + // A values-identical ordinary-expression control makes the evaluator's + // additional admission/error/mask cost visible. It is not a full-domain + // semantic baseline for checked numeric errors or a SQL rollout gate. + for (bool checked : {false, true}) { + auto type = checked ? sirius::logical_type::make_mo_decimal(narrow, false) + : sirius::logical_type::make_decimal(18, 0); + std::vector> arguments; + arguments.push_back(std::make_unique(sirius::ast::reference{0, type})); + arguments.push_back(std::make_unique(sirius::ast::reference{1, type})); + sirius::ast::node expression(sirius::ast::function_call{ + checked ? sirius::function_id::mo_decimal_add : sirius::function_id::add, + std::move(arguments), + type}); + sirius::expression_evaluator evaluator( + expression, mr, stream.view(), sirius::expression_evaluator_strategy::MATERIALIZE); + measure(checked ? "exact_expression_add64" : "ordinary_expression_add64", stream.view(), [&] { + return evaluator.evaluate(cudf::table_view{{a->view(), b->view()}}); + }); + } } diff --git a/cmake/sirius-embedding-tests.cmake b/cmake/sirius-embedding-tests.cmake index 3e8bae102a..7cb4f6105e 100644 --- a/cmake/sirius-embedding-tests.cmake +++ b/cmake/sirius-embedding-tests.cmake @@ -49,17 +49,24 @@ target_link_libraries(sirius_native_result_integration PRIVATE Sirius::embed add_executable( sirius_native_binding_unittest test/cpp/exec/execution_evidence_unittest.cpp - test/cpp/embedding/test_native_binding.cpp src/embedding/plan.cpp) + test/cpp/embedding/test_native_binding.cpp + test/cpp/embedding/test_decimal_import.cpp + test/cpp/expression/test_function_id.cpp + test/cpp/helper/test_logical_type.cpp + src/embedding/plan.cpp) target_compile_features(sirius_native_binding_unittest PRIVATE cxx_std_20) target_include_directories( sirius_native_binding_unittest PRIVATE "${CMAKE_SOURCE_DIR}/third_party/catch" + "${SIRIUS_SUBSTRAIT_DIR}/src/include" "${CMAKE_CURRENT_SOURCE_DIR}/src" "${CMAKE_CURRENT_SOURCE_DIR}/tae-scanner/include" "${SIRIUS_SUBSTRAIT_DIR}/third_party" "${SIRIUS_SUBSTRAIT_DIR}/third_party/substrait") -target_link_libraries(sirius_native_binding_unittest PRIVATE sirius_extension - Threads::Threads) +target_link_libraries( + sirius_native_binding_unittest + PRIVATE sirius_extension duckdb_static dummy_static_extension_loader + Threads::Threads) add_executable( sirius_native_gpu_unittest @@ -68,6 +75,7 @@ add_executable( test/cpp/embedding/test_native_admission.cpp 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_tae_gpu.cpp src/embedding/c_api.cpp) target_compile_features(sirius_native_gpu_unittest PRIVATE cxx_std_20) @@ -75,6 +83,9 @@ target_include_directories( sirius_native_gpu_unittest PRIVATE "${CMAKE_SOURCE_DIR}/third_party/catch" "${CMAKE_CURRENT_SOURCE_DIR}/test/cpp" + "${SIRIUS_SUBSTRAIT_DIR}/src/include" + "${SIRIUS_SUBSTRAIT_DIR}/third_party" + "${SIRIUS_SUBSTRAIT_DIR}/third_party/substrait" "${CMAKE_CURRENT_SOURCE_DIR}/tae-scanner/include" "${CMAKE_CURRENT_SOURCE_DIR}/src" "${CMAKE_CURRENT_SOURCE_DIR}/src/compression/simpatico_codegen/src") diff --git a/cmake/sirius-embedding.cmake b/cmake/sirius-embedding.cmake index 7209a1763e..0def528be9 100644 --- a/cmake/sirius-embedding.cmake +++ b/cmake/sirius-embedding.cmake @@ -26,9 +26,11 @@ target_link_libraries(sirius_c_smoke PRIVATE Sirius::embed) set_property( TARGET sirius_c_smoke APPEND - PROPERTY LINK_DEPENDS - "${CMAKE_CURRENT_SOURCE_DIR}/scripts/export_embed_link.py" - "${CMAKE_CURRENT_SOURCE_DIR}/src/sirius_c.h") + PROPERTY + LINK_DEPENDS + "${CMAKE_CURRENT_SOURCE_DIR}/scripts/export_embed_link.py" + "${CMAKE_CURRENT_SOURCE_DIR}/proto/matrixone/sirius/numeric/v1/exact_decimal.proto" + "${CMAKE_CURRENT_SOURCE_DIR}/src/sirius_c.h") add_custom_command( TARGET sirius_c_smoke POST_BUILD diff --git a/docs/mo-exact-decimal-import-kernels.csv b/docs/mo-exact-decimal-import-kernels.csv new file mode 100644 index 0000000000..24e6e5e6ce --- /dev/null +++ b/docs/mo-exact-decimal-import-kernels.csv @@ -0,0 +1,57 @@ +case,rows,repetition,wall_ms,rows_per_second +cudf_add64,262144,0,0.201209,1.30284e+09 +cudf_add64,262144,1,0.18722,1.40019e+09 +cudf_add64,262144,2,0.204459,1.28213e+09 +cudf_add64,262144,3,0.199659,1.31296e+09 +cudf_add64,262144,4,0.19854,1.32036e+09 +cudf_add64,262144,5,0.198959,1.31758e+09 +cudf_add64,262144,6,0.19894,1.3177e+09 +checked_add64,262144,0,0.02799,9.36563e+09 +checked_add64,262144,1,0.02695,9.72705e+09 +checked_add64,262144,2,0.02728,9.60938e+09 +checked_add64,262144,3,0.02825,9.27943e+09 +checked_add64,262144,4,0.033239,7.88664e+09 +checked_add64,262144,5,0.02721,9.63411e+09 +checked_add64,262144,6,0.03398,7.71466e+09 +cudf_widen_multiply64_to128,262144,0,0.510259,5.13747e+08 +cudf_widen_multiply64_to128,262144,1,0.512328,5.11672e+08 +cudf_widen_multiply64_to128,262144,2,0.509408,5.14605e+08 +cudf_widen_multiply64_to128,262144,3,0.505559,5.18523e+08 +cudf_widen_multiply64_to128,262144,4,0.505398,5.18688e+08 +cudf_widen_multiply64_to128,262144,5,0.503878,5.20253e+08 +cudf_widen_multiply64_to128,262144,6,0.505139,5.18954e+08 +checked_multiply64_to128,262144,0,0.510949,5.13053e+08 +checked_multiply64_to128,262144,1,0.503798,5.20336e+08 +checked_multiply64_to128,262144,2,0.510058,5.13949e+08 +checked_multiply64_to128,262144,3,0.503409,5.20738e+08 +checked_multiply64_to128,262144,4,0.507878,5.16155e+08 +checked_multiply64_to128,262144,5,0.504078,5.20047e+08 +checked_multiply64_to128,262144,6,0.504059,5.20066e+08 +checked_add256,262144,0,0.263589,9.94518e+08 +checked_add256,262144,1,0.260059,1.00802e+09 +checked_add256,262144,2,0.268119,9.77715e+08 +checked_add256,262144,3,0.259359,1.01074e+09 +checked_add256,262144,4,0.260209,1.00744e+09 +checked_add256,262144,5,0.26219,9.99825e+08 +checked_add256,262144,6,0.264499,9.91096e+08 +checked_divide256,262144,0,4.84093,5.41515e+07 +checked_divide256,262144,1,4.83929,5.41699e+07 +checked_divide256,262144,2,4.83929,5.41699e+07 +checked_divide256,262144,3,4.8423,5.41362e+07 +checked_divide256,262144,4,4.84289,5.41296e+07 +checked_divide256,262144,5,4.84123,5.41482e+07 +checked_divide256,262144,6,4.84392,5.41181e+07 +ordinary_expression_add64,262144,0,0.199989,1.31079e+09 +ordinary_expression_add64,262144,1,0.20326,1.2897e+09 +ordinary_expression_add64,262144,2,0.19919,1.31605e+09 +ordinary_expression_add64,262144,3,0.198689,1.31937e+09 +ordinary_expression_add64,262144,4,0.198569,1.32017e+09 +ordinary_expression_add64,262144,5,0.1986,1.31996e+09 +ordinary_expression_add64,262144,6,0.199959,1.31099e+09 +exact_expression_add64,262144,0,0.15899,1.64881e+09 +exact_expression_add64,262144,1,0.167109,1.5687e+09 +exact_expression_add64,262144,2,0.16494,1.58933e+09 +exact_expression_add64,262144,3,0.156899,1.67078e+09 +exact_expression_add64,262144,4,0.16549,1.58405e+09 +exact_expression_add64,262144,5,0.157249,1.66706e+09 +exact_expression_add64,262144,6,0.15738,1.66568e+09 diff --git a/docs/mo-exact-decimal-import.md b/docs/mo-exact-decimal-import.md new file mode 100644 index 0000000000..c739a59b55 --- /dev/null +++ b/docs/mo-exact-decimal-import.md @@ -0,0 +1,124 @@ +# Embedded MO exact-decimal import and scalar execution + +This is the first of the two native integration increments approved for +[MO #28968](https://github.com/matrixorigin/matrixone/issues/28968) and +[MO #28966](https://github.com/matrixorigin/matrixone/issues/28966). It follows +merged Sirius #25 and importer #4. The semantic contract remains the approved +MO #29449 document blob `42a89f09a1d168d02b9583cb3ea7b4de6dbb5634`. + +## Contract and ownership + +Physical coefficient width, declared precision, scale, and nullability survive +Substrait import, DuckDB binding, Sirius translation, cloning and GPU evaluation. +The first type owner is the validated import descriptor. Ordinary DuckDB +decimal inference cannot substitute a narrower carrier. Decimal64/128 use +tagged BIGINT/HUGEINT DuckDB carriers; Decimal256 uses the signed-high/unsigned +low limb STRUCT introduced by #25. Sirius retains a distinct MO decimal type. + +The private alias records the descriptor and is validated against the complete +physical carrier on conversion. It is internal binding metadata, not a new +wire encoding or registered SQL type. The public literal wire contract is the +approved `ExactDecimalLiteral` protobuf with one fixed-width coefficient field, +defined under `proto/matrixone/sirius/numeric/v1/`. +The build-tree SDK exports and fingerprints this schema beside its C header. +The pinned protobuf reader accepts standard unknown fields and length varints; +the final coefficient must still have exactly the declared physical width. + +One scoped importer validates exact-decimal identities and output type anchors +from its own bounded serialized plan. Private scalar markers bind immutable +input/output types and retain MO's declared result. They cannot execute on CPU. +The typed signature survives the DuckDB/Sirius round trip. Existing function +IDs retain their values; exact-decimal IDs are appended. + +GPU scalar execution uses the caller's task stream and resource. Checked +results retain the failing operation's internal error kind. CASE and COALESCE +propagate active rows, so an inactive arithmetic or narrowing failure does not +become a query failure. At most one conditional branch's temporary columns +remain owned at a time. Failure proves stream quiescence before cleanup; +unprovable owners transfer to the existing fatal process quarantine. +cuDF STRUCT selection may superimpose the parent bitmap on child columns. +After validating canonical operands and proving stream completion, conditional +selection removes only those redundant child masks without copying coefficient +data. A following exact operation verifies that parent-only validity survives. + +No process-wide type/function registry or GPU allocator is replaced. Query +metadata remains bounded by the existing plan/admission contract. Literal +broadcast and error reduction allocate only through the supplied resource; +there is no per-row host callback or numeric string conversion. + +## Exposure boundary + +The helper and direct-consumer tests can import and evaluate the scalar family. +Production embedded admission still rejects the new extension, and the exact +capability remains clear. Ordinary importer calls continue without the scoped +handler. Aggregates, keys, complete admission, public statuses 12/13 and +capability advertisement belong to the next native increment. + +The remaining approved sequence is: + +1. Sirius import/type/scalar plumbing, with capability disabled (this increment). +2. Sirius aggregates/keys, complete native admission and capability. +3. MO lowering, reconstruction, typed errors and query-local evidence. +4. Real native-MO/embedded-MO SF1/SF10 parity, resource and performance campaign. +5. MO embedded default and strict Flight configuration/recovery cutover. +6. Sidecar Flight retirement after the verified MO release artifact exists. + +The campaign requires streams 2 plus controls 1/4, keeps the 10% numeric +regression policy, and retains the 1.0 Flight+MO gate only on available equivalent +coverage. Native MO is the complete correctness oracle. Direct TAE and Docker +image/base changes remain outside this migration milestone. + +## Validation map + +| Closure | Risk | Required evidence | +| --- | --- | --- | +| Type/literal and binding descriptors | R2 | Independent byte/domain controls, malformed identity/width rejection, binding and clone/round-trip fidelity | +| GPU scalar and conditional owners | R3 | Actual GPU values, masks, scalar/cast error kinds, empty inputs, inactive failures, same-stream cleanup | +| Existing consumers and exposure | R2 | Ordinary type/function controls, native ownership suites, unchanged C ABI/capability | +| Build, schema and documentation | R0/R1 | Native build, registered tests and pinned static hooks | + +Full public SQL acceptance is the MO lowering/campaign boundary. This increment +does not claim all-22 execution, resource rollout acceptance or migration completion. + +The standalone decimal benchmark retains the primitive measurements and adds +ordinary/exact expression add controls on the same values and resource. Report +their raw repetitions separately: these controls measure evaluator cost, not +full-domain error equivalence or complete query performance. + +## Local evidence + +Validated on 2026-10-08 using the frozen `mo` Pixi environment: CUDA 13.3.73, +GCC 14.4.0, cuDF/RMM 26.08, an RTX 3070 and driver 615.71.09. + +| Check | Result | +| --- | --- | +| Native binding, logical types and ordinary function controls | 727 assertions / 75 cases | +| CPU exact-decimal foundation | 52,330 assertions / 5 cases | +| Native ownership and protocol controls | 594 assertions / 77 cases | +| Imported expressions, decimal kernels and native GPU/result controls | 31,459 assertions / 18 cases | +| Compute Sanitizer memcheck, imported expressions | 1,664 assertions / 7 cases; zero sanitizer errors | +| SDK export tests | 8 tests | +| Independent C ABI consumer and SDK schema fingerprint | Passed | + +All listed tests passed. The native targets built and pinned pre-commit hooks +passed. Memcheck covers invalid memory access; it does not establish the later +resource, concurrency or leak acceptance campaign. + +Reproduce the GPU and sanitizer groups from the Sirius source directory: + +```bash +pixi run --frozen -e mo build/mo/extension/sirius/sirius_native_gpu_unittest \ + '[decimal_import_gpu],[exact_decimal_gpu],[native_gpu],[native_result_gpu]' +pixi run --frozen -e mo compute-sanitizer --tool memcheck --error-exitcode 99 \ + build/mo/extension/sirius/sirius_native_gpu_unittest '[decimal_import_gpu]' +pixi run --frozen -e mo build/mo/extension/sirius/sirius_exact_decimal_benchmark +``` + +The benchmark uses two warmups; the +[raw scalar repetitions](mo-exact-decimal-import-kernels.csv) record seven +measured repetitions at 262,144 rows. Against the +recorded #25 primitive run, checked-operation median wall time changed by ++1.49% (add64), +0.77% (widened multiply), -0.05% (add256), and +3.83% +(divide256). This is a historical local comparison, not a concurrent SQL +acceptance run. The new ordinary/exact expression add controls measured +0.199190/0.158990 ms median on identical values and the same bounded pool. diff --git a/proto/matrixone/sirius/numeric/v1/exact_decimal.proto b/proto/matrixone/sirius/numeric/v1/exact_decimal.proto new file mode 100644 index 0000000000..ceb1e8fba0 --- /dev/null +++ b/proto/matrixone/sirius/numeric/v1/exact_decimal.proto @@ -0,0 +1,10 @@ +// Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 +syntax = "proto3"; +package matrixone.sirius.numeric.v1; +option go_package = "github.com/matrixorigin/matrixone/pkg/pb/siriusnumeric"; + +// Type parameters carry physical width, precision and scale; this payload is +// exactly the signed little-endian two's-complement coefficient for that width. +message ExactDecimalLiteral { + bytes coefficient_le = 1; +} diff --git a/scripts/export_embed_link.py b/scripts/export_embed_link.py index bcaba9600e..af3951b788 100644 --- a/scripts/export_embed_link.py +++ b/scripts/export_embed_link.py @@ -251,6 +251,15 @@ def main(): manifest["artifact_sha256"][str((args.output / "sirius_c.h").resolve())] = sha256( args.output / "sirius_c.h" ) + # The versioned literal schema is owned by Sirius, so downstream language + # bindings generate from the same pinned SDK contract instead of duplicating it. + relative_schema = Path("proto/matrixone/sirius/numeric/v1/exact_decimal.proto") + schema = Path(manifest["source_directory"]) / relative_schema + if schema.is_file(): + destination = args.output / relative_schema + destination.parent.mkdir(parents=True, exist_ok=True) + shutil.copyfile(schema, destination) + manifest["artifact_sha256"][str(destination.resolve())] = sha256(destination) (args.output / "link.json").write_text(json.dumps(manifest, indent=2) + "\n") # Reusing a build tree must not leave the retired GPU-provider manifest # beside a newly exported SDK that no longer records or validates it. diff --git a/scripts/tests/test_export_embed_link.py b/scripts/tests/test_export_embed_link.py index 2e6b7ff6c0..2ba74f5ce3 100644 --- a/scripts/tests/test_export_embed_link.py +++ b/scripts/tests/test_export_embed_link.py @@ -2,6 +2,7 @@ # SPDX-License-Identifier: Apache-2.0 import importlib.util +import hashlib import json from pathlib import Path import subprocess @@ -18,9 +19,16 @@ class ExportTest(unittest.TestCase): - def test_export_retires_stale_gpu_manifest(self): + def test_export_records_literal_schema_and_retires_stale_gpu_manifest(self): with tempfile.TemporaryDirectory() as directory: root = Path(directory) + relative_schema = Path( + "proto/matrixone/sirius/numeric/v1/exact_decimal.proto" + ) + schema = root / relative_schema + schema.parent.mkdir(parents=True) + payload = b'syntax = "proto3"; message ExactDecimalLiteral { bytes coefficient_le = 1; }\n' + schema.write_bytes(payload) output = root / "embedding-sdk" output.mkdir() stale = output / "toolchain.json" @@ -51,12 +59,21 @@ def test_export_retires_stale_gpu_manifest(self): ), patch.object( sdk, "c_compiler", return_value="/usr/bin/cc" ), patch.object( - sdk, "source_provenance", return_value={"source_revision": "test"} + sdk, + "source_provenance", + return_value={"source_revision": "test", "source_directory": str(root)}, ), patch.object( sdk, "artifact_hashes", return_value={} ): sdk.main() self.assertFalse(stale.exists()) + destination = output / relative_schema + self.assertEqual(destination.read_bytes(), payload) + manifest = json.loads((output / "link.json").read_text()) + self.assertEqual( + manifest["artifact_sha256"][str(destination.resolve())], + hashlib.sha256(payload).hexdigest(), + ) self.assertNotIn( "gpu_toolchain_manifest", json.loads((output / "link.json").read_text()) ) diff --git a/src/cudf/cudf_utils.hpp b/src/cudf/cudf_utils.hpp index b17a71dba5..ca335ea9de 100644 --- a/src/cudf/cudf_utils.hpp +++ b/src/cudf/cudf_utils.hpp @@ -20,6 +20,7 @@ #define CUDF_VERSION_NUM (CUDF_VERSION_MAJOR * 100 + CUDF_VERSION_MINOR) #include "helper/logical_type.hpp" +#include "numeric/exact_decimal_gpu.hpp" #include "sirius/exception.hpp" #include @@ -172,6 +173,13 @@ inline std::size_t estimate_referenced_column_bytes( */ inline cudf::data_type get_cudf_type(const logical_type& t) { + if (t.is_mo_decimal()) { + auto type = t.mo_decimal_type(); + return type.bits == 256 ? cudf::data_type{cudf::type_id::STRUCT} + : cudf::data_type{type.bits == 64 ? cudf::type_id::DECIMAL64 + : cudf::type_id::DECIMAL128, + -static_cast(type.scale)}; + } switch (t.id()) { case type_id::TINYINT: return cudf::data_type(cudf::type_id::INT8); case type_id::SMALLINT: return cudf::data_type(cudf::type_id::INT16); @@ -341,11 +349,19 @@ template * types). Used wherever an empty but schema-bearing table is needed — e.g. an all-pruned GPU-values * source, or the synthesized missing side of a join against an empty table. */ -inline std::unique_ptr make_empty_table(const duckdb::vector& types) +inline std::unique_ptr make_empty_table( + const duckdb::vector& types, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) { std::vector> columns; columns.reserve(types.size()); for (auto const& t : types) { + if (t.is_mo_decimal()) { + columns.push_back(mo_decimal::make_decimal_column( + t.mo_decimal_type(), 0, cudf::mask_state::UNALLOCATED, stream, mr)); + continue; + } columns.push_back(t.is_array() ? cudf::make_empty_lists_column(get_cudf_type(t.array_child())) : cudf::make_empty_column(get_cudf_type(t))); } diff --git a/src/expression/from_duckdb.cpp b/src/expression/from_duckdb.cpp index 05baa9b23b..1ee8bc53e5 100644 --- a/src/expression/from_duckdb.cpp +++ b/src/expression/from_duckdb.cpp @@ -16,6 +16,8 @@ #include "expression/ast/from_duckdb.hpp" +#include "numeric/decimal_functions.hpp" + // sirius #include "expression/aggregate_id.hpp" #include "expression/ast/aggregate.hpp" @@ -189,6 +191,7 @@ std::unique_ptr translate_function(duckdb::BoundFunctionExpression const& { auto func_id_opt = sirius::from_duckdb_function_name(expr.function.name); if (!func_id_opt.has_value()) { return nullptr; } + if (mo_decimal::is_decimal_function(*func_id_opt)) mo_decimal::validate_bound_scalar(expr); auto arguments = translate_children(expr.children); if (!arguments) { return nullptr; } auto return_type = sirius::from_duckdb(expr.return_type); diff --git a/src/expression/function_id.cpp b/src/expression/function_id.cpp index f764359788..5e621a9649 100644 --- a/src/expression/function_id.cpp +++ b/src/expression/function_id.cpp @@ -30,7 +30,7 @@ namespace { // Forward table: DuckDB function name -> Sirius function id. // Symbolic SQL operators and their Substrait spellings resolve to the same ids. // Linear scan; called once per BoundFunctionExpression at executor entry. -constexpr std::array, 35> kForwardTable = {{ +constexpr std::array, 49> kForwardTable = {{ {"+", function_id::add}, {"add", function_id::add}, {"-", function_id::sub}, @@ -66,24 +66,75 @@ constexpr std::array, 35> kForwardTable {"row", function_id::row}, {"struct_pack", function_id::struct_pack}, {"error", function_id::error}, + {"__sirius_mo_decimal_add", function_id::mo_decimal_add}, + {"__sirius_mo_decimal_subtract", function_id::mo_decimal_subtract}, + {"__sirius_mo_decimal_multiply", function_id::mo_decimal_multiply}, + {"__sirius_mo_decimal_divide", function_id::mo_decimal_divide}, + {"__sirius_mo_decimal_integer_divide", function_id::mo_decimal_integer_divide}, + {"__sirius_mo_decimal_modulo", function_id::mo_decimal_modulo}, + {"__sirius_mo_decimal_negate", function_id::mo_decimal_negate}, + {"__sirius_mo_decimal_cast", function_id::mo_decimal_cast}, + {"__sirius_mo_decimal_equal", function_id::mo_decimal_equal}, + {"__sirius_mo_decimal_not_equal", function_id::mo_decimal_not_equal}, + {"__sirius_mo_decimal_less", function_id::mo_decimal_less}, + {"__sirius_mo_decimal_less_equal", function_id::mo_decimal_less_equal}, + {"__sirius_mo_decimal_greater", function_id::mo_decimal_greater}, + {"__sirius_mo_decimal_greater_equal", function_id::mo_decimal_greater_equal}, }}; // Reverse table: Sirius function id -> canonical DuckDB function name. // Indexed directly by enum value; never searched. -constexpr std::array kReverseTable = { - "+", "-", "*", "/", "//", - "%", "substring", "~~", "!~~", "contains", - "prefix", "suffix", "strlen", "length", "regexp_replace", - "concat", "||", "year", "month", "day", - "hour", "minute", "second", "millisecond", "microsecond", - "date_trunc", "row", "struct_pack", "error", +constexpr std::array kReverseTable = { + "+", + "-", + "*", + "/", + "//", + "%", + "substring", + "~~", + "!~~", + "contains", + "prefix", + "suffix", + "strlen", + "length", + "regexp_replace", + "concat", + "||", + "year", + "month", + "day", + "hour", + "minute", + "second", + "millisecond", + "microsecond", + "date_trunc", + "row", + "struct_pack", + "error", + "__sirius_mo_decimal_add", + "__sirius_mo_decimal_subtract", + "__sirius_mo_decimal_multiply", + "__sirius_mo_decimal_divide", + "__sirius_mo_decimal_integer_divide", + "__sirius_mo_decimal_modulo", + "__sirius_mo_decimal_negate", + "__sirius_mo_decimal_cast", + "__sirius_mo_decimal_equal", + "__sirius_mo_decimal_not_equal", + "__sirius_mo_decimal_less", + "__sirius_mo_decimal_less_equal", + "__sirius_mo_decimal_greater", + "__sirius_mo_decimal_greater_equal", }; static_assert(static_cast(function_id::error) + 1 == 29, - "function_id::error must be the last entry; cardinality locked at 29."); -static_assert(kReverseTable.size() == 29, + "Ordinary function IDs must retain their original ABI values."); +static_assert(kReverseTable.size() == 43, "kReverseTable must have one slot per function_id value."); -static_assert(kForwardTable.size() == 35, +static_assert(kForwardTable.size() == 49, "kForwardTable includes SQL and Substrait aliases for supported function ids."); // Walks both tables to ensure every enum value has exactly one canonical diff --git a/src/expression/function_id.hpp b/src/expression/function_id.hpp index 5287edfe03..8f3af63cf2 100644 --- a/src/expression/function_id.hpp +++ b/src/expression/function_id.hpp @@ -31,7 +31,7 @@ namespace sirius { * test_function_id.cpp). The order is grouped by category for human * readability; integer values are part of the public ABI — new entries go * at the end of their category, never in the middle. Cardinality is - * exactly 29 (D-01). + * 29 ordinary functions plus the appended MO exact-decimal family. */ enum class function_id : uint16_t { // Arithmetic — 6 entries (also the contents of supported_ast_functions) @@ -70,6 +70,22 @@ enum class function_id : uint16_t { row, struct_pack, error, + + // Append only: embedded MO markers never enter ordinary function inference. + mo_decimal_add, + mo_decimal_subtract, + mo_decimal_multiply, + mo_decimal_divide, + mo_decimal_integer_divide, + mo_decimal_modulo, + mo_decimal_negate, + mo_decimal_cast, + mo_decimal_equal, + mo_decimal_not_equal, + mo_decimal_less, + mo_decimal_less_equal, + mo_decimal_greater, + mo_decimal_greater_equal, }; /** diff --git a/src/expression/to_duckdb.cpp b/src/expression/to_duckdb.cpp index 9ecf3850fa..f25ebb7bb2 100644 --- a/src/expression/to_duckdb.cpp +++ b/src/expression/to_duckdb.cpp @@ -16,6 +16,8 @@ #include "expression/ast/to_duckdb.hpp" +#include "numeric/decimal_functions.hpp" + // sirius #include "expression/ast/aggregate.hpp" #include "expression/ast/between.hpp" @@ -248,6 +250,10 @@ std::unique_ptr to_duckdb(function_call const& alt) children.push_back(to_duck_ptr(to_duckdb(*arg))); } + if (mo_decimal::is_decimal_function(alt.function())) + return from_duck_ptr(mo_decimal::bound_scalar( + mo_decimal::operation(alt.function()), std::move(children), alt.return_type())); + return from_duck_derived_ptr( duckdb::make_uniq(std::move(return_type), std::move(stub_fn), diff --git a/src/expression/value.cpp b/src/expression/value.cpp index f636e9ef37..907ef00040 100644 --- a/src/expression/value.cpp +++ b/src/expression/value.cpp @@ -17,6 +17,7 @@ #include "expression/value.hpp" #include "helper/type_conversions.hpp" +#include "numeric/decimal_types.hpp" #include "sirius/exception.hpp" // duckdb @@ -38,8 +39,13 @@ value from_duckdb(const duckdb::Value& v, const logical_type& type) // Typed NULL fidelity — every IsNull() input becomes null_value, regardless // of the value's logical type. The SQL type of the NULL is recovered later // via the supplied logical_type at the to_duckdb boundary. - if (v.IsNull()) { return value{null_value{}}; } + if (v.IsNull()) { + if (type.is_mo_decimal() && type.nullability() == 1) + throw std::invalid_argument("NULL MO decimal literal has a required type"); + return value{null_value{}}; + } + if (type.is_mo_decimal()) return value{mo_decimal::from_duckdb_value(v, type)}; switch (type.id()) { case type_id::BOOLEAN: return value{v.GetValue()}; case type_id::TINYINT: return value{v.GetValue()}; @@ -101,8 +107,14 @@ duckdb::Value to_duckdb(const value& v, const logical_type& type) // Typed NULL recovery — null_value plus the supplied type yields a // typed-null duckdb::Value via the `Value(LogicalType)` ctor. Reuses the // sirius::to_duckdb(logical_type) overload from helper/type_conversions.hpp. - if (std::holds_alternative(v)) { return duckdb::Value(to_duckdb(type)); } + if (std::holds_alternative(v)) { + if (type.is_mo_decimal() && type.nullability() == 1) + throw std::invalid_argument("NULL MO decimal literal has a required type"); + return duckdb::Value(to_duckdb(type)); + } + if (type.is_mo_decimal()) + return mo_decimal::duckdb_value(std::get(v), type); switch (type.id()) { case type_id::BOOLEAN: return duckdb::Value::BOOLEAN(std::get(v)); case type_id::TINYINT: return duckdb::Value::TINYINT(std::get(v)); diff --git a/src/expression/value.hpp b/src/expression/value.hpp index 9af7f0dac0..5b0d720848 100644 --- a/src/expression/value.hpp +++ b/src/expression/value.hpp @@ -149,7 +149,8 @@ using value = std::variant; + std::string, + mo_decimal::coefficient>; /** * @brief Convert a duckdb::Value to a sirius::value. diff --git a/src/expression_evaluator/CMakeLists.txt b/src/expression_evaluator/CMakeLists.txt index 4cb7116398..094aef6e1c 100644 --- a/src/expression_evaluator/CMakeLists.txt +++ b/src/expression_evaluator/CMakeLists.txt @@ -12,6 +12,8 @@ target_sources( specializations/conjunction.cpp specializations/constant.cpp specializations/function.cpp + specializations/mo_decimal.cpp + specializations/mo_conditional.cpp specializations/narrow_domain.cpp specializations/operator.cpp specializations/reference.cpp) diff --git a/src/expression_evaluator/expression_evaluator.cpp b/src/expression_evaluator/expression_evaluator.cpp index ad6fa0591d..9dc1183d40 100644 --- a/src/expression_evaluator/expression_evaluator.cpp +++ b/src/expression_evaluator/expression_evaluator.cpp @@ -15,7 +15,10 @@ */ // sirius +#include "pipeline/gpu_stream_quiescence_error.hpp" + #include +#include #include // sirius::ast::aggregate #include // sirius::ast::node alternatives @@ -266,86 +269,108 @@ void expression_evaluator::release_temporaries( std::unique_ptr expression_evaluator::evaluate(cudf::table_view input) { - _output_columns.clear(); - _output_columns.reserve(_ast_expressions.size()); - - // Reset AST state from any previous invocation so that the tree, temp scalars, temp columns, - // the restored-reference cache over them, and the per-call instrumentation do not carry stale - // state across calls. - _restored_reference_cache.clear(); - _restored_reference_cast_count = 0; - _narrow_domain_comparison_count = 0; - _ast_tree = cudf::ast::tree{}; - _temp_scalars.clear(); - _temp_columns.clear(); - - // Get the table_view from the input_batch - _input_table = std::move(input); - - // Per-result column post-processing. The node's kind and result type are read - // natively from the Sirius AST (no DuckDB round-trip). - auto post_process = [this](sirius::ast::node const& expr, evaluate_result result) { - if (expr.holds()) { - // A narrowed reference restores as a view of a column owned by the per-evaluation restoration - // cache; an unchanged reference stays a zero-copy view of the input. Copying the view (or - // releasing an owned column) gives the output table independent columns. - if (result.is_owned_column()) { - _output_columns.push_back(result.release_column()); - } else { - _output_columns.push_back( - std::make_unique(result.get_column_view(), _stream, _mr)); - } - } else { - // Cast the `result` from libcudf to the node's return type if `result` has a different - // type. E.g., `extract(year from col)` from libcudf returns int16_t but the SQL type is - // int64_t. Only use cudf::cast when both types are fixed-width (cast does not support - // STRING/LIST/STRUCT). - auto const cudf_return_type = sirius::get_cudf_type(expr.return_type()); - std::unique_ptr result_column; - if (result.is_scalar()) { - result_column = - cudf::make_column_from_scalar(result.get_scalar(), _input_table.num_rows(), _stream, _mr); - } else if (result.is_column_view()) { - result_column = std::make_unique(result.get_column_view(), _stream, _mr); + struct unsafe_temporaries { + std::vector> outputs, columns; + std::vector> scalars; + }; + auto quarantine = uses_mo_expressions() ? std::make_unique() : nullptr; + try { + _output_columns.clear(); + _output_columns.reserve(_ast_expressions.size()); + + // Reset AST state from any previous invocation so that the tree, temp scalars, temp columns, + // the restored-reference cache over them, and the per-call instrumentation do not carry stale + // state across calls. + _restored_reference_cache.clear(); + _restored_reference_cast_count = 0; + _narrow_domain_comparison_count = 0; + _ast_tree = cudf::ast::tree{}; + _temp_scalars.clear(); + _temp_columns.clear(); + + // Get the table_view from the input_batch + _input_table = std::move(input); + + // Per-result column post-processing. The node's kind and result type are read + // natively from the Sirius AST (no DuckDB round-trip). + auto post_process = [this](sirius::ast::node const& expr, evaluate_result result) { + if (expr.holds()) { + // A narrowed reference restores as a view of a column owned by the per-evaluation + // restoration cache; an unchanged reference stays a zero-copy view of the input. Copying + // the view (or releasing an owned column) gives the output table independent columns. + if (result.is_owned_column()) { + _output_columns.push_back(result.release_column()); + } else { + _output_columns.push_back( + std::make_unique(result.get_column_view(), _stream, _mr)); + } } else { - result_column = result.release_column(); - } - if (result_column->type() != cudf_return_type) { - // Cast is only valid for fixed-width types (no STRING/LIST/STRUCT/etc.). - if (IsFixedWidth(result_column->type()) && IsFixedWidth(cudf_return_type)) { - // Final result-type reconciliation is a physical schema restore. The declared output - // type provides the provenance needed to restore a narrowed DATE representation. - result_column = - sirius::cast_through_rep(result_column->view(), cudf_return_type, _stream, _mr); + // Cast the `result` from libcudf to the node's return type if `result` has a different + // type. E.g., `extract(year from col)` from libcudf returns int16_t but the SQL type is + // int64_t. Only use cudf::cast when both types are fixed-width (cast does not support + // STRING/LIST/STRUCT). + auto const cudf_return_type = sirius::get_cudf_type(expr.return_type()); + std::unique_ptr result_column; + if (result.is_scalar()) { + result_column = cudf::make_column_from_scalar( + result.get_scalar(), _input_table.num_rows(), _stream, _mr); + } else if (result.is_column_view()) { + result_column = std::make_unique(result.get_column_view(), _stream, _mr); } else { - throw internal_exception("[expression_evaluator] Unsupported type conversion: {} to {}", - cudf::type_to_name(result_column->type()), - cudf::type_to_name(cudf_return_type)); + result_column = result.release_column(); + } + if (expr.return_type().is_mo_decimal() && + !mo_decimal::decimal_column_matches(result_column->view(), + expr.return_type().mo_decimal_type())) + throw std::invalid_argument( + "MO exact-decimal result carrier does not match its descriptor"); + if (result_column->type() != cudf_return_type) { + // Cast is only valid for fixed-width types (no STRING/LIST/STRUCT/etc.). + if (IsFixedWidth(result_column->type()) && IsFixedWidth(cudf_return_type)) { + // Final result-type reconciliation is a physical schema restore. The declared output + // type provides the provenance needed to restore a narrowed DATE representation. + result_column = + sirius::cast_through_rep(result_column->view(), cudf_return_type, _stream, _mr); + } else { + throw internal_exception("[expression_evaluator] Unsupported type conversion: {} to {}", + cudf::type_to_name(result_column->type()), + cudf::type_to_name(cudf_return_type)); + } } + _output_columns.push_back(std::move(result_column)); } - _output_columns.push_back(std::move(result_column)); + }; + + // Iterate _ast_expressions and route through the std::visit dispatcher, then + // post-process each result using the node's native kind + return type. + for (auto const* ast_expr : _ast_expressions) { + if (!ast_expr) { + // This is a Sirius bug, not a user error: a null slot means from_duckdb + // declined an expression as unsupported, but the planner still routed it + // to the GPU executor. This hard-fails the query instead of falling back to CPU. + // TODO fix to ensure the planner never builds a GPU projection containing unsupported + // expressions. + throw internal_exception( + "[expression_evaluator] null expression in select list — " + "from_duckdb returned nullptr for an unsupported expression; " + "cannot evaluate on GPU"); + } + auto result = evaluate(*ast_expr, evaluation_mode::MATERIALIZE); + post_process(*ast_expr, std::move(result)); } - }; - // Iterate _ast_expressions and route through the std::visit dispatcher, then - // post-process each result using the node's native kind + return type. - for (auto const* ast_expr : _ast_expressions) { - if (!ast_expr) { - // This is a Sirius bug, not a user error: a null slot means from_duckdb - // declined an expression as unsupported, but the planner still routed it - // to the GPU executor. This hard-fails the query instead of falling back to CPU. - // TODO fix to ensure the planner never builds a GPU projection containing unsupported - // expressions. - throw internal_exception( - "[expression_evaluator] null expression in select list — " - "from_duckdb returned nullptr for an unsupported expression; " - "cannot evaluate on GPU"); + return std::make_unique(std::move(_output_columns)); + } catch (...) { + if (quarantine && cudaStreamSynchronize(_stream.get()) != cudaSuccess) { + quarantine->outputs = std::move(_output_columns); + quarantine->columns = std::move(_temp_columns); + quarantine->scalars = std::move(_temp_scalars); + (void)quarantine.release(); + throw pipeline::gpu_stream_quiescence_error( + "MO expression temporaries could not prove quiescence"); } - auto result = evaluate(*ast_expr, evaluation_mode::MATERIALIZE); - post_process(*ast_expr, std::move(result)); + throw; } - - return std::make_unique(std::move(_output_columns)); } std::unique_ptr expression_evaluator::compute_mask(cudf::table_view input) diff --git a/src/expression_evaluator/expression_evaluator.hpp b/src/expression_evaluator/expression_evaluator.hpp index 71ca4fadac..6b93b86825 100644 --- a/src/expression_evaluator/expression_evaluator.hpp +++ b/src/expression_evaluator/expression_evaluator.hpp @@ -550,6 +550,13 @@ class expression_evaluator { evaluate_result evaluate(sirius::ast::comparison const& expr, evaluation_mode mode); evaluate_result evaluate(sirius::ast::conjunction const& expr, evaluation_mode mode); evaluate_result evaluate(sirius::ast::function_call const& expr, evaluation_mode mode); + evaluate_result evaluate_decimal_function(sirius::ast::function_call const& expr, + evaluation_mode mode); + bool const* _mo_active_rows{nullptr}; + std::optional _has_mo_expressions; + evaluate_result evaluate_mo_case(ast::case_expr const& expression, evaluation_mode mode); + evaluate_result evaluate_mo_coalesce(ast::coalesce const& expression, evaluation_mode mode); + bool uses_mo_expressions(); evaluate_result evaluate(sirius::ast::unary_op const& expr, evaluation_mode mode); evaluate_result evaluate(sirius::ast::coalesce const& expr, evaluation_mode mode); evaluate_result evaluate(sirius::ast::in_list const& expr, evaluation_mode mode); diff --git a/src/expression_evaluator/specializations/case.cpp b/src/expression_evaluator/specializations/case.cpp index 66c3361629..01f5196744 100644 --- a/src/expression_evaluator/specializations/case.cpp +++ b/src/expression_evaluator/specializations/case.cpp @@ -15,6 +15,8 @@ */ // sirius +#include "numeric/decimal_functions.hpp" + #include #include #include @@ -27,6 +29,7 @@ #include // standard library +#include #include namespace sirius { @@ -35,6 +38,15 @@ using evaluate_result = expression_evaluator::evaluate_result; evaluate_result expression_evaluator::evaluate(sirius::ast::case_expr const& alt, evaluation_mode mode) { + if (_mo_active_rows || + (uses_mo_expressions() && + (alt.return_type().is_mo_decimal() || mo_decimal::contains_exact_expression(*alt.else_) || + std::any_of(alt.cases.begin(), alt.cases.end(), [](auto const& entry) { + return mo_decimal::contains_exact_expression(*entry.when_) || + mo_decimal::contains_exact_expression(*entry.then_); + })))) + return evaluate_mo_case(alt, mode); + //===----------MATERIALIZE (AST breaker)----------===// // CASE cannot be represented as a cudf AST operation, so we always materialize it. If the caller // requested AST mode, we materialize the result and wrap it as a temporary column that the diff --git a/src/expression_evaluator/specializations/constant.cpp b/src/expression_evaluator/specializations/constant.cpp index f34c40b757..fccd6dc3a5 100644 --- a/src/expression_evaluator/specializations/constant.cpp +++ b/src/expression_evaluator/specializations/constant.cpp @@ -15,6 +15,8 @@ */ // sirius +#include "numeric/exact_decimal_gpu.hpp" + #include #include #include @@ -51,6 +53,14 @@ using evaluate_result = expression_evaluator::evaluate_result; evaluate_result expression_evaluator::evaluate(sirius::ast::constant const& alt, evaluation_mode mode) { + if (alt.return_type().is_mo_decimal()) { + bool valid = !std::holds_alternative(alt.payload); + auto value = valid ? std::get(alt.payload) : mo_decimal::coefficient{}; + auto column = mo_decimal::make_decimal_literal( + alt.return_type().mo_decimal_type(), value, valid, _input_table.num_rows(), _stream, _mr); + return mode == evaluation_mode::AST ? materialize_as_ast_column(std::move(column)) + : evaluate_result(std::move(column)); + } auto const cudf_type = sirius::get_cudf_type(alt.return_type()); bool const is_valid = !std::holds_alternative(alt.payload); diff --git a/src/expression_evaluator/specializations/function.cpp b/src/expression_evaluator/specializations/function.cpp index 59e551ccb0..8ab28dcfff 100644 --- a/src/expression_evaluator/specializations/function.cpp +++ b/src/expression_evaluator/specializations/function.cpp @@ -15,6 +15,8 @@ */ // sirius +#include "numeric/decimal_functions.hpp" + #include #include #include @@ -72,6 +74,7 @@ like_multiliteral_cache::entry_ptr const& expression_evaluator::get_or_classify_ evaluate_result expression_evaluator::evaluate(sirius::ast::function_call const& alt, evaluation_mode mode) { + if (mo_decimal::is_decimal_function(alt.function())) return evaluate_decimal_function(alt, mode); auto const resolved_id = alt.function(); auto const& args = alt.arguments(); diff --git a/src/expression_evaluator/specializations/mo_conditional.cpp b/src/expression_evaluator/specializations/mo_conditional.cpp new file mode 100644 index 0000000000..80588b7644 --- /dev/null +++ b/src/expression_evaluator/specializations/mo_conditional.cpp @@ -0,0 +1,260 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "expression_evaluator/expression_evaluator.hpp" +#include "numeric/decimal_functions.hpp" +#include "numeric/exact_decimal_gpu.hpp" +#include "pipeline/gpu_stream_quiescence_error.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace sirius { +namespace { +using result = expression_evaluator::evaluate_result; +struct active_scope { + bool const*& target; + bool const* previous; + active_scope(bool const*& t, cudf::column_view mask) : target(t), previous(t) + { + target = mask.data(); + } + ~active_scope() { target = previous; } +}; +// At most one branch's temporaries are retained. On a failed synchronization, +// ownership transfers to process quarantine rather than unwinding live buffers. +struct conditional_owners { + std::unique_ptr output, remaining, condition, selected, inverse; + std::unique_ptr next_output, next_remaining, lifted; + std::unique_ptr null_scalar, truth, falsity, reduced; + std::optional predicate, value; + + void retire_branch(logical_type const& type) + { + if (next_output && type.is_mo_decimal() && type.mo_decimal_type().bits == 256) { + // cuDF selection superimposes parent validity on STRUCT children. Inputs + // were canonical, so these child masks contain no independent NULLs. + // The caller has synchronized the task stream before retiring them. + auto rows = next_output->size(), nulls = next_output->null_count(); + auto contents = next_output->release(); + for (auto& child : contents.children) { + auto child_type = child->type(); + auto count = child->size(); + auto data = child->release(); + child = + std::make_unique(child_type, + count, + data.data ? std::move(*data.data) : rmm::device_buffer{}, + rmm::device_buffer{}, + 0); + } + next_output = 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)); + } + if (next_output) output = std::move(next_output); + if (next_remaining) remaining = std::move(next_remaining); + predicate.reset(); + value.reset(); + lifted.reset(); + condition.reset(); + selected.reset(); + inverse.reset(); + reduced.reset(); + } +}; +void require_canonical(cudf::column_view value, logical_type const& type) +{ + if (type.is_mo_decimal() && !mo_decimal::decimal_column_matches(value, type.mo_decimal_type())) + throw std::invalid_argument("MO conditional input is not a canonical exact-decimal column"); +} +void initialize(conditional_owners& owner, + logical_type const& type, + cudf::size_type rows, + bool const* active, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + owner.truth = std::make_unique>(true, true, stream, mr); + owner.falsity = std::make_unique>(false, true, stream, mr); + if (type.is_mo_decimal()) + owner.output = + mo_decimal::make_decimal_literal(type.mo_decimal_type(), {}, false, rows, stream, mr); + else { + owner.null_scalar = cudf::make_default_constructed_scalar(get_cudf_type(type), stream, mr); + owner.null_scalar->set_valid_async(false, stream); + owner.output = cudf::make_column_from_scalar(*owner.null_scalar, rows, stream, mr); + } + if (active) + owner.remaining = std::make_unique( + cudf::column_view(cudf::data_type{cudf::type_id::BOOL8}, rows, active, nullptr, 0), + stream, + mr); + else + owner.remaining = cudf::make_column_from_scalar(*owner.truth, rows, stream, mr); +} +cudf::column_view column(result& value, + conditional_owners& owner, + cudf::size_type rows, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + if (!value.is_scalar()) return value.get_column_view(); + owner.lifted = cudf::make_column_from_scalar(value.get_scalar(), rows, stream, mr); + return owner.lifted->view(); +} +bool any(cudf::column_view mask, + conditional_owners& owner, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + if (!mask.size()) return false; + owner.reduced = cudf::reduce(mask, + *cudf::make_any_aggregation(), + cudf::data_type{cudf::type_id::BOOL8}, + stream, + mr); + return static_cast&>(*owner.reduced).value(stream); +} +void rethrow_after_quiescence(std::unique_ptr& owner, + rmm::cuda_stream_view stream) +{ + auto error = std::current_exception(); + if (cudaStreamSynchronize(stream.value()) != cudaSuccess) { + (void)owner.release(); + throw pipeline::gpu_stream_quiescence_error( + "MO conditional expression could not prove quiescence"); + } + std::rethrow_exception(error); +} +} // namespace + +bool expression_evaluator::uses_mo_expressions() +{ + if (!_has_mo_expressions) { + _has_mo_expressions = false; + for (auto expression : _ast_expressions) + if (expression && mo_decimal::contains_exact_expression(*expression)) { + _has_mo_expressions = true; + break; + } + } + return *_has_mo_expressions; +} +expression_evaluator::evaluate_result expression_evaluator::evaluate_mo_case( + ast::case_expr const& expression, evaluation_mode mode) +{ + if (expression.cases.empty() || !expression.else_) + throw std::invalid_argument("invalid MO numeric CASE expression"); + auto owner = std::make_unique(); + auto rows = _input_table.num_rows(); + try { + initialize(*owner, expression.return_type(), rows, _mo_active_rows, _stream, _mr); + for (auto const& branch : expression.cases) { + if (!branch.when_ || !branch.then_) + throw std::invalid_argument("missing MO numeric CASE branch"); + if (!any(owner->remaining->view(), *owner, _stream, _mr)) break; + { + active_scope active(_mo_active_rows, owner->remaining->view()); + owner->predicate.emplace(evaluate(*branch.when_, evaluation_mode::MATERIALIZE)); + } + auto predicate = column(*owner->predicate, *owner, rows, _stream, _mr); + owner->condition = cudf::replace_nulls(predicate, *owner->falsity, _stream, _mr); + owner->selected = cudf::binary_operation(owner->remaining->view(), + owner->condition->view(), + cudf::binary_operator::LOGICAL_AND, + cudf::data_type{cudf::type_id::BOOL8}, + _stream, + _mr); + // Complete any scalar lift before another branch can replace its owner. + CUDF_CUDA_TRY(cudaStreamSynchronize(_stream.get())); + owner->lifted.reset(); + if (any(owner->selected->view(), *owner, _stream, _mr)) { + { + active_scope active(_mo_active_rows, owner->selected->view()); + owner->value.emplace(evaluate(*branch.then_, evaluation_mode::MATERIALIZE)); + } + auto value = column(*owner->value, *owner, rows, _stream, _mr); + require_canonical(value, expression.return_type()); + owner->next_output = + cudf::copy_if_else(value, owner->output->view(), owner->selected->view(), _stream, _mr); + owner->inverse = + cudf::unary_operation(owner->condition->view(), cudf::unary_operator::NOT, _stream, _mr); + owner->next_remaining = cudf::binary_operation(owner->remaining->view(), + owner->inverse->view(), + cudf::binary_operator::LOGICAL_AND, + cudf::data_type{cudf::type_id::BOOL8}, + _stream, + _mr); + } + CUDF_CUDA_TRY(cudaStreamSynchronize(_stream.get())); + owner->retire_branch(expression.return_type()); + } + if (any(owner->remaining->view(), *owner, _stream, _mr)) { + { + active_scope active(_mo_active_rows, owner->remaining->view()); + owner->value.emplace(evaluate(*expression.else_, evaluation_mode::MATERIALIZE)); + } + auto value = column(*owner->value, *owner, rows, _stream, _mr); + require_canonical(value, expression.return_type()); + owner->next_output = + cudf::copy_if_else(value, owner->output->view(), owner->remaining->view(), _stream, _mr); + } + CUDF_CUDA_TRY(cudaStreamSynchronize(_stream.get())); + owner->retire_branch(expression.return_type()); + auto output = std::move(owner->output); + return mode == evaluation_mode::AST ? materialize_as_ast_column(std::move(output)) + : evaluate_result(std::move(output)); + } catch (...) { + rethrow_after_quiescence(owner, _stream); + throw; + } +} +expression_evaluator::evaluate_result expression_evaluator::evaluate_mo_coalesce( + ast::coalesce const& expression, evaluation_mode mode) +{ + if (expression.children.size() < 2) throw std::invalid_argument("invalid MO numeric COALESCE"); + auto owner = std::make_unique(); + auto rows = _input_table.num_rows(); + try { + initialize(*owner, expression.return_type(), rows, _mo_active_rows, _stream, _mr); + for (auto const& child : expression.children) { + if (!child) throw std::invalid_argument("missing MO numeric COALESCE argument"); + if (!any(owner->remaining->view(), *owner, _stream, _mr)) break; + { + active_scope active(_mo_active_rows, owner->remaining->view()); + owner->value.emplace(evaluate(*child, evaluation_mode::MATERIALIZE)); + } + auto value = column(*owner->value, *owner, rows, _stream, _mr); + require_canonical(value, expression.return_type()); + owner->next_output = + cudf::copy_if_else(value, owner->output->view(), owner->remaining->view(), _stream, _mr); + owner->condition = cudf::is_null(value, _stream, _mr); + owner->next_remaining = cudf::binary_operation(owner->remaining->view(), + owner->condition->view(), + cudf::binary_operator::LOGICAL_AND, + cudf::data_type{cudf::type_id::BOOL8}, + _stream, + _mr); + CUDF_CUDA_TRY(cudaStreamSynchronize(_stream.get())); + owner->retire_branch(expression.return_type()); + } + CUDF_CUDA_TRY(cudaStreamSynchronize(_stream.get())); + auto output = std::move(owner->output); + return mode == evaluation_mode::AST ? materialize_as_ast_column(std::move(output)) + : evaluate_result(std::move(output)); + } catch (...) { + rethrow_after_quiescence(owner, _stream); + throw; + } +} +} // namespace sirius diff --git a/src/expression_evaluator/specializations/mo_decimal.cpp b/src/expression_evaluator/specializations/mo_decimal.cpp new file mode 100644 index 0000000000..9944eb8d8f --- /dev/null +++ b/src/expression_evaluator/specializations/mo_decimal.cpp @@ -0,0 +1,93 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "expression_evaluator/expression_evaluator.hpp" +#include "numeric/decimal_error.hpp" +#include "numeric/decimal_functions.hpp" +#include "numeric/exact_decimal_gpu.hpp" +#include "pipeline/gpu_stream_quiescence_error.hpp" + +#include +#include +#include + +namespace sirius { +expression_evaluator::evaluate_result expression_evaluator::evaluate_decimal_function( + ast::function_call const& expression, evaluation_mode mode) +{ + using namespace mo_decimal; + auto op = operation(expression.function()); + auto const& arguments = expression.arguments(); + std::vector types; + for (auto const& argument : arguments) { + if (!argument) throw std::invalid_argument("missing MO decimal argument"); + types.push_back(argument->return_type()); + } + validate_signature(op, types, expression.return_type()); + struct owners { + std::vector arguments; + std::vector> lifted; + std::optional result; + }; + auto owner = std::make_unique(); + owner->arguments.reserve(arguments.size()); + owner->lifted.reserve(2 * arguments.size()); + try { + std::vector columns; + std::vector descriptors; + columns.reserve(arguments.size()); + descriptors.reserve(arguments.size()); + for (std::size_t i = 0; i < arguments.size(); ++i) { + owner->arguments.push_back(evaluate(*arguments[i], evaluation_mode::MATERIALIZE)); + auto& value = owner->arguments.back(); + if (value.is_scalar()) { + owner->lifted.push_back( + cudf::make_column_from_scalar(value.get_scalar(), _input_table.num_rows(), _stream, _mr)); + columns.push_back(owner->lifted.back()->view()); + } else + columns.push_back(value.get_column_view()); + if (types[i].is_mo_decimal()) + descriptors.push_back(types[i].mo_decimal_type()); + else { + if (columns.back().type().id() != cudf::type_id::INT64) { + owner->lifted.push_back( + cudf::cast(columns.back(), cudf::data_type{cudf::type_id::INT64}, _stream, _mr)); + columns.back() = owner->lifted.back()->view(); + } + auto column = columns.back(); + columns.back() = cudf::column_view(cudf::data_type{cudf::type_id::DECIMAL64, 0}, + column.size(), + column.head(), + column.null_mask(), + column.null_count(), + column.offset()); + descriptors.push_back({64, 18, 0}); + } + } + auto output = expression.return_type().is_mo_decimal() + ? expression.return_type().mo_decimal_type() + : decimal_type{64, 18, 0}; + auto right = columns.size() == 1 ? 0u : 1u; + owner->result.emplace(evaluate_decimal_columns(op, + columns[0], + descriptors[0], + columns[right], + descriptors[right], + output, + _mo_active_rows, + _stream, + _mr)); + auto error = column_error(*owner->result, _stream, _mr); + if (error != decimal_error::none) throw numeric_error(error); + auto result = std::move(owner->result->values); + return mode == evaluation_mode::AST ? materialize_as_ast_column(std::move(result)) + : evaluate_result(std::move(result)); + } catch (...) { + auto error = std::current_exception(); + if (cudaStreamSynchronize(_stream.get()) != cudaSuccess) { + (void)owner.release(); + throw pipeline::gpu_stream_quiescence_error( + "MO decimal expression could not prove quiescence"); + } + std::rethrow_exception(error); + } +} +} // namespace sirius diff --git a/src/expression_evaluator/specializations/operator.cpp b/src/expression_evaluator/specializations/operator.cpp index 4814b5f171..bd51a7780b 100644 --- a/src/expression_evaluator/specializations/operator.cpp +++ b/src/expression_evaluator/specializations/operator.cpp @@ -15,6 +15,8 @@ */ // sirius +#include "numeric/decimal_functions.hpp" + #include #include #include @@ -534,6 +536,14 @@ evaluate_result expression_evaluator::evaluate(sirius::ast::in_list const& alt, evaluate_result expression_evaluator::evaluate(sirius::ast::coalesce const& alt, evaluation_mode mode) { + if (_mo_active_rows || + (uses_mo_expressions() && + (alt.return_type().is_mo_decimal() || + std::any_of(alt.children.begin(), alt.children.end(), [](auto const& child) { + return mo_decimal::contains_exact_expression(*child); + })))) + return evaluate_mo_coalesce(alt, mode); + D_ASSERT(alt.children.size() > 1); auto current_result = evaluate(*alt.children[0], evaluation_mode::MATERIALIZE); diff --git a/src/expression_evaluator/specializations/reference.cpp b/src/expression_evaluator/specializations/reference.cpp index 09b98c4cd8..4b91e122c6 100644 --- a/src/expression_evaluator/specializations/reference.cpp +++ b/src/expression_evaluator/specializations/reference.cpp @@ -72,8 +72,14 @@ evaluate_result expression_evaluator::get_or_create_restored_reference(std::uint evaluate_result expression_evaluator::evaluate(sirius::ast::reference const& alt, evaluation_mode mode) { - auto const source = _input_table.column(alt.column_index); auto const& logical = alt.return_type(); + 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); + 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"); if (is_narrowable_numeric_type(logical)) { auto const native = get_cudf_type(logical); if (can_restore_to(source.type(), native)) { diff --git a/src/helper/logical_type.hpp b/src/helper/logical_type.hpp index b44b5e4840..f7ae609928 100644 --- a/src/helper/logical_type.hpp +++ b/src/helper/logical_type.hpp @@ -16,6 +16,8 @@ #pragma once +#include "numeric/decimal_type.hpp" + #include #include #include @@ -60,6 +62,7 @@ enum class type_id : uint8_t { LIST, ARRAY, // fixed-sized list type DECIMAL, + MO_DECIMAL, ///< Embedded MO carrier width is independent of declared precision. }; //===----------------------------------------------------------------------===// @@ -105,7 +108,12 @@ class logical_type { * @brief Construct a non-DECIMAL logical type. * @param id The SQL type identifier (must not be DECIMAL). */ - static logical_type make(type_id id) { return logical_type(id, 0, 0); } + static logical_type make(type_id id) + { + if (id == type_id::MO_DECIMAL) + throw std::invalid_argument("MO decimal types require an explicit physical descriptor"); + return logical_type(id, 0, 0); + } /** * @brief Construct a DECIMAL logical type with full precision and scale. @@ -117,6 +125,30 @@ class logical_type { return logical_type(type_id::DECIMAL, precision, scale); } + static logical_type make_mo_decimal(mo_decimal::decimal_type type, bool nullable) + { + if (!type.valid()) throw std::invalid_argument("invalid MO exact-decimal type"); + logical_type result(type_id::MO_DECIMAL, type.precision, type.scale); + result._decimal_bits = type.bits; + result._nullability = nullable ? 2 : 1; + return result; + } + + bool is_mo_decimal() const noexcept { return _id == type_id::MO_DECIMAL; } + mo_decimal::decimal_type mo_decimal_type() const + { + if (!is_mo_decimal()) throw std::invalid_argument("not an MO exact-decimal type"); + return {_decimal_bits, _precision, _scale}; + } + // Zero is unspecified for ordinary types; exact bindings preserve required/nullable. + uint8_t nullability() const noexcept { return _nullability; } + logical_type with_nullability(bool nullable) const + { + auto result = *this; + result._nullability = nullable ? 2 : 1; + return result; + } + /** * @brief Construct an ARRAY logical type with child element type and size. * @param child The type of elements in the array. @@ -168,7 +200,7 @@ class logical_type { bool is_numeric() const noexcept { return is_integer() || _id == type_id::FLOAT || _id == type_id::DOUBLE || - _id == type_id::DECIMAL; + _id == type_id::DECIMAL || _id == type_id::MO_DECIMAL; } /// Returns true if this is a date or timestamp type (any precision). @@ -256,6 +288,7 @@ class logical_type { case type_id::TIMESTAMP_MS: case type_id::TIMESTAMP: case type_id::TIMESTAMP_NS: return 8; + case type_id::MO_DECIMAL: return _decimal_bits / 8; case type_id::DECIMAL: if (_precision <= decimal_max_precision_int32) return 4; // DECIMAL32 if (_precision <= decimal_max_precision_int64) return 8; // DECIMAL64 @@ -302,6 +335,9 @@ class logical_type { std::string child = _child ? _child->to_string() : "?"; return _array_size == 0 ? child + "[ANY]" : child + "[" + std::to_string(_array_size) + "]"; } + case type_id::MO_DECIMAL: + return "MO_DECIMAL" + std::to_string(_decimal_bits) + "(" + std::to_string(_precision) + + "," + std::to_string(_scale) + ")"; case type_id::DECIMAL: return "DECIMAL(" + std::to_string(_precision) + "," + std::to_string(_scale) + ")"; default: return "UNKNOWN"; @@ -316,7 +352,8 @@ class logical_type { { // base field bool base_eq = _id == other._id && _precision == other._precision && _scale == other._scale && - _array_size == other._array_size; + _array_size == other._array_size && _decimal_bits == other._decimal_bits && + _nullability == other._nullability; if (!base_eq) return false; // recursively compare types for child @@ -336,6 +373,9 @@ class logical_type { uint8_t _precision{0}; ///< Meaningful only for DECIMAL: total significant digits (1–38) uint8_t _scale{0}; ///< Meaningful only for DECIMAL: fractional digits (0–precision) + uint16_t _decimal_bits{0}; + uint8_t _nullability{0}; + // array uint32_t _array_size{0}; ///< Meaningful only for ARRAY: fixed size (0 = any size) std::shared_ptr _child; ///< Child type for ARRAY/LIST (nullptr if not nested) diff --git a/src/helper/type_conversions.cpp b/src/helper/type_conversions.cpp index 34c6ba1442..d8153ebfd2 100644 --- a/src/helper/type_conversions.cpp +++ b/src/helper/type_conversions.cpp @@ -17,6 +17,7 @@ #include "helper/type_conversions.hpp" #include "duckdb/common/optional_idx.hpp" +#include "numeric/decimal_types.hpp" #include #include @@ -26,6 +27,7 @@ namespace sirius { logical_type from_duckdb(const duckdb::LogicalType& t) { + if (auto exact = mo_decimal::from_duckdb_type(t)) return *exact; using duckdb::LogicalTypeId; switch (t.id()) { case LogicalTypeId::SQLNULL: return logical_type::make(type_id::SQLNULL); @@ -68,6 +70,9 @@ logical_type from_duckdb(const duckdb::LogicalType& t) duckdb::LogicalType to_duckdb(const logical_type& t) { + if (t.is_mo_decimal() || + (t.nullability() && (t.id() == type_id::BOOLEAN || t.id() == type_id::BIGINT))) + return mo_decimal::duckdb_type(t); using duckdb::LogicalType; switch (t.id()) { case type_id::SQLNULL: return LogicalType::SQLNULL; diff --git a/src/numeric/CMakeLists.txt b/src/numeric/CMakeLists.txt index c88cba83f6..fdfd831856 100644 --- a/src/numeric/CMakeLists.txt +++ b/src/numeric/CMakeLists.txt @@ -1 +1,2 @@ -target_sources(sirius_objects PRIVATE exact_decimal_gpu.cu) +target_sources(sirius_objects PRIVATE exact_decimal_gpu.cu decimal_types.cpp + decimal_functions.cpp decimal_import.cpp) diff --git a/src/numeric/decimal_error.hpp b/src/numeric/decimal_error.hpp new file mode 100644 index 0000000000..1f96771c5e --- /dev/null +++ b/src/numeric/decimal_error.hpp @@ -0,0 +1,21 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include "numeric/decimal_type.hpp" + +#include +namespace sirius::mo_decimal { +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 arithmetic is out of range"), + code_(code) + { + } + decimal_error code() const noexcept { return code_; } + + private: + decimal_error code_; +}; +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_functions.cpp b/src/numeric/decimal_functions.cpp new file mode 100644 index 0000000000..8eb2425c9c --- /dev/null +++ b/src/numeric/decimal_functions.cpp @@ -0,0 +1,244 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "numeric/decimal_functions.hpp" + +#include "expression/ast/node.hpp" +#include "helper/type_conversions.hpp" +#include "numeric/exact_decimal.hpp" + +#include +#include +#include +#include +#include + +namespace sirius::mo_decimal { +bool is_decimal_function(function_id id) noexcept +{ + return id >= function_id::mo_decimal_add && id <= function_id::mo_decimal_greater_equal; +} +decimal_op operation(function_id id) +{ + if (!is_decimal_function(id)) throw std::invalid_argument("not an MO decimal operation"); + return static_cast(static_cast(id) - + static_cast(function_id::mo_decimal_add)); +} +function_id function(decimal_op op) +{ + if (op > decimal_op::greater_equal) throw std::invalid_argument("unknown MO decimal operation"); + return static_cast(static_cast(function_id::mo_decimal_add) + + static_cast(op)); +} +bool contains_exact_expression(ast::node const& expression) +{ + if (expression.return_type().is_mo_decimal()) return true; + return std::visit( + [](auto const& node) { + auto child = [](auto const& p) { return p && contains_exact_expression(*p); }; + if constexpr (requires { + node.function(); + node.arguments(); + }) { + if constexpr (std::is_same_v, ast::function_call>) + if (is_decimal_function(node.function())) return true; + for (auto const& argument : node.arguments()) + if (child(argument)) return true; + } else if constexpr (requires { + node.cases; + node.else_; + }) { + if (child(node.else_)) return true; + for (auto const& entry : node.cases) + if (child(entry.when_) || child(entry.then_)) return true; + } else if constexpr (requires { node.children; }) { + for (auto const& argument : node.children) + if (child(argument)) return true; + } else if constexpr (requires { node.child; }) + return child(node.child); + else if constexpr (requires { + node.left; + node.right; + }) + return child(node.left) || child(node.right); + else if constexpr (requires { + node.input; + node.lower; + node.upper; + }) + return child(node.input) || child(node.lower) || child(node.upper); + else if constexpr (requires { + node.probe; + node.values; + }) { + if (child(node.probe)) return true; + for (auto const& value : node.values) + if (child(value)) return true; + } + return false; + }, + expression.v); +} +namespace { +bool signed_integer_cast_fits(logical_type const& input, logical_type const& output) +{ + unsigned bits; + switch (input.id()) { + case type_id::TINYINT: bits = 8; break; + case type_id::SMALLINT: bits = 16; break; + case type_id::INTEGER: bits = 32; break; + case type_id::BIGINT: bits = 64; break; + default: return false; + } + uint64_t endpoints[]{~uint64_t(0) << (bits - 1), (uint64_t(1) << (bits - 1)) - 1}; + for (auto value : endpoints) { + auto coefficient = load_coefficient(reinterpret_cast(&value), 8); + if (evaluate_decimal(decimal_op::cast, + coefficient, + {64, 18, 0}, + true, + {}, + {64, 18, 0}, + true, + output.mo_decimal_type()) + .error != decimal_error::none) + return false; + } + return true; +} +struct scalar_data final : duckdb::FunctionData { + decimal_op op; + duckdb::vector inputs; + duckdb::LogicalType output; + scalar_data(decimal_op o, duckdb::vector i, duckdb::LogicalType r) + : op(o), inputs(std::move(i)), output(std::move(r)) + { + } + duckdb::unique_ptr Copy() const override + { + return duckdb::make_uniq(op, inputs, output); + } + bool Equals(duckdb::FunctionData const& other) const override + { + auto data = dynamic_cast(&other); + return data && op == data->op && inputs == data->inputs && output == data->output; + } +}; +void never_execute(duckdb::DataChunk&, duckdb::ExpressionState&, duckdb::Vector&) +{ + throw duckdb::NotImplementedException("MO exact-decimal marker cannot execute in DuckDB"); +} +duckdb::unique_ptr bind_scalar( + duckdb::ClientContext&, + duckdb::ScalarFunction& fn, + duckdb::vector>& arguments) +{ + auto id = from_duckdb_function_name(fn.name); + if (!id || !is_decimal_function(*id) || arguments.empty() || + arguments.back()->GetExpressionClass() != duckdb::ExpressionClass::BOUND_CONSTANT) + throw duckdb::InvalidInputException("invalid MO decimal binding marker"); + auto const& marker = arguments.back()->Cast().value; + auto output = from_duckdb_type(marker.type()); + if (!marker.IsNull() || !output) + throw duckdb::InvalidInputException("MO decimal binding needs its declared output type"); + std::vector inputs; + duckdb::vector carriers; + for (std::size_t i = 0; i + 1 < arguments.size(); ++i) { + inputs.push_back(sirius::from_duckdb(arguments[i]->return_type)); + carriers.push_back(arguments[i]->return_type); + } + validate_signature(operation(*id), inputs, *output); + fn.return_type = marker.type(); + arguments.pop_back(); + fn.arguments = carriers; + return duckdb::make_uniq(operation(*id), std::move(carriers), fn.return_type); +} +duckdb::ScalarFunction scalar_function(decimal_op op) +{ + auto arity = op == decimal_op::cast || op == decimal_op::negate ? 1u : 2u; + duckdb::vector arguments(arity + 1, duckdb::LogicalType::ANY); + return duckdb::ScalarFunction(std::string(to_duckdb_function_name(function(op))), + std::move(arguments), + duckdb::LogicalType::ANY, + never_execute, + bind_scalar, + nullptr, + nullptr, + nullptr, + duckdb::LogicalType::INVALID, + duckdb::FunctionStability::VOLATILE, + duckdb::FunctionNullHandling::SPECIAL_HANDLING); +} +} // namespace + +void validate_signature(decimal_op op, + std::vector const& inputs, + logical_type const& output) +{ + bool unary = op == decimal_op::cast || op == decimal_op::negate; + if (op > decimal_op::greater_equal || inputs.size() != (unary ? 1u : 2u)) + throw std::invalid_argument("invalid MO decimal operation arity"); + if (op >= decimal_op::equal || op == decimal_op::integer_divide) { + if (output.id() != (op >= decimal_op::equal ? type_id::BOOLEAN : type_id::BIGINT) || + !output.nullability()) + throw std::invalid_argument("invalid MO decimal comparison or DIV result"); + } else if (!output.is_mo_decimal()) + throw std::invalid_argument("MO decimal operation requires an exact result descriptor"); + if (op == decimal_op::cast && !inputs[0].is_mo_decimal()) { + if (!signed_integer_cast_fits(inputs[0], output)) + throw std::invalid_argument("MO decimal integer cast does not contain the full input domain"); + return; + } + for (auto const& input : inputs) + if (!input.is_mo_decimal()) + throw std::invalid_argument("MO decimal operands must retain exact type descriptors"); + if (op == decimal_op::negate) { + auto a = inputs[0].mo_decimal_type(), r = output.mo_decimal_type(); + if (r.bits < a.bits || r.precision < a.precision || r.scale != a.scale) + throw std::invalid_argument("MO decimal negation is not proven domain-preserving"); + } +} +void register_scalar_functions(duckdb::DatabaseInstance& instance) +{ + auto transaction = duckdb::CatalogTransaction::GetSystemTransaction(instance); + auto& catalog = duckdb::Catalog::GetSystemCatalog(instance); + for (unsigned i = 0; i <= static_cast(decimal_op::greater_equal); ++i) { + duckdb::CreateScalarFunctionInfo info(scalar_function(static_cast(i))); + catalog.CreateFunction(transaction, info); + } +} +void validate_bound_scalar(duckdb::BoundFunctionExpression const& expression) +{ + auto data = dynamic_cast(expression.bind_info.get()); + auto id = from_duckdb_function_name(expression.function.name); + if (!data || !id || !is_decimal_function(*id) || data->op != operation(*id) || + data->output != expression.return_type || data->inputs.size() != expression.children.size() || + expression.function.GetBindCallback() != bind_scalar) + throw std::invalid_argument("MO decimal bound function lost its checked signature"); + std::vector inputs; + for (std::size_t i = 0; i < data->inputs.size(); ++i) { + if (data->inputs[i] != expression.children[i]->return_type) + throw std::invalid_argument("MO decimal bound operand type mismatch"); + inputs.push_back(sirius::from_duckdb(data->inputs[i])); + } + validate_signature(data->op, inputs, sirius::from_duckdb(data->output)); +} +duckdb::unique_ptr bound_scalar( + decimal_op op, + duckdb::vector> children, + logical_type const& output) +{ + duckdb::vector carriers; + std::vector inputs; + for (auto const& child : children) { + carriers.push_back(child->return_type); + inputs.push_back(sirius::from_duckdb(child->return_type)); + } + validate_signature(op, inputs, output); + auto fn = scalar_function(op); + fn.arguments = carriers; + fn.return_type = sirius::to_duckdb(output); + auto data = duckdb::make_uniq(op, std::move(carriers), fn.return_type); + auto result_type = fn.return_type; + return duckdb::make_uniq( + std::move(result_type), std::move(fn), std::move(children), std::move(data)); +} +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_functions.hpp b/src/numeric/decimal_functions.hpp new file mode 100644 index 0000000000..f9660384d8 --- /dev/null +++ b/src/numeric/decimal_functions.hpp @@ -0,0 +1,31 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include "expression/function_id.hpp" +#include "numeric/decimal_types.hpp" + +#include +#include + +namespace duckdb { +class DatabaseInstance; +class BoundFunctionExpression; +} // namespace duckdb +namespace sirius::ast { +struct node; +} + +namespace sirius::mo_decimal { +bool is_decimal_function(function_id id) noexcept; +decimal_op operation(function_id id); +function_id function(decimal_op op); +bool contains_exact_expression(ast::node const& expression); +void validate_signature(decimal_op op, + std::vector const& inputs, + logical_type const& output); +void register_scalar_functions(duckdb::DatabaseInstance& instance); +void validate_bound_scalar(duckdb::BoundFunctionExpression const& expression); +duckdb::unique_ptr bound_scalar( + decimal_op op, + duckdb::vector> children, + logical_type const& output); +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_import.cpp b/src/numeric/decimal_import.cpp new file mode 100644 index 0000000000..88a1604017 --- /dev/null +++ b/src/numeric/decimal_import.cpp @@ -0,0 +1,167 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "numeric/decimal_import.hpp" + +#include "from_substrait.hpp" +#include "numeric/decimal_functions.hpp" +#include "numeric/exact_decimal.hpp" + +#include +#include +#include +#include + +#include +#include + +namespace sirius::mo_decimal { +namespace { +bool nullable(substrait::Type_Nullability value) +{ + if (value != substrait::Type::NULLABILITY_NULLABLE && + value != substrait::Type::NULLABILITY_REQUIRED) + throw std::invalid_argument("MO numeric type requires explicit nullability"); + return value == substrait::Type::NULLABILITY_NULLABLE; +} +template +decimal_type descriptor(UserType const& user) +{ + if (user.type_parameters_size() != 3) + throw std::invalid_argument("MO numeric type requires width precision and scale"); + int64_t parameters[3]{}; + for (int i = 0; i < 3; ++i) { + if (!user.type_parameters(i).has_integer()) + throw std::invalid_argument("MO numeric type parameters must be integers"); + parameters[i] = user.type_parameters(i).integer(); + } + if ((parameters[0] != 64 && parameters[0] != 128 && parameters[0] != 256) || parameters[1] < 1 || + parameters[1] > 76 || parameters[2] < 0 || parameters[2] > 76) + throw std::invalid_argument("invalid MO numeric type domain"); + decimal_type result{static_cast(parameters[0]), + static_cast(parameters[1]), + static_cast(parameters[2])}; + if (!result.valid()) throw std::invalid_argument("invalid MO numeric precision or scale"); + return result; +} +logical_type result_type(substrait::Type const& type, std::vector const& anchors) +{ + if (type.has_user_defined()) { + if (!std::binary_search(anchors.begin(), anchors.end(), type.user_defined().type_reference())) + throw std::invalid_argument("MO numeric result references another type identity"); + return logical_type::make_mo_decimal(descriptor(type.user_defined()), + nullable(type.user_defined().nullability())); + } + if (type.has_bool_()) + return logical_type::make(type_id::BOOLEAN) + .with_nullability(nullable(type.bool_().nullability())); + if (type.has_i64()) + return logical_type::make(type_id::BIGINT).with_nullability(nullable(type.i64().nullability())); + throw std::invalid_argument("unsupported MO decimal result type"); +} +class importer final : public duckdb::SubstraitExtensionHandler { + std::vector anchors_; + + public: + explicit importer(std::vector anchors) : anchors_(std::move(anchors)) {} + bool Handles(duckdb::SubstraitExtensionIdentity const& identity) const override + { + return identity.urn == extension_uri; + } + duckdb::LogicalType Type(duckdb::ClientContext&, + duckdb::SubstraitExtensionIdentity const& identity, + substrait::Type const& type) const override + { + if (!Handles(identity) || identity.name != "mo_exact_decimal" || !type.has_user_defined()) + throw std::invalid_argument("unhandled MO numeric type identity"); + return duckdb_type(result_type(type, anchors_)); + } + duckdb::unique_ptr Literal( + duckdb::ClientContext&, + duckdb::SubstraitExtensionIdentity const& identity, + substrait::Expression_Literal const& literal) const override + { + if (!Handles(identity) || identity.name != "mo_exact_decimal" || !literal.has_user_defined()) + throw std::invalid_argument("unhandled MO numeric literal identity"); + auto const& user = literal.user_defined(); + auto type = descriptor(user); + auto const& any = user.value(); + auto const& payload = any.value(); + if (any.type_url() != literal_type_url) + 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 + // is retained, in bounded stack storage; no numeric strings are parsed. + duckdb::google::protobuf::io::CodedInputStream input( + reinterpret_cast(payload.data()), static_cast(payload.size())); + std::array coefficient{}; + bool found = false, width_matches = false; + while (auto tag = input.ReadTag()) { + if (!(tag >> 3)) throw std::invalid_argument("invalid MO decimal literal field"); + if (tag == 10) { + uint32_t bytes; + if (!input.ReadVarint32(&bytes) || bytes > payload.size()) + throw std::invalid_argument("invalid MO decimal literal length"); + found = true; + width_matches = bytes == type.bytes(); + if (!(width_matches ? input.ReadRaw(coefficient.data(), bytes) + : input.Skip(static_cast(bytes)))) + throw std::invalid_argument("truncated MO decimal literal coefficient"); + } else if (!duckdb::google::protobuf::internal::WireFormatLite::SkipField(&input, tag)) + throw std::invalid_argument("invalid MO decimal literal payload"); + } + if (!found || !width_matches || !input.ConsumedEntireMessage()) + throw std::invalid_argument("MO decimal literal coefficient width mismatch"); + auto value = load_coefficient(coefficient.data(), type.bytes()); + return duckdb::make_uniq( + duckdb_value(value, logical_type::make_mo_decimal(type, literal.nullable()))); + } + duckdb::unique_ptr Scalar( + duckdb::ClientContext&, + duckdb::SubstraitExtensionIdentity const& identity, + substrait::Expression_ScalarFunction const& scalar, + duckdb::vector> children) const override + { + if (!Handles(identity) || !scalar.has_output_type() || + scalar.arguments_size() != static_cast(children.size()) || scalar.options_size()) + throw std::invalid_argument("invalid MO exact-decimal scalar signature"); + auto id = from_duckdb_function_name("__sirius_" + identity.name); + if (!id || !is_decimal_function(*id)) + throw std::invalid_argument("unknown MO exact-decimal scalar identity"); + auto op = operation(*id); + if (children.size() != (op == decimal_op::cast || op == decimal_op::negate ? 1u : 2u)) + throw std::invalid_argument("invalid MO exact-decimal scalar arity"); + auto output = result_type(scalar.output_type(), anchors_); + children.push_back( + duckdb::make_uniq(duckdb::Value(duckdb_type(output)))); + return duckdb::make_uniq(std::string(to_duckdb_function_name(*id)), + std::move(children)); + } +}; +} // namespace +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"); + substrait::Plan plan; + if (!plan.ParseFromArray(bytes.data(), static_cast(bytes.size()))) + throw std::invalid_argument("invalid MO numeric plan"); + // Retain only this extension's anchors, not copies of unrelated URI strings. + // The converter still validates all URI/type/function declarations. + std::vector urns; + for (auto const& entry : plan.extension_urns()) + if (entry.urn() == extension_uri) urns.push_back(entry.extension_urn_anchor()); + std::sort(urns.begin(), urns.end()); + if (std::adjacent_find(urns.begin(), urns.end()) != urns.end()) + throw std::invalid_argument("duplicate MO numeric extension URI anchor"); + std::vector anchors; + for (auto const& extension : plan.extensions()) { + if (!extension.has_extension_type()) continue; + auto const& type = extension.extension_type(); + if (std::binary_search(urns.begin(), urns.end(), type.extension_urn_reference()) && + type.name() == "mo_exact_decimal") + anchors.push_back(type.type_anchor()); + } + std::sort(anchors.begin(), anchors.end()); + if (std::adjacent_find(anchors.begin(), anchors.end()) != anchors.end()) + throw std::invalid_argument("duplicate MO numeric type anchor"); + return duckdb::make_shared_ptr(std::move(anchors)); +} +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_import.hpp b/src/numeric/decimal_import.hpp new file mode 100644 index 0000000000..32da20f14c --- /dev/null +++ b/src/numeric/decimal_import.hpp @@ -0,0 +1,14 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include + +#include +namespace duckdb { +class SubstraitExtensionHandler; +} +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); +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_type.hpp b/src/numeric/decimal_type.hpp new file mode 100644 index 0000000000..31b31dd0da --- /dev/null +++ b/src/numeric/decimal_type.hpp @@ -0,0 +1,51 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include +#ifdef __CUDACC__ +#define SIRIUS_DECIMAL_TYPE_HD __host__ __device__ +#else +#define SIRIUS_DECIMAL_TYPE_HD +#endif +namespace sirius::mo_decimal { +// MO physical width is independent of SQL precision. This descriptor does not +// change ordinary DuckDB/cuDF decimal type inference. +struct decimal_type { + uint16_t bits; + uint8_t precision, scale; + SIRIUS_DECIMAL_TYPE_HD bool valid() const + { + auto limit = bits == 64 ? 18 : bits == 128 ? 38 : bits == 256 ? 76 : 0; + return precision > 0 && precision <= limit && scale <= precision; + } + SIRIUS_DECIMAL_TYPE_HD uint32_t bytes() const { return bits / 8; } +}; + +struct coefficient { + uint32_t words[8]{}; // Little-endian, signed two's complement, sign extended. + SIRIUS_DECIMAL_TYPE_HD bool negative() const { return (words[7] >> 31) != 0; } +}; +enum class decimal_error : uint8_t { none, out_of_range, invalid_input }; +enum class decimal_op : uint8_t { + add, + subtract, + multiply, + divide, + integer_divide, + modulo, + negate, + cast, + equal, + not_equal, + less, + less_equal, + greater, + greater_equal +}; +struct decimal_result { + coefficient value{}; + decimal_error error{decimal_error::none}; + bool valid{false}; +}; + +} // namespace sirius::mo_decimal +#undef SIRIUS_DECIMAL_TYPE_HD diff --git a/src/numeric/decimal_types.cpp b/src/numeric/decimal_types.cpp new file mode 100644 index 0000000000..d96c0b89e4 --- /dev/null +++ b/src/numeric/decimal_types.cpp @@ -0,0 +1,139 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "numeric/decimal_types.hpp" + +#include "numeric/exact_decimal.hpp" + +#include +#include +#include + +namespace sirius::mo_decimal { +namespace { +constexpr std::string_view prefix = "__sirius_mo_exact_v1_"; +std::string alias(logical_type const& type) +{ + auto descriptor = type.is_mo_decimal() ? type.mo_decimal_type() : decimal_type{}; + auto tag = type.is_mo_decimal() ? "d" : type.id() == type_id::BOOLEAN ? "b" : "i"; + return std::string(prefix) + tag + "_" + std::to_string(descriptor.bits) + "_" + + std::to_string(descriptor.precision) + "_" + std::to_string(descriptor.scale) + "_" + + std::to_string(type.nullability()); +} +duckdb::LogicalType carrier(logical_type const& type) +{ + if (type.id() == type_id::BOOLEAN) return duckdb::LogicalType::BOOLEAN; + if (type.id() == type_id::BIGINT) return duckdb::LogicalType::BIGINT; + auto bits = type.mo_decimal_type().bits; + if (bits == 64) return duckdb::LogicalType::BIGINT; + if (bits == 128) return duckdb::LogicalType::HUGEINT; + return duckdb::LogicalType::STRUCT({{"high", duckdb::LogicalType::BIGINT}, + {"mid_high", duckdb::LogicalType::UBIGINT}, + {"mid_low", duckdb::LogicalType::UBIGINT}, + {"low", duckdb::LogicalType::UBIGINT}}); +} +} // namespace + +duckdb::LogicalType duckdb_type(logical_type const& type) +{ + if (!type.nullability() || + (!type.is_mo_decimal() && type.id() != type_id::BOOLEAN && type.id() != type_id::BIGINT)) + throw std::invalid_argument("invalid MO exact-decimal carrier"); + auto result = carrier(type); + result.SetAlias(alias(type)); + return result; +} +std::optional from_duckdb_type(duckdb::LogicalType const& type) +{ + auto const& name = type.GetAlias(); + if (!std::string_view(name).starts_with(prefix)) return std::nullopt; + auto tail = std::string_view(name).substr(prefix.size()); + if (tail.size() < 3 || tail[1] != '_') + throw std::invalid_argument("invalid MO exact-decimal carrier alias"); + auto tag = tail[0]; + tail.remove_prefix(2); + uint32_t fields[4]{}; + for (int i = 0; i < 4; ++i) { + auto end = i == 3 ? tail.size() : tail.find('_'); + if (end == std::string_view::npos || end == 0) + throw std::invalid_argument("invalid MO exact-decimal carrier parameters"); + auto text = tail.substr(0, end); + auto [last, error] = std::from_chars(text.data(), text.data() + text.size(), fields[i]); + if (error != std::errc{} || last != text.data() + text.size()) + throw std::invalid_argument("invalid MO exact-decimal carrier parameter"); + tail.remove_prefix(end + (i == 3 ? 0 : 1)); + } + if (!tail.empty() || (fields[3] != 1 && fields[3] != 2) || fields[0] > 256 || fields[1] > 76 || + fields[2] > 76) + throw std::invalid_argument("invalid MO exact-decimal carrier domain"); + logical_type result; + if (tag == 'd') + result = logical_type::make_mo_decimal({static_cast(fields[0]), + static_cast(fields[1]), + static_cast(fields[2])}, + fields[3] == 2); + else if ((tag == 'b' || tag == 'i') && !fields[0] && !fields[1] && !fields[2]) + result = logical_type::make(tag == 'b' ? type_id::BOOLEAN : type_id::BIGINT) + .with_nullability(fields[3] == 2); + else + throw std::invalid_argument("invalid MO exact-decimal carrier kind"); + if (duckdb_type(result) != type) + throw std::invalid_argument("noncanonical MO exact-decimal carrier"); + return result; +} +duckdb::Value duckdb_value(coefficient const& value, logical_type const& type) +{ + auto descriptor = type.mo_decimal_type(); + auto checked = + evaluate_decimal(decimal_op::cast, value, descriptor, true, {}, descriptor, true, descriptor); + if (checked.error != decimal_error::none) + throw std::invalid_argument("MO exact-decimal literal exceeds its declared domain"); + uint64_t limbs[4]{}; + for (int i = 0; i < 4; ++i) + limbs[i] = uint64_t(value.words[2 * i]) | uint64_t(value.words[2 * i + 1]) << 32; + duckdb::Value result; + if (descriptor.bits == 64) { + int64_t signed_value; + std::memcpy(&signed_value, limbs, 8); + result = duckdb::Value::BIGINT(signed_value); + } else if (descriptor.bits == 128) { + int64_t high; + std::memcpy(&high, limbs + 1, 8); + result = duckdb::Value::HUGEINT(duckdb::hugeint_t{high, limbs[0]}); + } else { + int64_t high; + std::memcpy(&high, limbs + 3, 8); + result = duckdb::Value::STRUCT(carrier(type), + {duckdb::Value::BIGINT(high), + duckdb::Value::UBIGINT(limbs[2]), + duckdb::Value::UBIGINT(limbs[1]), + duckdb::Value::UBIGINT(limbs[0])}); + } + result.Reinterpret(duckdb_type(type)); + return result; +} +coefficient from_duckdb_value(duckdb::Value const& value, logical_type const& type) +{ + if (value.IsNull() || value.type() != duckdb_type(type)) + throw std::invalid_argument("MO exact-decimal literal carrier mismatch"); + auto descriptor = type.mo_decimal_type(); + uint64_t limbs[4]{}; + if (descriptor.bits == 64) { + auto signed_value = value.GetValueUnsafe(); + std::memcpy(limbs, &signed_value, 8); + } else if (descriptor.bits == 128) { + auto big = value.GetValueUnsafe(); + limbs[0] = big.lower; + std::memcpy(limbs + 1, &big.upper, 8); + } else { + auto const& children = duckdb::StructValue::GetChildren(value); + for (int i = 0; i < 4; ++i) { + if (children[i].IsNull()) throw std::invalid_argument("NULL MO decimal coefficient limb"); + if (i == 0) { + auto high = children[i].GetValueUnsafe(); + std::memcpy(limbs + 3, &high, 8); + } else + limbs[3 - i] = children[i].GetValueUnsafe(); + } + } + return load_coefficient(reinterpret_cast(limbs), descriptor.bytes()); +} +} // namespace sirius::mo_decimal diff --git a/src/numeric/decimal_types.hpp b/src/numeric/decimal_types.hpp new file mode 100644 index 0000000000..9043fd4f96 --- /dev/null +++ b/src/numeric/decimal_types.hpp @@ -0,0 +1,17 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include "helper/logical_type.hpp" + +#include +#include + +#include + +namespace sirius::mo_decimal { +// Private DuckDB carriers preserve the complete MO descriptor. These are not +// registered SQL types and never participate in ordinary decimal inference. +duckdb::LogicalType duckdb_type(logical_type const& type); +std::optional from_duckdb_type(duckdb::LogicalType const& type); +duckdb::Value duckdb_value(coefficient const& value, logical_type const& type); +coefficient from_duckdb_value(duckdb::Value const& value, logical_type const& type); +} // namespace sirius::mo_decimal diff --git a/src/numeric/exact_decimal.hpp b/src/numeric/exact_decimal.hpp index 9e31c3a2c2..5049b43bc0 100644 --- a/src/numeric/exact_decimal.hpp +++ b/src/numeric/exact_decimal.hpp @@ -1,7 +1,7 @@ /* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ #pragma once -#include +#include "numeric/decimal_type.hpp" #ifdef __CUDACC__ #define SIRIUS_DECIMAL_HD __host__ __device__ @@ -10,46 +10,6 @@ #endif namespace sirius::mo_decimal { -// MO physical width is independent of SQL precision. This descriptor does not -// change ordinary DuckDB/cuDF decimal type inference. -struct decimal_type { - uint16_t bits; - uint8_t precision, scale; - SIRIUS_DECIMAL_HD bool valid() const - { - auto limit = bits == 64 ? 18 : bits == 128 ? 38 : bits == 256 ? 76 : 0; - return precision > 0 && precision <= limit && scale <= precision; - } - SIRIUS_DECIMAL_HD uint32_t bytes() const { return bits / 8; } -}; - -struct coefficient { - uint32_t words[8]{}; // Little-endian, signed two's complement, sign extended. - SIRIUS_DECIMAL_HD bool negative() const { return (words[7] >> 31) != 0; } -}; -enum class decimal_error : uint8_t { none, out_of_range, invalid_input }; -enum class decimal_op : uint8_t { - add, - subtract, - multiply, - divide, - integer_divide, - modulo, - negate, - cast, - equal, - not_equal, - less, - less_equal, - greater, - greater_equal -}; -struct decimal_result { - coefficient value{}; - decimal_error error{decimal_error::none}; - bool valid{false}; -}; - // Same-scale 64-bit operations need no wide scratch. Unsigned arithmetic makes // physical overflow checks defined, including signed-minimum negation. SIRIUS_DECIMAL_HD inline decimal_result evaluate_same_scale64(decimal_op op, diff --git a/src/numeric/exact_decimal_gpu.cu b/src/numeric/exact_decimal_gpu.cu index f636f9b483..b97417703b 100644 --- a/src/numeric/exact_decimal_gpu.cu +++ b/src/numeric/exact_decimal_gpu.cu @@ -4,6 +4,8 @@ #include #include +#include +#include #include #include @@ -12,6 +14,26 @@ namespace sirius::mo_decimal { namespace { +__global__ void broadcast_literal(coefficient value, + uint32_t bytes, + cudf::size_type rows, + uint8_t* data, + uint64_t* high, + uint64_t* mid_high, + uint64_t* mid_low, + uint64_t* low) +{ + for (int64_t row = blockIdx.x * blockDim.x + threadIdx.x; row < rows; + row += blockDim.x * gridDim.x) { + if (bytes != 32) + store_coefficient(value, data + row * bytes, bytes); + else { + uint64_t* limbs[]{low, mid_low, mid_high, high}; + for (int i = 0; i < 4; ++i) + limbs[i][row] = uint64_t(value.words[2 * i]) | uint64_t(value.words[2 * i + 1]) << 32; + } + } +} struct operand { uint8_t const* data{}; uint64_t const* limbs[4]{}; @@ -125,6 +147,73 @@ __global__ void evaluate(decimal_op op, } } // namespace +std::unique_ptr make_decimal_literal(decimal_type type, + coefficient value, + bool valid, + cudf::size_type rows, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + if (!type.valid() || rows < 0) + throw std::invalid_argument("invalid MO exact-decimal literal descriptor"); + if (valid && evaluate_decimal(decimal_op::cast, value, type, true, {}, type, true, type).error != + decimal_error::none) + throw std::invalid_argument("MO exact-decimal literal exceeds declared precision"); + auto owner = std::make_unique>(); + try { + *owner = make_decimal_column( + type, rows, valid ? cudf::mask_state::ALL_VALID : cudf::mask_state::ALL_NULL, stream, mr); + if (rows) { + auto view = (*owner)->mutable_view(); + broadcast_literal<<<128, 256, 0, stream.value()>>>( + valid ? value : coefficient{}, + type.bytes(), + rows, + type.bits == 256 ? nullptr : view.data(), + type.bits == 256 ? view.child(0).data() : nullptr, + type.bits == 256 ? view.child(1).data() : nullptr, + type.bits == 256 ? view.child(2).data() : nullptr, + type.bits == 256 ? view.child(3).data() : nullptr); + CUDF_CUDA_TRY(cudaGetLastError()); + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + } + return std::move(*owner); + } catch (...) { + auto error = std::current_exception(); + if (cudaStreamSynchronize(stream.value()) != cudaSuccess) { + (void)owner.release(); + throw pipeline::gpu_stream_quiescence_error("MO decimal literal could not prove quiescence"); + } + std::rethrow_exception(error); + } +} +decimal_error column_error(decimal_column_result const& result, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + if (!result.errors.size()) return decimal_error::none; + cudf::column_view errors( + cudf::data_type{cudf::type_id::UINT8}, result.errors.size(), result.errors.data(), nullptr, 0); + auto owner = std::make_unique>(); + try { + *owner = cudf::reduce(errors, + *cudf::make_max_aggregation(), + cudf::data_type{cudf::type_id::UINT8}, + stream, + mr); + auto code = static_cast&>(**owner).value(stream); + return static_cast(code); + } catch (...) { + auto error = std::current_exception(); + if (cudaStreamSynchronize(stream.value()) != cudaSuccess) { + (void)owner.release(); + throw pipeline::gpu_stream_quiescence_error( + "MO decimal error reduction could not prove quiescence"); + } + std::rethrow_exception(error); + } +} + std::unique_ptr make_decimal_column(decimal_type type, cudf::size_type rows, cudf::mask_state mask, diff --git a/src/numeric/exact_decimal_gpu.hpp b/src/numeric/exact_decimal_gpu.hpp index 80bb3769b3..61e58dba25 100644 --- a/src/numeric/exact_decimal_gpu.hpp +++ b/src/numeric/exact_decimal_gpu.hpp @@ -24,6 +24,15 @@ struct decimal_column_result { std::unique_ptr values; rmm::device_uvector errors; }; +std::unique_ptr make_decimal_literal(decimal_type type, + coefficient value, + bool valid, + cudf::size_type rows, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); +decimal_error column_error(decimal_column_result const& result, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); // Caller retains operands/mask and the returned buffers until its task stream // is quiescent. All allocations use that stream's reservation-aware resource. decimal_column_result evaluate_decimal_columns(decimal_op op, diff --git a/substrait b/substrait index dd4cab14b8..95d9ce8d78 160000 --- a/substrait +++ b/substrait @@ -1 +1 @@ -Subproject commit dd4cab14b82754ca919633913436cd530f496e00 +Subproject commit 95d9ce8d78490db3991ab6145653716aa3ec42c9 diff --git a/test/cpp/embedding/decimal_import_fixture.hpp b/test/cpp/embedding/decimal_import_fixture.hpp new file mode 100644 index 0000000000..b1ab492304 --- /dev/null +++ b/test/cpp/embedding/decimal_import_fixture.hpp @@ -0,0 +1,147 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include "expression/ast/from_duckdb.hpp" +#include "from_substrait.hpp" +#include "numeric/decimal_functions.hpp" +#include "numeric/decimal_import.hpp" +#include "numeric/exact_decimal.hpp" + +#include +#include +#include +#include +#include +#include + +namespace decimal_fixture { +inline sirius::mo_decimal::coefficient small(int64_t value) +{ + return sirius::mo_decimal::load_coefficient(reinterpret_cast(&value), 8); +} +template +void parameters(User* value, sirius::mo_decimal::decimal_type type) +{ + for (auto item : {int(type.bits), int(type.precision), int(type.scale)}) + value->add_type_parameters()->set_integer(item); +} +inline substrait::Type type(sirius::mo_decimal::decimal_type descriptor, bool nullable = false) +{ + substrait::Type result; + auto user = result.mutable_user_defined(); + user->set_type_reference(11); + user->set_nullability(nullable ? substrait::Type::NULLABILITY_NULLABLE + : substrait::Type::NULLABILITY_REQUIRED); + parameters(user, descriptor); + return result; +} +inline substrait::Expression literal(sirius::mo_decimal::decimal_type type, + sirius::mo_decimal::coefficient const& value) +{ + substrait::Expression result; + auto user = result.mutable_literal()->mutable_user_defined(); + user->set_type_reference(11); + parameters(user, type); + user->mutable_value()->set_type_url(sirius::mo_decimal::literal_type_url); + std::string payload(type.bytes() + 2, '\0'); + payload[0] = 10; + payload[1] = type.bytes(); + sirius::mo_decimal::store_coefficient( + value, reinterpret_cast(payload.data() + 2), type.bytes()); + user->mutable_value()->set_value(payload); + return result; +} +inline substrait::Expression literal(sirius::mo_decimal::decimal_type type, int64_t value) +{ + return literal(type, small(value)); +} +inline substrait::Plan plan(std::string name, + substrait::Type output, + std::vector arguments) +{ + substrait::Plan result; + result.mutable_version()->set_minor_number(78); + auto urn = result.add_extension_urns(); + urn->set_extension_urn_anchor(1); + urn->set_urn(sirius::mo_decimal::extension_uri); + auto type = result.add_extensions()->mutable_extension_type(); + type->set_type_anchor(11); + type->set_extension_urn_reference(1); + type->set_name("mo_exact_decimal"); + auto fn = result.add_extensions()->mutable_extension_function(); + fn->set_function_anchor(21); + fn->set_extension_urn_reference(1); + fn->set_name(name); + auto root = result.add_relations()->mutable_root(); + root->add_names("exact"); + auto project = root->mutable_input()->mutable_project(); + project->mutable_common()->mutable_emit()->add_output_mapping(1); + auto read = project->mutable_input()->mutable_read(); + read->mutable_base_schema()->add_names("seed"); + read->mutable_base_schema()->mutable_struct_()->add_types()->mutable_i64()->set_nullability( + substrait::Type::NULLABILITY_REQUIRED); + read->mutable_virtual_table()->add_expressions()->add_fields()->mutable_literal()->set_i64(1); + auto scalar = project->add_expressions()->mutable_scalar_function(); + scalar->set_function_reference(21); + *scalar->mutable_output_type() = std::move(output); + for (auto& argument : arguments) + *scalar->add_arguments()->mutable_value() = std::move(argument); + return result; +} +class importer { + static std::unique_ptr config() + { + auto result = std::make_unique(); + result->options.load_extensions = false; + return result; + } + std::unique_ptr config_{config()}; + + public: + duckdb::DuckDB db{nullptr, config_.get()}; + duckdb::Connection connection{db}; + importer() + { + db.LoadStaticExtension(); + sirius::mo_decimal::register_scalar_functions(*db.instance); + } + std::unique_ptr bind(substrait::Plan const& plan) + { + auto bytes = plan.SerializeAsString(); + connection.BeginTransaction(); + try { + duckdb::SubstraitToDuckDB converter( + connection.context, bytes, false, false, sirius::mo_decimal::make_import_handler(bytes)); + auto relation = converter.TransformPlan(); + duckdb::Planner planner(*connection.context); + planner.CreatePlan(duckdb::make_uniq(relation)); + std::unique_ptr found; + std::function expression; + expression = [&](duckdb::Expression const& value) { + if (found) return; + if (value.GetExpressionClass() == duckdb::ExpressionClass::BOUND_FUNCTION) { + auto const& function = value.Cast(); + auto id = sirius::from_duckdb_function_name(function.function.name); + if (id && sirius::mo_decimal::is_decimal_function(*id)) { + found = sirius::ast::from_duckdb(value); + return; + } + } + duckdb::ExpressionIterator::EnumerateChildren(value, expression); + }; + std::function visit = [&](auto const& op) { + for (auto const& value : op.expressions) + expression(*value); + for (auto const& child : op.children) + visit(*child); + }; + visit(*planner.plan); + if (!found) throw std::runtime_error("numeric function disappeared during import"); + connection.Rollback(); + return found; + } catch (...) { + if (connection.HasActiveTransaction()) connection.Rollback(); + throw; + } + } +}; +} // namespace decimal_fixture diff --git a/test/cpp/embedding/test_decimal_expression_gpu.cpp b/test/cpp/embedding/test_decimal_expression_gpu.cpp new file mode 100644 index 0000000000..72a917a2ec --- /dev/null +++ b/test/cpp/embedding/test_decimal_expression_gpu.cpp @@ -0,0 +1,407 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "decimal_import_fixture.hpp" +#include "expression/ast/utils.hpp" +#include "expression_evaluator/expression_evaluator.hpp" +#include "numeric/decimal_error.hpp" +#include "numeric/exact_decimal_gpu.hpp" +#include "sirius_c.h" + +#include + +#include + +#include + +using namespace sirius; +using namespace sirius::mo_decimal; +namespace { +std::vector download(cudf::column_view column, + uint32_t bytes, + rmm::cuda_stream_view stream) +{ + std::vector result(column.size()); + if (bytes == 32) { + std::vector limb(column.size()); + for (int i = 0; i < 4; ++i) { + REQUIRE(cudaMemcpyAsync(limb.data(), + column.child(3 - i).data(), + limb.size() * 8, + cudaMemcpyDeviceToHost, + stream.value()) == cudaSuccess); + stream.synchronize(); + for (std::size_t row = 0; row < result.size(); ++row) { + result[row].words[2 * i] = limb[row]; + result[row].words[2 * i + 1] = limb[row] >> 32; + } + } + } else { + std::vector data(result.size() * bytes); + REQUIRE( + cudaMemcpyAsync( + data.data(), column.data(), data.size(), cudaMemcpyDeviceToHost, stream.value()) == + cudaSuccess); + stream.synchronize(); + for (std::size_t row = 0; row < result.size(); ++row) { + if (bytes == 1) + result[row].words[0] = data[row]; + else + result[row] = load_coefficient(data.data() + row * bytes, bytes); + } + } + return result; +} +void expect(cudf::column_view column, + uint32_t bytes, + std::vector const& expected, + rmm::cuda_stream_view stream) +{ + REQUIRE(column.size() == static_cast(expected.size())); + REQUIRE(column.null_count() == 0); + auto values = download(column, bytes, stream); + for (std::size_t row = 0; row < values.size(); ++row) { + auto reference = decimal_fixture::small(expected[row]); + for (int word = 0; word < 8; ++word) + CHECK(values[row].words[word] == reference.words[word]); + } +} +std::unique_ptr constant(int64_t value, logical_type const& type) +{ + return std::make_unique(ast::constant{decimal_fixture::small(value), type}); +} +std::unique_ptr reference(uint32_t index, logical_type const& type) +{ + return std::make_unique(ast::reference{index, type}); +} +std::unique_ptr binary(decimal_op op, + std::unique_ptr left, + std::unique_ptr right, + logical_type const& output) +{ + std::vector> arguments; + arguments.push_back(std::move(left)); + arguments.push_back(std::move(right)); + return std::make_unique( + ast::function_call{function(op), std::move(arguments), output}); +} +std::unique_ptr values(decimal_type type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr, + bool second_null = false, + int64_t second_value = 125) +{ + auto column = make_decimal_literal(type, decimal_fixture::small(999), true, 2, stream, mr); + auto view = column->mutable_view(); + if (type.bits != 256) { + uint8_t coefficient[16]{}; + store_coefficient(decimal_fixture::small(second_value), coefficient, type.bytes()); + REQUIRE(cudaMemcpyAsync(view.data() + type.bytes(), + coefficient, + type.bytes(), + cudaMemcpyHostToDevice, + stream.value()) == cudaSuccess); + stream.synchronize(); + } else { + for (int i = 0; i < 4; ++i) { + auto coefficient = decimal_fixture::small(second_value); + auto index = (3 - i) * 2; + uint64_t value = uint64_t(coefficient.words[index]) | uint64_t(coefficient.words[index + 1]) + << 32; + REQUIRE( + cudaMemcpyAsync( + view.child(i).data() + 1, &value, 8, cudaMemcpyHostToDevice, stream.value()) == + cudaSuccess); + stream.synchronize(); + } + } + if (second_null) { + uint32_t mask = 1; + REQUIRE(cudaMemcpyAsync(view.null_mask(), &mask, 4, cudaMemcpyHostToDevice, stream.value()) == + cudaSuccess); + stream.synchronize(); + column->set_null_count(1); + } + return column; +} +} // namespace + +TEST_CASE("MO imported scalar signatures execute on the GPU without coercion", + "[decimal_import_gpu]") +{ + decimal_fixture::importer importer; + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + char const* names[]{"mo_decimal_add", + "mo_decimal_subtract", + "mo_decimal_multiply", + "mo_decimal_divide", + "mo_decimal_integer_divide", + "mo_decimal_modulo", + "mo_decimal_negate", + "mo_decimal_cast", + "mo_decimal_equal", + "mo_decimal_not_equal", + "mo_decimal_less", + "mo_decimal_less_equal", + "mo_decimal_greater", + "mo_decimal_greater_equal"}; + int64_t expected[]{225, 25, 125, 125, 1, 25, -125, 13, 0, 1, 0, 0, 1, 1}; + for (uint16_t bits : {64, 128, 256}) { + decimal_type input{bits, 9, 2}; + auto dummy = make_decimal_literal(input, {}, true, 3, stream, mr); + for (unsigned i = 0; i < 14; ++i) { + INFO("bits=" << bits << " function=" << names[i]); + auto op = static_cast(i); + bool unary = op == decimal_op::cast || op == decimal_op::negate; + decimal_type output{bits, 15, static_cast(op == decimal_op::cast ? 1 : 2)}; + auto result_type = decimal_fixture::type(output); + if (op >= decimal_op::equal) + result_type.mutable_bool_()->set_nullability(substrait::Type::NULLABILITY_REQUIRED); + else if (op == decimal_op::integer_divide) + result_type.mutable_i64()->set_nullability(substrait::Type::NULLABILITY_NULLABLE); + std::vector arguments{decimal_fixture::literal(input, 125)}; + if (!unary) arguments.push_back(decimal_fixture::literal(input, 100)); + auto plan = decimal_fixture::plan(names[i], result_type, std::move(arguments)); + auto expression = importer.bind(plan); + expression_evaluator evaluator(*expression, mr, stream.view()); + auto result = evaluator.evaluate(cudf::table_view{{dummy->view()}}); + auto bytes = op >= decimal_op::equal ? 1u + : op == decimal_op::integer_divide ? 8u + : input.bytes(); + expect( + result->view().column(0), bytes, {expected[i], expected[i], expected[i]}, stream.view()); + if (op < decimal_op::equal && op != decimal_op::integer_divide) + CHECK(decimal_column_matches(result->view().column(0), output)); + } + } + CHECK((sirius_capabilities() & 16u) == 0); +} + +TEST_CASE("MO imported literals preserve distinct high limbs and signs", "[decimal_import_gpu]") +{ + decimal_fixture::importer importer; + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + for (uint16_t bits : {64, 128, 256}) { + decimal_type type{bits, static_cast(bits == 64 ? 18 : bits == 128 ? 38 : 65), 4}; + coefficient value; + value.words[0] = 0x89abcdef; + value.words[1] = 0x01234567; + if (bits >= 128) { + value.words[2] = 0xaabbccdd; + value.words[3] = 0x01234567; + } + if (bits == 256) { + value.words[4] = 0x12345678; + value.words[5] = 0x9abcdef0; + value.words[6] = 5; + } + auto dummy = make_decimal_literal(type, {}, true, 2, stream, mr); + for (bool negative : {false, true}) { + auto expected = value; + if (negative) { + uint64_t carry = 1; + for (auto& word : expected.words) { + auto next = uint64_t(~word) + carry; + word = next; + carry = next >> 32; + } + } + auto plan = decimal_fixture::plan( + "mo_decimal_cast", decimal_fixture::type(type), {decimal_fixture::literal(type, expected)}); + auto expression = importer.bind(plan); + expression_evaluator evaluator(*expression, mr, stream.view()); + auto output = evaluator.evaluate(cudf::table_view{{dummy->view()}}); + auto observed = download(output->view().column(0), type.bytes(), stream.view()); + REQUIRE(output->view().column(0).null_count() == 0); + for (auto const& row : observed) + for (int word = 0; word < 8; ++word) + CHECK(row.words[word] == expected.words[word]); + } + } +} + +TEST_CASE("MO integer casts preserve the full signed source domain", "[decimal_import_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + auto input = cudf::make_numeric_column( + cudf::data_type{cudf::type_id::INT64}, 2, cudf::mask_state::UNALLOCATED, stream, mr); + int64_t endpoints[]{INT64_MIN, INT64_MAX}; + REQUIRE(cudaMemcpyAsync(input->mutable_view().data(), + endpoints, + sizeof(endpoints), + cudaMemcpyHostToDevice, + stream.value()) == cudaSuccess); + stream.synchronize(); + auto type = logical_type::make_mo_decimal({128, 19, 0}, false); + std::vector> arguments; + arguments.push_back(reference(0, logical_type::make(type_id::BIGINT))); + ast::node cast(ast::function_call{function(decimal_op::cast), std::move(arguments), type}); + expression_evaluator evaluator(cast, mr, stream.view()); + auto output = evaluator.evaluate(cudf::table_view{{input->view()}}); + expect(output->view().column(0), 16, {INT64_MIN, INT64_MAX}, stream.view()); +} + +TEST_CASE("MO conditional masks suppress inactive arithmetic and cast errors", + "[decimal_import_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + for (uint16_t bits : {64, 128, 256}) { + INFO("conditional bits=" << bits); + auto type = logical_type::make_mo_decimal({bits, 3, 0}, false); + auto input = values(type.mo_decimal_type(), stream, mr); + auto nullable_input = values(type.mo_decimal_type(), stream, mr, true); + auto predicate = cudf::make_numeric_column( + cudf::data_type{cudf::type_id::BOOL8}, 2, cudf::mask_state::UNALLOCATED, stream, mr); + bool selected[]{false, true}; + REQUIRE(cudaMemcpyAsync(predicate->mutable_view().data(), + selected, + sizeof(selected), + cudaMemcpyHostToDevice, + stream.value()) == cudaSuccess); + stream.synchronize(); + std::vector branches; + branches.push_back({reference(1, logical_type::make(type_id::BOOLEAN)), + binary(decimal_op::add, reference(0, type), constant(1, type), type)}); + auto expression = + std::make_unique(ast::case_expr{std::move(branches), constant(9, type), type}); + expression_evaluator case_only(*expression, mr, stream.view()); + auto case_result = case_only.evaluate(cudf::table_view{{input->view(), predicate->view()}}); + INFO("CASE parent nullable=" << case_result->view().column(0).nullable()); + if (bits == 256) { + INFO("CASE high limb nullable=" << case_result->view().column(0).child(0).nullable()); + REQUIRE(decimal_column_matches(case_result->view().column(0), type.mo_decimal_type())); + } + // A following exact operation also checks that conditional output retained + // the canonical parent-validity-only Decimal256 representation. + expression = binary(decimal_op::add, std::move(expression), constant(1, type), type); + expression_evaluator evaluator(*expression, mr, stream.view()); + auto result = evaluator.evaluate(cudf::table_view{{input->view(), predicate->view()}}); + expect(result->view().column(0), type.fixed_width_byte_size(), {10, 127}, stream.view()); + CHECK(decimal_column_matches(result->view().column(0), type.mo_decimal_type())); + selected[0] = true; + REQUIRE(cudaMemcpyAsync(predicate->mutable_view().data(), + selected, + sizeof(selected), + cudaMemcpyHostToDevice, + stream.value()) == cudaSuccess); + stream.synchronize(); + try { + evaluator.evaluate(cudf::table_view{{input->view(), predicate->view()}}); + FAIL("active CASE overflow was ignored"); + } catch (numeric_error const& error) { + CHECK(error.code() == decimal_error::out_of_range); + } + selected[0] = false; + REQUIRE(cudaMemcpyAsync(predicate->mutable_view().data(), + selected, + sizeof(selected), + cudaMemcpyHostToDevice, + stream.value()) == cudaSuccess); + stream.synchronize(); + auto recovered = evaluator.evaluate(cudf::table_view{{input->view(), predicate->view()}}); + expect(recovered->view().column(0), type.fixed_width_byte_size(), {10, 127}, stream.view()); + + auto narrow_type = logical_type::make_mo_decimal({bits, 2, 0}, false); + auto cast_input = values(type.mo_decimal_type(), stream, mr, false, 12); + std::vector> masked_cast_arguments; + masked_cast_arguments.push_back(reference(0, type)); + auto masked_cast = std::make_unique(ast::function_call{ + function(decimal_op::cast), std::move(masked_cast_arguments), narrow_type}); + std::vector cast_branches; + cast_branches.push_back( + {reference(1, logical_type::make(type_id::BOOLEAN)), std::move(masked_cast)}); + ast::node cast_case( + ast::case_expr{std::move(cast_branches), constant(9, narrow_type), narrow_type}); + expression_evaluator cast_selector(cast_case, mr, stream.view()); + auto cast_result = + cast_selector.evaluate(cudf::table_view{{cast_input->view(), predicate->view()}}); + expect(cast_result->view().column(0), type.fixed_width_byte_size(), {9, 12}, stream.view()); + + auto nullable_type = logical_type::make_mo_decimal({bits, 3, 0}, true); + std::vector> alternatives; + alternatives.push_back(reference(1, nullable_type)); + alternatives.push_back(binary(decimal_op::add, reference(0, type), constant(1, type), type)); + ast::node coalesce(ast::coalesce{std::move(alternatives), nullable_type}); + expression_evaluator coalescer(coalesce, mr, stream.view()); + auto combined = coalescer.evaluate(cudf::table_view{{input->view(), nullable_input->view()}}); + expect(combined->view().column(0), type.fixed_width_byte_size(), {999, 126}, stream.view()); + auto active = binary(decimal_op::add, reference(0, type), constant(1, type), type); + expression_evaluator failing(*active, mr, stream.view()); + try { + failing.evaluate(cudf::table_view{{input->view()}}); + FAIL("active arithmetic overflow was ignored"); + } catch (numeric_error const& error) { + CHECK(error.code() == decimal_error::out_of_range); + } + std::vector> cast_arguments; + cast_arguments.push_back(reference(0, type)); + ast::node cast(ast::function_call{function(decimal_op::cast), + std::move(cast_arguments), + logical_type::make_mo_decimal({bits, 2, 0}, false)}); + expression_evaluator narrowing(cast, mr, stream.view()); + try { + narrowing.evaluate(cudf::table_view{{input->view()}}); + FAIL("active cast overflow was ignored"); + } catch (numeric_error const& error) { + CHECK(error.code() == decimal_error::invalid_input); + } + } +} + +TEST_CASE("MO SELECT zero divisors produce NULL without numeric failure", "[decimal_import_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + auto input_type = logical_type::make_mo_decimal({64, 9, 0}, false); + auto output_type = logical_type::make_mo_decimal({128, 19, 2}, true); + auto dummy = make_decimal_literal(input_type.mo_decimal_type(), {}, true, 2, stream, mr); + for (int64_t divisor : {0, 2}) { + auto expression = binary( + decimal_op::divide, constant(1, input_type), constant(divisor, input_type), output_type); + expression_evaluator evaluator(*expression, mr, stream.view()); + auto result = evaluator.evaluate(cudf::table_view{{dummy->view()}}); + if (divisor == 0) + CHECK(result->view().column(0).null_count() == 2); + else + expect(result->view().column(0), 16, {50, 50}, stream.view()); + } +} + +TEST_CASE("MO references reject width restoration while ordinary decimals retain it", + "[decimal_import_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + auto compressed = cudf::make_fixed_width_column( + cudf::data_type{cudf::type_id::DECIMAL32, 0}, 2, cudf::mask_state::UNALLOCATED, stream, mr); + REQUIRE(cudaMemsetAsync(compressed->mutable_view().data(), 0, 8, stream.value()) == + cudaSuccess); + stream.synchronize(); + auto exact = reference(0, logical_type::make_mo_decimal({64, 18, 0}, false)); + expression_evaluator strict(*exact, mr, stream.view()); + REQUIRE_THROWS_AS(strict.evaluate(cudf::table_view{{compressed->view()}}), std::invalid_argument); + auto ordinary = reference(0, logical_type::make_decimal(18, 0)); + expression_evaluator restoring(*ordinary, mr, stream.view()); + auto output = restoring.evaluate(cudf::table_view{{compressed->view()}}); + CHECK(output->view().column(0).type() == cudf::data_type{cudf::type_id::DECIMAL64, 0}); + expect(output->view().column(0), 8, {0, 0}, stream.view()); +} + +TEST_CASE("MO exact GPU expressions retain empty schemas and evaluator reuse", + "[decimal_import_gpu]") +{ + rmm::cuda_stream stream; + auto mr = cudf::get_current_device_resource_ref(); + auto type = logical_type::make_mo_decimal({256, 15, 2}, false); + auto expression = binary(decimal_op::add, constant(125, type), constant(-25, type), type); + expression_evaluator evaluator(*expression, mr, stream.view()); + for (int rows : {0, 2, 0}) { + auto input = make_decimal_literal(type.mo_decimal_type(), {}, true, rows, stream, mr); + auto output = evaluator.evaluate(cudf::table_view{{input->view()}}); + CHECK(output->num_rows() == rows); + CHECK(decimal_column_matches(output->view().column(0), type.mo_decimal_type())); + if (rows) expect(output->view().column(0), 32, {100, 100}, stream.view()); + } +} diff --git a/test/cpp/embedding/test_decimal_import.cpp b/test/cpp/embedding/test_decimal_import.cpp new file mode 100644 index 0000000000..9d20d0b843 --- /dev/null +++ b/test/cpp/embedding/test_decimal_import.cpp @@ -0,0 +1,165 @@ +/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */ +#include "decimal_import_fixture.hpp" +#include "expression/ast/node.hpp" +#include "expression/ast/to_duckdb.hpp" +#include "expression/ast/utils.hpp" +#include "expression/value.hpp" +#include "helper/type_conversions.hpp" +#include "numeric/decimal_types.hpp" +#include "numeric/exact_decimal.hpp" + +#include + +using namespace sirius; + +TEST_CASE("MO import carriers retain width precision scale and nullability", "[decimal_import]") +{ + for (uint16_t bits : {64, 128, 256}) { + for (bool nullable : {false, true}) { + auto type = logical_type::make_mo_decimal({bits, 9, 2}, nullable); + CHECK(type.fixed_width_byte_size() == bits / 8); + auto carrier = to_duckdb(type); + CHECK(from_duckdb(carrier) == type); + CHECK(carrier.id() == (bits == 64 ? duckdb::LogicalTypeId::BIGINT + : bits == 128 ? duckdb::LogicalTypeId::HUGEINT + : duckdb::LogicalTypeId::STRUCT)); + for (int64_t number : {0, 125, -125}) { + auto coefficient = + mo_decimal::load_coefficient(reinterpret_cast(&number), 8); + auto value = mo_decimal::duckdb_value(coefficient, type); + auto observed = mo_decimal::from_duckdb_value(value, type); + for (int i = 0; i < 8; ++i) + CHECK(observed.words[i] == coefficient.words[i]); + auto literal = ast::node(ast::constant{from_duckdb(value, type), type}); + auto copied = ast::clone(literal); + CHECK(copied->return_type() == type); + CHECK( + to_duckdb(std::get(copied->get().payload), type) + .type() == carrier); + } + if (nullable) { + auto null = from_duckdb(duckdb::Value(carrier), type); + CHECK(std::holds_alternative(null)); + CHECK(to_duckdb(null, type).type() == carrier); + CHECK(to_duckdb(null, type).IsNull()); + } else + REQUIRE_THROWS(from_duckdb(duckdb::Value(carrier), type)); + } + } + auto ordinary = from_duckdb(duckdb::LogicalType::DECIMAL(9, 2)); + CHECK(ordinary.id() == type_id::DECIMAL); + CHECK(ordinary.fixed_width_byte_size() == 4); + CHECK(to_duckdb(ordinary) == duckdb::LogicalType::DECIMAL(9, 2)); +} + +TEST_CASE("MO import rejects malformed and forged physical carriers", "[decimal_import]") +{ + REQUIRE_THROWS(logical_type::make_mo_decimal({64, 19, 0}, false)); + REQUIRE_THROWS(logical_type::make_mo_decimal({256, 77, 0}, true)); + auto real = to_duckdb(logical_type::make_mo_decimal({256, 15, 2}, true)); + duckdb::LogicalType forged = duckdb::LogicalType::BIGINT; + forged.SetAlias(real.GetAlias()); + REQUIRE_THROWS(from_duckdb(forged)); + for (auto alias : {"__sirius_mo_exact_v1_d_256_15_2_0", + "__sirius_mo_exact_v1_d_256_015_2_2", + "__sirius_mo_exact_v1_d_256_15_16_2", + "__sirius_mo_exact_v1_d_256_15_2_2_extra"}) { + auto invalid = real; + invalid.SetAlias(alias); + REQUIRE_THROWS(from_duckdb(invalid)); + } + int64_t value = 1000; + auto coefficient = mo_decimal::load_coefficient(reinterpret_cast(&value), 8); + REQUIRE_THROWS( + mo_decimal::duckdb_value(coefficient, logical_type::make_mo_decimal({64, 3, 0}, false))); +} + +TEST_CASE("MO scalar import binds exact types without numeric coercion", "[decimal_import]") +{ + decimal_fixture::importer importer; + for (uint16_t bits : {64, 128, 256}) { + mo_decimal::decimal_type input{bits, 9, 2}, output{bits, 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 expression = importer.bind(plan); + REQUIRE(expression->is_function_call()); + CHECK(expression->as_function_call().function() == function_id::mo_decimal_add); + CHECK(expression->return_type() == logical_type::make_mo_decimal(output, false)); + REQUIRE(expression->as_function_call().arguments().size() == 2); + for (auto const& argument : expression->as_function_call().arguments()) + CHECK(argument->return_type() == logical_type::make_mo_decimal(input, false)); + auto copied = ast::clone(*expression); + auto reconstructed = ast::to_duckdb(*copied); + auto roundtrip = ast::from_duckdb(*reconstructed); + CHECK(roundtrip->return_type() == expression->return_type()); + CHECK(roundtrip->as_function_call().function() == function_id::mo_decimal_add); + } +} + +TEST_CASE("MO scalar import rejects unrelated identities and malformed literals", + "[decimal_import]") +{ + decimal_fixture::importer importer; + mo_decimal::decimal_type input{256, 15, 2}; + auto original = decimal_fixture::plan( + "mo_decimal_cast", decimal_fixture::type(input), {decimal_fixture::literal(input, 125)}); + REQUIRE_NOTHROW(importer.bind(original)); + auto bad = original; + auto scalar = bad.mutable_relations(0) + ->mutable_root() + ->mutable_input() + ->mutable_project() + ->mutable_expressions(0) + ->mutable_scalar_function(); + scalar->mutable_output_type()->mutable_user_defined()->set_type_reference(999); + REQUIRE_THROWS(importer.bind(bad)); + bad = original; + scalar = bad.mutable_relations(0) + ->mutable_root() + ->mutable_input() + ->mutable_project() + ->mutable_expressions(0) + ->mutable_scalar_function(); + scalar->mutable_arguments(0) + ->mutable_value() + ->mutable_literal() + ->mutable_user_defined() + ->mutable_value() + ->set_value(std::string(32, '\0')); + REQUIRE_THROWS(importer.bind(bad)); + bad = original; + bad.mutable_extensions(1)->mutable_extension_function()->set_name("mo_decimal_cast:unproven"); + REQUIRE_THROWS(importer.bind(bad)); + REQUIRE_NOTHROW(importer.bind(original)); + bad = original; + scalar = bad.mutable_relations(0) + ->mutable_root() + ->mutable_input() + ->mutable_project() + ->mutable_expressions(0) + ->mutable_scalar_function(); + auto payload = scalar->mutable_arguments(0) + ->mutable_value() + ->mutable_literal() + ->mutable_user_defined() + ->mutable_value(); + auto extended = payload->value(); + extended.append("\x10\x11", 2); // Unknown protobuf field 2, varint 17. + payload->set_value(extended); + REQUIRE_NOTHROW(importer.bind(bad)); + extended = payload->value(); + extended[1] = static_cast(static_cast(extended[1]) | 0x80); + extended.insert(extended.begin() + 2, '\0'); // Valid nonminimal length varint. + payload->set_value(extended); + REQUIRE_NOTHROW(importer.bind(bad)); + payload->set_type_url("type.googleapis.com/unrelated.Numeric"); + REQUIRE_THROWS(importer.bind(bad)); + auto type = logical_type::make_mo_decimal({128, 19, 0}, false); + REQUIRE_NOTHROW(mo_decimal::validate_signature( + mo_decimal::decimal_op::cast, {logical_type::make(type_id::BIGINT)}, type)); + REQUIRE_THROWS(mo_decimal::validate_signature(mo_decimal::decimal_op::cast, + {logical_type::make(type_id::BIGINT)}, + logical_type::make_mo_decimal({64, 18, 0}, false))); +}