From 7e7de2b2bf239c580d6711a74150f2b2ebe2745a Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Wed, 6 Jul 2022 15:35:20 -0700 Subject: [PATCH 1/2] debug first pass --- cpp/src/binaryop/compiled/binary_ops.cu | 24 ++++++++++++------------ cpp/tests/CMakeLists.txt | 4 ++-- cpp/tests/binaryop/assert-binops.h | 20 ++++++++++++-------- 3 files changed, 26 insertions(+), 22 deletions(-) diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu index ee9fe840fd6..0b86657c9df 100644 --- a/cpp/src/binaryop/compiled/binary_ops.cu +++ b/cpp/src/binaryop/compiled/binary_ops.cu @@ -314,14 +314,14 @@ case binary_operator::FLOOR_DIV: apply_binary_op(out, case binary_operator::MOD: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::PYMOD: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::POW: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; -case binary_operator::EQUAL: -case binary_operator::NOT_EQUAL: -if(out.type().id() != type_id::BOOL8) CUDF_FAIL("Output type of Comparison operator should be bool type"); -dispatch_equality_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, op, stream); break; -case binary_operator::LESS: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; -case binary_operator::GREATER: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; -case binary_operator::LESS_EQUAL: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; -case binary_operator::GREATER_EQUAL: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; +// case binary_operator::EQUAL: +// case binary_operator::NOT_EQUAL: +// if(out.type().id() != type_id::BOOL8) CUDF_FAIL("Output type of Comparison operator should be bool type"); +// dispatch_equality_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, op, stream); break; +// case binary_operator::LESS: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; +// case binary_operator::GREATER: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; +// case binary_operator::LESS_EQUAL: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; +// case binary_operator::GREATER_EQUAL: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::BITWISE_AND: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::BITWISE_OR: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::BITWISE_XOR: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; @@ -362,8 +362,8 @@ void binary_operation(mutable_column_view& out, binary_operator op, rmm::cuda_stream_view stream) { - auto [lhsv, aux] = scalar_to_column_view(lhs, stream); - operator_dispatcher(out, lhsv, rhs, true, false, op, stream); + auto lhsc = make_column_from_scalar(lhs, 1, stream); + operator_dispatcher(out, lhsc->view(), rhs, true, false, op, stream); } // vector_scalar void binary_operation(mutable_column_view& out, @@ -372,8 +372,8 @@ void binary_operation(mutable_column_view& out, binary_operator op, rmm::cuda_stream_view stream) { - auto [rhsv, aux] = scalar_to_column_view(rhs, stream); - operator_dispatcher(out, lhs, rhsv, false, true, op, stream); + auto rhsc = make_column_from_scalar(rhs, 1, stream); + operator_dispatcher(out, lhs, rhsc->view(), false, true, op, stream); } } // namespace compiled } // namespace binops diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 816c5a1c59c..2a2b689801e 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -191,8 +191,8 @@ ConfigureTest( binaryop/binop-verify-input-test.cpp binaryop/binop-null-test.cpp binaryop/binop-compiled-test.cpp - binaryop/binop-compiled-fixed_point-test.cpp - binaryop/binop-generic-ptx-test.cpp + # binaryop/binop-compiled-fixed_point-test.cpp + # binaryop/binop-generic-ptx-test.cpp ) # ################################################################################################## diff --git a/cpp/tests/binaryop/assert-binops.h b/cpp/tests/binaryop/assert-binops.h index b257ca21dd7..b9fbfac9878 100644 --- a/cpp/tests/binaryop/assert-binops.h +++ b/cpp/tests/binaryop/assert-binops.h @@ -91,10 +91,12 @@ void ASSERT_BINOP(column_view const& out, auto out_data = out_h.first; ASSERT_EQ(out_data.size(), rhs_data.size()); - for (size_t i = 0; i < out_data.size(); ++i) { - auto lhs = out_data[i]; - auto rhs = (TypeOut)(op(lhs_h, rhs_data[i])); - ASSERT_TRUE(value_comparator(lhs, rhs)) << stringify_out_values{}(i, lhs, rhs); + if(lhs.is_valid()) { + for (size_t i = 0; i < out_data.size(); ++i) { + auto lhs = out_data[i]; + auto rhs = (TypeOut)(op(lhs_h, rhs_data[i])); + ASSERT_TRUE(value_comparator(lhs, rhs)) << stringify_out_values{}(i, lhs, rhs); + } } if (rhs.nullable()) { @@ -138,10 +140,12 @@ void ASSERT_BINOP(column_view const& out, auto out_data = out_h.first; ASSERT_EQ(out_data.size(), lhs_data.size()); - for (size_t i = 0; i < out_data.size(); ++i) { - auto lhs = out_data[i]; - auto rhs = (TypeOut)(op(lhs_data[i], rhs_h)); - ASSERT_TRUE(value_comparator(lhs, rhs)) << stringify_out_values{}(i, lhs, rhs); + if(rhs.is_valid()) { + for (size_t i = 0; i < out_data.size(); ++i) { + auto lhs = out_data[i]; + auto rhs = (TypeOut)(op(lhs_data[i], rhs_h)); + ASSERT_TRUE(value_comparator(lhs, rhs)) << stringify_out_values{}(i, lhs, rhs); + } } if (lhs.nullable()) { From e0095ad0beed24a3ed48d1902f1498dac14c2d9d Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 7 Jul 2022 16:49:40 -0700 Subject: [PATCH 2/2] line of failure pinpointed --- cpp/src/binaryop/compiled/binary_ops.cu | 140 +++++++++++----------- cpp/src/binaryop/compiled/binary_ops.cuh | 4 +- cpp/src/binaryop/compiled/equality_ops.cu | 3 + 3 files changed, 75 insertions(+), 72 deletions(-) diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu index 0b86657c9df..6595574d7cb 100644 --- a/cpp/src/binaryop/compiled/binary_ops.cu +++ b/cpp/src/binaryop/compiled/binary_ops.cu @@ -36,71 +36,71 @@ namespace binops { namespace compiled { namespace { -/** - * @brief Converts scalar to column_view with single element. - * - * @return pair with column_view and column containing any auxiliary data to create column_view from - * scalar - */ -struct scalar_as_column_view { - using return_type = typename std::pair>; - template ())>* = nullptr> - return_type operator()(scalar const& s, rmm::cuda_stream_view, rmm::mr::device_memory_resource*) - { - auto& h_scalar_type_view = static_cast&>(const_cast(s)); - auto col_v = - column_view(s.type(), 1, h_scalar_type_view.data(), (bitmask_type const*)s.validity_data()); - return std::pair{col_v, std::unique_ptr(nullptr)}; - } - template ())>* = nullptr> - return_type operator()(scalar const&, rmm::cuda_stream_view, rmm::mr::device_memory_resource*) - { - CUDF_FAIL("Unsupported type"); - } -}; -// specialization for cudf::string_view -template <> -scalar_as_column_view::return_type scalar_as_column_view::operator()( - scalar const& s, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) -{ - using T = cudf::string_view; - auto& h_scalar_type_view = static_cast&>(const_cast(s)); +// /** +// * @brief Converts scalar to column_view with single element. +// * +// * @return pair with column_view and column containing any auxiliary data to create column_view from +// * scalar +// */ +// struct scalar_as_column_view { +// using return_type = typename std::pair>; +// template ())>* = nullptr> +// return_type operator()(scalar const& s, rmm::cuda_stream_view, rmm::mr::device_memory_resource*) +// { +// auto& h_scalar_type_view = static_cast&>(const_cast(s)); +// auto col_v = +// column_view(s.type(), 1, h_scalar_type_view.data(), (bitmask_type const*)s.validity_data()); +// return std::pair{col_v, std::unique_ptr(nullptr)}; +// } +// template ())>* = nullptr> +// return_type operator()(scalar const&, rmm::cuda_stream_view, rmm::mr::device_memory_resource*) +// { +// CUDF_FAIL("Unsupported type"); +// } +// }; +// // specialization for cudf::string_view +// template <> +// scalar_as_column_view::return_type scalar_as_column_view::operator()( +// scalar const& s, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) +// { +// using T = cudf::string_view; +// auto& h_scalar_type_view = static_cast&>(const_cast(s)); - // build offsets column from the string size - auto offsets_transformer_itr = - thrust::make_constant_iterator(h_scalar_type_view.size()); - auto offsets_column = strings::detail::make_offsets_child_column( - offsets_transformer_itr, offsets_transformer_itr + 1, stream, mr); +// // build offsets column from the string size +// auto offsets_transformer_itr = +// thrust::make_constant_iterator(h_scalar_type_view.size()); +// auto offsets_column = strings::detail::make_offsets_child_column( +// offsets_transformer_itr, offsets_transformer_itr + 1, stream, mr); - auto chars_column_v = - column_view(data_type{type_id::INT8}, h_scalar_type_view.size(), h_scalar_type_view.data()); - // Construct string column_view - auto col_v = column_view(s.type(), - 1, - nullptr, - (bitmask_type const*)s.validity_data(), - cudf::UNKNOWN_NULL_COUNT, - 0, - {offsets_column->view(), chars_column_v}); - return std::pair{col_v, std::move(offsets_column)}; -} +// auto chars_column_v = +// column_view(data_type{type_id::INT8}, h_scalar_type_view.size(), h_scalar_type_view.data()); +// // Construct string column_view +// auto col_v = column_view(s.type(), +// 1, +// nullptr, +// (bitmask_type const*)s.validity_data(), +// cudf::UNKNOWN_NULL_COUNT, +// 0, +// {offsets_column->view(), chars_column_v}); +// return std::pair{col_v, std::move(offsets_column)}; +// } -/** - * @brief Converts scalar to column_view with single element. - * - * @param scal scalar to convert - * @param stream CUDA stream used for device memory operations and kernel launches. - * @param mr Device memory resource used to allocate the returned column's device memory - * @return pair with column_view and column containing any auxiliary data to create - * column_view from scalar - */ -auto scalar_to_column_view( - scalar const& scal, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) -{ - return type_dispatcher(scal.type(), scalar_as_column_view{}, scal, stream, mr); -} +// /** +// * @brief Converts scalar to column_view with single element. +// * +// * @param scal scalar to convert +// * @param stream CUDA stream used for device memory operations and kernel launches. +// * @param mr Device memory resource used to allocate the returned column's device memory +// * @return pair with column_view and column containing any auxiliary data to create +// * column_view from scalar +// */ +// auto scalar_to_column_view( +// scalar const& scal, +// rmm::cuda_stream_view stream, +// rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +// { +// return type_dispatcher(scal.type(), scalar_as_column_view{}, scal, stream, mr); +// } // This functor does the actual comparison between string column value and a scalar string // or between two string column values using a comparator @@ -314,14 +314,14 @@ case binary_operator::FLOOR_DIV: apply_binary_op(out, case binary_operator::MOD: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::PYMOD: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::POW: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; -// case binary_operator::EQUAL: -// case binary_operator::NOT_EQUAL: -// if(out.type().id() != type_id::BOOL8) CUDF_FAIL("Output type of Comparison operator should be bool type"); -// dispatch_equality_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, op, stream); break; -// case binary_operator::LESS: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; -// case binary_operator::GREATER: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; -// case binary_operator::LESS_EQUAL: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; -// case binary_operator::GREATER_EQUAL: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; +case binary_operator::EQUAL: +case binary_operator::NOT_EQUAL: +if(out.type().id() != type_id::BOOL8) CUDF_FAIL("Output type of Comparison operator should be bool type"); +dispatch_equality_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, op, stream); break; +case binary_operator::LESS: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; +case binary_operator::GREATER: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; +case binary_operator::LESS_EQUAL: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; +case binary_operator::GREATER_EQUAL: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::BITWISE_AND: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::BITWISE_OR: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::BITWISE_XOR: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; diff --git a/cpp/src/binaryop/compiled/binary_ops.cuh b/cpp/src/binaryop/compiled/binary_ops.cuh index d88d2be2499..02ea6392e4a 100644 --- a/cpp/src/binaryop/compiled/binary_ops.cuh +++ b/cpp/src/binaryop/compiled/binary_ops.cuh @@ -123,8 +123,8 @@ struct ops_wrapper { // To suppress nvcc warning return std::invoke_result_t{}; }(); - if constexpr (is_bool_result()) - out.element(i) = result; + if constexpr (is_bool_result()){} + // out.element(i) = result; else type_dispatcher(out.type(), typed_casted_writer{}, i, out, result); } diff --git a/cpp/src/binaryop/compiled/equality_ops.cu b/cpp/src/binaryop/compiled/equality_ops.cu index 61f02252a26..17fcd7e3360 100644 --- a/cpp/src/binaryop/compiled/equality_ops.cu +++ b/cpp/src/binaryop/compiled/equality_ops.cu @@ -33,6 +33,9 @@ void dispatch_equality_op(mutable_column_view& out, auto rhsd = column_device_view::create(rhs, stream); if (common_dtype) { if (op == binary_operator::EQUAL) { + // cudf::test::print(out); + // cudf::test::print(lhs); + // cudf::test::print(rhs); for_each(stream, out.size(), binary_op_device_dispatcher{