feat(numeric): add checked MO exact-decimal primitives

bench/mo_exact_decimal.cpp

@@ -0,0 +1,120 @@
+/* Copyright 2026 Sirius Contributors. SPDX-License-Identifier: Apache-2.0 */
+#include "numeric/exact_decimal_gpu.hpp"
+
+#include <cudf/binaryop.hpp>
+#include <cudf/unary.hpp>
+#include <cudf/utilities/error.hpp>
+
+#include <rmm/cuda_stream.hpp>
+#include <rmm/mr/cuda_memory_resource.hpp>
+#include <rmm/mr/pool_memory_resource.hpp>
+
+#include <chrono>
+#include <iostream>
+#include <vector>
+
+using namespace sirius::mo_decimal;
+namespace {
+constexpr cudf::size_type rows = 1 << 18;
+
+std::unique_ptr<cudf::column> input(decimal_type type,
+                                    rmm::cuda_stream_view stream,
+                                    rmm::device_async_resource_ref mr)
+{
+  auto result = make_decimal_column(type, rows, cudf::mask_state::ALL_VALID, stream, mr);
+  auto view   = result->mutable_view();
+  std::vector<uint64_t> ones(rows, 1);
+  if (type.bits == 64) {
+    CUDF_CUDA_TRY(cudaMemcpyAsync(view.data<uint64_t>(),
+                                  ones.data(),
+                                  rows * sizeof(uint64_t),
+                                  cudaMemcpyHostToDevice,
+                                  stream.value()));
+  } else {
+    for (int child = 0; child < 3; ++child)
+      CUDF_CUDA_TRY(cudaMemsetAsync(
+        view.child(child).data<uint64_t>(), 0, rows * sizeof(uint64_t), stream.value()));
+    CUDF_CUDA_TRY(cudaMemcpyAsync(view.child(3).data<uint64_t>(),
+                                  ones.data(),
+                                  rows * sizeof(uint64_t),
+                                  cudaMemcpyHostToDevice,
+                                  stream.value()));
+  }
+  stream.synchronize();
+  return result;
+}
+
+template <class Evaluate>
+void measure(char const* name, rmm::cuda_stream_view stream, Evaluate evaluate)
+{
+  // Include allocations, masks, error storage and required stream quiescence.
+  // Exclude input upload, warm-up and output destruction for every case.
+  for (int repetition = -2; repetition < 7; ++repetition) {
+    auto start  = std::chrono::steady_clock::now();
+    auto result = evaluate();
+    stream.synchronize();
+    auto ms =
+      std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now() - start).count();
+    if (repetition >= 0)
+      std::cout << name << ',' << rows << ',' << repetition << ',' << ms << ','
+                << rows * 1000.0 / ms << '\n';
+  }
+}
+}  // namespace
+
+int main()
+{
+  rmm::cuda_stream stream;
+  // Match the pooled allocation model of a task resource, and use the same
+  // bounded caller-owned pool for cuDF controls and checked operations.
+  rmm::mr::pool_memory_resource pool(
+    cuda::mr::any_resource<cuda::mr::device_accessible>{rmm::mr::cuda_memory_resource{}},
+    64u << 20,
+    128u << 20);
+  auto mr = rmm::to_device_async_resource_ref_checked(&pool);
+  decimal_type narrow{64, 18, 0}, wide{256, 65, 0};
+  auto a = input(narrow, stream.view(), mr);
+  auto b = input(narrow, stream.view(), mr);
+  auto x = input(wide, stream.view(), mr);
+  auto y = input(wide, stream.view(), mr);
+  std::cout << "case,rows,repetition,wall_ms,rows_per_second\n";
+  for (auto operation : {decimal_op::add, decimal_op::multiply}) {
+    auto out = operation == decimal_op::add ? narrow : decimal_type{128, 38, 0};
+    auto cudf_op =
+      operation == decimal_op::add ? cudf::binary_operator::ADD : cudf::binary_operator::MUL;
+    measure(operation == decimal_op::add ? "cudf_add64" : "cudf_widen_multiply64_to128",
+            stream.view(),
+            [&] {
+              auto output_type = cudf::data_type{
+                operation == decimal_op::add ? cudf::type_id::DECIMAL64 : cudf::type_id::DECIMAL128,
+                0};
+              // cuDF arithmetic runs at the operand width. Widen both operands for
+              // an equivalent multiply control; changing only the output type wraps.
+              std::unique_ptr<cudf::column> left, right;
+              if (operation == decimal_op::multiply) {
+                left  = cudf::cast(a->view(), output_type, stream.view(), mr);
+                right = cudf::cast(b->view(), output_type, stream.view(), mr);
+              }
+              auto result = cudf::binary_operation(left ? left->view() : a->view(),
+                                                   right ? right->view() : b->view(),
+                                                   cudf_op,
+                                                   output_type,
+                                                   stream.view(),
+                                                   mr);
+              stream.synchronize();
+              return result;
+            });
+    measure(operation == decimal_op::add ? "checked_add64" : "checked_multiply64_to128",
+            stream.view(),
+            [&] {
+              return evaluate_decimal_columns(
+                operation, a->view(), narrow, b->view(), narrow, out, nullptr, stream.view(), mr);
+            });
+  }
+  for (auto operation : {decimal_op::add, decimal_op::divide})
+    measure(
+      operation == decimal_op::add ? "checked_add256" : "checked_divide256", stream.view(), [&] {
+        return evaluate_decimal_columns(
+          operation, x->view(), wide, y->view(), wide, wide, nullptr, stream.view(), mr);
+      });
+}

cmake/sirius-embedding-tests.cmake

@@ -1,5 +1,20 @@
 # Embedding control tests need neither a GPU nor a DuckDB runtime. Keep them
 # independently runnable from the linked ABI and GPU integration tests.
+add_executable(
+  sirius_exact_decimal_unittest test/cpp/exec/execution_evidence_unittest.cpp
+                                test/cpp/embedding/test_exact_decimal.cpp)
+target_compile_features(sirius_exact_decimal_unittest PRIVATE cxx_std_20)
+target_include_directories(
+  sirius_exact_decimal_unittest PRIVATE "${CMAKE_SOURCE_DIR}/third_party/catch"
+                                        "${CMAKE_CURRENT_SOURCE_DIR}/src")
+
+add_executable(sirius_exact_decimal_benchmark EXCLUDE_FROM_ALL
+               bench/mo_exact_decimal.cpp)
+target_compile_features(sirius_exact_decimal_benchmark PRIVATE cxx_std_20)
+target_include_directories(sirius_exact_decimal_benchmark
+                           PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/src")
+target_link_libraries(sirius_exact_decimal_benchmark PRIVATE sirius_extension)
+
 add_executable(
   sirius_native_control_unittest
   test/cpp/exec/execution_evidence_unittest.cpp
@@ -52,6 +67,7 @@ add_executable(
   test/cpp/embedding/test_native_gpu.cpp
   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_tae_gpu.cpp
   src/embedding/c_api.cpp)
 target_compile_features(sirius_native_gpu_unittest PRIVATE cxx_std_20)

docs/mo-exact-decimal-kernels.csv

@@ -0,0 +1,43 @@
+case,rows,repetition,wall_ms,rows_per_second
+cudf_add64,262144,0,0.20445,1.28219e+09
+cudf_add64,262144,1,0.189379,1.38423e+09
+cudf_add64,262144,2,0.197139,1.32974e+09
+cudf_add64,262144,3,0.195489,1.34097e+09
+cudf_add64,262144,4,0.19587,1.33836e+09
+cudf_add64,262144,5,0.195839,1.33857e+09
+cudf_add64,262144,6,0.204669,1.28082e+09
+checked_add64,262144,0,0.02729,9.60586e+09
+checked_add64,262144,1,0.02719,9.64119e+09
+checked_add64,262144,2,0.02758,9.50486e+09
+checked_add64,262144,3,0.02637,9.94099e+09
+checked_add64,262144,4,0.033949,7.7217e+09
+checked_add64,262144,5,0.03364,7.79263e+09
+checked_add64,262144,6,0.03487,7.51775e+09
+cudf_widen_multiply64_to128,262144,0,0.502128,5.22066e+08
+cudf_widen_multiply64_to128,262144,1,0.495698,5.28838e+08
+cudf_widen_multiply64_to128,262144,2,0.499908,5.24384e+08
+cudf_widen_multiply64_to128,262144,3,0.497848,5.26554e+08
+cudf_widen_multiply64_to128,262144,4,0.499428,5.24888e+08
+cudf_widen_multiply64_to128,262144,5,0.499538,5.24773e+08
+cudf_widen_multiply64_to128,262144,6,0.507368,5.16674e+08
+checked_multiply64_to128,262144,0,0.505538,5.18545e+08
+checked_multiply64_to128,262144,1,0.500578,5.23683e+08
+checked_multiply64_to128,262144,2,0.500228,5.24049e+08
+checked_multiply64_to128,262144,3,0.495228,5.2934e+08
+checked_multiply64_to128,262144,4,0.512808,5.11193e+08
+checked_multiply64_to128,262144,5,0.498808,5.25541e+08
+checked_multiply64_to128,262144,6,0.499539,5.24772e+08
+checked_add256,262144,0,0.263789,9.93764e+08
+checked_add256,262144,1,0.262309,9.99371e+08
+checked_add256,262144,2,0.258849,1.01273e+09
+checked_add256,262144,3,0.262029,1.00044e+09
+checked_add256,262144,4,0.262509,9.9861e+08
+checked_add256,262144,5,0.262309,9.99371e+08
+checked_add256,262144,6,0.263589,9.94518e+08
+checked_divide256,262144,0,4.67064,5.61259e+07
+checked_divide256,262144,1,4.66343,5.62127e+07
+checked_divide256,262144,2,4.66613,5.61802e+07
+checked_divide256,262144,3,4.66244,5.62246e+07
+checked_divide256,262144,4,4.65689,5.62916e+07
+checked_divide256,262144,5,4.65525,5.63115e+07
+checked_divide256,262144,6,4.65287,5.63403e+07

docs/mo-exact-decimal-primitives.md

@@ -0,0 +1,118 @@
+# MO exact-decimal primitives
+
+This is the first native implementation increment for
+[MO #28968](https://github.com/matrixorigin/matrixone/issues/28968), following
+the approved numeric design revision
+`42a89f09a1d168d02b9583cb3ea7b4de6dbb5634` in MO #29449.
+
+It preserves the existing Decimal64/128 carriers and input/result ownership.
+It does not register the MO exact-decimal Substrait family, advertise its
+capability, change ordinary DuckDB decimal execution, or admit new MO queries.
+The importer is pinned to merged duckdb-substrait #3; its scoped handler is
+installed by the following native integration increment.
+
+## Representation and arithmetic
+
+`sirius::mo_decimal::decimal_type` stores physical coefficient width, precision
+and scale independently. The native wire remains little-endian signed
+two's-complement coefficients: 8, 16 and 32 bytes, with separate validity.
+Working Decimal256 precision reaches 76; public result precision stops at 65.
+In particular, Decimal256(15,2) remains physically 32 bytes.
+
+Decimal64/128 use existing cuDF fixed-point columns. Decimal256 uses a STRUCT
+of signed high64 followed by three unsigned limbs in descending significance;
+validity belongs to the parent, not to individual limbs. Input decoding and
+result interleaving reuse the existing admitted storage and scratch slab.
+
+The checked host/device core supports arithmetic, comparisons, proven-safe
+negation and checked normal rescaling. It uses fixed 512-bit scratch, never
+per-row allocation, floating point or numeric strings. Multiplication rescales
+and rounds before checking the final physical and declared domains. Division
+consumes the supplied result scale; it does not derive MO types or session
+settings. Scale reduction and division round half away from zero; DIV and
+modulo truncate. NULL, an inactive execution mask and SELECT zero divisors
+produce NULL without arithmetic failure.
+
+Scalar failure is `decimal_error::out_of_range`; checked cast failure is
+`decimal_error::invalid_input`. These are internal errors here, not new public
+C statuses. The native integration increment maps the operation-owned status
+without parsing messages or inferring from the enclosing aggregate.
+
+## Ownership and use
+
+`evaluate_decimal_columns` takes the caller's task stream and reservation-aware
+RMM resource. Operand and mask owners must outlive that stream's work. Result
+values and the bounded per-row error vector are owned together. Failure first
+proves stream quiescence; an unprovable owner is retained and reported through
+the existing fatal GPU path. No device-wide synchronization or independent
+allocator/pool is installed. The pinned cuDF API requires an explicit result
+null count, computed on the supplied stream.
+
+## Local validation
+
+Use the frozen Sirius MO Pixi environment in the existing checkout:
+
+```sh
+pixi run --frozen -e mo mo-configure-embedding-sdk
+pixi run --frozen -e mo cmake --build build/mo --target \
+  sirius_exact_decimal_unittest sirius_native_control_unittest \
+  sirius_native_gpu_unittest -j8
+pixi run --frozen -e mo build/mo/extension/sirius/sirius_exact_decimal_unittest
+pixi run --frozen -e mo build/mo/extension/sirius/sirius_native_control_unittest
+pixi run --frozen -e mo build/mo/extension/sirius/sirius_native_gpu_unittest \
+  '[exact_decimal_gpu],[native_gpu],[native_result_gpu]'
+```
+
+The CPU suite includes independent integer/rational boundary vectors, not
+expectations computed by the native implementation. GPU tests check those same
+independent full-width results, masks, 64/128-bit controls, sliced wide output,
+and constant/NULL/retry input decoding. The existing native suite retains its
+1/2/4-worker checks. Small native tests are not SF10/all-22 correctness or
+performance evidence. Full capability, MO lowering and public/resource/timing
+gates remain required before default cutover.
+
+## Kernel measurements
+
+The standalone benchmark uses 262,144 rows, two warm-ups and seven measured
+repetitions per operation. One caller-owned RMM pool starts at 64 MiB and is
+capped at 128 MiB. Both the cuDF controls and checked kernels use that pool
+and one task stream. Timing includes output/error/mask allocation and stream
+quiescence, and excludes input upload, warm-ups and output destruction.
+
+```sh
+pixi run --frozen -e mo cmake --build build/mo \
+  --target sirius_exact_decimal_benchmark -j8
+pixi run --frozen -e mo build/mo/extension/sirius/sirius_exact_decimal_benchmark \
+  > decimal-kernels.csv
+```
+
+The cuDF multiply control widens both operands to Decimal128 first. Merely
+requesting a Decimal128 output still computes at Decimal64 width in the
+pinned cuDF and can wrap; the signed physical-domain test uses an independent
+host `__int128` oracle. The add control measures cuDF's unchecked arithmetic
+on the same inputs. Decimal256 has no equivalent cuDF decimal carrier.
+
+The CSV retains every sample. These kernel measurements establish the cost
+of these primitives; all-22 numeric query timing and the rollout regression
+gate belong to the native plan integration and public campaign.
+
+### Local snapshot
+
+[Raw repetitions](mo-exact-decimal-kernels.csv), measured on 2026-10-08 from
+this increment based on Sirius `e2e2f08f9fd1eaa1253297493df0ab11a6078664`:
+RTX 3070, NVIDIA driver 615.71.09, CUDA 13.3.73, GCC 14.4.0, frozen Pixi `mo`
+environment, DuckDB `069cc9f9b5be802405797faecc284961b07c70ef`, importer
+`dd4cab14b82754ca919633913436cd530f496e00`. Pixi lock SHA-256:
+`ed808ec91232d916b49e11abb14aba0dcdd82fa13d908342bc970ce584ce4fef`.
+
+| Operation | Median wall time (ms) |
+| --- | ---: |
+| cuDF add64, unchecked | 0.195870 |
+| Checked add64 | 0.027580 |
+| cuDF widened multiply64-to128 | 0.499538 |
+| Checked multiply64-to128 | 0.500228 |
+| Checked add256 | 0.262309 |
+| Checked divide256 | 4.662440 |
+
+The checked multiply differs by approximately 0.14% in this sample. Wide
+operations report absolute throughput because cuDF has no equivalent carrier.

src/CMakeLists.txt

@@ -27,6 +27,7 @@ add_subdirectory(io)
 add_subdirectory(late_mat)
 add_subdirectory(log)
 add_subdirectory(memory)
+add_subdirectory(numeric)
 add_subdirectory(op)
 add_subdirectory(parallel)
 add_subdirectory(pipeline)