From 32269fb619033eeeffb32bffff29c35d54e5ed80 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Mon, 20 May 2019 18:16:30 +0530 Subject: [PATCH 01/14] Implemented logical AND and logical OR in libcudf and exposed in cuDF --- cpp/include/cudf/types.h | 2 ++ cpp/src/binary/jit/code/operation.cpp | 18 ++++++++++++++++++ cpp/src/binary/jit/util/type.cpp | 4 ++++ cpp/src/unary/math_ops.cu | 1 + python/cudf/bindings/binops.pxd | 2 ++ python/cudf/bindings/binops.pyx | 3 +++ python/cudf/dataframe/numerical.py | 6 +++++- python/cudf/dataframe/series.py | 18 ++++++++++++++++++ 8 files changed, 53 insertions(+), 1 deletion(-) diff --git a/cpp/include/cudf/types.h b/cpp/include/cudf/types.h index 0f5833be8fb..627fcc2898a 100644 --- a/cpp/include/cudf/types.h +++ b/cpp/include/cudf/types.h @@ -223,6 +223,8 @@ typedef enum { GDF_BITWISE_AND, ///< operator & GDF_BITWISE_OR, ///< operator | GDF_BITWISE_XOR, ///< operator ^ + GDF_LOGICAL_AND, ///< operator && + GDF_LOGICAL_OR, ///< operator || GDF_COALESCE, ///< operator x,y x is null ? y : x GDF_INVALID_BINARY ///< invalid operation } gdf_binary_operator; diff --git a/cpp/src/binary/jit/code/operation.cpp b/cpp/src/binary/jit/code/operation.cpp index 1cf0cca7513..90fb1849f51 100644 --- a/cpp/src/binary/jit/code/operation.cpp +++ b/cpp/src/binary/jit/code/operation.cpp @@ -268,6 +268,24 @@ R"***( using RBitwiseXor = BitwiseXor; + struct LogicalAnd { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (x && y); + } + }; + + using RLogicalAnd = LogicalAnd; + + struct LogicalOr { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (x || y); + } + }; + + using RLogicalOr = LogicalOr; + )***"; } // namespace code diff --git a/cpp/src/binary/jit/util/type.cpp b/cpp/src/binary/jit/util/type.cpp index 0a43b7c8890..dad7d517905 100644 --- a/cpp/src/binary/jit/util/type.cpp +++ b/cpp/src/binary/jit/util/type.cpp @@ -115,6 +115,10 @@ namespace jit { return "BitwiseOr"; case GDF_BITWISE_XOR: return "BitwiseXor"; + case GDF_LOGICAL_AND: + return "LogicalAnd"; + case GDF_LOGICAL_OR: + return "LogicalOr"; default: return "None"; } diff --git a/cpp/src/unary/math_ops.cu b/cpp/src/unary/math_ops.cu index 3bc884812ba..8d42aed6af0 100644 --- a/cpp/src/unary/math_ops.cu +++ b/cpp/src/unary/math_ops.cu @@ -127,6 +127,7 @@ struct DeviceAbs { // bitwise op struct DeviceInvert { + // TODO: maybe sfinae overload this for cudf::bool8 template __device__ T apply(T data) { diff --git a/python/cudf/bindings/binops.pxd b/python/cudf/bindings/binops.pxd index f1c764e1cb6..a5b8d1eb7d5 100644 --- a/python/cudf/bindings/binops.pxd +++ b/python/cudf/bindings/binops.pxd @@ -28,6 +28,8 @@ cdef extern from "cudf.h" nogil: GDF_BITWISE_AND, GDF_BITWISE_OR, GDF_BITWISE_XOR, + GDF_LOGICAL_AND, + GDF_LOGICAL_OR, cdef gdf_error gdf_binary_operation_s_v(gdf_column* out, gdf_scalar* lhs, gdf_column* rhs, gdf_binary_operator ope) except + cdef gdf_error gdf_binary_operation_v_s(gdf_column* out, gdf_column* lhs, gdf_scalar* rhs, gdf_binary_operator ope) except + diff --git a/python/cudf/bindings/binops.pyx b/python/cudf/bindings/binops.pyx index b25bc347159..a8c66987307 100644 --- a/python/cudf/bindings/binops.pyx +++ b/python/cudf/bindings/binops.pyx @@ -19,6 +19,7 @@ _COMPILED_OPS = [ 'le', 'ge', 'and', 'or', 'xor' ] +# TODO: convert to single declaration of dictionary _BINARY_OP = {} _BINARY_OP['add'] = GDF_ADD _BINARY_OP['sub'] = GDF_SUB @@ -37,6 +38,8 @@ _BINARY_OP['ge'] = GDF_GREATER_EQUAL _BINARY_OP['and'] = GDF_BITWISE_AND _BINARY_OP['or'] = GDF_BITWISE_OR _BINARY_OP['xor'] = GDF_BITWISE_XOR +_BINARY_OP['l_and'] = GDF_LOGICAL_AND +_BINARY_OP['l_or'] = GDF_LOGICAL_OR cdef apply_jit_op(gdf_column* c_lhs, gdf_column* c_rhs, gdf_column* c_out, op): diff --git a/python/cudf/dataframe/numerical.py b/python/cudf/dataframe/numerical.py index 2a5495b4f0d..01a8ae5ab66 100644 --- a/python/cudf/dataframe/numerical.py +++ b/python/cudf/dataframe/numerical.py @@ -64,7 +64,11 @@ def deserialize(cls, deserialize, header, frames): def binary_operator(self, binop, rhs): if isinstance(rhs, NumericalColumn): - out_dtype = np.result_type(self.dtype, rhs.dtype) + if (np.issubdtype(self.dtype, np.bool_) or + np.issubdtype(rhs.dtype, np.bool_)): + out_dtype = np.bool_ + else: + out_dtype = np.result_type(self.dtype, rhs.dtype) return numeric_column_binop(lhs=self, rhs=rhs, op=binop, out_dtype=out_dtype) else: diff --git a/python/cudf/dataframe/series.py b/python/cudf/dataframe/series.py index c1ff7332ca7..0f782e780cd 100644 --- a/python/cudf/dataframe/series.py +++ b/python/cudf/dataframe/series.py @@ -417,6 +417,8 @@ def _binaryop(self, other, fn): """ from cudf import DataFrame if isinstance(other, DataFrame): + # i removed _binaryop from dataframe in #1292 (bitwise binary ops) + # that should've broken this. This needs a fix and a test return other._binaryop(self, fn) nvtx_range_push("CUDF_BINARY_OP", "orange") other = self._normalize_binop_value(other) @@ -502,9 +504,13 @@ def __rtruediv__(self, other): __div__ = __truediv__ def _bitwise_binop(self, other, op): + # TODO: replace self.dtype.type with self.dtype if (np.issubdtype(self.dtype.type, np.integer) and np.issubdtype(other.dtype.type, np.integer)): + # this broke Series (op) DataFrame because dataframe doesn't have dtype return self._binaryop(other, op) + # (1) Allow bool and check what happens. + # It worked but I'm not sure if it's right else: raise TypeError( f"Operation 'bitwise {op}' not supported between " @@ -515,12 +521,22 @@ def __and__(self, other): """Performs vectorized bitwise and (&) on corresponding elements of two series. """ + # Check if bool and pass logical-and in arg (requires logical-and in libcudf) + # Do this after (1) fails + # TODO: replace self.dtype.type with self.dtype + if (np.issubdtype(self.dtype.type, np.bool_) or + np.issubdtype(other.dtype.type, np.bool_)): + return self._binaryop(other, 'l_and') return self._bitwise_binop(other, 'and') def __or__(self, other): """Performs vectorized bitwise or (|) on corresponding elements of two series. """ + # Check if bool and pass logical-or in arg (requires logical-or in libcudf) + if (np.issubdtype(self.dtype.type, np.bool_) or + np.issubdtype(other.dtype.type, np.bool_)): + return self._binaryop(other, 'l_or') return self._bitwise_binop(other, 'or') def __xor__(self, other): @@ -584,6 +600,8 @@ def __invert__(self): """ if np.issubdtype(self.dtype.type, np.integer): return self._unaryop('not') + # Allow bool and check what happens. 0000 0001 will become 1111 1110 + # or maybe because of proper casting in C++, we'll get the correct result regardless else: raise TypeError( f"Operation `~` not supported on {self.dtype.type.__name__}" From d3278503a2b191d138cc726d312552faf3ce4386 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Mon, 20 May 2019 19:01:50 +0530 Subject: [PATCH 02/14] small cleanups in python binops --- python/cudf/bindings/binops.pyx | 42 ++++++++++++++++----------------- python/cudf/dataframe/series.py | 14 +++++------ 2 files changed, 27 insertions(+), 29 deletions(-) diff --git a/python/cudf/bindings/binops.pyx b/python/cudf/bindings/binops.pyx index a8c66987307..bb23f7f0d3d 100644 --- a/python/cudf/bindings/binops.pyx +++ b/python/cudf/bindings/binops.pyx @@ -20,27 +20,27 @@ _COMPILED_OPS = [ ] # TODO: convert to single declaration of dictionary -_BINARY_OP = {} -_BINARY_OP['add'] = GDF_ADD -_BINARY_OP['sub'] = GDF_SUB -_BINARY_OP['mul'] = GDF_MUL -_BINARY_OP['div'] = GDF_DIV -_BINARY_OP['truediv'] = GDF_TRUE_DIV -_BINARY_OP['floordiv'] = GDF_FLOOR_DIV -_BINARY_OP['mod'] = GDF_MOD -_BINARY_OP['pow'] = GDF_POW -_BINARY_OP['eq'] = GDF_EQUAL -_BINARY_OP['ne'] = GDF_NOT_EQUAL -_BINARY_OP['lt'] = GDF_LESS -_BINARY_OP['gt'] = GDF_GREATER -_BINARY_OP['le'] = GDF_LESS_EQUAL -_BINARY_OP['ge'] = GDF_GREATER_EQUAL -_BINARY_OP['and'] = GDF_BITWISE_AND -_BINARY_OP['or'] = GDF_BITWISE_OR -_BINARY_OP['xor'] = GDF_BITWISE_XOR -_BINARY_OP['l_and'] = GDF_LOGICAL_AND -_BINARY_OP['l_or'] = GDF_LOGICAL_OR - +_BINARY_OP = { + 'add' : GDF_ADD, + 'sub' : GDF_SUB, + 'mul' : GDF_MUL, + 'div' : GDF_DIV, + 'truediv' : GDF_TRUE_DIV, + 'floordiv' : GDF_FLOOR_DIV, + 'mod' : GDF_MOD, + 'pow' : GDF_POW, + 'eq' : GDF_EQUAL, + 'ne' : GDF_NOT_EQUAL, + 'lt' : GDF_LESS, + 'gt' : GDF_GREATER, + 'le' : GDF_LESS_EQUAL, + 'ge' : GDF_GREATER_EQUAL, + 'and' : GDF_BITWISE_AND, + 'or' : GDF_BITWISE_OR, + 'xor' : GDF_BITWISE_XOR, + 'l_and' : GDF_LOGICAL_AND, + 'l_or' : GDF_LOGICAL_OR, +} cdef apply_jit_op(gdf_column* c_lhs, gdf_column* c_rhs, gdf_column* c_out, op): """ diff --git a/python/cudf/dataframe/series.py b/python/cudf/dataframe/series.py index 0f782e780cd..331fe95ecdf 100644 --- a/python/cudf/dataframe/series.py +++ b/python/cudf/dataframe/series.py @@ -504,9 +504,8 @@ def __rtruediv__(self, other): __div__ = __truediv__ def _bitwise_binop(self, other, op): - # TODO: replace self.dtype.type with self.dtype - if (np.issubdtype(self.dtype.type, np.integer) and - np.issubdtype(other.dtype.type, np.integer)): + if (np.issubdtype(self.dtype, np.integer) and + np.issubdtype(other.dtype, np.integer)): # this broke Series (op) DataFrame because dataframe doesn't have dtype return self._binaryop(other, op) # (1) Allow bool and check what happens. @@ -523,9 +522,8 @@ def __and__(self, other): """ # Check if bool and pass logical-and in arg (requires logical-and in libcudf) # Do this after (1) fails - # TODO: replace self.dtype.type with self.dtype - if (np.issubdtype(self.dtype.type, np.bool_) or - np.issubdtype(other.dtype.type, np.bool_)): + if (np.issubdtype(self.dtype, np.bool_) or + np.issubdtype(other.dtype, np.bool_)): return self._binaryop(other, 'l_and') return self._bitwise_binop(other, 'and') @@ -534,8 +532,8 @@ def __or__(self, other): series. """ # Check if bool and pass logical-or in arg (requires logical-or in libcudf) - if (np.issubdtype(self.dtype.type, np.bool_) or - np.issubdtype(other.dtype.type, np.bool_)): + if (np.issubdtype(self.dtype, np.bool_) or + np.issubdtype(other.dtype, np.bool_)): return self._binaryop(other, 'l_or') return self._bitwise_binop(other, 'or') From f1e2c57bde3535d93967acb20878a0446bca8336 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Tue, 21 May 2019 16:54:20 +0530 Subject: [PATCH 03/14] logical binop pytest and doc update --- python/cudf/dataframe/series.py | 9 ++------- python/cudf/tests/test_binops.py | 30 ++++++++++++++++++++++++++++++ 2 files changed, 32 insertions(+), 7 deletions(-) diff --git a/python/cudf/dataframe/series.py b/python/cudf/dataframe/series.py index 331fe95ecdf..5eb84505cd7 100644 --- a/python/cudf/dataframe/series.py +++ b/python/cudf/dataframe/series.py @@ -508,8 +508,6 @@ def _bitwise_binop(self, other, op): np.issubdtype(other.dtype, np.integer)): # this broke Series (op) DataFrame because dataframe doesn't have dtype return self._binaryop(other, op) - # (1) Allow bool and check what happens. - # It worked but I'm not sure if it's right else: raise TypeError( f"Operation 'bitwise {op}' not supported between " @@ -518,10 +516,8 @@ def _bitwise_binop(self, other, op): def __and__(self, other): """Performs vectorized bitwise and (&) on corresponding elements of two - series. + series. Performs logical AND if one of the series is of bool dtype. """ - # Check if bool and pass logical-and in arg (requires logical-and in libcudf) - # Do this after (1) fails if (np.issubdtype(self.dtype, np.bool_) or np.issubdtype(other.dtype, np.bool_)): return self._binaryop(other, 'l_and') @@ -529,9 +525,8 @@ def __and__(self, other): def __or__(self, other): """Performs vectorized bitwise or (|) on corresponding elements of two - series. + series. Performs logical OR if one of the series is of bool dtype. """ - # Check if bool and pass logical-or in arg (requires logical-or in libcudf) if (np.issubdtype(self.dtype, np.bool_) or np.issubdtype(other.dtype, np.bool_)): return self._binaryop(other, 'l_or') diff --git a/python/cudf/tests/test_binops.py b/python/cudf/tests/test_binops.py index 74bd229b277..02d2d483a4d 100644 --- a/python/cudf/tests/test_binops.py +++ b/python/cudf/tests/test_binops.py @@ -98,6 +98,36 @@ def test_series_bitwise_binop(binop, obj_class, lhs_dtype, rhs_dtype): np.testing.assert_almost_equal(result.to_array(), binop(arr1, arr2)) +_logical_binops = [ + operator.and_, + operator.or_, +] + + +@pytest.mark.parametrize('obj_class', ['Series', 'Index']) +@pytest.mark.parametrize('binop', _logical_binops) +def test_series_logical_binop(binop, obj_class): + import pandas as pd + + arr1 = pd.Series(np.random.choice([True, False], 100)) + sr1 = Series(arr1) + + arr2 = pd.Series(np.random.choice([True, False], 100)) + sr2 = Series(arr2) + + if obj_class == 'Index': + sr1 = as_index(sr1) + sr2 = as_index(sr2) + + result = binop(sr1, sr2) + expect = binop(arr1, arr2) + + if obj_class == 'Index': + result = Series(result) + + np.testing.assert_almost_equal(result.to_array(), expect.to_numpy()) + + _cmpops = [ operator.lt, operator.gt, From fa97550b77aa0ae0803198aab0b061f5263c4a26 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Wed, 22 May 2019 00:27:43 +0530 Subject: [PATCH 04/14] Added logical NOT (!) to libcudf along with gtest --- cpp/include/cudf/types.h | 1 + cpp/src/unary/math_ops.cu | 41 +++++++++++++++++++++++++++++ cpp/src/utilities/wrapper_types.hpp | 7 +++++ cpp/tests/unary/unary_ops_test.cu | 33 +++++++++++++++++++++++ 4 files changed, 82 insertions(+) diff --git a/cpp/include/cudf/types.h b/cpp/include/cudf/types.h index 627fcc2898a..dbfdb76cac2 100644 --- a/cpp/include/cudf/types.h +++ b/cpp/include/cudf/types.h @@ -247,6 +247,7 @@ typedef enum { GDF_FLOOR, ///< largest integer value not greater than arg GDF_ABS, ///< Absolute value GDF_BIT_INVERT, ///< Bitwise Not (~) + GDF_NOT, ///< Logical Not (!) } gdf_unary_math_op; diff --git a/cpp/src/unary/math_ops.cu b/cpp/src/unary/math_ops.cu index 8d42aed6af0..21060a1f3e6 100644 --- a/cpp/src/unary/math_ops.cu +++ b/cpp/src/unary/math_ops.cu @@ -135,6 +135,16 @@ struct DeviceInvert { } }; +// logical op + +struct DeviceNot { + template + __device__ + cudf::bool8 apply(T data) { + return static_cast( !data ); + } +}; + template static gdf_error launch(gdf_column *input, gdf_column *output) { @@ -174,6 +184,33 @@ struct BitwiseOpDispatcher { }; +template +struct LogicalOpDispatcher { +private: + template + static constexpr bool is_supported() { + return std::is_arithmetic::value || + std::is_same::value; + + // TODO: try using member detector + // std::is_member_function_pointer::value; + } + +public: + template + typename std::enable_if_t(), gdf_error> + operator()(gdf_column *input, gdf_column *output) { + return cudf::unary::Launcher::launch(input, output); + } + + template + typename std::enable_if_t(), gdf_error> + operator()(gdf_column *input, gdf_column *output) { + return GDF_UNSUPPORTED_DTYPE; + } +}; + + gdf_error gdf_unary_math(gdf_column *input, gdf_column *output, gdf_unary_math_op op) { cudf::unary::handleChecksAndValidity(input, output); @@ -230,6 +267,10 @@ gdf_error gdf_unary_math(gdf_column *input, gdf_column *output, gdf_unary_math_o return cudf::type_dispatcher(input->dtype, BitwiseOpDispatcher{}, input, output); + case GDF_NOT: + return cudf::type_dispatcher(input->dtype, + LogicalOpDispatcher{}, + input, output); default: return GDF_INVALID_API_CALL; } diff --git a/cpp/src/utilities/wrapper_types.hpp b/cpp/src/utilities/wrapper_types.hpp index b955e38484a..3846751057b 100644 --- a/cpp/src/utilities/wrapper_types.hpp +++ b/cpp/src/utilities/wrapper_types.hpp @@ -460,6 +460,13 @@ cudf::bool8& operator/=(cudf::bool8 &lhs, cudf::bool8 const &rhs) return lhs; } +template +CUDA_HOST_DEVICE_CALLABLE +cudf::bool8 operator!(wrapper const& me) +{ + return static_cast( ! static_cast(me.value) ); +} + } // namespace detail } // namespace cudf diff --git a/cpp/tests/unary/unary_ops_test.cu b/cpp/tests/unary/unary_ops_test.cu index a775333db3e..86f821282c8 100644 --- a/cpp/tests/unary/unary_ops_test.cu +++ b/cpp/tests/unary/unary_ops_test.cu @@ -2438,3 +2438,36 @@ TEST_F(gdf_timestamp_casting_TEST, timestamp_to_timestamp) { EXPECT_TRUE( outputCol == expectOut ); } } + +template +struct gdf_logical_test : public GdfTest {}; + +using type_list = ::testing::Types< + int8_t, int16_t, int32_t, int64_t, float, double, cudf::bool8>; + +TYPED_TEST_CASE(gdf_logical_test, type_list); + +TYPED_TEST(gdf_logical_test, LogicalNot) { + const int colSize = 1000; + + // init input vector + std::vector h_input_v(colSize); + initialize_vector(h_input_v, colSize, 10, false); + + auto inputCol = cudf::test::column_wrapper{h_input_v}; + auto outputCol = cudf::test::column_wrapper{colSize}; + + std::vector h_expect_v{colSize}; + + // compute NOT + for (gdf_size_type i = 0; i < colSize; ++i) + h_expect_v[i] = static_cast( !h_input_v[i] ); + + // Use vector to build expected output + auto expectCol = cudf::test::column_wrapper{h_expect_v}; + + auto error = gdf_unary_math(inputCol, outputCol, GDF_NOT); + EXPECT_EQ(error, GDF_SUCCESS); + + EXPECT_EQ(expectCol, outputCol); +} \ No newline at end of file From 09c67b4eb1a3a99b5d05f6375d3c5ed458fef5b6 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Wed, 22 May 2019 10:36:50 +0530 Subject: [PATCH 05/14] Add python bindings and cuDF methods for logical NOT --- python/cudf/bindings/unaryops.pxd | 1 + python/cudf/bindings/unaryops.pyx | 3 ++- python/cudf/dataframe/series.py | 9 +++++---- python/cudf/tests/test_unaops.py | 10 +++++++++- 4 files changed, 17 insertions(+), 6 deletions(-) diff --git a/python/cudf/bindings/unaryops.pxd b/python/cudf/bindings/unaryops.pxd index c9c12d09a61..f4d2534442b 100644 --- a/python/cudf/bindings/unaryops.pxd +++ b/python/cudf/bindings/unaryops.pxd @@ -24,6 +24,7 @@ cdef extern from "cudf.h" nogil: GDF_FLOOR, GDF_ABS, GDF_BIT_INVERT, + GDF_NOT, cdef gdf_error gdf_unary_math(gdf_column *input, gdf_column *output, gdf_unary_math_op op) except + diff --git a/python/cudf/bindings/unaryops.pyx b/python/cudf/bindings/unaryops.pyx index 6b2983756be..0fbdb3269d8 100644 --- a/python/cudf/bindings/unaryops.pyx +++ b/python/cudf/bindings/unaryops.pyx @@ -26,7 +26,8 @@ _MATH_OP = { 'ceil' : GDF_CEIL, 'floor' : GDF_FLOOR, 'abs' : GDF_ABS, - 'not' : GDF_BIT_INVERT, + 'invert': GDF_BIT_INVERT, + 'not' : GDF_NOT, } def apply_math_op(incol, outcol, op): diff --git a/python/cudf/dataframe/series.py b/python/cudf/dataframe/series.py index 5eb84505cd7..4718e7ed77a 100644 --- a/python/cudf/dataframe/series.py +++ b/python/cudf/dataframe/series.py @@ -587,14 +587,15 @@ def __ge__(self, other): return self._ordered_compare(other, 'ge') def __invert__(self): - """Bitwise invert (~)/(not) for each element + """Bitwise invert (~) for each element. + Logical NOT if dtype is bool Returns a new Series. """ - if np.issubdtype(self.dtype.type, np.integer): + if np.issubdtype(self.dtype, np.integer): + return self._unaryop('invert') + elif np.issubdtype(self.dtype, np.bool_): return self._unaryop('not') - # Allow bool and check what happens. 0000 0001 will become 1111 1110 - # or maybe because of proper casting in C++, we'll get the correct result regardless else: raise TypeError( f"Operation `~` not supported on {self.dtype.type.__name__}" diff --git a/python/cudf/tests/test_unaops.py b/python/cudf/tests/test_unaops.py index 04619bf10a8..5f72d2167a9 100644 --- a/python/cudf/tests/test_unaops.py +++ b/python/cudf/tests/test_unaops.py @@ -19,13 +19,21 @@ def test_series_abs(dtype): @pytest.mark.parametrize('dtype', [np.int8, np.int16, np.int32, np.int64]) -def test_series_not(dtype): +def test_series_invert(dtype): arr = (np.random.random(1000) * 100).astype(dtype) sr = Series(arr) np.testing.assert_equal((~sr).to_array(), np.invert(arr)) np.testing.assert_equal((~sr).to_array(), ~arr) +def test_series_not(): + import pandas as pd + arr = pd.Series(np.random.choice([True, False], 1000)) + sr = Series(arr) + np.testing.assert_equal((~sr).to_array(), np.logical_not(arr)) + np.testing.assert_equal((~sr).to_array(), ~arr) + + def test_series_neg(): arr = np.random.random(100) * 100 sr = Series(arr) From 3580d6cbbb6293dc28337785e13f903ab06af88c Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Wed, 22 May 2019 18:40:10 +0530 Subject: [PATCH 06/14] gtests for logical binops --- .../binary-operation-integration-test.cpp | 51 ++++++++++++++++--- cpp/tests/binary/util/operation.h | 14 +++++ 2 files changed, 57 insertions(+), 8 deletions(-) diff --git a/cpp/tests/binary/integration/binary-operation-integration-test.cpp b/cpp/tests/binary/integration/binary-operation-integration-test.cpp index b8a7b652413..3865a17bc42 100644 --- a/cpp/tests/binary/integration/binary-operation-integration-test.cpp +++ b/cpp/tests/binary/integration/binary-operation-integration-test.cpp @@ -269,13 +269,13 @@ TEST_F(BinaryOperationIntegrationTest, Pow_Vector_Vector_SI64) { TEST_F(BinaryOperationIntegrationTest, And_Vector_Vector_SI16_SI64_SI32) { using AND = cudf::library::operation::BitwiseAnd; - auto lhs = cudf::test::column_wrapper{500, + auto lhs = cudf::test::column_wrapper{500, [](gdf_size_type row) {return row;}, [](gdf_size_type row) {return (row % 6 > 0);}}; - auto rhs = cudf::test::column_wrapper{500, - [](gdf_size_type row) {return 2;}, + auto rhs = cudf::test::column_wrapper{500, + [](gdf_size_type row) {return 2 * row + 3;}, [](gdf_size_type row) {return (row % 4 > 0);}}; - auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_BITWISE_AND); ASSERT_TRUE(result == GDF_SUCCESS); @@ -287,13 +287,13 @@ TEST_F(BinaryOperationIntegrationTest, And_Vector_Vector_SI16_SI64_SI32) { TEST_F(BinaryOperationIntegrationTest, Or_Vector_Vector_SI64_SI16_SI32) { using OR = cudf::library::operation::BitwiseOr; - auto lhs = cudf::test::column_wrapper{500, + auto lhs = cudf::test::column_wrapper{500, [](gdf_size_type row) {return row;}, [](gdf_size_type row) {return (row % 6 > 0);}}; - auto rhs = cudf::test::column_wrapper{500, - [](gdf_size_type row) {return 2;}, + auto rhs = cudf::test::column_wrapper{500, + [](gdf_size_type row) {return 2 * row + 3;}, [](gdf_size_type row) {return (row % 4 > 0);}}; - auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_BITWISE_OR); ASSERT_TRUE(result == GDF_SUCCESS); @@ -319,6 +319,41 @@ TEST_F(BinaryOperationIntegrationTest, Xor_Vector_Vector_SI32_SI16_SI64) { ASSERT_BINOP(out, lhs, rhs, XOR()); } + +TEST_F(BinaryOperationIntegrationTest, Logical_And_Vector_Vector_SI16_FP64_B8) { + using AND = cudf::library::operation::LogicalAnd; + + auto lhs = cudf::test::column_wrapper{500, + [](gdf_size_type row) {return (row % 5);}, + [](gdf_size_type row) {return (row % 6 > 0);}}; + auto rhs = cudf::test::column_wrapper{500, + [](gdf_size_type row) {return cudf::bool8{row % 3 > 0};}, + [](gdf_size_type row) {return (row % 4 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_LOGICAL_AND); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, AND()); +} + +TEST_F(BinaryOperationIntegrationTest, Logical_Or_Vector_Vector_B8_SI16_FP32) { + using OR = cudf::library::operation::LogicalOr; + + auto lhs = cudf::test::column_wrapper{500, + [](gdf_size_type row) {return (row % 5);}, + [](gdf_size_type row) {return (row % 6 > 0);}}; + auto rhs = cudf::test::column_wrapper{500, + [](gdf_size_type row) {return (row % 3 > 0);}, + [](gdf_size_type row) {return (row % 4 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_LOGICAL_OR); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, OR()); +} + } // namespace binop } // namespace test } // namespace cudf diff --git a/cpp/tests/binary/util/operation.h b/cpp/tests/binary/util/operation.h index f6f5a16c8ba..b1b14cfc17d 100644 --- a/cpp/tests/binary/util/operation.h +++ b/cpp/tests/binary/util/operation.h @@ -129,6 +129,20 @@ namespace operation { } }; + template + struct LogicalAnd { + TypeOut operator()(TypeLhs lhs, TypeRhs rhs) { + return TypeOut(lhs && rhs); + } + }; + + template + struct LogicalOr { + TypeOut operator()(TypeLhs lhs, TypeRhs rhs) { + return TypeOut(lhs || rhs); + } + }; + } // namespace operation } // namespace library } // namespace cudf From 39cde1a8abf48a47c8a7cc5515275f6e198e09e1 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Wed, 22 May 2019 22:20:57 +0530 Subject: [PATCH 07/14] Changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 6e679e2d492..da42e4459e3 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -9,6 +9,7 @@ - PR #1310 Implemented the slice/split functionality. - PR #1630 Add Python layer to the GPU-accelerated JSON reader - PR #1745 Add rounding of numeric columns via Numba +- PR #1745 Logical operators (AND, OR, NOT) for libcudf and cuDF ## Improvements From 79babfa23de72c5e17045a5e5fb4ff2f1dcc0118 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Thu, 23 May 2019 15:41:44 +0530 Subject: [PATCH 08/14] Fix CI issues, implement PR review changes. --- cpp/include/cudf/types.h | 1 + python/cudf/bindings/binops.pxd | 1 + python/cudf/bindings/unaryops.pxd | 1 + python/cudf/dataframe/series.py | 6 +++--- python/cudf/tests/test_binops.py | 2 +- 5 files changed, 7 insertions(+), 4 deletions(-) diff --git a/cpp/include/cudf/types.h b/cpp/include/cudf/types.h index 99845f666d9..60b6c4c351f 100644 --- a/cpp/include/cudf/types.h +++ b/cpp/include/cudf/types.h @@ -248,6 +248,7 @@ typedef enum { GDF_ABS, ///< Absolute value GDF_BIT_INVERT, ///< Bitwise Not (~) GDF_NOT, ///< Logical Not (!) + GDF_INVALID_UNARY ///< invalid operation } gdf_unary_math_op; diff --git a/python/cudf/bindings/binops.pxd b/python/cudf/bindings/binops.pxd index a5b8d1eb7d5..76404bff828 100644 --- a/python/cudf/bindings/binops.pxd +++ b/python/cudf/bindings/binops.pxd @@ -30,6 +30,7 @@ cdef extern from "cudf.h" nogil: GDF_BITWISE_XOR, GDF_LOGICAL_AND, GDF_LOGICAL_OR, + GDF_INVALID_BINARY cdef gdf_error gdf_binary_operation_s_v(gdf_column* out, gdf_scalar* lhs, gdf_column* rhs, gdf_binary_operator ope) except + cdef gdf_error gdf_binary_operation_v_s(gdf_column* out, gdf_column* lhs, gdf_scalar* rhs, gdf_binary_operator ope) except + diff --git a/python/cudf/bindings/unaryops.pxd b/python/cudf/bindings/unaryops.pxd index f4d2534442b..36e00b27768 100644 --- a/python/cudf/bindings/unaryops.pxd +++ b/python/cudf/bindings/unaryops.pxd @@ -25,6 +25,7 @@ cdef extern from "cudf.h" nogil: GDF_ABS, GDF_BIT_INVERT, GDF_NOT, + GDF_INVALID_UNARY cdef gdf_error gdf_unary_math(gdf_column *input, gdf_column *output, gdf_unary_math_op op) except + diff --git a/python/cudf/dataframe/series.py b/python/cudf/dataframe/series.py index 2db0128664a..be54c5da224 100644 --- a/python/cudf/dataframe/series.py +++ b/python/cudf/dataframe/series.py @@ -430,9 +430,9 @@ def _binaryop(self, other, fn): """ from cudf import DataFrame if isinstance(other, DataFrame): - # i removed _binaryop from dataframe in #1292 (bitwise binary ops) - # that should've broken this. This needs a fix and a test - return other._binaryop(self, fn) + # TODO: fn is not the same as arg expected by _apply_op + # e.g. for fn = 'and', _apply_op equivalent is '__and__' + return other._apply_op(self, fn) nvtx_range_push("CUDF_BINARY_OP", "orange") other = self._normalize_binop_value(other) outcol = self._column.binary_operator(fn, other._column) diff --git a/python/cudf/tests/test_binops.py b/python/cudf/tests/test_binops.py index 02d2d483a4d..39c8a612545 100644 --- a/python/cudf/tests/test_binops.py +++ b/python/cudf/tests/test_binops.py @@ -125,7 +125,7 @@ def test_series_logical_binop(binop, obj_class): if obj_class == 'Index': result = Series(result) - np.testing.assert_almost_equal(result.to_array(), expect.to_numpy()) + np.testing.assert_almost_equal(result.to_array(), expect) _cmpops = [ From 2553ce53c78cbaa6fce6b9fd9e3b89c20be4b557 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Thu, 23 May 2019 15:43:57 +0530 Subject: [PATCH 09/14] style fix --- python/cudf/dataframe/series.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/python/cudf/dataframe/series.py b/python/cudf/dataframe/series.py index be54c5da224..dfe5fc4eb9b 100644 --- a/python/cudf/dataframe/series.py +++ b/python/cudf/dataframe/series.py @@ -519,7 +519,8 @@ def __rtruediv__(self, other): def _bitwise_binop(self, other, op): if (np.issubdtype(self.dtype, np.integer) and np.issubdtype(other.dtype, np.integer)): - # this broke Series (op) DataFrame because dataframe doesn't have dtype + # TODO: This doesn't work on Series (op) DataFrame + # because dataframe doesn't have dtype return self._binaryop(other, op) else: raise TypeError( From 6f5c0c31e82e6a6e6a17b8733ece58c4cbf8c4b7 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Sat, 25 May 2019 01:36:10 +0530 Subject: [PATCH 10/14] Make ops more closely resemble pandas (& is bitwise) and add ufuncs. --- python/cudf/__init__.py | 3 ++- python/cudf/dataframe/numerical.py | 10 ++++----- python/cudf/dataframe/series.py | 35 ++++++++++++++++++++++-------- python/cudf/ops.py | 21 ++++++++++++++++++ python/cudf/tests/test_binops.py | 31 +++++++++++++------------- python/cudf/tests/test_unaops.py | 14 +++++++++--- 6 files changed, 81 insertions(+), 33 deletions(-) diff --git a/python/cudf/__init__.py b/python/cudf/__init__.py index 1e1a29e47a3..5d216f80d61 100644 --- a/python/cudf/__init__.py +++ b/python/cudf/__init__.py @@ -10,7 +10,8 @@ read_hdf, read_orc, from_dlpack) from cudf.settings import set_options from cudf.reshape import melt -from cudf.ops import (sqrt, sin, cos, tan, arcsin, arccos, arctan, exp, log) +from cudf.ops import (sqrt, sin, cos, tan, arcsin, arccos, arctan, exp, log, + logical_not, logical_and, logical_or) from librmm_cffi import librmm as rmm diff --git a/python/cudf/dataframe/numerical.py b/python/cudf/dataframe/numerical.py index a6dfbea4b80..a869810f3f8 100644 --- a/python/cudf/dataframe/numerical.py +++ b/python/cudf/dataframe/numerical.py @@ -64,11 +64,7 @@ def deserialize(cls, deserialize, header, frames): def binary_operator(self, binop, rhs): if isinstance(rhs, NumericalColumn): - if (np.issubdtype(self.dtype, np.bool_) or - np.issubdtype(rhs.dtype, np.bool_)): - out_dtype = np.bool_ - else: - out_dtype = np.result_type(self.dtype, rhs.dtype) + out_dtype = np.result_type(self.dtype, rhs.dtype) return numeric_column_binop(lhs=self, rhs=rhs, op=binop, out_dtype=out_dtype) else: @@ -79,6 +75,10 @@ def unary_operator(self, unaryop): return numeric_column_unaryop(self, op=unaryop, out_dtype=self.dtype) + def unary_logic_op(self, unaryop): + return numeric_column_unaryop(self, op=unaryop, + out_dtype=np.bool_) + def unordered_compare(self, cmpop, rhs): return numeric_column_compare(self, rhs, op=cmpop) diff --git a/python/cudf/dataframe/series.py b/python/cudf/dataframe/series.py index dfe5fc4eb9b..6eef278b985 100644 --- a/python/cudf/dataframe/series.py +++ b/python/cudf/dataframe/series.py @@ -517,11 +517,22 @@ def __rtruediv__(self, other): __div__ = __truediv__ def _bitwise_binop(self, other, op): - if (np.issubdtype(self.dtype, np.integer) and - np.issubdtype(other.dtype, np.integer)): + if ( + np.issubdtype(self.dtype, np.bool_) + or np.issubdtype(self.dtype, np.integer) + ) and ( + np.issubdtype(other.dtype, np.bool_) + or np.issubdtype(other.dtype, np.integer) + ): # TODO: This doesn't work on Series (op) DataFrame # because dataframe doesn't have dtype - return self._binaryop(other, op) + ser = self._binaryop(other, op) + if ( + np.issubdtype(self.dtype, np.bool_) + or np.issubdtype(other.dtype, np.bool_) + ): + ser = ser.astype(np.bool_) + return ser else: raise TypeError( f"Operation 'bitwise {op}' not supported between " @@ -532,18 +543,12 @@ def __and__(self, other): """Performs vectorized bitwise and (&) on corresponding elements of two series. Performs logical AND if one of the series is of bool dtype. """ - if (np.issubdtype(self.dtype, np.bool_) or - np.issubdtype(other.dtype, np.bool_)): - return self._binaryop(other, 'l_and') return self._bitwise_binop(other, 'and') def __or__(self, other): """Performs vectorized bitwise or (|) on corresponding elements of two series. Performs logical OR if one of the series is of bool dtype. """ - if (np.issubdtype(self.dtype, np.bool_) or - np.issubdtype(other.dtype, np.bool_)): - return self._binaryop(other, 'l_or') return self._bitwise_binop(other, 'or') def __xor__(self, other): @@ -552,6 +557,18 @@ def __xor__(self, other): """ return self._bitwise_binop(other, 'xor') + def logical_and(self, other): + ser = self._binaryop(other, 'l_and') + return ser.astype(np.bool_) + + def logical_or(self, other): + ser = self._binaryop(other, 'l_or') + return ser.astype(np.bool_) + + def logical_not(self): + outcol = self._column.unary_logic_op('not') + return self._copy_construct(data=outcol) + def _normalize_binop_value(self, other): if isinstance(other, Series): return other diff --git a/python/cudf/ops.py b/python/cudf/ops.py index 5e83e6b59bf..3e48a4e8886 100644 --- a/python/cudf/ops.py +++ b/python/cudf/ops.py @@ -68,3 +68,24 @@ def sqrt(arbitrary): return np.sqrt(arbitrary) else: return getattr(arbitrary, 'sqrt')() + + +def logical_not(arbitrary): + if isinstance(arbitrary, Number): + return np.logical_not(arbitrary) + else: + return getattr(arbitrary, 'logical_not')() + + +def logical_and(lhs, rhs): + if isinstance(lhs, Number) and isinstance(rhs, Number): + return np.logical_and(lhs, rhs) + else: + return getattr(lhs, 'logical_and')(rhs) + + +def logical_or(lhs, rhs): + if isinstance(lhs, Number) and isinstance(rhs, Number): + return np.logical_or(lhs, rhs) + else: + return getattr(lhs, 'logical_or')(rhs) diff --git a/python/cudf/tests/test_binops.py b/python/cudf/tests/test_binops.py index 39c8a612545..749b493bf3c 100644 --- a/python/cudf/tests/test_binops.py +++ b/python/cudf/tests/test_binops.py @@ -8,6 +8,7 @@ import pytest import numpy as np +import cudf from cudf.dataframe import Series from cudf.dataframe.index import as_index @@ -99,32 +100,32 @@ def test_series_bitwise_binop(binop, obj_class, lhs_dtype, rhs_dtype): _logical_binops = [ - operator.and_, - operator.or_, + (operator.and_, operator.and_), + (operator.or_, operator.or_), + (np.logical_and, cudf.logical_and), + (np.logical_or, cudf.logical_or), ] -@pytest.mark.parametrize('obj_class', ['Series', 'Index']) -@pytest.mark.parametrize('binop', _logical_binops) -def test_series_logical_binop(binop, obj_class): +@pytest.mark.parametrize('lhstype', _int_types + [np.bool_]) +@pytest.mark.parametrize('rhstype', _int_types + [np.bool_]) +@pytest.mark.parametrize('binop,cubinop', _logical_binops) +def test_series_logical_binop(lhstype, rhstype, binop, cubinop): import pandas as pd - arr1 = pd.Series(np.random.choice([True, False], 100)) + arr1 = pd.Series(np.random.choice([True, False], 10)) + if lhstype is not np.bool_: + arr1 = arr1 * (np.random.random(10) * 100).astype(lhstype) sr1 = Series(arr1) - arr2 = pd.Series(np.random.choice([True, False], 100)) + arr2 = pd.Series(np.random.choice([True, False], 10)) + if rhstype is not np.bool_: + arr2 = arr2 * (np.random.random(10) * 100).astype(rhstype) sr2 = Series(arr2) - if obj_class == 'Index': - sr1 = as_index(sr1) - sr2 = as_index(sr2) - - result = binop(sr1, sr2) + result = cubinop(sr1, sr2) expect = binop(arr1, arr2) - if obj_class == 'Index': - result = Series(result) - np.testing.assert_almost_equal(result.to_array(), expect) diff --git a/python/cudf/tests/test_unaops.py b/python/cudf/tests/test_unaops.py index 5f72d2167a9..0110c597595 100644 --- a/python/cudf/tests/test_unaops.py +++ b/python/cudf/tests/test_unaops.py @@ -4,6 +4,7 @@ import numpy as np import pandas as pd +import cudf from cudf.dataframe import Series @@ -26,11 +27,18 @@ def test_series_invert(dtype): np.testing.assert_equal((~sr).to_array(), ~arr) -def test_series_not(): +@pytest.mark.parametrize('dtype', + [np.int8, np.int16, np.int32, np.int64, np.bool_]) +def test_series_not(dtype): import pandas as pd - arr = pd.Series(np.random.choice([True, False], 1000)) + arr = pd.Series(np.random.choice([True, False], 1000)).astype(dtype) + if dtype is not np.bool_: + arr = arr * (np.random.random(1000) * 100).astype(dtype) sr = Series(arr) - np.testing.assert_equal((~sr).to_array(), np.logical_not(arr)) + + result = cudf.logical_not(sr).to_array() + expect = np.logical_not(arr) + np.testing.assert_equal(result, expect) np.testing.assert_equal((~sr).to_array(), ~arr) From 0a4f0b1361c6c2c98dfa624ade1303095caaf577 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Mon, 27 May 2019 16:57:00 +0530 Subject: [PATCH 11/14] Replaced braces with round brackets in constructor to prevent initialiser list confusion. --- .../binary-operation-integration-test.cpp | 20 +++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/cpp/tests/binary/integration/binary-operation-integration-test.cpp b/cpp/tests/binary/integration/binary-operation-integration-test.cpp index 3865a17bc42..7f795d2f6cf 100644 --- a/cpp/tests/binary/integration/binary-operation-integration-test.cpp +++ b/cpp/tests/binary/integration/binary-operation-integration-test.cpp @@ -323,13 +323,13 @@ TEST_F(BinaryOperationIntegrationTest, Xor_Vector_Vector_SI32_SI16_SI64) { TEST_F(BinaryOperationIntegrationTest, Logical_And_Vector_Vector_SI16_FP64_B8) { using AND = cudf::library::operation::LogicalAnd; - auto lhs = cudf::test::column_wrapper{500, + auto lhs = cudf::test::column_wrapper(500, [](gdf_size_type row) {return (row % 5);}, - [](gdf_size_type row) {return (row % 6 > 0);}}; - auto rhs = cudf::test::column_wrapper{500, + [](gdf_size_type row) {return (row % 6 > 0);}); + auto rhs = cudf::test::column_wrapper(500, [](gdf_size_type row) {return cudf::bool8{row % 3 > 0};}, - [](gdf_size_type row) {return (row % 4 > 0);}}; - auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + [](gdf_size_type row) {return (row % 4 > 0);}); + auto out = cudf::test::column_wrapper(lhs.get()->size, true); auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_LOGICAL_AND); ASSERT_TRUE(result == GDF_SUCCESS); @@ -340,13 +340,13 @@ TEST_F(BinaryOperationIntegrationTest, Logical_And_Vector_Vector_SI16_FP64_B8) { TEST_F(BinaryOperationIntegrationTest, Logical_Or_Vector_Vector_B8_SI16_FP32) { using OR = cudf::library::operation::LogicalOr; - auto lhs = cudf::test::column_wrapper{500, + auto lhs = cudf::test::column_wrapper(500, [](gdf_size_type row) {return (row % 5);}, - [](gdf_size_type row) {return (row % 6 > 0);}}; - auto rhs = cudf::test::column_wrapper{500, + [](gdf_size_type row) {return (row % 6 > 0);}); + auto rhs = cudf::test::column_wrapper(500, [](gdf_size_type row) {return (row % 3 > 0);}, - [](gdf_size_type row) {return (row % 4 > 0);}}; - auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + [](gdf_size_type row) {return (row % 4 > 0);}); + auto out = cudf::test::column_wrapper(lhs.get()->size, true); auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_LOGICAL_OR); ASSERT_TRUE(result == GDF_SUCCESS); From 974a2daca911311061e0597da31e8ea6fbfa6ac5 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Mon, 27 May 2019 20:16:09 +0530 Subject: [PATCH 12/14] revert python methods' doc. Testing column_wrapper without using wrapper_types --- .../integration/binary-operation-integration-test.cpp | 8 ++++---- python/cudf/dataframe/series.py | 4 ++-- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/cpp/tests/binary/integration/binary-operation-integration-test.cpp b/cpp/tests/binary/integration/binary-operation-integration-test.cpp index 7f795d2f6cf..19d55c84ab4 100644 --- a/cpp/tests/binary/integration/binary-operation-integration-test.cpp +++ b/cpp/tests/binary/integration/binary-operation-integration-test.cpp @@ -320,14 +320,14 @@ TEST_F(BinaryOperationIntegrationTest, Xor_Vector_Vector_SI32_SI16_SI64) { } -TEST_F(BinaryOperationIntegrationTest, Logical_And_Vector_Vector_SI16_FP64_B8) { - using AND = cudf::library::operation::LogicalAnd; +TEST_F(BinaryOperationIntegrationTest, Logical_And_Vector_Vector_SI16_FP64_SI8) { + using AND = cudf::library::operation::LogicalAnd; auto lhs = cudf::test::column_wrapper(500, [](gdf_size_type row) {return (row % 5);}, [](gdf_size_type row) {return (row % 6 > 0);}); - auto rhs = cudf::test::column_wrapper(500, - [](gdf_size_type row) {return cudf::bool8{row % 3 > 0};}, + auto rhs = cudf::test::column_wrapper(500, + [](gdf_size_type row) {return (row % 3 > 0);}, [](gdf_size_type row) {return (row % 4 > 0);}); auto out = cudf::test::column_wrapper(lhs.get()->size, true); diff --git a/python/cudf/dataframe/series.py b/python/cudf/dataframe/series.py index 6eef278b985..c004a8aecf8 100644 --- a/python/cudf/dataframe/series.py +++ b/python/cudf/dataframe/series.py @@ -541,13 +541,13 @@ def _bitwise_binop(self, other, op): def __and__(self, other): """Performs vectorized bitwise and (&) on corresponding elements of two - series. Performs logical AND if one of the series is of bool dtype. + series. """ return self._bitwise_binop(other, 'and') def __or__(self, other): """Performs vectorized bitwise or (|) on corresponding elements of two - series. Performs logical OR if one of the series is of bool dtype. + series. """ return self._bitwise_binop(other, 'or') From e90086c63b66e808fe989900fa1c80643ad47c18 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Mon, 27 May 2019 20:40:43 +0530 Subject: [PATCH 13/14] change the remaining column_wrapper to primitive type to fix centos error. change bintop pytests to use pandas testing --- .../binary/integration/binary-operation-integration-test.cpp | 2 +- python/cudf/tests/test_binops.py | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/tests/binary/integration/binary-operation-integration-test.cpp b/cpp/tests/binary/integration/binary-operation-integration-test.cpp index 19d55c84ab4..1b440d50503 100644 --- a/cpp/tests/binary/integration/binary-operation-integration-test.cpp +++ b/cpp/tests/binary/integration/binary-operation-integration-test.cpp @@ -346,7 +346,7 @@ TEST_F(BinaryOperationIntegrationTest, Logical_Or_Vector_Vector_B8_SI16_FP32) { auto rhs = cudf::test::column_wrapper(500, [](gdf_size_type row) {return (row % 3 > 0);}, [](gdf_size_type row) {return (row % 4 > 0);}); - auto out = cudf::test::column_wrapper(lhs.get()->size, true); + auto out = cudf::test::column_wrapper(lhs.get()->size, true); auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_LOGICAL_OR); ASSERT_TRUE(result == GDF_SUCCESS); diff --git a/python/cudf/tests/test_binops.py b/python/cudf/tests/test_binops.py index 24307f946ee..beb81e94c34 100644 --- a/python/cudf/tests/test_binops.py +++ b/python/cudf/tests/test_binops.py @@ -126,7 +126,7 @@ def test_series_logical_binop(lhstype, rhstype, binop, cubinop): result = cubinop(sr1, sr2) expect = binop(arr1, arr2) - np.testing.assert_almost_equal(result.to_array(), expect) + assert_eq(result, expect) _cmpops = [ From 5893d0cbd7e724d92fd893e5fdc9e5deff3c32ed Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Mon, 27 May 2019 20:50:30 +0530 Subject: [PATCH 14/14] fix binop pytest bug. forgot to rebuild test --- python/cudf/tests/test_binops.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/cudf/tests/test_binops.py b/python/cudf/tests/test_binops.py index beb81e94c34..facaa86907f 100644 --- a/python/cudf/tests/test_binops.py +++ b/python/cudf/tests/test_binops.py @@ -126,7 +126,7 @@ def test_series_logical_binop(lhstype, rhstype, binop, cubinop): result = cubinop(sr1, sr2) expect = binop(arr1, arr2) - assert_eq(result, expect) + utils.assert_eq(result, expect) _cmpops = [