[FEA] ANSI SQL Operator JIT Support (2) : Implement Operator Library - #22514
Conversation
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
📝 WalkthroughWalkthroughAdds a device-side operator library under cudf::ops: error codes, type utilities, promotion helpers, and many CUDA device operator headers implementing arithmetic, ANSI-arithmetic, bitwise, comparison, logic, math, casts, trigonometric, and null-handling functions with optional/null semantics. ChangesDevice Operator Library
🎯 4 (Complex) | ⏱️ ~60 minutes
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
There was a problem hiding this comment.
Actionable comments posted: 6
🧹 Nitpick comments (6)
cpp/include/cudf/ast/expressions.hpp (1)
517-527: 💤 Low valueMinor documentation style inconsistency.
The Doxygen comment at line 517 uses
///style without a proper@brieftag, while other class documentation in this file uses multi-line/** ... */format with explicit tags. For consistency with the rest of the file (e.g.,generic_scalar_device_view,literal), consider using the standard format.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/include/cudf/ast/expressions.hpp` around lines 517 - 527, Replace the triple-slash doc comment for the predicate class with the project's standard Doxygen block style: convert the leading `///` comment into a `/** ... */` block containing an explicit `@brief` (and any other relevant tags such as `@internal` or a short note that this is for internal filter operations) so it matches the formatting used by other classes like `generic_scalar_device_view` and `literal`; update the comment above the `class predicate : public expression` declaration and keep the existing description and param notes (reference symbol: predicate, constructor predicate(expression const& source)).cpp/src/jit/row_ir.hpp (1)
9-11: ⚖️ Poor tradeoffConsider moving factory includes to implementation file.
Including
<cudf/column/column_factories.hpp>and<cudf/scalar/scalar_factories.hpp>in this header adds compilation overhead for all translation units that includerow_ir.hpp. Since these are only used byinstance_context::add_input(scalar const&)which is defined inline at lines 139-143, consider either:
- Moving that method implementation to the .cpp file, or
- Forward-declaring the necessary types if possible
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/src/jit/row_ir.hpp` around lines 9 - 11, The header currently pulls in heavy factory headers for cudf used only by the inline method instance_context::add_input(scalar const&); to reduce compile overhead, move the implementation of instance_context::add_input(scalar const&) out of the header into the corresponding .cpp and add `#include` <cudf/column/column_factories.hpp> and `#include` <cudf/scalar/scalar_factories.hpp> to that .cpp (or alternatively forward-declare the minimal cudf types if feasible); update the header to keep only the declaration of instance_context::add_input(scalar const&) and ensure the .cpp defines the method with the required includes.cpp/src/stream_compaction/filter/filter.cu (1)
63-86: 💤 Low valueMissing
CUDF_FUNC_RANGE()in public function.The public
cudf::filterfunction should includeCUDF_FUNC_RANGE()before delegating to detail functions for proper NVTX profiling, consistent withfilter_extendedat line 99.Suggested fix
std::unique_ptr<table> filter(table_view const& predicate_table, ast::expression const& predicate_expr, table_view const& filter_table, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { + CUDF_FUNC_RANGE(); auto args = cudf::detail::row_ir::ast_converter::filter(cudf::detail::row_ir::target::CUDA,🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/src/stream_compaction/filter/filter.cu` around lines 63 - 86, The public filter function is missing a CUDF_FUNC_RANGE() NVTX scope; add a CUDF_FUNC_RANGE() call at the start of the filter(table_view const& predicate_table, ast::expression const& predicate_expr, table_view const& filter_table, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) function (same placement as in filter_extended) before performing the ast_converter call and delegating to detail::filter so profiling/stack traces include this public API call.cpp/src/join/filter_join_indices_jit.cu (1)
51-54: 💤 Low valueParameter should be
std::spanof const elements for const-correctness.The
table_sourcesparameter is not modified by these functions, so it should bestd::span<std::optional<int32_t> const>rather thanstd::span<std::optional<int32_t>>.Suggested fix
jitify2::StringVec build_join_filter_template_params( std::span<transform_input const> inputs, - std::span<std::optional<int32_t>> table_sources, + std::span<std::optional<int32_t> const> table_sources, null_aware is_null_aware)jitify2::ConfiguredKernel build_join_filter_kernel(std::string const& predicate_code, std::span<transform_input const> inputs, - std::span<std::optional<int32_t>> table_sources, + std::span<std::optional<int32_t> const> table_sources, bool is_ptx,Also applies to: 91-94
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/src/join/filter_join_indices_jit.cu` around lines 51 - 54, Change the parameter type for table_sources to a span of const optionals to preserve const-correctness: update the signature of build_join_filter_template_params to take std::span<std::optional<int32_t> const> (instead of std::span<std::optional<int32_t>>) and do the same for the other overload/variant in this file that accepts table_sources; adjust any related variable names or usages to compile with the new const-qualified span but do not modify the function body logic.cpp/include/cudf/operators/error.hpp (1)
13-13: ⚡ Quick winAdd Doxygen documentation for the public enum.
The
errcenum is a public API type but lacks Doxygen documentation. As per coding guidelines, all public API entities incpp/include/cudf/should have Doxygen documentation tags.📝 Suggested documentation
+/** + * `@brief` Error codes returned by device operators. + */ enum errc : int { OK = 0, OVERFLOW = 1, DIVISION_BY_ZERO = 2 };As per coding guidelines: "Add Doxygen documentation tags (
@param,@return,@throw,@tparam,@brief) on all public API functions" (applies to public types as well).🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/include/cudf/operators/error.hpp` at line 13, The public enum errc (OK, OVERFLOW, DIVISION_BY_ZERO) is missing Doxygen documentation; add a Doxygen comment block (/** ... */) immediately above the enum declaration that includes a `@brief` describing the purpose of errc and short descriptions for each enumerator (OK, OVERFLOW, DIVISION_BY_ZERO) so tools and users can understand the API; ensure the block uses standard tags (e.g., `@brief` and either `@enum` or inline descriptions for each enumerator) and follows the project's Doxygen style for public API types.cpp/include/cudf/operators/casts.cuh (1)
471-471: Address the TODO comment or track as an issue.The TODO indicates that casting from int and float to decimal32 is not yet implemented.
Do you want me to open a new issue to track this missing functionality, or is it already being handled elsewhere?
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/include/cudf/operators/casts.cuh` at line 471, Resolve the TODO for CAST_TO_DEC32: either implement casting logic from int and float to decimal32 inside the CAST_TO_DEC32 specialization in cudf/operators/casts.cuh (add conversion rules, handle scale/precision, overflow/rounding behavior consistent with other decimal casts, and add unit tests) or remove the TODO and open a tracked issue referencing CAST_TO_DEC32 and this file describing the missing int/float→decimal32 work and expected behavior; update the comment to point to the issue if you choose the latter.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@cpp/include/cudf/operators/ansi_arithmetic.cuh`:
- Around line 622-636: Add a doc comment for the optional overload of ansi_abs
to match the other functions in the file: document the template<typename T>
__device__ inline errc ansi_abs(optional<T>* out, optional<T> const* a)
behavior, parameters (out and a), and return value (errc::OK or propagated
error) and note that nullopt is written to out when input has no value or an
error occurs; follow the same style/format used by the other optional overloads
in this file.
In `@cpp/include/cudf/operators/casts.cuh`:
- Around line 610-623: The optional overload of rescale is calling rescale(&r,
&a->value(), new_scale->value()) with new_scale->value() (int32_t const&) but
the called rescale expects an int32_t const*; fix by passing the address of the
scale value instead (i.e., pass &new_scale->value()) so the call matches the
signature; update the call in the rescale(optional<decimal<R>>*,
optional<decimal<R>> const*, optional<int32_t> const*) function that currently
invokes rescale(&r, &a->value(), new_scale->value()) to use &new_scale->value()
as the third argument.
In `@cpp/include/cudf/operators/logic.cuh`:
- Around line 244-255: The if_else(optional<T>* out, ...) implementation invokes
out->value() before out is engaged, causing undefined behavior; fix by creating
a local T result variable (e.g., T result), call the internal
if_else<T>(&result, &true_value->value(), &false_value->value(), &pred->value())
when the optionals are engaged, then assign the result into out via
out->emplace(result) or *out = result; keep the else branch setting *out =
nullopt and return errc::OK.
In `@cpp/include/cudf/operators/math.cuh`:
- Around line 96-109: The ceil and floor implementations use detail::ipow10 with
a->scale(), but fixed_point::scale() is negative for decimal places; update both
functions (template <typename R> __device__ inline errc ceil(decimal<R>* out,
decimal<R> const* a) and the corresponding floor implementation) to pass the
negated scale to ipow10 (i.e., use detail::ipow10(-a->scale())) so the
multiplier factor is computed correctly for positive exponents and
division/remainder logic remains valid for decimal<R>.
In `@cpp/include/cudf/operators/types.cuh`:
- Around line 41-57: Document that ipow10 requires a non-negative exponent by
adding a clear precondition comment above the template for ipow10(T exponent)
stating exponent must be >= 0 (or change the API to take an unsigned type). Also
consider adding a runtime check/assert (e.g., assert(exponent >= 0)) inside
ipow10 or switching the template parameter to an unsigned integer to enforce the
precondition; reference the ipow10 function to locate the change.
In `@cpp/src/join/jit/filter_join_kernel.cuh`:
- Around line 21-35: The Doxygen comment above the template function
filter_join_kernel is out of sync with its signature: replace or update the
`@param` entries that mention left_table, right_table, and scalars with correct
descriptions for the actual parameters (num_rows, left_indices, right_indices,
columns, predicate_results, and user_data) and keep any notes about the template
parameters (has_user_data, is_null_aware, Accessors) if relevant; specifically,
document num_rows as the number of rows to process, left_indices/right_indices
as device spans of indices, columns as the device column view array,
predicate_results as the output boolean array, and user_data as optional user
data for the predicate. Ensure parameter names in the comment exactly match the
function parameter names (filter_join_kernel signature) to avoid confusion.
---
Nitpick comments:
In `@cpp/include/cudf/ast/expressions.hpp`:
- Around line 517-527: Replace the triple-slash doc comment for the predicate
class with the project's standard Doxygen block style: convert the leading `///`
comment into a `/** ... */` block containing an explicit `@brief` (and any other
relevant tags such as `@internal` or a short note that this is for internal
filter operations) so it matches the formatting used by other classes like
`generic_scalar_device_view` and `literal`; update the comment above the `class
predicate : public expression` declaration and keep the existing description and
param notes (reference symbol: predicate, constructor predicate(expression
const& source)).
In `@cpp/include/cudf/operators/casts.cuh`:
- Line 471: Resolve the TODO for CAST_TO_DEC32: either implement casting logic
from int and float to decimal32 inside the CAST_TO_DEC32 specialization in
cudf/operators/casts.cuh (add conversion rules, handle scale/precision,
overflow/rounding behavior consistent with other decimal casts, and add unit
tests) or remove the TODO and open a tracked issue referencing CAST_TO_DEC32 and
this file describing the missing int/float→decimal32 work and expected behavior;
update the comment to point to the issue if you choose the latter.
In `@cpp/include/cudf/operators/error.hpp`:
- Line 13: The public enum errc (OK, OVERFLOW, DIVISION_BY_ZERO) is missing
Doxygen documentation; add a Doxygen comment block (/** ... */) immediately
above the enum declaration that includes a `@brief` describing the purpose of errc
and short descriptions for each enumerator (OK, OVERFLOW, DIVISION_BY_ZERO) so
tools and users can understand the API; ensure the block uses standard tags
(e.g., `@brief` and either `@enum` or inline descriptions for each enumerator) and
follows the project's Doxygen style for public API types.
In `@cpp/src/jit/row_ir.hpp`:
- Around line 9-11: The header currently pulls in heavy factory headers for cudf
used only by the inline method instance_context::add_input(scalar const&); to
reduce compile overhead, move the implementation of
instance_context::add_input(scalar const&) out of the header into the
corresponding .cpp and add `#include` <cudf/column/column_factories.hpp> and
`#include` <cudf/scalar/scalar_factories.hpp> to that .cpp (or alternatively
forward-declare the minimal cudf types if feasible); update the header to keep
only the declaration of instance_context::add_input(scalar const&) and ensure
the .cpp defines the method with the required includes.
In `@cpp/src/join/filter_join_indices_jit.cu`:
- Around line 51-54: Change the parameter type for table_sources to a span of
const optionals to preserve const-correctness: update the signature of
build_join_filter_template_params to take std::span<std::optional<int32_t>
const> (instead of std::span<std::optional<int32_t>>) and do the same for the
other overload/variant in this file that accepts table_sources; adjust any
related variable names or usages to compile with the new const-qualified span
but do not modify the function body logic.
In `@cpp/src/stream_compaction/filter/filter.cu`:
- Around line 63-86: The public filter function is missing a CUDF_FUNC_RANGE()
NVTX scope; add a CUDF_FUNC_RANGE() call at the start of the filter(table_view
const& predicate_table, ast::expression const& predicate_expr, table_view const&
filter_table, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr)
function (same placement as in filter_extended) before performing the
ast_converter call and delegating to detail::filter so profiling/stack traces
include this public API call.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 6cb67a67-631a-4678-88d6-2acbe2289ca7
📒 Files selected for processing (25)
cpp/include/cudf/ast/detail/operator_functor.cuhcpp/include/cudf/ast/expressions.hppcpp/include/cudf/operators/ansi_arithmetic.cuhcpp/include/cudf/operators/arithmetic.cuhcpp/include/cudf/operators/bitwise.cuhcpp/include/cudf/operators/casts.cuhcpp/include/cudf/operators/comparison.cuhcpp/include/cudf/operators/detail/promote.cuhcpp/include/cudf/operators/error.hppcpp/include/cudf/operators/logic.cuhcpp/include/cudf/operators/math.cuhcpp/include/cudf/operators/null_handling.cuhcpp/include/cudf/operators/trigonometric.cuhcpp/include/cudf/operators/types.cuhcpp/src/ast/expressions.cppcpp/src/jit/column_accessor.cuhcpp/src/jit/join_column_accessor.cuhcpp/src/jit/row_ir.cppcpp/src/jit/row_ir.hppcpp/src/join/filter_join_indices_jit.cucpp/src/join/jit/filter_join_kernel.cucpp/src/join/jit/filter_join_kernel.cuhcpp/src/stream_compaction/filter/filter.cucpp/src/transform/transform.cucpp/tests/jit/row_ir.cpp
💤 Files with no reviewable changes (1)
- cpp/src/jit/join_column_accessor.cuh
…rrr/cudf into ansi-jit-1--refactor-row-ir
Co-authored-by: Bradley Dice <bdice@bradleydice.com>
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@cpp/include/cudf/operators/ansi_arithmetic.cuh`:
- Around line 622-629: The Doxygen block for the ANSI absolute-value operator
uses backticks around commands which breaks parsing; edit the comment above the
function that "Computes absolute value for optional inputs with ANSI overflow
checks" (the docblock containing `@brief`, `@tparam`, `@param`, `@return`) and
remove all surrounding backticks so the tags are plain `@brief`, `@tparam`, `@param`,
and `@return` to match the rest of the file.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: b457725b-4b7e-40f5-9d0e-a14c47f85dd3
📒 Files selected for processing (4)
cpp/include/cudf/operators/ansi_arithmetic.cuhcpp/include/cudf/operators/casts.cuhcpp/include/cudf/operators/logic.cuhcpp/include/cudf/operators/types.cuh
It duplicates the operators in AST cc: @PointKernel @wence- |
…ison requirements
…ved type flexibility and remove redundant comments
…mpatibility, and add value_or_null method for improved error handling
There was a problem hiding this comment.
Pull request overview
Implements an initial reusable device-side operator library (cudf::ops) and refactors the AST/JIT operator functor plumbing to delegate operator implementations to that library, enabling reuse by JIT codegen paths.
Changes:
- Introduces
cudf::opsoperator headers (arithmetic, casts, comparisons, logic, math, trigonometric, null-handling) plus an error/result type for operators. - Refactors
cudf::ast::detail::operator_functorto map AST operators tocudf::ops::*implementations and updates null-handling behavior accordingly. - Updates ROW IR codegen and tests to use the new
operator_functor<op>form and the newcudf::ops::predicate.
Reviewed changes
Copilot reviewed 21 out of 21 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
| cpp/tests/jit/row_ir.cpp | Updates expected generated CUDA code strings for new operator functor instantiation and predicate symbol. |
| cpp/src/jit/row_ir.cpp | Emits cudf::ops::predicate and switches generated operator calls to operator_functor<op> (no has_nulls template arg). |
| cpp/src/ast/operators.cpp | Updates AST operator arity/type dispatch to use the new operator_functor<op> template signature. |
| cpp/include/cudf/operators/types.cuh | Adds cudf::ops::optional alias for device-side optional values. |
| cpp/include/cudf/operators/{arithmetic,ansi_arithmetic,bitwise,casts,comparison,concepts,error,identity,logic,math,null_handling,trigonometric}.cuh/.hpp | Adds the operator library implementations, concepts, and error/result types. |
| cpp/include/cudf/operators/all.cuh | Aggregates the operator library includes. |
| cpp/include/cudf/fixed_point/fixed_point.hpp | Adds numeric::decimal<Rep> alias for base-10 fixed_point types. |
| cpp/include/cudf/ast/detail/operators.cuh | Updates operator validity traits to use requires-based checks. |
| cpp/include/cudf/ast/detail/operator_functor.cuh | Replaces prior per-operator implementations with a mapping layer to cudf::ops functions plus optional/null propagation handling. |
| cpp/include/cudf/ast/detail/expression_evaluator.cuh | Switches evaluator dispatch/type deduction to the new operator_functor<op> form. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
PointKernel
left a comment
There was a problem hiding this comment.
Some final pass cleanups.
This is a really nice simplification of operator_functor.cuh by eliminating the per-operator specializations. Great work!
As a follow-up, @lamarrr could you open an issue to track some additional cleanup work I'd like to land? The same operator logic still exists in two other places within cudf:
FloorDiv,PMod, comparisons,NullLogicalAnd/Or, shifts, bitwise, e.g. https://github.com/rapidsai/cudf/blob/b8a14e2ce87308f7e271bd05ef570e32b6f50dfb/cpp/src/binaryop/compiled/operation.cuh#L85- Basically all operators under anonymous namespace in
src/unary, e.g. https://github.com/rapidsai/cudf/blob/b8a14e2ce87308f7e271bd05ef570e32b6f50dfb/cpp/src/unary/math_ops.cu#L31
I believe most of them can be replaced with the corresponding cudf::ops operators.
One other question is about the location of these new operators. I may be mistaken, but my understanding is that they don't necessarily need to live in the public namespace. libcudf has generally tried to avoid exposing device APIs unless there's a clear need to do so. Would it be possible to move these operators into a detail namespace instead?
|
@PointKernel I've added a tracking issue as requested: #22742 |
PointKernel
left a comment
There was a problem hiding this comment.
What a journey!
I'm really happy with how this PR turned out. @lamarrr I especially appreciate how thoughtfully you engaged with the review feedback throughout the process and how willing you were to revisit ideas and refine the design as the discussion evolved.
This is a significant addition to cudf and I think we've landed in a very clean and maintainable place. The final result is much stronger because of the effort you put into iterating on it. Thanks for all the hard work, and congratulations on getting this across the finish line. Great work! 🔥 🔥 🔥
|
/merge |
…22836) Follows up #22514. Implements refactoring and bug fixes for the operator library: - renames the `bit_*` operators to `bitwise_*` for consistency with AST operator names - renames `cast_to_b8` to `cast_to_bool8` and other `cast_to_*` functions for consistency with existing AST operator names - renames `ansi_*` arithmetic to use `_overflow` and `ansi_precision_check` to `precision_check` for consistency with CUDF's operators - updates binary operators to use double template types (to make it usable with AST's double dispatch mechanism) - removes invalid `.overflow` checks in the ANSI operators - fixes error returns in ANSI operators (return `std::unexpected`) - removes `cuda::std::convertible_to` concept checks, which are not available in CUDA 12.2 - properly constrains `if_else` operator to avoid ambiguous argument deduction - moves the error code enum to a public header - adds `to_string` function for the error code enum --- to be used in exceptions - refactors the transform kernel - removes `nullify_if` as it is not necessary for a first PoC for ANSI support - Fixes call to `detail::decimal_cast` in `cast_to_decimal_*` functions. Authors: - Basit Ayantunde (https://github.com/lamarrr) Approvers: - Igor Peshansky (https://github.com/igorpeshansky) - Yunsong Wang (https://github.com/PointKernel) URL: #22836
…VIDIA#22836) Follows up NVIDIA#22514. Implements refactoring and bug fixes for the operator library: - renames the `bit_*` operators to `bitwise_*` for consistency with AST operator names - renames `cast_to_b8` to `cast_to_bool8` and other `cast_to_*` functions for consistency with existing AST operator names - renames `ansi_*` arithmetic to use `_overflow` and `ansi_precision_check` to `precision_check` for consistency with CUDF's operators - updates binary operators to use double template types (to make it usable with AST's double dispatch mechanism) - removes invalid `.overflow` checks in the ANSI operators - fixes error returns in ANSI operators (return `std::unexpected`) - removes `cuda::std::convertible_to` concept checks, which are not available in CUDA 12.2 - properly constrains `if_else` operator to avoid ambiguous argument deduction - moves the error code enum to a public header - adds `to_string` function for the error code enum --- to be used in exceptions - refactors the transform kernel - removes `nullify_if` as it is not necessary for a first PoC for ANSI support - Fixes call to `detail::decimal_cast` in `cast_to_decimal_*` functions. Authors: - Basit Ayantunde (https://github.com/lamarrr) Approvers: - Igor Peshansky (https://github.com/igorpeshansky) - Yunsong Wang (https://github.com/PointKernel) URL: NVIDIA#22836
Description
Split from #22224
Preceded by #22511
Story: #22598
This Pull request:
coalesceoperatorresulttype to allow error returns from operatorsoperator_functorsto use the operator libraryChecklist