Use column_device_view instead of raw pointers in Decimal Division Expression Functors

[simoneves]

Implemented column_device_view / mutable_column_device_view in the divide functors and moved null handling into the kernel.

• GPU (DecimalExpressionKernelsGpu.cu)

• Functors now take cudf::column_device_view / cudf::mutable_column_device_view instead of raw pointers.
• Device views are created on the host via ::create() before launch (same pattern as DecimalAggregationDevice.cu).
• In-kernel null rules:
  • Column–column: null if lhs or rhs is null, or rhs == 0.
  • Column–scalar: null if lhs is null, or scalar divisor is 0.
  • Scalar–column: null if rhs is null, or rhs == 0.
• Valid rows use element<T>(idx); null rows use out.set_null(idx).
• Removed the denom == 0 early return from decimalDivideImpl — callers skip zero divisors before calling it.

• Host (DecimalExpressionKernels.cpp)

• Removed host-side null-mask prep (bitmask_and, copy_bitmask) and scatterNullsAtZeroDivisor.
• Output columns are created with mask_state::ALL_VALID so the kernel can set null bits.
• Added finalizeDivideOutputNullCount() to compute null_count via cudf::null_count() after the kernel.

• Header cleanup

• Removed the public scatterNullsAtZeroDivisor API.
• Updated docs to describe in-kernel null propagation.

This drops the extra binary_operation + copy_if_else pass for zero divisors. Scalar divide-by-zero now also produces null (previously only column–column did, via the scatter step).

Does this change ensure that garbage values (for null rows) will never reach decimalDivideImpl and therefore never cause markDecimalOverflow to be called incorrectly?

Yes — for rows that should produce null output, garbage payload values no longer reach decimalDivideImpl, so they cannot spuriously set overflowFlag.

Each functor returns before calling decimalDivideImpl when:

Functor	Early-return conditions
DivideFunctor	lhs.is_null(idx), rhs.is_null(idx), or rhs == 0
DivideRhsScalarFunctor	lhs.is_null(idx) or rhsValue == 0
DivideLhsScalarFunctor	rhs.is_null(idx) or rhs == 0

Because of short-circuit evaluation, element() is not called on null inputs. For example in DivideFunctor:

  __device__ void operator()(cudf::size_type idx) const {
    if (lhs.is_null(idx) || rhs.is_null(idx) || rhs.element<InT>(idx) == 0) {
      out.set_null(idx);
      return;
    }

• Null lhs → return before reading rhs.element
• Null rhs → return before reading rhs.element
• Non-null rhs with value 0 → rhs.element is read intentionally (zero-divisor case)

So garbage stored in null slots is never fed into decimalDivideImpl.

Compared to the old approach

Previously, the host pre-merged null masks and the kernel still ran on every row via raw pointers (lhs[idx], rhs[idx]). Null rows could contain arbitrary payloads, and those values were passed into decimalDivideImpl, which could call markDecimalOverflow on garbage. The new in-kernel null guards fix that.

Caveats worth knowing

1. Valid rows with extreme values can still correctly set overflowFlag — that's intended.
2. Null output rows may have uninitialized data in the value buffer; only the null bit matters. set_null(idx) is called without writing out.element<OutT>(idx).
3. overflowFlag is global — one real overflow on any valid row fails the whole divide, but null rows no longer contribute false positives.

Bottom line: null-input and zero-divisor rows are filtered before decimalDivideImpl, so garbage in null slots cannot cause incorrect overflow reporting.

[github-actions]

Selective Build Plan

Linux release with adapters is running a full build (changes touch velox/experimental/ or velox/external/). See the CI workflows README for what this means.

---

Selective build plan