From a10533343ed497d97ec2762eb2c2f1d1cbeb8932 Mon Sep 17 00:00:00 2001 From: Simon Eves Date: Tue, 7 Jul 2026 15:23:08 -0700 Subject: [PATCH] Implemented `column_device_view` / `mutable_column_device_view` in the divide functors and moved null handling into the kernel. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * 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(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). --- .../expression/DecimalExpressionKernels.cpp | 83 ++++--------------- .../expression/DecimalExpressionKernels.h | 34 ++------ .../expression/DecimalExpressionKernelsGpu.cu | 71 +++++++++------- .../expression/DecimalExpressionKernelsGpu.h | 4 +- 4 files changed, 65 insertions(+), 127 deletions(-) diff --git a/velox/experimental/cudf/expression/DecimalExpressionKernels.cpp b/velox/experimental/cudf/expression/DecimalExpressionKernels.cpp index ed9ff82f94c..591a759915b 100644 --- a/velox/experimental/cudf/expression/DecimalExpressionKernels.cpp +++ b/velox/experimental/cudf/expression/DecimalExpressionKernels.cpp @@ -21,13 +21,10 @@ #include "velox/type/DecimalUtil.h" #include "velox/type/Type.h" -#include #include -#include #include #include #include -#include namespace facebook::velox::cudf_velox { namespace { @@ -77,52 +74,19 @@ void checkDecimalDivideTypes( } } -} // namespace - -// Scatters null values to positions where the divisor is zero. -// Returns a new column with nulls at zero-divisor positions. -std::unique_ptr scatterNullsAtZeroDivisor( - std::unique_ptr result, - const cudf::column_view& divisor, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) { - // Create zero scalar for comparison and null scalar for scattering. - std::unique_ptr zeroScalar; - std::unique_ptr nullScalar; - auto divisorScale = numeric::scale_type{divisor.type().scale()}; - auto outputScale = numeric::scale_type{result->type().scale()}; - - if (divisor.type().id() == cudf::type_id::DECIMAL64) { - zeroScalar = cudf::make_fixed_point_scalar( - 0, divisorScale, stream, mr); - nullScalar = cudf::make_fixed_point_scalar( - 0, outputScale, stream, mr); - } else if (divisor.type().id() == cudf::type_id::DECIMAL128) { - zeroScalar = cudf::make_fixed_point_scalar( - 0, divisorScale, stream, mr); - nullScalar = cudf::make_fixed_point_scalar( - 0, outputScale, stream, mr); - } else { - VELOX_FAIL( - "Unsupported decimal type {} for division", - cudf::ast::type_id_to_string(divisor.type().id())); +void finalizeDivideOutputNullCount( + cudf::column& out, + rmm::cuda_stream_view stream) { + if (out.size() == 0) { + return; } - nullScalar->set_valid_async(false, stream); - - // Create boolean column: TRUE where divisor == 0, FALSE otherwise. - auto divisorIsZero = cudf::binary_operation( - divisor, - *zeroScalar, - cudf::binary_operator::EQUAL, - cudf::data_type{cudf::type_id::BOOL8}, - stream, - mr); - - // Scatter nulls where divisor is zero. - return cudf::copy_if_else( - *nullScalar, *result, divisorIsZero->view(), stream, mr); + auto const nullCount = + cudf::null_count(out.view().null_mask(), 0, out.size(), stream); + out.set_null_count(nullCount); } +} // namespace + std::unique_ptr decimalDivide( const cudf::column_view& lhs, const cudf::column_view& rhs, @@ -147,13 +111,8 @@ std::unique_ptr decimalDivide( const auto outType = outputType.id(); checkDecimalDivideTypes(inType, outType); - // Combine input null masks (lhs and rhs nulls). - auto [nullMask, nullCount] = - cudf::bitmask_and(cudf::table_view({lhs, rhs}), stream, mr); - - // Create output column with input null mask and perform division. auto out = cudf::make_fixed_width_column( - outputType, lhs.size(), std::move(nullMask), nullCount, stream, mr); + outputType, lhs.size(), cudf::mask_state::ALL_VALID, stream, mr); const __int128_t rescaleFactor = DecimalUtil::kPowersOfTen[aRescale]; VELOX_USER_CHECK( @@ -167,8 +126,8 @@ std::unique_ptr decimalDivide( stream), "Decimal overflow"); - // Scatter nulls where divisor is zero. - return scatterNullsAtZeroDivisor(std::move(out), rhs, stream, mr); + finalizeDivideOutputNullCount(*out, stream); + return out; } std::unique_ptr decimalDivide( @@ -189,10 +148,8 @@ std::unique_ptr decimalDivide( return makeAllNullDecimalColumn(outputType, lhs.size(), stream, mr); } - auto nullMask = cudf::copy_bitmask(lhs, stream, mr); - auto nullCount = lhs.null_count(); auto out = cudf::make_fixed_width_column( - outputType, lhs.size(), std::move(nullMask), nullCount, stream, mr); + outputType, lhs.size(), cudf::mask_state::ALL_VALID, stream, mr); auto rhsValue = getDecimalScalarValue(rhs, stream); @@ -211,6 +168,7 @@ std::unique_ptr decimalDivide( stream), "Decimal overflow"); + finalizeDivideOutputNullCount(*out, stream); return out; } @@ -232,13 +190,8 @@ std::unique_ptr decimalDivide( return makeAllNullDecimalColumn(outputType, rhs.size(), stream, mr); } - // Copy rhs null mask. - auto nullMask = cudf::copy_bitmask(rhs, stream, mr); - auto nullCount = rhs.null_count(); - - // Create output column and perform division. auto out = cudf::make_fixed_width_column( - outputType, rhs.size(), std::move(nullMask), nullCount, stream, mr); + outputType, rhs.size(), cudf::mask_state::ALL_VALID, stream, mr); auto lhsValue = getDecimalScalarValue(lhs, stream); @@ -258,8 +211,8 @@ std::unique_ptr decimalDivide( stream), "Decimal overflow"); - // Scatter nulls where divisor is zero. - return scatterNullsAtZeroDivisor(std::move(out), rhs, stream, mr); + finalizeDivideOutputNullCount(*out, stream); + return out; } } // namespace facebook::velox::cudf_velox diff --git a/velox/experimental/cudf/expression/DecimalExpressionKernels.h b/velox/experimental/cudf/expression/DecimalExpressionKernels.h index 0fe37053f99..f38d620a09b 100644 --- a/velox/experimental/cudf/expression/DecimalExpressionKernels.h +++ b/velox/experimental/cudf/expression/DecimalExpressionKernels.h @@ -29,9 +29,8 @@ namespace facebook::velox::cudf_velox { /** * @brief Element-wise decimal division of two columns. * - * Builds the output null mask as the bitwise AND of lhs and rhs validity, runs - * the GPU divide into outputType, and applies scatterNullsAtZeroDivisor so - * rows with a zero divisor are null. + * Runs the GPU divide into outputType. The kernel propagates input nulls and + * nulls rows where the divisor is zero. * * @param lhs Left-hand decimal operand column (DECIMAL64 or DECIMAL128). * @param rhs Right-hand decimal operand column (same type as lhs). @@ -54,8 +53,8 @@ std::unique_ptr decimalDivide( * @brief Element-wise decimal division of a column by a scalar. * * If the scalar is invalid, returns an all-null column of outputType; - * otherwise copies lhs nulls and divides without zero-divisor scattering (rhs - * is not per-row). + * otherwise the kernel propagates lhs nulls and nulls rows when dividing by + * zero. * * @param lhs Left-hand decimal operand column. * @param rhs Right-hand decimal operand scalar. @@ -77,9 +76,8 @@ std::unique_ptr decimalDivide( /** * @brief Element-wise decimal division of a scalar by a column. * - * Invalid lhs yields all-null output; otherwise rhs nulls are propagated, then - * divide and scatterNullsAtZeroDivisor on rhs so division-by-zero rows are - * null. + * Invalid lhs yields all-null output; otherwise the kernel propagates rhs + * nulls and nulls rows where the divisor is zero. * * @param lhs Left-hand decimal operand scalar. * @param rhs Right-hand decimal operand column. @@ -98,24 +96,4 @@ std::unique_ptr decimalDivide( rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); -/** - * @brief Nulls output rows where the divisor column equals zero. - * - * After a decimal divide, forces output rows to null where the divisor column - * compares equal to zero (DECIMAL64 or DECIMAL128), using copy_if_else. Kept in - * the .cpp translation unit so it can use Velox checks alongside cuDF APIs - * without pulling those into CUDA compilation units. - * - * @param result Column produced by decimal division. - * @param divisor Per-row divisor column used to detect division by zero. - * @param stream CUDA stream for GPU execution. - * @param mr Memory resource for output allocation. - * @return Column with nulls scattered at zero-divisor rows. - */ -std::unique_ptr scatterNullsAtZeroDivisor( - std::unique_ptr result, - const cudf::column_view& divisor, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr); - } // namespace facebook::velox::cudf_velox diff --git a/velox/experimental/cudf/expression/DecimalExpressionKernelsGpu.cu b/velox/experimental/cudf/expression/DecimalExpressionKernelsGpu.cu index 95ba18f8643..2537442b943 100644 --- a/velox/experimental/cudf/expression/DecimalExpressionKernelsGpu.cu +++ b/velox/experimental/cudf/expression/DecimalExpressionKernelsGpu.cu @@ -16,6 +16,7 @@ #include "velox/experimental/cudf/expression/DecimalExpressionKernelsGpu.h" +#include #include #include @@ -101,10 +102,6 @@ __device__ OutT decimalDivideImpl( __int128_t denom, __int128_t rescaleFactor, int32_t* overflowFlag) { - if (denom == 0) { - return OutT{0}; - } - bool negative = false; unsigned __int128 const uNum = absToUnsigned(numerator, negative); unsigned __int128 const uDenom = absToUnsigned(denom, negative); @@ -151,43 +148,58 @@ __device__ OutT decimalDivideImpl( template struct DivideFunctor { - const InT* lhs; - const InT* rhs; - OutT* out; + cudf::column_device_view lhs; + cudf::column_device_view rhs; + cudf::mutable_column_device_view out; __int128_t rescaleFactor; int32_t* overflowFlag; __device__ void operator()(cudf::size_type idx) const { - out[idx] = decimalDivideImpl( - lhs[idx], rhs[idx], rescaleFactor, overflowFlag); + if (lhs.is_null(idx) || rhs.is_null(idx) || rhs.element(idx) == 0) { + out.set_null(idx); + return; + } + out.element(idx) = decimalDivideImpl( + lhs.element(idx), + rhs.element(idx), + rescaleFactor, + overflowFlag); } }; template struct DivideLhsScalarFunctor { __int128_t lhsValue; - const InColT* rhs; - OutT* out; + cudf::column_device_view rhs; + cudf::mutable_column_device_view out; __int128_t rescaleFactor; int32_t* overflowFlag; __device__ void operator()(cudf::size_type idx) const { - out[idx] = decimalDivideImpl( - lhsValue, rhs[idx], rescaleFactor, overflowFlag); + if (rhs.is_null(idx) || rhs.element(idx) == 0) { + out.set_null(idx); + return; + } + out.element(idx) = decimalDivideImpl( + lhsValue, rhs.element(idx), rescaleFactor, overflowFlag); } }; template struct DivideRhsScalarFunctor { - const InColT* lhs; + cudf::column_device_view lhs; __int128_t rhsValue; - OutT* out; + cudf::mutable_column_device_view out; __int128_t rescaleFactor; int32_t* overflowFlag; __device__ void operator()(cudf::size_type idx) const { - out[idx] = decimalDivideImpl( - lhs[idx], rhsValue, rescaleFactor, overflowFlag); + if (lhs.is_null(idx) || rhsValue == 0) { + out.set_null(idx); + return; + } + out.element(idx) = decimalDivideImpl( + lhs.element(idx), rhsValue, rescaleFactor, overflowFlag); } }; @@ -228,15 +240,14 @@ struct divideColumnColumnKernel { template requires ValidDecimalDivideStorageTypes bool operator()() const { + auto lhsDev = cudf::column_device_view::create(lhs, stream); + auto rhsDev = cudf::column_device_view::create(rhs, stream); + auto outDev = cudf::mutable_column_device_view::create(out, stream); return launchDecimalDivide( lhs.size(), [&](int32_t* overflowFlag) { return DivideFunctor{ - lhs.data(), - rhs.data(), - out.data(), - rescaleFactor, - overflowFlag}; + *lhsDev, *rhsDev, *outDev, rescaleFactor, overflowFlag}; }, stream); } @@ -259,15 +270,13 @@ struct divideColumnScalarKernel { template requires ValidDecimalDivideStorageTypes bool operator()() const { + auto lhsDev = cudf::column_device_view::create(lhs, stream); + auto outDev = cudf::mutable_column_device_view::create(out, stream); return launchDecimalDivide( lhs.size(), [&](int32_t* overflowFlag) { return DivideRhsScalarFunctor{ - lhs.data(), - rhsValue, - out.data(), - rescaleFactor, - overflowFlag}; + *lhsDev, rhsValue, *outDev, rescaleFactor, overflowFlag}; }, stream); } @@ -290,15 +299,13 @@ struct divideScalarColumnKernel { template requires ValidDecimalDivideStorageTypes bool operator()() const { + auto rhsDev = cudf::column_device_view::create(rhs, stream); + auto outDev = cudf::mutable_column_device_view::create(out, stream); return launchDecimalDivide( rhs.size(), [&](int32_t* overflowFlag) { return DivideLhsScalarFunctor{ - lhsValue, - rhs.data(), - out.data(), - rescaleFactor, - overflowFlag}; + lhsValue, *rhsDev, *outDev, rescaleFactor, overflowFlag}; }, stream); } diff --git a/velox/experimental/cudf/expression/DecimalExpressionKernelsGpu.h b/velox/experimental/cudf/expression/DecimalExpressionKernelsGpu.h index 6f72f8cf598..d0b1aa6b93f 100644 --- a/velox/experimental/cudf/expression/DecimalExpressionKernelsGpu.h +++ b/velox/experimental/cudf/expression/DecimalExpressionKernelsGpu.h @@ -28,8 +28,8 @@ namespace facebook::velox::cudf_velox::detail { * @brief Dispatches a per-row device loop for fixed-point decimal division. * * Computes (lhs * rescaleFactor) / rhs with half-away-from-zero rounding on the - * remainder, writing into out. Zero divisors produce a numeric zero in out - * (callers patch nulls). + * remainder, writing into out. Input nulls and zero divisors are written as + * null in the output null mask. * * @param inType DECIMAL64 or DECIMAL128 input storage width selector. * @param outType DECIMAL64 or DECIMAL128 output storage width selector.