From c18ba0187ceb9004bbfc57d9bbfc0129e8199930 Mon Sep 17 00:00:00 2001 From: "Robert (Bobby) Evans" Date: Tue, 11 Jan 2022 09:20:45 -0600 Subject: [PATCH 1/7] Add in support for NULL_LOGICAL_AND and NULL_LOGICAL_OR binops --- cpp/CMakeLists.txt | 4 +- cpp/include/cudf/binaryop.hpp | 12 ++++- cpp/src/binaryop/binaryop.cpp | 5 +- cpp/src/binaryop/compiled/binary_ops.cu | 4 +- cpp/src/binaryop/compiled/binary_ops.cuh | 6 ++- cpp/src/binaryop/compiled/operation.cuh | 52 ++++++++++++++++++- cpp/src/binaryop/compiled/util.cpp | 9 +++- cpp/tests/binaryop/binop-compiled-test.cpp | 36 ++++++++++++- cpp/tests/binaryop/util/operation.h | 44 +++++++++++++++- .../main/java/ai/rapids/cudf/BinaryOp.java | 6 ++- .../java/ai/rapids/cudf/BinaryOpTest.java | 21 +++++++- 11 files changed, 184 insertions(+), 15 deletions(-) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 84e486c7e18..3dce777d77d 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -1,5 +1,5 @@ # ============================================================================= -# Copyright (c) 2018-2021, NVIDIA CORPORATION. +# Copyright (c) 2018-2022, NVIDIA CORPORATION. # # Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except # in compliance with the License. You may obtain a copy of the License at @@ -186,6 +186,8 @@ add_library( src/binaryop/compiled/Mod.cu src/binaryop/compiled/Mul.cu src/binaryop/compiled/NullEquals.cu + src/binaryop/compiled/NullLogicalOr.cu + src/binaryop/compiled/NullLogicalAnd.cu src/binaryop/compiled/NullMax.cu src/binaryop/compiled/NullMin.cu src/binaryop/compiled/PMod.cu diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index a514010c1f0..a6d37c948bc 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021, NVIDIA CORPORATION. + * Copyright (c) 2019-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -72,7 +72,15 @@ enum class binary_operator : int32_t { ///< operand when one is null; or invalid when both are null GENERIC_BINARY, ///< generic binary operator to be generated with input ///< ptx code - INVALID_BINARY ///< invalid operation + NULL_LOGICAL_AND, ///< operator && with Spark rules: NULL_LOGICAL_AND(null, null) is null, + ///< NULL_LOGICAL_AND(null, true) is + ///< null, NULL_LOGICAL_AND(null, false) is false, and NULL_LOGICAL_AND(valid, + ///< valid) == LOGICAL_AND(valid, valid) + NULL_LOGICAL_OR, ///< operator || with Spark rules: NULL_LOGICAL_OR(null, null) is null, + ///< NULL_LOGICAL_OR(null, true) is true, + ///< NULL_LOGICAL_OR(null, false) is null, and NULL_LOGICAL_OR(valid, valid) == + ///< LOGICAL_OR(valid, valid) + INVALID_BINARY ///< invalid operation }; /** * @brief Performs a binary operation between a scalar and a column. diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 7087b71a84e..5f9ff2574e3 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021, NVIDIA CORPORATION. + * Copyright (c) 2019-2022, NVIDIA CORPORATION. * * Copyright 2018-2019 BlazingDB, Inc. * Copyright 2018 Christian Noboa Mardini @@ -74,7 +74,8 @@ rmm::device_buffer scalar_col_valid_mask_and(column_view const& col, inline bool is_null_dependent(binary_operator op) { return op == binary_operator::NULL_EQUALS || op == binary_operator::NULL_MIN || - op == binary_operator::NULL_MAX; + op == binary_operator::NULL_MAX || op == binary_operator::NULL_LOGICAL_AND || + op == binary_operator::NULL_LOGICAL_OR; } /** diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu index 71d9b615153..d38fcab00c4 100644 --- a/cpp/src/binaryop/compiled/binary_ops.cu +++ b/cpp/src/binaryop/compiled/binary_ops.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021, NVIDIA CORPORATION. + * Copyright (c) 2019-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -339,6 +339,8 @@ case binary_operator::PMOD: apply_binary_op(out, lhs, case binary_operator::NULL_EQUALS: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::NULL_MAX: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; case binary_operator::NULL_MIN: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; +case binary_operator::NULL_LOGICAL_AND: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; +case binary_operator::NULL_LOGICAL_OR: apply_binary_op(out, lhs, rhs, is_lhs_scalar, is_rhs_scalar, stream); break; default:; } // clang-format on diff --git a/cpp/src/binaryop/compiled/binary_ops.cuh b/cpp/src/binaryop/compiled/binary_ops.cuh index 5bd639dd295..7f5e204f8e8 100644 --- a/cpp/src/binaryop/compiled/binary_ops.cuh +++ b/cpp/src/binaryop/compiled/binary_ops.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -103,6 +103,8 @@ struct ops_wrapper { type_dispatcher(rhs.type(), type_casted_accessor{}, i, rhs, is_rhs_scalar); auto result = [&]() { if constexpr (std::is_same_v or + std::is_same_v or + std::is_same_v or std::is_same_v or std::is_same_v) { bool output_valid = false; @@ -150,6 +152,8 @@ struct ops2_wrapper { TypeRhs y = rhs.element(is_rhs_scalar ? 0 : i); auto result = [&]() { if constexpr (std::is_same_v or + std::is_same_v or + std::is_same_v or std::is_same_v or std::is_same_v) { bool output_valid = false; diff --git a/cpp/src/binaryop/compiled/operation.cuh b/cpp/src/binaryop/compiled/operation.cuh index 86645e2cb8a..51b476794fc 100644 --- a/cpp/src/binaryop/compiled/operation.cuh +++ b/cpp/src/binaryop/compiled/operation.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -415,6 +415,56 @@ struct NullMin { -> decltype(static_cast(static_cast(x) < static_cast(y) ? x : y)); }; +struct NullLogicalAnd { + template + CUDA_DEVICE_CALLABLE auto operator()( + TypeLhs x, TypeRhs y, bool lhs_valid, bool rhs_valid, bool& output_valid) -> decltype(x && y) + { + if (lhs_valid && !x) { + output_valid = true; + return false; + } + if (rhs_valid && !y) { + output_valid = true; + return false; + } + if (lhs_valid && rhs_valid) { + output_valid = true; + return true; + } + output_valid = false; + return false; + } + // To allow std::is_invocable_v = true + template + CUDA_DEVICE_CALLABLE auto operator()(TypeLhs x, TypeRhs y) -> decltype(x && y); +}; + +struct NullLogicalOr { + template + CUDA_DEVICE_CALLABLE auto operator()( + TypeLhs x, TypeRhs y, bool lhs_valid, bool rhs_valid, bool& output_valid) -> decltype(x || y) + { + if (lhs_valid && x) { + output_valid = true; + return true; + } + if (rhs_valid && y) { + output_valid = true; + return true; + } + if (lhs_valid && rhs_valid) { + output_valid = true; + return false; + } + output_valid = false; + return false; + } + // To allow std::is_invocable_v = true + template + CUDA_DEVICE_CALLABLE auto operator()(TypeLhs x, TypeRhs y) -> decltype(x || y); +}; + } // namespace ops } // namespace compiled } // namespace binops diff --git a/cpp/src/binaryop/compiled/util.cpp b/cpp/src/binaryop/compiled/util.cpp index f89941a3d68..146e53aae59 100644 --- a/cpp/src/binaryop/compiled/util.cpp +++ b/cpp/src/binaryop/compiled/util.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -71,8 +71,9 @@ struct is_binary_operation_supported { if constexpr (has_common_type_v) { using common_t = std::common_type_t; return std::is_invocable_v; - } else + } else { return std::is_invocable_v; + } } else { return false; } @@ -166,6 +167,10 @@ struct is_supported_operation_functor { case binary_operator::LESS_EQUAL: return bool_op(out); case binary_operator::GREATER_EQUAL: return bool_op(out); case binary_operator::NULL_EQUALS: return bool_op(out); + case binary_operator::NULL_LOGICAL_AND: + return bool_op(out); + case binary_operator::NULL_LOGICAL_OR: + return bool_op(out); default: return type_dispatcher(out, nested_support_functor{}, op); } return false; diff --git a/cpp/tests/binaryop/binop-compiled-test.cpp b/cpp/tests/binaryop/binop-compiled-test.cpp index 0339d52dda9..bf47f01970e 100644 --- a/cpp/tests/binaryop/binop-compiled-test.cpp +++ b/cpp/tests/binaryop/binop-compiled-test.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -580,6 +580,40 @@ auto NullOp_Result(column_view lhs, column_view rhs) return column_wrapper(result.cbegin(), result.cend(), result_mask.cbegin()); } +// Yes this is ugly but these logical tests need some things from nullops to work, and this +// kept the changes smallest +TYPED_TEST(BinaryOperationCompiledTest_Logical, NullLogicalAnd_Vector_Vector) +{ + using TypeOut = bool; + using TypeLhs = typename TestFixture::TypeLhs; + using TypeRhs = typename TestFixture::TypeRhs; + using NULL_AND = cudf::library::operation::NullLogicalAnd; + + auto lhs = lhs_random_column(col_size); + auto rhs = rhs_random_column(col_size); + auto const expected = NullOp_Result(lhs, rhs); + + auto const result = cudf::binary_operation( + lhs, rhs, cudf::binary_operator::NULL_LOGICAL_AND, data_type(type_to_id())); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + +TYPED_TEST(BinaryOperationCompiledTest_Logical, NullLogicalOr_Vector_Vector) +{ + using TypeOut = bool; + using TypeLhs = typename TestFixture::TypeLhs; + using TypeRhs = typename TestFixture::TypeRhs; + using NULL_OR = cudf::library::operation::NullLogicalOr; + + auto lhs = lhs_random_column(col_size); + auto rhs = rhs_random_column(col_size); + auto const expected = NullOp_Result(lhs, rhs); + + auto const result = cudf::binary_operation( + lhs, rhs, cudf::binary_operator::NULL_LOGICAL_OR, data_type(type_to_id())); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + TYPED_TEST(BinaryOperationCompiledTest_NullOps, NullEquals_Vector_Vector) { using TypeOut = bool; diff --git a/cpp/tests/binaryop/util/operation.h b/cpp/tests/binaryop/util/operation.h index 481e5cfd4a9..3b0ba2e38d0 100644 --- a/cpp/tests/binaryop/util/operation.h +++ b/cpp/tests/binaryop/util/operation.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021, NVIDIA CORPORATION. + * Copyright (c) 2019-2022, NVIDIA CORPORATION. * * Copyright 2018-2019 BlazingDB, Inc. * Copyright 2018 Christian Noboa Mardini @@ -323,6 +323,48 @@ struct PyMod { } }; +template +struct NullLogicalAnd { + TypeOut operator()(TypeLhs x, TypeRhs y, bool lhs_valid, bool rhs_valid, bool& output_valid) const + { + if (lhs_valid && !x) { + output_valid = true; + return false; + } + if (rhs_valid && !y) { + output_valid = true; + return false; + } + if (lhs_valid && rhs_valid) { + output_valid = true; + return true; + } + output_valid = false; + return false; + } +}; + +template +struct NullLogicalOr { + TypeOut operator()(TypeLhs x, TypeRhs y, bool lhs_valid, bool rhs_valid, bool& output_valid) const + { + if (lhs_valid && x) { + output_valid = true; + return true; + } + if (rhs_valid && y) { + output_valid = true; + return true; + } + if (lhs_valid && rhs_valid) { + output_valid = true; + return false; + } + output_valid = false; + return false; + } +}; + template struct NullEquals { TypeOut operator()(TypeLhs x, TypeRhs y, bool lhs_valid, bool rhs_valid, bool& output_valid) const diff --git a/java/src/main/java/ai/rapids/cudf/BinaryOp.java b/java/src/main/java/ai/rapids/cudf/BinaryOp.java index 8b58d8383b4..15b8d32d6da 100644 --- a/java/src/main/java/ai/rapids/cudf/BinaryOp.java +++ b/java/src/main/java/ai/rapids/cudf/BinaryOp.java @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2020,2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -49,8 +49,10 @@ public enum BinaryOp { GREATER_EQUAL(25), // >= NULL_EQUALS(26), // like EQUAL but NULL == NULL is TRUE and NULL == not NULL is FALSE NULL_MAX(27), // MAX but NULL < not NULL - NULL_MIN(28); // MIN but NULL > not NULL + NULL_MIN(28), // MIN but NULL > not NULL //NOT IMPLEMENTED YET GENERIC_BINARY(29); + NULL_LOGICAL_AND(30), + NULL_LOGICAL_OR(31); static final EnumSet COMPARISON = EnumSet.of( diff --git a/java/src/test/java/ai/rapids/cudf/BinaryOpTest.java b/java/src/test/java/ai/rapids/cudf/BinaryOpTest.java index 0ca997d3c80..f94919fa25d 100644 --- a/java/src/test/java/ai/rapids/cudf/BinaryOpTest.java +++ b/java/src/test/java/ai/rapids/cudf/BinaryOpTest.java @@ -1,6 +1,6 @@ /* * - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -1363,6 +1363,25 @@ public void testBitXor() { } } + @Test + public void testNullAnd() { + try (ColumnVector icv1 = ColumnVector.fromBoxedBooleans(BOOLEANS_1); + ColumnVector icv2 = ColumnVector.fromBoxedBooleans(BOOLEANS_2)) { + try (ColumnVector answer = icv1.binaryOp(BinaryOp.NULL_LOGICAL_AND, icv2, DType.BOOL8); + ColumnVector expected = forEach(DType.BOOL8, icv1, icv2, + (b, l, r, i) -> b.append(l.getBoolean(i) && r.getBoolean(i)))) { + assertColumnsAreEqual(expected, answer, "boolean AND boolean"); + } + + try (Scalar s = Scalar.fromBool(true); + ColumnVector answer = icv1.binaryOp(BinaryOp.NULL_LOGICAL_AND, s, DType.BOOL8); + ColumnVector expected = forEachS(DType.BOOL8, icv1, true, + (b, l, r, i) -> b.append(l.getBoolean(i) && r))) { + assertColumnsAreEqual(expected, answer, "boolean AND true"); + } + } + } + @Test public void testAnd() { try (ColumnVector icv1 = ColumnVector.fromBoxedBooleans(BOOLEANS_1); From f4fdbd3d247c6310513c80f09effab7a22852307 Mon Sep 17 00:00:00 2001 From: "Robert (Bobby) Evans" Date: Tue, 11 Jan 2022 09:55:53 -0600 Subject: [PATCH 2/7] Addressed review comments and style checks --- cpp/include/cudf/binaryop.hpp | 14 +++---- cpp/tests/binaryop/binop-compiled-test.cpp | 12 +++--- cpp/tests/binaryop/util/operation.h | 8 ++-- .../java/ai/rapids/cudf/BinaryOpTest.java | 41 ++++++++++++++----- 4 files changed, 46 insertions(+), 29 deletions(-) diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index a6d37c948bc..39deaff4301 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -72,15 +72,11 @@ enum class binary_operator : int32_t { ///< operand when one is null; or invalid when both are null GENERIC_BINARY, ///< generic binary operator to be generated with input ///< ptx code - NULL_LOGICAL_AND, ///< operator && with Spark rules: NULL_LOGICAL_AND(null, null) is null, - ///< NULL_LOGICAL_AND(null, true) is - ///< null, NULL_LOGICAL_AND(null, false) is false, and NULL_LOGICAL_AND(valid, - ///< valid) == LOGICAL_AND(valid, valid) - NULL_LOGICAL_OR, ///< operator || with Spark rules: NULL_LOGICAL_OR(null, null) is null, - ///< NULL_LOGICAL_OR(null, true) is true, - ///< NULL_LOGICAL_OR(null, false) is null, and NULL_LOGICAL_OR(valid, valid) == - ///< LOGICAL_OR(valid, valid) - INVALID_BINARY ///< invalid operation + NULL_LOGICAL_AND, ///< operator && with Spark rules: (null, null) is null, (null, true) is null, + ///< (null, false)is false, and (valid, valid) == LOGICAL_AND(valid, valid) + NULL_LOGICAL_OR, ///< operator || with Spark rules: (null, null) is null, (null, true) is true, + ///< (null, false) is null, and (valid, valid) == LOGICAL_OR(valid, valid) + INVALID_BINARY ///< invalid operation }; /** * @brief Performs a binary operation between a scalar and a column. diff --git a/cpp/tests/binaryop/binop-compiled-test.cpp b/cpp/tests/binaryop/binop-compiled-test.cpp index bf47f01970e..57766495091 100644 --- a/cpp/tests/binaryop/binop-compiled-test.cpp +++ b/cpp/tests/binaryop/binop-compiled-test.cpp @@ -584,9 +584,9 @@ auto NullOp_Result(column_view lhs, column_view rhs) // kept the changes smallest TYPED_TEST(BinaryOperationCompiledTest_Logical, NullLogicalAnd_Vector_Vector) { - using TypeOut = bool; - using TypeLhs = typename TestFixture::TypeLhs; - using TypeRhs = typename TestFixture::TypeRhs; + using TypeOut = bool; + using TypeLhs = typename TestFixture::TypeLhs; + using TypeRhs = typename TestFixture::TypeRhs; using NULL_AND = cudf::library::operation::NullLogicalAnd; auto lhs = lhs_random_column(col_size); @@ -600,9 +600,9 @@ TYPED_TEST(BinaryOperationCompiledTest_Logical, NullLogicalAnd_Vector_Vector) TYPED_TEST(BinaryOperationCompiledTest_Logical, NullLogicalOr_Vector_Vector) { - using TypeOut = bool; - using TypeLhs = typename TestFixture::TypeLhs; - using TypeRhs = typename TestFixture::TypeRhs; + using TypeOut = bool; + using TypeLhs = typename TestFixture::TypeLhs; + using TypeRhs = typename TestFixture::TypeRhs; using NULL_OR = cudf::library::operation::NullLogicalOr; auto lhs = lhs_random_column(col_size); diff --git a/cpp/tests/binaryop/util/operation.h b/cpp/tests/binaryop/util/operation.h index 3b0ba2e38d0..22802580cd0 100644 --- a/cpp/tests/binaryop/util/operation.h +++ b/cpp/tests/binaryop/util/operation.h @@ -333,11 +333,11 @@ struct NullLogicalAnd { } if (rhs_valid && !y) { output_valid = true; - return false; + return false; } if (lhs_valid && rhs_valid) { output_valid = true; - return true; + return true; } output_valid = false; return false; @@ -354,11 +354,11 @@ struct NullLogicalOr { } if (rhs_valid && y) { output_valid = true; - return true; + return true; } if (lhs_valid && rhs_valid) { output_valid = true; - return false; + return false; } output_valid = false; return false; diff --git a/java/src/test/java/ai/rapids/cudf/BinaryOpTest.java b/java/src/test/java/ai/rapids/cudf/BinaryOpTest.java index f94919fa25d..862f3860d3d 100644 --- a/java/src/test/java/ai/rapids/cudf/BinaryOpTest.java +++ b/java/src/test/java/ai/rapids/cudf/BinaryOpTest.java @@ -1365,19 +1365,40 @@ public void testBitXor() { @Test public void testNullAnd() { - try (ColumnVector icv1 = ColumnVector.fromBoxedBooleans(BOOLEANS_1); - ColumnVector icv2 = ColumnVector.fromBoxedBooleans(BOOLEANS_2)) { + try (ColumnVector icv1 = ColumnVector.fromBoxedBooleans( + true, true, true, + false, false, false, + null, null, null); + ColumnVector icv2 = ColumnVector.fromBoxedBooleans( + true, false, null, + true, false, null, + true, false, null)) { try (ColumnVector answer = icv1.binaryOp(BinaryOp.NULL_LOGICAL_AND, icv2, DType.BOOL8); - ColumnVector expected = forEach(DType.BOOL8, icv1, icv2, - (b, l, r, i) -> b.append(l.getBoolean(i) && r.getBoolean(i)))) { - assertColumnsAreEqual(expected, answer, "boolean AND boolean"); + ColumnVector expected = ColumnVector.fromBoxedBooleans( + true, false, null, + false, false, false, + null, false, null)) { + assertColumnsAreEqual(expected, answer, "boolean NULL AND boolean"); } + } + } - try (Scalar s = Scalar.fromBool(true); - ColumnVector answer = icv1.binaryOp(BinaryOp.NULL_LOGICAL_AND, s, DType.BOOL8); - ColumnVector expected = forEachS(DType.BOOL8, icv1, true, - (b, l, r, i) -> b.append(l.getBoolean(i) && r))) { - assertColumnsAreEqual(expected, answer, "boolean AND true"); + @Test + public void testNullOr() { + try (ColumnVector icv1 = ColumnVector.fromBoxedBooleans( + true, true, true, + false, false, false, + null, null, null); + ColumnVector icv2 = ColumnVector.fromBoxedBooleans( + true, false, null, + true, false, null, + true, false, null)) { + try (ColumnVector answer = icv1.binaryOp(BinaryOp.NULL_LOGICAL_OR, icv2, DType.BOOL8); + ColumnVector expected = ColumnVector.fromBoxedBooleans( + true, true, true, + true, false, null, + true, null, null)) { + assertColumnsAreEqual(expected, answer, "boolean NULL OR boolean"); } } } From 7ec2e64e707a20ebd0bb33b1861a9dd8c16bddce Mon Sep 17 00:00:00 2001 From: "Robert (Bobby) Evans" Date: Tue, 11 Jan 2022 10:27:25 -0600 Subject: [PATCH 3/7] Missed some files --- cpp/src/binaryop/compiled/NullLogicalAnd.cu | 26 +++++++++++++++++++++ cpp/src/binaryop/compiled/NullLogicalOr.cu | 26 +++++++++++++++++++++ 2 files changed, 52 insertions(+) create mode 100644 cpp/src/binaryop/compiled/NullLogicalAnd.cu create mode 100644 cpp/src/binaryop/compiled/NullLogicalOr.cu diff --git a/cpp/src/binaryop/compiled/NullLogicalAnd.cu b/cpp/src/binaryop/compiled/NullLogicalAnd.cu new file mode 100644 index 00000000000..48ae125bc93 --- /dev/null +++ b/cpp/src/binaryop/compiled/NullLogicalAnd.cu @@ -0,0 +1,26 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "binary_ops.cuh" + +namespace cudf::binops::compiled { +template void apply_binary_op(mutable_column_device_view&, + column_device_view const&, + column_device_view const&, + bool is_lhs_scalar, + bool is_rhs_scalar, + rmm::cuda_stream_view); +} // namespace cudf::binops::compiled diff --git a/cpp/src/binaryop/compiled/NullLogicalOr.cu b/cpp/src/binaryop/compiled/NullLogicalOr.cu new file mode 100644 index 00000000000..e0ea95ac3ee --- /dev/null +++ b/cpp/src/binaryop/compiled/NullLogicalOr.cu @@ -0,0 +1,26 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "binary_ops.cuh" + +namespace cudf::binops::compiled { +template void apply_binary_op(mutable_column_device_view&, + column_device_view const&, + column_device_view const&, + bool is_lhs_scalar, + bool is_rhs_scalar, + rmm::cuda_stream_view); +} // namespace cudf::binops::compiled From 6a6312525418a2de2918b4b79a4d6059685da427 Mon Sep 17 00:00:00 2001 From: "Robert (Bobby) Evans" Date: Tue, 11 Jan 2022 13:10:39 -0600 Subject: [PATCH 4/7] Fix merge conflict --- cpp/src/binaryop/compiled/operation.cuh | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/src/binaryop/compiled/operation.cuh b/cpp/src/binaryop/compiled/operation.cuh index bc5ee522c47..538efaeea04 100644 --- a/cpp/src/binaryop/compiled/operation.cuh +++ b/cpp/src/binaryop/compiled/operation.cuh @@ -417,7 +417,7 @@ struct NullMin { struct NullLogicalAnd { template - CUDA_DEVICE_CALLABLE auto operator()( + __device__ inline auto operator()( TypeLhs x, TypeRhs y, bool lhs_valid, bool rhs_valid, bool& output_valid) -> decltype(x && y) { if (lhs_valid && !x) { @@ -437,12 +437,12 @@ struct NullLogicalAnd { } // To allow std::is_invocable_v = true template - CUDA_DEVICE_CALLABLE auto operator()(TypeLhs x, TypeRhs y) -> decltype(x && y); + __device__ inline auto operator()(TypeLhs x, TypeRhs y) -> decltype(x && y); }; struct NullLogicalOr { template - CUDA_DEVICE_CALLABLE auto operator()( + __device__ inline auto operator()( TypeLhs x, TypeRhs y, bool lhs_valid, bool rhs_valid, bool& output_valid) -> decltype(x || y) { if (lhs_valid && x) { @@ -462,7 +462,7 @@ struct NullLogicalOr { } // To allow std::is_invocable_v = true template - CUDA_DEVICE_CALLABLE auto operator()(TypeLhs x, TypeRhs y) -> decltype(x || y); + __device__ inline auto operator()(TypeLhs x, TypeRhs y) -> decltype(x || y); }; } // namespace ops From a273f0005bddc6e03699706f41f5a508a420b95d Mon Sep 17 00:00:00 2001 From: "Robert (Bobby) Evans" Date: Wed, 12 Jan 2022 11:12:43 -0600 Subject: [PATCH 5/7] Addressed review comments --- cpp/include/cudf/binaryop.hpp | 2 +- cpp/tests/binaryop/binop-compiled-test.cpp | 118 ++++++++++----------- 2 files changed, 59 insertions(+), 61 deletions(-) diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index 39deaff4301..daf55c0befe 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -73,7 +73,7 @@ enum class binary_operator : int32_t { GENERIC_BINARY, ///< generic binary operator to be generated with input ///< ptx code NULL_LOGICAL_AND, ///< operator && with Spark rules: (null, null) is null, (null, true) is null, - ///< (null, false)is false, and (valid, valid) == LOGICAL_AND(valid, valid) + ///< (null, false) is false, and (valid, valid) == LOGICAL_AND(valid, valid) NULL_LOGICAL_OR, ///< operator || with Spark rules: (null, null) is null, (null, true) is true, ///< (null, false) is null, and (valid, valid) == LOGICAL_OR(valid, valid) INVALID_BINARY ///< invalid operation diff --git a/cpp/tests/binaryop/binop-compiled-test.cpp b/cpp/tests/binaryop/binop-compiled-test.cpp index 57766495091..00408741653 100644 --- a/cpp/tests/binaryop/binop-compiled-test.cpp +++ b/cpp/tests/binaryop/binop-compiled-test.cpp @@ -475,6 +475,64 @@ TYPED_TEST(BinaryOperationCompiledTest_Logical, LogicalOr_Vector_Vector) this->template test(cudf::binary_operator::LOGICAL_OR); } +template +using column_wrapper = std::conditional_t, + cudf::test::strings_column_wrapper, + cudf::test::fixed_width_column_wrapper>; + +template +auto NullOp_Result(column_view lhs, column_view rhs) +{ + auto [lhs_data, lhs_mask] = cudf::test::to_host(lhs); + auto [rhs_data, rhs_mask] = cudf::test::to_host(rhs); + std::vector result(lhs.size()); + std::vector result_mask; + std::transform(thrust::make_counting_iterator(0), + thrust::make_counting_iterator(lhs.size()), + result.begin(), + [&lhs_data, &lhs_mask, &rhs_data, &rhs_mask, &result_mask](auto i) -> TypeOut { + auto lhs_valid = lhs_mask.data() and cudf::bit_is_set(lhs_mask.data(), i); + auto rhs_valid = rhs_mask.data() and cudf::bit_is_set(rhs_mask.data(), i); + bool output_valid = lhs_valid or rhs_valid; + auto result = OP{}(lhs_data[i], rhs_data[i], lhs_valid, rhs_valid, output_valid); + result_mask.push_back(output_valid); + return result; + }); + return column_wrapper(result.cbegin(), result.cend(), result_mask.cbegin()); +} + +TYPED_TEST(BinaryOperationCompiledTest_Logical, NullLogicalAnd_Vector_Vector) +{ + using TypeOut = bool; + using TypeLhs = typename TestFixture::TypeLhs; + using TypeRhs = typename TestFixture::TypeRhs; + using NULL_AND = cudf::library::operation::NullLogicalAnd; + + auto lhs = lhs_random_column(col_size); + auto rhs = rhs_random_column(col_size); + auto const expected = NullOp_Result(lhs, rhs); + + auto const result = cudf::binary_operation( + lhs, rhs, cudf::binary_operator::NULL_LOGICAL_AND, data_type(type_to_id())); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + +TYPED_TEST(BinaryOperationCompiledTest_Logical, NullLogicalOr_Vector_Vector) +{ + using TypeOut = bool; + using TypeLhs = typename TestFixture::TypeLhs; + using TypeRhs = typename TestFixture::TypeRhs; + using NULL_OR = cudf::library::operation::NullLogicalOr; + + auto lhs = lhs_random_column(col_size); + auto rhs = rhs_random_column(col_size); + auto const expected = NullOp_Result(lhs, rhs); + + auto const result = cudf::binary_operation( + lhs, rhs, cudf::binary_operator::NULL_LOGICAL_OR, data_type(type_to_id())); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + // Comparison Operations ==, !=, <, >, <=, >= // nn, tt, dd, ss, dcdc using Comparison_types = cudf::test::Types, @@ -554,66 +612,6 @@ struct BinaryOperationCompiledTest_NullOps : public BinaryOperationCompiledTest< }; TYPED_TEST_SUITE(BinaryOperationCompiledTest_NullOps, Null_types); -template -using column_wrapper = std::conditional_t, - cudf::test::strings_column_wrapper, - cudf::test::fixed_width_column_wrapper>; - -template -auto NullOp_Result(column_view lhs, column_view rhs) -{ - auto [lhs_data, lhs_mask] = cudf::test::to_host(lhs); - auto [rhs_data, rhs_mask] = cudf::test::to_host(rhs); - std::vector result(lhs.size()); - std::vector result_mask; - std::transform(thrust::make_counting_iterator(0), - thrust::make_counting_iterator(lhs.size()), - result.begin(), - [&lhs_data, &lhs_mask, &rhs_data, &rhs_mask, &result_mask](auto i) -> TypeOut { - auto lhs_valid = lhs_mask.data() and cudf::bit_is_set(lhs_mask.data(), i); - auto rhs_valid = rhs_mask.data() and cudf::bit_is_set(rhs_mask.data(), i); - bool output_valid = lhs_valid or rhs_valid; - auto result = OP{}(lhs_data[i], rhs_data[i], lhs_valid, rhs_valid, output_valid); - result_mask.push_back(output_valid); - return result; - }); - return column_wrapper(result.cbegin(), result.cend(), result_mask.cbegin()); -} - -// Yes this is ugly but these logical tests need some things from nullops to work, and this -// kept the changes smallest -TYPED_TEST(BinaryOperationCompiledTest_Logical, NullLogicalAnd_Vector_Vector) -{ - using TypeOut = bool; - using TypeLhs = typename TestFixture::TypeLhs; - using TypeRhs = typename TestFixture::TypeRhs; - using NULL_AND = cudf::library::operation::NullLogicalAnd; - - auto lhs = lhs_random_column(col_size); - auto rhs = rhs_random_column(col_size); - auto const expected = NullOp_Result(lhs, rhs); - - auto const result = cudf::binary_operation( - lhs, rhs, cudf::binary_operator::NULL_LOGICAL_AND, data_type(type_to_id())); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); -} - -TYPED_TEST(BinaryOperationCompiledTest_Logical, NullLogicalOr_Vector_Vector) -{ - using TypeOut = bool; - using TypeLhs = typename TestFixture::TypeLhs; - using TypeRhs = typename TestFixture::TypeRhs; - using NULL_OR = cudf::library::operation::NullLogicalOr; - - auto lhs = lhs_random_column(col_size); - auto rhs = rhs_random_column(col_size); - auto const expected = NullOp_Result(lhs, rhs); - - auto const result = cudf::binary_operation( - lhs, rhs, cudf::binary_operator::NULL_LOGICAL_OR, data_type(type_to_id())); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); -} - TYPED_TEST(BinaryOperationCompiledTest_NullOps, NullEquals_Vector_Vector) { using TypeOut = bool; From b4c7487e211d000f407c4d3b4516e47a572395fc Mon Sep 17 00:00:00 2001 From: "Robert (Bobby) Evans" Date: Thu, 13 Jan 2022 12:57:02 -0600 Subject: [PATCH 6/7] Performance boost --- cpp/src/binaryop/compiled/operation.cuh | 38 +++++++------------------ 1 file changed, 10 insertions(+), 28 deletions(-) diff --git a/cpp/src/binaryop/compiled/operation.cuh b/cpp/src/binaryop/compiled/operation.cuh index 538efaeea04..c42cd66194f 100644 --- a/cpp/src/binaryop/compiled/operation.cuh +++ b/cpp/src/binaryop/compiled/operation.cuh @@ -420,20 +420,11 @@ struct NullLogicalAnd { __device__ inline auto operator()( TypeLhs x, TypeRhs y, bool lhs_valid, bool rhs_valid, bool& output_valid) -> decltype(x && y) { - if (lhs_valid && !x) { - output_valid = true; - return false; - } - if (rhs_valid && !y) { - output_valid = true; - return false; - } - if (lhs_valid && rhs_valid) { - output_valid = true; - return true; - } - output_valid = false; - return false; + bool lhs_false = lhs_valid && !x; + bool rhs_false = rhs_valid && !y; + bool both_valid = lhs_valid && rhs_valid; + output_valid = lhs_false || rhs_false || both_valid; + return both_valid && !lhs_false && !rhs_false; } // To allow std::is_invocable_v = true template @@ -445,20 +436,11 @@ struct NullLogicalOr { __device__ inline auto operator()( TypeLhs x, TypeRhs y, bool lhs_valid, bool rhs_valid, bool& output_valid) -> decltype(x || y) { - if (lhs_valid && x) { - output_valid = true; - return true; - } - if (rhs_valid && y) { - output_valid = true; - return true; - } - if (lhs_valid && rhs_valid) { - output_valid = true; - return false; - } - output_valid = false; - return false; + bool lhs_true = lhs_valid && x; + bool rhs_true = rhs_valid && y; + bool both_valid = lhs_valid && rhs_valid; + output_valid = lhs_true || rhs_true || both_valid; + return lhs_true || rhs_true; } // To allow std::is_invocable_v = true template From e1c6d68578a006824b1097c69b07ec1202c823eb Mon Sep 17 00:00:00 2001 From: "Robert (Bobby) Evans" Date: Thu, 13 Jan 2022 13:11:12 -0600 Subject: [PATCH 7/7] fix style again --- cpp/src/binaryop/compiled/operation.cuh | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/cpp/src/binaryop/compiled/operation.cuh b/cpp/src/binaryop/compiled/operation.cuh index c42cd66194f..313fc34567d 100644 --- a/cpp/src/binaryop/compiled/operation.cuh +++ b/cpp/src/binaryop/compiled/operation.cuh @@ -420,10 +420,10 @@ struct NullLogicalAnd { __device__ inline auto operator()( TypeLhs x, TypeRhs y, bool lhs_valid, bool rhs_valid, bool& output_valid) -> decltype(x && y) { - bool lhs_false = lhs_valid && !x; - bool rhs_false = rhs_valid && !y; + bool lhs_false = lhs_valid && !x; + bool rhs_false = rhs_valid && !y; bool both_valid = lhs_valid && rhs_valid; - output_valid = lhs_false || rhs_false || both_valid; + output_valid = lhs_false || rhs_false || both_valid; return both_valid && !lhs_false && !rhs_false; } // To allow std::is_invocable_v = true @@ -436,10 +436,10 @@ struct NullLogicalOr { __device__ inline auto operator()( TypeLhs x, TypeRhs y, bool lhs_valid, bool rhs_valid, bool& output_valid) -> decltype(x || y) { - bool lhs_true = lhs_valid && x; - bool rhs_true = rhs_valid && y; + bool lhs_true = lhs_valid && x; + bool rhs_true = rhs_valid && y; bool both_valid = lhs_valid && rhs_valid; - output_valid = lhs_true || rhs_true || both_valid; + output_valid = lhs_true || rhs_true || both_valid; return lhs_true || rhs_true; } // To allow std::is_invocable_v = true