From e0a24ef192f1fefc82a730c94a2a6f674b0d028d Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 28 Sep 2021 12:42:49 -0400 Subject: [PATCH 1/9] add benchmarks for copy-if-else and replace-nans --- cpp/benchmarks/CMakeLists.txt | 7 +- .../copying/copy_if_else_benchmark.cpp | 69 +++++++++++++++++++ cpp/benchmarks/replace/clamp_benchmark.cpp | 20 +++--- cpp/benchmarks/replace/nans_benchmark.cpp | 63 +++++++++++++++++ 4 files changed, 148 insertions(+), 11 deletions(-) create mode 100644 cpp/benchmarks/copying/copy_if_else_benchmark.cpp create mode 100644 cpp/benchmarks/replace/nans_benchmark.cpp diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt index b3b92003573..aeaee60086d 100644 --- a/cpp/benchmarks/CMakeLists.txt +++ b/cpp/benchmarks/CMakeLists.txt @@ -88,6 +88,10 @@ ConfigureBench(CONTIGUOUS_SPLIT_BENCH copying/contiguous_split_benchmark.cu) # - shift benchmark ------------------------------------------------------------------------------- ConfigureBench(SHIFT_BENCH copying/shift_benchmark.cu) +################################################################################################### +# - copy-if-else benchmark ----------------------------------------------------------------------------- +ConfigureBench(COPY_IF_ELSE_BENCH copying/copy_if_else_benchmark.cpp) + ################################################################################################### # - transpose benchmark --------------------------------------------------------------------------- ConfigureBench(TRANSPOSE_BENCH transpose/transpose_benchmark.cu) @@ -141,7 +145,8 @@ ConfigureBench(REDUCTION_BENCH ################################################################################################### # - reduction benchmark --------------------------------------------------------------------------- ConfigureBench(REPLACE_BENCH - replace/clamp_benchmark.cpp) + replace/clamp_benchmark.cpp + replace/nans_benchmark.cpp) ################################################################################################### # - filling benchmark ----------------------------------------------------------------------------- diff --git a/cpp/benchmarks/copying/copy_if_else_benchmark.cpp b/cpp/benchmarks/copying/copy_if_else_benchmark.cpp new file mode 100644 index 00000000000..0a5c425e601 --- /dev/null +++ b/cpp/benchmarks/copying/copy_if_else_benchmark.cpp @@ -0,0 +1,69 @@ +/* + * Copyright (c) 2021, 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 +#include +#include +#include + +#include + +#include + +// #include + +class CopyIfElse : public cudf::benchmark { +}; + +template +static void BM_copy_if_else(benchmark::State& state, bool nulls) +{ + cudf::size_type const n_rows{(cudf::size_type)state.range(0)}; + auto input_type = cudf::type_to_id(); + auto bool_type = cudf::type_id::BOOL8; + auto const input = create_random_table({input_type, input_type, bool_type}, 3, row_count{n_rows}); + + if (!nulls) { + input->get_column(2).set_null_mask(rmm::device_buffer{}, 0); + input->get_column(1).set_null_mask(rmm::device_buffer{}, 0); + input->get_column(0).set_null_mask(rmm::device_buffer{}, 0); + } + + cudf::column_view decision(input->view().column(2)); + cudf::column_view rhs(input->view().column(1)); + cudf::column_view lhs(input->view().column(0)); + + for (auto _ : state) { + cuda_event_timer raii(state, true, rmm::cuda_stream_default); + cudf::copy_if_else(lhs, rhs, decision); + } +} + +#define COPY_BENCHMARK_DEFINE(name, type, b) \ + BENCHMARK_DEFINE_F(CopyIfElse, name) \ + (::benchmark::State & st) { BM_copy_if_else(st, b); } \ + BENCHMARK_REGISTER_F(CopyIfElse, name) \ + ->RangeMultiplier(8) \ + ->Ranges({{1 << 12, 1 << 27}}) \ + ->UseManualTime() \ + ->Unit(benchmark::kMillisecond); + +COPY_BENCHMARK_DEFINE(int16, int16_t, true) +COPY_BENCHMARK_DEFINE(uint32, uint32_t, true) +COPY_BENCHMARK_DEFINE(float64, double, true) +COPY_BENCHMARK_DEFINE(int16_no_nulls, int16_t, false) +COPY_BENCHMARK_DEFINE(uint32_no_nulls, uint32_t, false) +COPY_BENCHMARK_DEFINE(float64_no_nulls, double, false) diff --git a/cpp/benchmarks/replace/clamp_benchmark.cpp b/cpp/benchmarks/replace/clamp_benchmark.cpp index f897b9d82cc..4d9da4aca6d 100644 --- a/cpp/benchmarks/replace/clamp_benchmark.cpp +++ b/cpp/benchmarks/replace/clamp_benchmark.cpp @@ -30,7 +30,7 @@ class ReplaceClamp : public cudf::benchmark { }; template -static void BM_reduction_scan(benchmark::State& state, bool include_nulls) +static void BM_clamp(benchmark::State& state, bool include_nulls) { cudf::size_type const n_rows{(cudf::size_type)state.range(0)}; auto const dtype = cudf::type_to_id(); @@ -58,15 +58,15 @@ static void BM_reduction_scan(benchmark::State& state, bool include_nulls) } } -#define CLAMP_BENCHMARK_DEFINE(name, type, nulls) \ - BENCHMARK_DEFINE_F(ReplaceClamp, name) \ - (::benchmark::State & state) { BM_reduction_scan(state, nulls); } \ - BENCHMARK_REGISTER_F(ReplaceClamp, name) \ - ->UseManualTime() \ - ->Arg(10000) /* 10k */ \ - ->Arg(100000) /* 100k */ \ - ->Arg(1000000) /* 1M */ \ - ->Arg(10000000) /* 10M */ \ +#define CLAMP_BENCHMARK_DEFINE(name, type, nulls) \ + BENCHMARK_DEFINE_F(ReplaceClamp, name) \ + (::benchmark::State & state) { BM_clamp(state, nulls); } \ + BENCHMARK_REGISTER_F(ReplaceClamp, name) \ + ->UseManualTime() \ + ->Arg(10000) /* 10k */ \ + ->Arg(100000) /* 100k */ \ + ->Arg(1000000) /* 1M */ \ + ->Arg(10000000) /* 10M */ \ ->Arg(100000000); /* 100M */ CLAMP_BENCHMARK_DEFINE(int8_no_nulls, int8_t, false); diff --git a/cpp/benchmarks/replace/nans_benchmark.cpp b/cpp/benchmarks/replace/nans_benchmark.cpp new file mode 100644 index 00000000000..a337ae5e7ad --- /dev/null +++ b/cpp/benchmarks/replace/nans_benchmark.cpp @@ -0,0 +1,63 @@ +/* + * Copyright (c) 2021, 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 +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +class ReplaceNans : public cudf::benchmark { +}; + +template +static void BM_replace_nans(benchmark::State& state, bool include_nulls) +{ + cudf::size_type const n_rows{(cudf::size_type)state.range(0)}; + auto const dtype = cudf::type_to_id(); + auto const table = create_random_table({dtype}, 1, row_count{n_rows}); + if (!include_nulls) { table->get_column(0).set_null_mask(rmm::device_buffer{}, 0); } + cudf::column_view input(table->view().column(0)); + + auto zero = cudf::make_fixed_width_scalar(0); + + for (auto _ : state) { + cuda_event_timer timer(state, true); + auto result = cudf::replace_nans(input, *zero); + } +} + +#define NANS_BENCHMARK_DEFINE(name, type, nulls) \ + BENCHMARK_DEFINE_F(ReplaceNans, name) \ + (::benchmark::State & state) { BM_replace_nans(state, nulls); } \ + BENCHMARK_REGISTER_F(ReplaceNans, name) \ + ->UseManualTime() \ + ->Arg(10000) /* 10k */ \ + ->Arg(100000) /* 100k */ \ + ->Arg(1000000) /* 1M */ \ + ->Arg(10000000) /* 10M */ \ + ->Arg(100000000); /* 100M */ + +NANS_BENCHMARK_DEFINE(float32_nulls, float, true); +NANS_BENCHMARK_DEFINE(float64_nulls, double, true); +NANS_BENCHMARK_DEFINE(float32_no_nulls, float, false); +NANS_BENCHMARK_DEFINE(float64_no_nulls, double, false); From 8f5d225893d33813750991b23848aa25136cabb9 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 28 Sep 2021 12:44:16 -0400 Subject: [PATCH 2/9] split up copy_tests into cu and cpp --- cpp/tests/CMakeLists.txt | 3 +- cpp/tests/copying/copy_tests.cpp | 614 +++++++++++++++++++++++++ cpp/tests/copying/detail_copy_tests.cu | 122 +++++ 3 files changed, 738 insertions(+), 1 deletion(-) create mode 100644 cpp/tests/copying/copy_tests.cpp create mode 100644 cpp/tests/copying/detail_copy_tests.cu diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index d9553d463ab..06d5864f3f1 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -214,7 +214,8 @@ ConfigureTest(COPYING_TEST copying/concatenate_tests.cu copying/copy_if_else_nested_tests.cpp copying/copy_range_tests.cpp - copying/copy_tests.cu + copying/copy_tests.cpp + copying/detail_copy_tests.cu copying/detail_gather_tests.cu copying/gather_list_tests.cpp copying/gather_str_tests.cpp diff --git a/cpp/tests/copying/copy_tests.cpp b/cpp/tests/copying/copy_tests.cpp new file mode 100644 index 00000000000..9a6a241d5c2 --- /dev/null +++ b/cpp/tests/copying/copy_tests.cpp @@ -0,0 +1,614 @@ +/* + * Copyright (c) 2019-2021, 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 +#include +#include +#include +#include + +#include +#include +#include +#include + +template +struct CopyTest : public cudf::test::BaseFixture { +}; + +TYPED_TEST_CASE(CopyTest, cudf::test::FixedWidthTypesWithoutFixedPoint); + +#define wrapper cudf::test::fixed_width_column_wrapper + +TYPED_TEST(CopyTest, CopyIfElseTestShort) +{ + using T = TypeParam; + + cudf::test::fixed_width_column_wrapper mask_w{1, 0, 0, 0}; + + wrapper lhs_w({5, 5, 5, 5}, {1, 1, 1, 1}); + wrapper rhs_w({6, 6, 6, 6}, {1, 1, 1, 1}); + wrapper expected_w({5, 6, 6, 6}); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTest, CopyIfElseTestManyNulls) +{ + using T = TypeParam; + + cudf::test::fixed_width_column_wrapper mask_w{{1, 0, 0, 0, 0, 0, 1}, {1, 1, 1, 1, 1, 1, 0}}; + + wrapper lhs_w({5, 5, 5, 5, 5, 5, 5}, {1, 1, 1, 1, 1, 1, 1}); + wrapper rhs_w({6, 6, 6, 6, 6, 6, 6}, {1, 0, 0, 0, 0, 0, 1}); + wrapper expected_w({5, 6, 6, 6, 6, 6, 6}, {1, 0, 0, 0, 0, 0, 1}); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTest, CopyIfElseTestLong) +{ + using T = TypeParam; + + // make sure we span at least 2 warps + int num_els = 64; + + bool mask[] = {1, 0, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, + 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); + + bool lhs_v[] = {1, 1, 1, 1, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; + wrapper lhs_w({5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, + 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, + 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5}, + lhs_v); + + bool rhs_v[] = {1, 1, 1, 1, 1, 1, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; + wrapper rhs_w({6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, + 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, + 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6}, + rhs_v); + + bool exp_v[] = {1, 1, 1, 1, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; + wrapper expected_w({5, 6, 5, 6, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 6, 6, 6, + 6, 5, 5, 5, 5, 5, 5, 5, 5, 5, 6, 6, 6, 6, 5, 5, 5, 5, 5, 5, 5, 5, + 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5}, + exp_v); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTest, CopyIfElseTestEmptyInputs) +{ + using T = TypeParam; + + cudf::test::fixed_width_column_wrapper mask_w{}; + + wrapper lhs_w{}; + wrapper rhs_w{}; + wrapper expected_w{}; + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTest, CopyIfElseMixedInputValidity) +{ + using T = TypeParam; + + cudf::test::fixed_width_column_wrapper mask_w{1, 0, 1, 1}; + + wrapper lhs_w({5, 5, 5, 5}, {1, 1, 1, 0}); + wrapper rhs_w({6, 6, 6, 6}, {1, 0, 1, 1}); + wrapper expected_w({5, 6, 5, 5}, {1, 0, 1, 0}); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTest, CopyIfElseMixedInputValidity2) +{ + using T = TypeParam; + + cudf::test::fixed_width_column_wrapper mask_w{1, 0, 1, 1}; + + wrapper lhs_w({5, 5, 5, 5}, {1, 1, 1, 0}); + wrapper rhs_w({6, 6, 6, 6}); + wrapper expected_w({5, 6, 5, 5}, {1, 1, 1, 0}); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTest, CopyIfElseMixedInputValidity3) +{ + using T = TypeParam; + + cudf::test::fixed_width_column_wrapper mask_w{1, 0, 1, 1}; + + wrapper lhs_w({5, 5, 5, 5}); + wrapper rhs_w({6, 6, 6, 6}, {1, 0, 1, 1}); + wrapper expected_w({5, 6, 5, 5}, {1, 0, 1, 1}); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTest, CopyIfElseMixedInputValidity4) +{ + using T = TypeParam; + + cudf::test::fixed_width_column_wrapper mask_w{1, 0, 1, 1}; + + wrapper lhs_w({5, 5, 5, 5}); + wrapper rhs_w({6, 6, 6, 6}); + wrapper expected_w({5, 6, 5, 5}); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTest, CopyIfElseBadInputLength) +{ + using T = TypeParam; + + // mask length mismatch + { + cudf::test::fixed_width_column_wrapper mask_w{1, 1, 1}; + + wrapper lhs_w({5, 5, 5, 5}); + wrapper rhs_w({6, 6, 6, 6}); + + EXPECT_THROW(cudf::copy_if_else(lhs_w, rhs_w, mask_w), cudf::logic_error); + } + + // column length mismatch + { + cudf::test::fixed_width_column_wrapper mask_w{1, 1, 1, 1}; + + wrapper lhs_w({5, 5, 5}); + wrapper rhs_w({6, 6, 6, 6}); + + EXPECT_THROW(cudf::copy_if_else(lhs_w, rhs_w, mask_w), cudf::logic_error); + } +} + +struct CopyEmptyNested : public cudf::test::BaseFixture { +}; + +TEST_F(CopyEmptyNested, CopyIfElseTestEmptyNestedColumns) +{ + // lists + { + cudf::test::lists_column_wrapper col{{{"abc", "def"}, {"xyz"}}}; + auto lhs = cudf::empty_like(col); + auto rhs = cudf::empty_like(col); + cudf::test::fixed_width_column_wrapper mask{}; + + auto expected = empty_like(col); + + auto out = cudf::copy_if_else(*lhs, *rhs, mask); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), *expected); + } + + // structs + { + cudf::test::lists_column_wrapper _col0{{{"abc", "def"}, {"xyz"}}}; + auto col0 = cudf::empty_like(_col0); + cudf::test::fixed_width_column_wrapper col1; + + std::vector> cols; + cols.push_back(std::move(col0)); + cols.push_back(col1.release()); + cudf::test::structs_column_wrapper struct_col(std::move(cols)); + auto lhs = cudf::empty_like(struct_col); + auto rhs = cudf::empty_like(struct_col); + + cudf::test::fixed_width_column_wrapper mask{}; + + auto expected = cudf::empty_like(struct_col); + + auto out = cudf::copy_if_else(*lhs, *rhs, mask); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), *expected); + } +} + +TEST_F(CopyEmptyNested, CopyIfElseTestEmptyNestedScalars) +{ + // lists + { + cudf::test::lists_column_wrapper _col{{{"abc", "def"}, {"xyz"}}}; + std::unique_ptr lhs = cudf::get_element(_col, 0); + std::unique_ptr rhs = cudf::get_element(_col, 0); + + cudf::test::fixed_width_column_wrapper mask{}; + + auto expected = empty_like(_col); + + auto out = cudf::copy_if_else(*lhs, *rhs, mask); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), *expected); + } + + // structs + { + cudf::test::lists_column_wrapper col0{{{"abc", "def"}, {"xyz"}}}; + cudf::test::fixed_width_column_wrapper col1{1}; + + cudf::table_view tbl({col0, col1}); + cudf::struct_scalar lhs(tbl); + cudf::struct_scalar rhs(tbl); + + std::vector> cols; + cols.push_back(col0.release()); + cols.push_back(col1.release()); + cudf::test::structs_column_wrapper struct_col(std::move(cols)); + + cudf::test::fixed_width_column_wrapper mask{}; + + auto expected = cudf::empty_like(struct_col); + + auto out = cudf::copy_if_else(lhs, rhs, mask); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), *expected); + } +} + +template +struct CopyTestNumeric : public cudf::test::BaseFixture { +}; +TYPED_TEST_CASE(CopyTestNumeric, cudf::test::NumericTypes); + +TYPED_TEST(CopyTestNumeric, CopyIfElseTestScalarColumn) +{ + using T = TypeParam; + + int num_els = 4; + + bool mask[] = {1, 0, 0, 1}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); + + cudf::numeric_scalar lhs_w(5); + + const auto rhs = cudf::test::make_type_param_vector({6, 6, 6, 6}); + bool rhs_v[] = {1, 0, 1, 1}; + wrapper rhs_w(rhs.begin(), rhs.end(), rhs_v); + + const auto expected = cudf::test::make_type_param_vector({5, 6, 6, 5}); + wrapper expected_w(expected.begin(), expected.end(), rhs_v); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTestNumeric, CopyIfElseTestColumnScalar) +{ + using T = TypeParam; + + int num_els = 4; + + bool mask[] = {1, 0, 0, 1}; + bool mask_v[] = {1, 1, 1, 0}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els, mask_v); + + const auto lhs = cudf::test::make_type_param_vector({5, 5, 5, 5}); + bool lhs_v[] = {0, 1, 1, 1}; + wrapper lhs_w(lhs.begin(), lhs.end(), lhs_v); + + cudf::numeric_scalar rhs_w(6); + + const auto expected = cudf::test::make_type_param_vector({5, 6, 6, 6}); + wrapper expected_w(expected.begin(), expected.end(), lhs_v); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTestNumeric, CopyIfElseTestScalarScalar) +{ + using T = TypeParam; + + int num_els = 4; + + bool mask[] = {1, 0, 0, 1}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); + + cudf::numeric_scalar lhs_w(5); + cudf::numeric_scalar rhs_w(6, false); + + const auto expected = cudf::test::make_type_param_vector({5, 6, 6, 5}); + wrapper expected_w(expected.begin(), expected.end(), mask); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +template +struct create_chrono_scalar { + template + typename std::enable_if_t< + std::is_same::type, std::true_type>::value, + cudf::timestamp_scalar> + operator()(Args&&... args) const + { + return cudf::timestamp_scalar(std::forward(args)...); + } + + template + typename std::enable_if_t< + std::is_same::type, std::true_type>::value, + cudf::duration_scalar> + operator()(Args&&... args) const + { + return cudf::duration_scalar(std::forward(args)...); + } +}; + +template +struct CopyTestChrono : public cudf::test::BaseFixture { +}; +TYPED_TEST_CASE(CopyTestChrono, cudf::test::ChronoTypes); + +TYPED_TEST(CopyTestChrono, CopyIfElseTestScalarColumn) +{ + using T = TypeParam; + + int num_els = 4; + + bool mask[] = {1, 0, 0, 1}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); + + auto lhs_w = create_chrono_scalar{}(cudf::test::make_type_param_scalar(5), true); + + bool rhs_v[] = {1, 0, 1, 1}; + wrapper rhs_w({6, 6, 6, 6}, rhs_v); + + wrapper expected_w({5, 6, 6, 5}, rhs_v); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTestChrono, CopyIfElseTestColumnScalar) +{ + using T = TypeParam; + + int num_els = 4; + + bool mask[] = {1, 0, 0, 1}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); + + bool lhs_v[] = {0, 1, 1, 1}; + wrapper lhs_w({5, 5, 5, 5}, lhs_v); + + auto rhs_w = create_chrono_scalar{}(cudf::test::make_type_param_scalar(6), true); + + wrapper expected_w({5, 6, 6, 5}, lhs_v); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +TYPED_TEST(CopyTestChrono, CopyIfElseTestScalarScalar) +{ + using T = TypeParam; + + int num_els = 4; + + bool mask[] = {1, 0, 0, 1}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); + + auto lhs_w = create_chrono_scalar{}(cudf::test::make_type_param_scalar(5), true); + auto rhs_w = create_chrono_scalar{}(cudf::test::make_type_param_scalar(6), false); + + wrapper expected_w({5, 6, 6, 5}, mask); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} + +struct CopyTestUntyped : public cudf::test::BaseFixture { +}; + +TEST_F(CopyTestUntyped, CopyIfElseTypeMismatch) +{ + cudf::test::fixed_width_column_wrapper mask_w{1, 1, 1, 1}; + + wrapper lhs_w{5, 5, 5, 5}; + wrapper rhs_w{6, 6, 6, 6}; + + EXPECT_THROW(cudf::copy_if_else(lhs_w, rhs_w, mask_w), cudf::logic_error); +} + +struct StringsCopyIfElseTest : public cudf::test::BaseFixture { +}; + +TEST_F(StringsCopyIfElseTest, CopyIfElse) +{ + auto valids = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return i % 2 == 0 ? true : false; }); + + std::vector h_strings1{"eee", "bb", "", "aa", "bbb", "ééé"}; + cudf::test::strings_column_wrapper strings1(h_strings1.begin(), h_strings1.end(), valids); + std::vector h_strings2{"zz", "", "yyy", "w", "ééé", "ooo"}; + cudf::test::strings_column_wrapper strings2(h_strings2.begin(), h_strings2.end(), valids); + + bool mask[] = {1, 1, 0, 1, 0, 1}; + bool mask_v[] = {1, 1, 1, 1, 1, 0}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + 6, mask_v); + + auto results = cudf::copy_if_else(strings1, strings2, mask_w); + + std::vector h_expected; + for (cudf::size_type idx = 0; idx < static_cast(h_strings1.size()); ++idx) { + if (mask[idx] and mask_v[idx]) + h_expected.push_back(h_strings1[idx]); + else + h_expected.push_back(h_strings2[idx]); + } + cudf::test::strings_column_wrapper expected(h_expected.begin(), h_expected.end(), valids); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); +} + +TEST_F(StringsCopyIfElseTest, CopyIfElseScalarColumn) +{ + auto valids = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return i % 2 == 0 ? true : false; }); + + std::vector h_string1{"eee"}; + cudf::string_scalar strings1{h_string1[0]}; + std::vector h_strings2{"zz", "", "yyy", "w", "ééé", "ooo"}; + cudf::test::strings_column_wrapper strings2(h_strings2.begin(), h_strings2.end(), valids); + + bool mask[] = {1, 0, 1, 0, 1, 0}; + bool mask_v[] = {1, 1, 1, 1, 1, 0}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + 6, mask_v); + + auto results = cudf::copy_if_else(strings1, strings2, mask_w); + + std::vector h_expected; + for (cudf::size_type idx = 0; idx < static_cast(h_strings2.size()); ++idx) { + if (mask[idx] and mask_v[idx]) { + h_expected.push_back(h_string1[0]); + } else { + h_expected.push_back(h_strings2[idx]); + } + } + cudf::test::strings_column_wrapper expected(h_expected.begin(), h_expected.end(), valids); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); +} + +TEST_F(StringsCopyIfElseTest, CopyIfElseColumnScalar) +{ + auto valids = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return i % 2 == 0 ? true : false; }); + + std::vector h_string1{"eee"}; + cudf::string_scalar strings1{h_string1[0]}; + std::vector h_strings2{"zz", "", "yyy", "w", "ééé", "ooo"}; + cudf::test::strings_column_wrapper strings2(h_strings2.begin(), h_strings2.end(), valids); + + bool mask[] = {0, 1, 1, 1, 0, 1}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + 6); + + auto results = cudf::copy_if_else(strings2, strings1, mask_w); + + std::vector h_expected; + for (cudf::size_type idx = 0; idx < static_cast(h_strings2.size()); ++idx) { + if (mask[idx]) { + h_expected.push_back(h_strings2[idx]); + } else { + h_expected.push_back(h_string1[0]); + } + } + cudf::test::strings_column_wrapper expected(h_expected.begin(), h_expected.end(), valids); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); +} + +TEST_F(StringsCopyIfElseTest, CopyIfElseScalarScalar) +{ + auto valids = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return i % 2 == 0 ? true : false; }); + + std::vector h_string1{"eee"}; + cudf::string_scalar string1{h_string1[0]}; + std::vector h_string2{"aaa"}; + cudf::string_scalar string2{h_string2[0], false}; + + constexpr cudf::size_type mask_size = 6; + bool mask[] = {1, 0, 1, 0, 1, 0}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + mask_size); + + auto results = cudf::copy_if_else(string1, string2, mask_w); + + std::vector h_expected; + for (cudf::size_type idx = 0; idx < static_cast(mask_size); ++idx) { + if (mask[idx]) { + h_expected.push_back(h_string1[0]); + } else { + h_expected.push_back(h_string2[0]); + } + } + cudf::test::strings_column_wrapper expected(h_expected.begin(), h_expected.end(), valids); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); +} + +template +struct FixedPointTypes : public cudf::test::BaseFixture { +}; + +TYPED_TEST_CASE(FixedPointTypes, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTypes, FixedPointSimple) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + + auto const mask = cudf::test::fixed_width_column_wrapper{0, 1, 1, 1, 0, 0}; + auto const a = fp_wrapper{{110, 220, 330, 440, 550, 660}, scale_type{-2}}; + auto const b = fp_wrapper{{0, 0, 0, 0, 0, 0}, scale_type{-2}}; + auto const expected = fp_wrapper{{0, 220, 330, 440, 0, 0}, scale_type{-2}}; + auto const result = cudf::copy_if_else(a, b, mask); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + +TYPED_TEST(FixedPointTypes, FixedPointLarge) +{ + using namespace numeric; + using namespace cudf::test; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + + auto a = thrust::make_counting_iterator(-1000); + auto b = thrust::make_constant_iterator(0); + auto m = cudf::detail::make_counting_transform_iterator(-1000, [](int i) { return i > 0; }); + auto e = + cudf::detail::make_counting_transform_iterator(-1000, [](int i) { return std::max(0, i); }); + + auto const mask = cudf::test::fixed_width_column_wrapper(m, m + 2000); + auto const A = fp_wrapper{a, a + 2000, scale_type{-3}}; + auto const B = fp_wrapper{b, b + 2000, scale_type{-3}}; + auto const expected = fp_wrapper{e, e + 2000, scale_type{-3}}; + auto const result = cudf::copy_if_else(A, B, mask); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + +TYPED_TEST(FixedPointTypes, FixedPointScaleMismatch) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + + auto const mask = cudf::test::fixed_width_column_wrapper{0, 1, 1, 1, 0, 0}; + auto const a = fp_wrapper{{110, 220, 330, 440, 550, 660}, scale_type{-2}}; + auto const b = fp_wrapper{{0, 0, 0, 0, 0, 0}, scale_type{-1}}; + + EXPECT_THROW(cudf::copy_if_else(a, b, mask), cudf::logic_error); +} diff --git a/cpp/tests/copying/detail_copy_tests.cu b/cpp/tests/copying/detail_copy_tests.cu new file mode 100644 index 00000000000..cd45c6169c7 --- /dev/null +++ b/cpp/tests/copying/detail_copy_tests.cu @@ -0,0 +1,122 @@ +/* + * Copyright (c) 2021, 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 +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +#include + +template +struct CopyDetailTest : public cudf::test::BaseFixture { +}; + +TYPED_TEST_CASE(CopyDetailTest, cudf::test::FixedWidthTypesWithoutFixedPoint); + +struct copy_if_else_tiny_grid_functor { + template ())> + std::unique_ptr operator()(cudf::column_view const& lhs, + cudf::column_view const& rhs, + Filter filter, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + // output + std::unique_ptr out = + cudf::allocate_like(lhs, lhs.size(), cudf::mask_allocation_policy::RETAIN, mr); + + // device views + auto lhs_view = cudf::column_device_view::create(lhs); + auto rhs_view = cudf::column_device_view::create(rhs); + auto lhs_iter = cudf::detail::make_optional_iterator(*lhs_view, cudf::contains_nulls::NO{}); + auto rhs_iter = cudf::detail::make_optional_iterator(*rhs_view, cudf::contains_nulls::NO{}); + auto out_dv = cudf::mutable_column_device_view::create(*out); + + // call the kernel with an artificially small grid + cudf::detail::copy_if_else_kernel<32, T, decltype(lhs_iter), decltype(rhs_iter), Filter, false> + <<<1, 32, 0, stream.value()>>>(lhs_iter, rhs_iter, filter, *out_dv, nullptr); + + return out; + } + + template ())> + std::unique_ptr operator()(cudf::column_view const& lhs, + cudf::column_view const& rhs, + Filter filter, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + CUDF_FAIL("Unexpected test execution"); + } +}; + +std::unique_ptr tiny_grid_launch(cudf::column_view const& lhs, + cudf::column_view const& rhs, + cudf::column_view const& boolean_mask) +{ + auto bool_mask_device_p = cudf::column_device_view::create(boolean_mask); + cudf::column_device_view bool_mask_device = *bool_mask_device_p; + auto filter = [bool_mask_device] __device__(cudf::size_type i) { + return bool_mask_device.element(i); + }; + return cudf::type_dispatcher(lhs.type(), + copy_if_else_tiny_grid_functor{}, + lhs, + rhs, + filter, + rmm::cuda_stream_default, + rmm::mr::get_current_device_resource()); +} + +TYPED_TEST(CopyDetailTest, CopyIfElseTestTinyGrid) +{ + using T = TypeParam; + + // make sure we span at least 2 warps + int num_els = 64; + + bool mask[] = {1, 0, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, + 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; + cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); + + cudf::test::fixed_width_column_wrapper lhs_w( + {5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, + 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, + 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5}); + + cudf::test::fixed_width_column_wrapper rhs_w( + {6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, + 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, + 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6}); + + cudf::test::fixed_width_column_wrapper expected_w( + {5, 6, 5, 6, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 6, 6, 6, + 6, 5, 5, 5, 5, 5, 5, 5, 5, 5, 6, 6, 6, 6, 5, 5, 5, 5, 5, 5, 5, 5, + 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5}); + + auto out = tiny_grid_launch(lhs_w, rhs_w, mask_w); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); +} From 45eff13f673d1203f8c9d27787d5644f46e2d3ea Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 28 Sep 2021 12:44:56 -0400 Subject: [PATCH 3/9] Use optional-iterator for copy-if-else kernel --- cpp/include/cudf/detail/copy_if_else.cuh | 31 +- cpp/include/cudf/detail/indexalator.cuh | 91 ++- cpp/src/copying/copy.cu | 119 ++-- cpp/src/copying/segmented_shift.cu | 161 ++---- cpp/src/dictionary/replace.cu | 15 +- cpp/src/replace/nans.cu | 60 +- cpp/tests/copying/copy_tests.cu | 703 ----------------------- 7 files changed, 218 insertions(+), 962 deletions(-) delete mode 100644 cpp/tests/copying/copy_tests.cu diff --git a/cpp/include/cudf/detail/copy_if_else.cuh b/cpp/include/cudf/detail/copy_if_else.cuh index 74a94f34ad8..596e119fd2f 100644 --- a/cpp/include/cudf/detail/copy_if_else.cuh +++ b/cpp/include/cudf/detail/copy_if_else.cuh @@ -18,19 +18,12 @@ #include #include -#include +#include #include #include -#include -#include -#include -#include -#include #include -#include - namespace cudf { namespace detail { namespace { // anonymous @@ -71,23 +64,17 @@ __launch_bounds__(block_size) __global__ size_type warp_cur = warp_begin + warp_id; size_type index = tid; while (warp_cur <= warp_end) { - bool in_range = (index >= begin && index < end); - - bool valid = true; - if (has_validity) { - valid = in_range && (filter(index) ? thrust::get<1>(lhs[index]) : thrust::get<1>(rhs[index])); - } - - // do the copy if-else - if (in_range) { - out.element(index) = filter(index) ? static_cast(thrust::get<0>(lhs[index])) - : static_cast(thrust::get<0>(rhs[index])); + bool valid = false; + if (index >= begin && index < end) { + auto value = filter(index) ? lhs[index] : rhs[index]; + valid = !has_validity || value.has_value(); + if (valid) { out.element(index) = static_cast(value.value()); } } // update validity if (has_validity) { // the final validity mask for this warp - int warp_mask = __ballot_sync(0xFFFF'FFFF, valid && in_range); + int warp_mask = __ballot_sync(0xFFFF'FFFF, valid); // only one guy in the warp needs to update the mask and count if (lane_id == 0) { out.set_mask_word(warp_cur, warp_mask); @@ -168,8 +155,8 @@ std::unique_ptr copy_if_else( rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { - using Element = - typename thrust::tuple_element<0, typename thrust::iterator_traits::value_type>::type; + // This is the type of the thrust::optional element in the passed iterators + using Element = typename thrust::iterator_traits::value_type::value_type; size_type size = std::distance(lhs_begin, lhs_end); size_type num_els = cudf::util::round_up_safe(size, warp_size); diff --git a/cpp/include/cudf/detail/indexalator.cuh b/cpp/include/cudf/detail/indexalator.cuh index d546162fc7a..df861f4e16b 100644 --- a/cpp/include/cudf/detail/indexalator.cuh +++ b/cpp/include/cudf/detail/indexalator.cuh @@ -484,7 +484,7 @@ struct indexalator_factory { /** * @brief An index accessor that returns a validity flag along with the index value. * - * This is suitable as a `pair_iterator` for calling functions like `copy_if_else`. + * This is suitable as a `pair_iterator`. */ struct nullable_index_accessor { input_indexalator iter; @@ -502,17 +502,32 @@ struct indexalator_factory { iter = make_input_iterator(col); } + __device__ thrust::pair operator()(size_type i) const + { + return {iter[i], (has_nulls ? bit_is_set(null_mask, i + offset) : true)}; + } + }; + + /** + * @brief An index accessor that returns a validity flag along with the index value. + * + * This is suitable as a `pair_iterator`. + */ + struct scalar_nullable_index_accessor { + input_indexalator iter; + bool const is_null; + /** * @brief Create an accessor from a scalar. */ - nullable_index_accessor(scalar const& input) : has_nulls{!input.is_valid()} + scalar_nullable_index_accessor(scalar const& input) : is_null{!input.is_valid()} { iter = indexalator_factory::make_input_iterator(input); } __device__ thrust::pair operator()(size_type i) const { - return {iter[i], (has_nulls ? bit_is_set(null_mask, i + offset) : true)}; + return {*iter, is_null}; } }; @@ -530,7 +545,75 @@ struct indexalator_factory { static auto make_input_pair_iterator(scalar const& input) { return thrust::make_transform_iterator(thrust::make_constant_iterator(0), - nullable_index_accessor{input}); + scalar_nullable_index_accessor{input}); + } + + /** + * @brief An index accessor that returns an index value if corresponding validity flag is true. + * + * This is suitable as an `optional_iterator`. + */ + struct optional_index_accessor { + input_indexalator iter; + bitmask_type const* null_mask{}; + size_type const offset{}; + bool const has_nulls{}; + + /** + * @brief Create an accessor from a column_view. + */ + optional_index_accessor(column_view const& col, bool has_nulls = false) + : null_mask{col.null_mask()}, offset{col.offset()}, has_nulls{has_nulls} + { + if (has_nulls) { CUDF_EXPECTS(col.nullable(), "Unexpected non-nullable column."); } + iter = make_input_iterator(col); + } + + __device__ thrust::optional operator()(size_type i) const + { + return has_nulls && !bit_is_set(null_mask, i + offset) ? thrust::nullopt + : thrust::make_optional(iter[i]); + } + }; + + /** + * @brief An index accessor that returns an index value if corresponding validity flag is true. + * + * This is suitable as an `optional_iterator`. + */ + struct scalar_optional_index_accessor { + input_indexalator iter; + bool const is_null; + + /** + * @brief Create an accessor from a scalar. + */ + scalar_optional_index_accessor(scalar const& input) : is_null{!input.is_valid()} + { + iter = indexalator_factory::make_input_iterator(input); + } + + __device__ thrust::optional operator()(size_type i) const + { + return is_null ? thrust::nullopt : thrust::make_optional(*iter); + } + }; + + /** + * @brief Create an index iterator with a nullable index accessor. + */ + static auto make_input_optional_iterator(column_view const& col) + { + return make_counting_transform_iterator(0, optional_index_accessor{col, col.has_nulls()}); + } + + /** + * @brief Create an index iterator with a nullable index accessor for a scalar. + */ + static auto make_input_optional_iterator(scalar const& input) + { + return thrust::make_transform_iterator(thrust::make_constant_iterator(0), + scalar_optional_index_accessor{input}); } }; diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu index df1ce6c0e67..72e5582f9aa 100644 --- a/cpp/src/copying/copy.cu +++ b/cpp/src/copying/copy.cu @@ -14,13 +14,14 @@ * limitations under the License. */ -#include #include #include #include #include #include #include +#include +#include #include #include @@ -73,28 +74,18 @@ struct copy_if_else_functor_impl auto const& lhs = *p_lhs; auto const& rhs = *p_rhs; - if (left_nullable) { - if (right_nullable) { - auto lhs_iter = cudf::detail::make_pair_iterator(lhs); - auto rhs_iter = cudf::detail::make_pair_iterator(rhs); - return detail::copy_if_else( - true, lhs_iter, lhs_iter + size, rhs_iter, filter, lhs.type(), stream, mr); - } - auto lhs_iter = cudf::detail::make_pair_iterator(lhs); - auto rhs_iter = cudf::detail::make_pair_iterator(rhs); - return detail::copy_if_else( - true, lhs_iter, lhs_iter + size, rhs_iter, filter, lhs.type(), stream, mr); - } - if (right_nullable) { - auto lhs_iter = cudf::detail::make_pair_iterator(lhs); - auto rhs_iter = cudf::detail::make_pair_iterator(rhs); - return detail::copy_if_else( - true, lhs_iter, lhs_iter + size, rhs_iter, filter, lhs.type(), stream, mr); - } - auto lhs_iter = cudf::detail::make_pair_iterator(lhs); - auto rhs_iter = cudf::detail::make_pair_iterator(rhs); - return detail::copy_if_else( - false, lhs_iter, lhs_iter + size, rhs_iter, filter, lhs.type(), stream, mr); + auto lhs_iter = + cudf::detail::make_optional_iterator(lhs, contains_nulls::DYNAMIC{}, left_nullable); + auto rhs_iter = + cudf::detail::make_optional_iterator(rhs, contains_nulls::DYNAMIC{}, right_nullable); + return detail::copy_if_else(left_nullable || right_nullable, + lhs_iter, + lhs_iter + size, + rhs_iter, + filter, + lhs.type(), + stream, + mr); } }; @@ -256,6 +247,38 @@ std::unique_ptr scatter_gather_based_if_else(cudf::scalar const& lhs, return scatter_gather_based_if_else(lhs, rhs_col->view(), size, is_left, stream, mr); } +template <> +struct copy_if_else_functor_impl { + template + std::unique_ptr operator()(Left const& lhs, + Right const& rhs, + size_type size, + bool, + bool, + Filter filter, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + return scatter_gather_based_if_else(lhs, rhs, size, filter, stream, mr); + } +}; + +template <> +struct copy_if_else_functor_impl { + template + std::unique_ptr operator()(Left const& lhs, + Right const& rhs, + size_type size, + bool, + bool, + Filter filter, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + return scatter_gather_based_if_else(lhs, rhs, size, filter, stream, mr); + } +}; + /** * @brief Functor called by the `type_dispatcher` to invoke copy_if_else on combinations * of column_view and scalar @@ -271,12 +294,6 @@ struct copy_if_else_functor { rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - if constexpr (std::is_same_v or std::is_same_v) { - (void)left_nullable; - (void)right_nullable; - return scatter_gather_based_if_else(lhs, rhs, size, filter, stream, mr); - } - copy_if_else_functor_impl copier{}; return copier(lhs, rhs, size, left_nullable, right_nullable, filter, stream, mr); } @@ -301,35 +318,21 @@ std::unique_ptr copy_if_else(Left const& lhs, auto bool_mask_device_p = column_device_view::create(boolean_mask); column_device_view bool_mask_device = *bool_mask_device_p; - if (boolean_mask.has_nulls()) { - auto filter = [bool_mask_device] __device__(cudf::size_type i) { - return bool_mask_device.is_valid_nocheck(i) and bool_mask_device.element(i); - }; - return cudf::type_dispatcher(lhs.type(), - copy_if_else_functor{}, - lhs, - rhs, - boolean_mask.size(), - left_nullable, - right_nullable, - filter, - stream, - mr); - } else { - auto filter = [bool_mask_device] __device__(cudf::size_type i) { - return bool_mask_device.element(i); - }; - return cudf::type_dispatcher(lhs.type(), - copy_if_else_functor{}, - lhs, - rhs, - boolean_mask.size(), - left_nullable, - right_nullable, - filter, - stream, - mr); - } + auto const has_nulls = boolean_mask.has_nulls(); + auto filter = [bool_mask_device, has_nulls] __device__(cudf::size_type i) { + return (!has_nulls || bool_mask_device.is_valid_nocheck(i)) and + bool_mask_device.element(i); + }; + return cudf::type_dispatcher(lhs.type(), + copy_if_else_functor{}, + lhs, + rhs, + boolean_mask.size(), + left_nullable, + right_nullable, + filter, + stream, + mr); } }; // namespace diff --git a/cpp/src/copying/segmented_shift.cu b/cpp/src/copying/segmented_shift.cu index 6fc785a61c6..ada5703073a 100644 --- a/cpp/src/copying/segmented_shift.cu +++ b/cpp/src/copying/segmented_shift.cu @@ -18,6 +18,7 @@ #include #include #include +#include #include #include #include @@ -33,88 +34,24 @@ namespace detail { namespace { /** - * @brief Helper function to invoke general `copy_if_else` + * @brief Common filter function to convert index values into copy-if-else left/right result. + * + * The offset position is used to identify which segment to copy from. */ -template -std::unique_ptr segmented_shift_rep_impl(PairIterator input_pair_iterator, - ScalarIterator fill_pair_iterator, - bool nullable, - size_type offset, - device_span segment_offsets, - data_type value_type, - size_type column_size, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - if (offset > 0) { - auto filter = [segment_offsets, offset] __device__(auto const& i) { - auto segment_bound_idx = - thrust::upper_bound(thrust::seq, segment_offsets.begin(), segment_offsets.end(), i) - 1; - return not(*segment_bound_idx <= i and i < *segment_bound_idx + offset); - }; - return copy_if_else(nullable, - input_pair_iterator, - input_pair_iterator + column_size, - fill_pair_iterator, - filter, - value_type, - stream, - mr); - } else { - auto filter = [segment_offsets, offset] __device__(auto const& i) { - auto segment_bound_idx = - thrust::upper_bound(thrust::seq, segment_offsets.begin(), segment_offsets.end(), i); - return not(*segment_bound_idx + offset <= i and i < *segment_bound_idx); - }; - return copy_if_else(nullable, - input_pair_iterator, - input_pair_iterator + column_size, - fill_pair_iterator, - filter, - value_type, - stream, - mr); - } -} +struct segmented_shift_filter { + device_span const segment_offsets; + size_type const offset; -/** - * @brief Helper function to invoke string specialization of `copy_if_else` - */ -template -std::unique_ptr segmented_shift_string_impl(PairIterator input_pair_iterator, - ScalarIterator fill_pair_iterator, - size_type offset, - device_span segment_offsets, - size_type column_size, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - if (offset > 0) { - auto filter = [segment_offsets, offset] __device__(auto const& i) { - auto segment_bound_idx = - thrust::upper_bound(thrust::seq, segment_offsets.begin(), segment_offsets.end(), i) - 1; - return not(*segment_bound_idx <= i and i < *segment_bound_idx + offset); - }; - return strings::detail::copy_if_else(input_pair_iterator, - input_pair_iterator + column_size, - fill_pair_iterator, - filter, - stream, - mr); - } else { - auto filter = [segment_offsets, offset] __device__(auto const& i) { - auto segment_bound_idx = - thrust::upper_bound(thrust::seq, segment_offsets.begin(), segment_offsets.end(), i); - return not(*segment_bound_idx + offset <= i and i < *segment_bound_idx); - }; - return strings::detail::copy_if_else(input_pair_iterator, - input_pair_iterator + column_size, - fill_pair_iterator, - filter, - stream, - mr); - } -} + __device__ bool operator()(size_type const i) const + { + auto const segment_bound_idx = + thrust::upper_bound(thrust::seq, segment_offsets.begin(), segment_offsets.end(), i) - + (offset > 0); + auto const left_idx = *segment_bound_idx + (offset < 0 ? offset : 0); + auto const right_idx = *segment_bound_idx + (offset > 0 ? offset : 0); + return not(left_idx <= i and i < right_idx); + }; +}; template struct segmented_shift_functor { @@ -138,32 +75,20 @@ struct segmented_shift_functor() rmm::mr::device_memory_resource* mr) { auto values_device_view = column_device_view::create(segmented_values, stream); - auto fill_pair_iterator = make_pair_iterator(fill_value); bool nullable = not fill_value.is_valid() or segmented_values.nullable(); - - if (segmented_values.has_nulls()) { - auto input_pair_iterator = make_pair_iterator(*values_device_view) - offset; - return segmented_shift_rep_impl(input_pair_iterator, - fill_pair_iterator, - nullable, - offset, - segment_offsets, - segmented_values.type(), - segmented_values.size(), - stream, - mr); - } else { - auto input_pair_iterator = make_pair_iterator(*values_device_view) - offset; - return segmented_shift_rep_impl(input_pair_iterator, - fill_pair_iterator, - nullable, - offset, - segment_offsets, - segmented_values.type(), - segmented_values.size(), - stream, - mr); - } + auto input_iterator = + cudf::detail::make_optional_iterator( + *values_device_view, contains_nulls::DYNAMIC{}, segmented_values.has_nulls()) - + offset; + auto fill_iterator = cudf::detail::make_optional_iterator(fill_value, contains_nulls::YES{}); + return copy_if_else(nullable, + input_iterator, + input_iterator + segmented_values.size(), + fill_iterator, + segmented_shift_filter{segment_offsets, offset}, + segmented_values.type(), + stream, + mr); } }; @@ -185,22 +110,20 @@ struct segmented_shift_functor { auto fill_pair_iterator = make_pair_iterator(fill_value); if (segmented_values.has_nulls()) { auto input_pair_iterator = make_pair_iterator(*values_device_view) - offset; - return segmented_shift_string_impl(input_pair_iterator, - fill_pair_iterator, - offset, - segment_offsets, - segmented_values.size(), - stream, - mr); + return strings::detail::copy_if_else(input_pair_iterator, + input_pair_iterator + segmented_values.size(), + fill_pair_iterator, + segmented_shift_filter{segment_offsets, offset}, + stream, + mr); } else { auto input_pair_iterator = make_pair_iterator(*values_device_view) - offset; - return segmented_shift_string_impl(input_pair_iterator, - fill_pair_iterator, - offset, - segment_offsets, - segmented_values.size(), - stream, - mr); + return strings::detail::copy_if_else(input_pair_iterator, + input_pair_iterator + segmented_values.size(), + fill_pair_iterator, + segmented_shift_filter{segment_offsets, offset}, + stream, + mr); } } }; diff --git a/cpp/src/dictionary/replace.cu b/cpp/src/dictionary/replace.cu index 37118779248..add56f45d74 100644 --- a/cpp/src/dictionary/replace.cu +++ b/cpp/src/dictionary/replace.cu @@ -58,17 +58,14 @@ std::unique_ptr replace_indices(column_view const& input, auto const d_input = *input_view; auto predicate = [d_input] __device__(auto i) { return d_input.is_valid(i); }; - using Element = typename thrust:: - tuple_element<0, typename thrust::iterator_traits::value_type>::type; - - auto input_pair_iterator = cudf::detail::indexalator_factory::make_input_pair_iterator(input); + auto input_iterator = cudf::detail::indexalator_factory::make_input_optional_iterator(input); return cudf::detail::copy_if_else(true, - input_pair_iterator, - input_pair_iterator + input.size(), + input_iterator, + input_iterator + input.size(), replacement_iter, predicate, - data_type{type_to_id()}, + data_type{type_to_id()}, stream, mr); } @@ -100,7 +97,7 @@ std::unique_ptr replace_nulls(dictionary_column_view const& input, auto new_indices = replace_indices(input_indices, - cudf::detail::indexalator_factory::make_input_pair_iterator(repl_indices), + cudf::detail::indexalator_factory::make_input_optional_iterator(repl_indices), stream, mr); @@ -133,7 +130,7 @@ std::unique_ptr replace_nulls(dictionary_column_view const& input, auto const input_indices = input_view.get_indices_annotated(); auto new_indices = replace_indices(input_indices, - cudf::detail::indexalator_factory::make_input_pair_iterator(*scalar_index), + cudf::detail::indexalator_factory::make_input_optional_iterator(*scalar_index), stream, mr); new_indices->set_null_mask(rmm::device_buffer{0, stream, mr}, 0); diff --git a/cpp/src/replace/nans.cu b/cpp/src/replace/nans.cu index 65ddb342ad7..a08a90f12af 100644 --- a/cpp/src/replace/nans.cu +++ b/cpp/src/replace/nans.cu @@ -26,6 +26,7 @@ #include #include +#include #include @@ -54,53 +55,18 @@ struct replace_nans_functor { return dinput.is_null(i) or !std::isnan(dinput.element(i)); }; - if (input.has_nulls()) { - auto input_pair_iterator = make_pair_iterator(*input_device_view); - if (replacement_nullable) { - auto replacement_pair_iterator = make_pair_iterator(replacement); - return copy_if_else(true, - input_pair_iterator, - input_pair_iterator + size, - replacement_pair_iterator, - predicate, - input.type(), - stream, - mr); - } else { - auto replacement_pair_iterator = make_pair_iterator(replacement); - return copy_if_else(true, - input_pair_iterator, - input_pair_iterator + size, - replacement_pair_iterator, - predicate, - input.type(), - stream, - mr); - } - } else { - auto input_pair_iterator = make_pair_iterator(*input_device_view); - if (replacement_nullable) { - auto replacement_pair_iterator = make_pair_iterator(replacement); - return copy_if_else(true, - input_pair_iterator, - input_pair_iterator + size, - replacement_pair_iterator, - predicate, - input.type(), - stream, - mr); - } else { - auto replacement_pair_iterator = make_pair_iterator(replacement); - return copy_if_else(false, - input_pair_iterator, - input_pair_iterator + size, - replacement_pair_iterator, - predicate, - input.type(), - stream, - mr); - } - } + auto input_iterator = + make_optional_iterator(*input_device_view, contains_nulls::DYNAMIC{}, input.has_nulls()); + auto replacement_iterator = + make_optional_iterator(replacement, contains_nulls::DYNAMIC{}, replacement_nullable); + return copy_if_else(input.has_nulls() or replacement_nullable, + input_iterator, + input_iterator + size, + replacement_iterator, + predicate, + input.type(), + stream, + mr); } template diff --git a/cpp/tests/copying/copy_tests.cu b/cpp/tests/copying/copy_tests.cu deleted file mode 100644 index 03869c37adf..00000000000 --- a/cpp/tests/copying/copy_tests.cu +++ /dev/null @@ -1,703 +0,0 @@ -/* - * Copyright (c) 2019-2020, 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 -#include -#include -#include -#include - -#include -#include -#include -#include -#include -#include -#include - -#include - -template -struct CopyTest : public cudf::test::BaseFixture { -}; - -TYPED_TEST_CASE(CopyTest, cudf::test::FixedWidthTypesWithoutFixedPoint); - -#define wrapper cudf::test::fixed_width_column_wrapper - -TYPED_TEST(CopyTest, CopyIfElseTestShort) -{ - using T = TypeParam; - - cudf::test::fixed_width_column_wrapper mask_w{1, 0, 0, 0}; - - wrapper lhs_w({5, 5, 5, 5}, {1, 1, 1, 1}); - wrapper rhs_w({6, 6, 6, 6}, {1, 1, 1, 1}); - wrapper expected_w({5, 6, 6, 6}); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTest, CopyIfElseTestManyNulls) -{ - using T = TypeParam; - - cudf::test::fixed_width_column_wrapper mask_w{{1, 0, 0, 0, 0, 0, 1}, {1, 1, 1, 1, 1, 1, 0}}; - - wrapper lhs_w({5, 5, 5, 5, 5, 5, 5}, {1, 1, 1, 1, 1, 1, 1}); - wrapper rhs_w({6, 6, 6, 6, 6, 6, 6}, {1, 0, 0, 0, 0, 0, 1}); - wrapper expected_w({5, 6, 6, 6, 6, 6, 6}, {1, 0, 0, 0, 0, 0, 1}); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -struct copy_if_else_tiny_grid_functor { - template ())> - std::unique_ptr operator()(cudf::column_view const& lhs, - cudf::column_view const& rhs, - Filter filter, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) - { - // output - std::unique_ptr out = - cudf::allocate_like(lhs, lhs.size(), cudf::mask_allocation_policy::RETAIN, mr); - - // device views - auto lhs_view = cudf::column_device_view::create(lhs); - auto rhs_view = cudf::column_device_view::create(rhs); - auto lhs_iter = cudf::detail::make_pair_iterator(*lhs_view); - auto rhs_iter = cudf::detail::make_pair_iterator(*rhs_view); - auto out_dv = cudf::mutable_column_device_view::create(*out); - - // call the kernel with an artificially small grid - cudf::detail::copy_if_else_kernel<32, T, decltype(lhs_iter), decltype(rhs_iter), Filter, false> - <<<1, 32, 0, stream.value()>>>(lhs_iter, rhs_iter, filter, *out_dv, nullptr); - - return out; - } - - template ())> - std::unique_ptr operator()(cudf::column_view const& lhs, - cudf::column_view const& rhs, - Filter filter, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) - { - CUDF_FAIL("Unexpected test execution"); - } -}; - -std::unique_ptr tiny_grid_launch(cudf::column_view const& lhs, - cudf::column_view const& rhs, - cudf::column_view const& boolean_mask) -{ - auto bool_mask_device_p = cudf::column_device_view::create(boolean_mask); - cudf::column_device_view bool_mask_device = *bool_mask_device_p; - auto filter = [bool_mask_device] __device__(cudf::size_type i) { - return bool_mask_device.element(i); - }; - return cudf::type_dispatcher(lhs.type(), - copy_if_else_tiny_grid_functor{}, - lhs, - rhs, - filter, - rmm::cuda_stream_default, - rmm::mr::get_current_device_resource()); -} - -TYPED_TEST(CopyTest, CopyIfElseTestTinyGrid) -{ - using T = TypeParam; - - // make sure we span at least 2 warps - int num_els = 64; - - bool mask[] = {1, 0, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, - 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); - - wrapper lhs_w({5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, - 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, - 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5}); - - wrapper rhs_w({6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, - 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, - 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6}); - - wrapper expected_w({5, 6, 5, 6, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 6, 6, 6, - 6, 5, 5, 5, 5, 5, 5, 5, 5, 5, 6, 6, 6, 6, 5, 5, 5, 5, 5, 5, 5, 5, - 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5}); - - auto out = tiny_grid_launch(lhs_w, rhs_w, mask_w); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTest, CopyIfElseTestLong) -{ - using T = TypeParam; - - // make sure we span at least 2 warps - int num_els = 64; - - bool mask[] = {1, 0, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, - 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); - - bool lhs_v[] = {1, 1, 1, 1, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; - wrapper lhs_w({5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, - 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, - 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5}, - lhs_v); - - bool rhs_v[] = {1, 1, 1, 1, 1, 1, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; - wrapper rhs_w({6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, - 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, - 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6}, - rhs_v); - - bool exp_v[] = {1, 1, 1, 1, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; - wrapper expected_w({5, 6, 5, 6, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 6, 6, 6, - 6, 5, 5, 5, 5, 5, 5, 5, 5, 5, 6, 6, 6, 6, 5, 5, 5, 5, 5, 5, 5, 5, - 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5}, - exp_v); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTest, CopyIfElseTestEmptyInputs) -{ - using T = TypeParam; - - cudf::test::fixed_width_column_wrapper mask_w{}; - - wrapper lhs_w{}; - wrapper rhs_w{}; - wrapper expected_w{}; - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTest, CopyIfElseMixedInputValidity) -{ - using T = TypeParam; - - cudf::test::fixed_width_column_wrapper mask_w{1, 0, 1, 1}; - - wrapper lhs_w({5, 5, 5, 5}, {1, 1, 1, 0}); - wrapper rhs_w({6, 6, 6, 6}, {1, 0, 1, 1}); - wrapper expected_w({5, 6, 5, 5}, {1, 0, 1, 0}); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTest, CopyIfElseMixedInputValidity2) -{ - using T = TypeParam; - - cudf::test::fixed_width_column_wrapper mask_w{1, 0, 1, 1}; - - wrapper lhs_w({5, 5, 5, 5}, {1, 1, 1, 0}); - wrapper rhs_w({6, 6, 6, 6}); - wrapper expected_w({5, 6, 5, 5}, {1, 1, 1, 0}); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTest, CopyIfElseMixedInputValidity3) -{ - using T = TypeParam; - - cudf::test::fixed_width_column_wrapper mask_w{1, 0, 1, 1}; - - wrapper lhs_w({5, 5, 5, 5}); - wrapper rhs_w({6, 6, 6, 6}, {1, 0, 1, 1}); - wrapper expected_w({5, 6, 5, 5}, {1, 0, 1, 1}); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTest, CopyIfElseMixedInputValidity4) -{ - using T = TypeParam; - - cudf::test::fixed_width_column_wrapper mask_w{1, 0, 1, 1}; - - wrapper lhs_w({5, 5, 5, 5}); - wrapper rhs_w({6, 6, 6, 6}); - wrapper expected_w({5, 6, 5, 5}); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTest, CopyIfElseBadInputLength) -{ - using T = TypeParam; - - // mask length mismatch - { - cudf::test::fixed_width_column_wrapper mask_w{1, 1, 1}; - - wrapper lhs_w({5, 5, 5, 5}); - wrapper rhs_w({6, 6, 6, 6}); - - EXPECT_THROW(cudf::copy_if_else(lhs_w, rhs_w, mask_w), cudf::logic_error); - } - - // column length mismatch - { - cudf::test::fixed_width_column_wrapper mask_w{1, 1, 1, 1}; - - wrapper lhs_w({5, 5, 5}); - wrapper rhs_w({6, 6, 6, 6}); - - EXPECT_THROW(cudf::copy_if_else(lhs_w, rhs_w, mask_w), cudf::logic_error); - } -} - -struct CopyEmptyNested : public cudf::test::BaseFixture { -}; - -TEST_F(CopyEmptyNested, CopyIfElseTestEmptyNestedColumns) -{ - // lists - { - cudf::test::lists_column_wrapper col{{{"abc", "def"}, {"xyz"}}}; - auto lhs = cudf::empty_like(col); - auto rhs = cudf::empty_like(col); - cudf::test::fixed_width_column_wrapper mask{}; - - auto expected = empty_like(col); - - auto out = cudf::copy_if_else(*lhs, *rhs, mask); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), *expected); - } - - // structs - { - cudf::test::lists_column_wrapper _col0{{{"abc", "def"}, {"xyz"}}}; - auto col0 = cudf::empty_like(_col0); - cudf::test::fixed_width_column_wrapper col1; - - std::vector> cols; - cols.push_back(std::move(col0)); - cols.push_back(col1.release()); - cudf::test::structs_column_wrapper struct_col(std::move(cols)); - auto lhs = cudf::empty_like(struct_col); - auto rhs = cudf::empty_like(struct_col); - - cudf::test::fixed_width_column_wrapper mask{}; - - auto expected = cudf::empty_like(struct_col); - - auto out = cudf::copy_if_else(*lhs, *rhs, mask); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), *expected); - } -} - -TEST_F(CopyEmptyNested, CopyIfElseTestEmptyNestedScalars) -{ - // lists - { - cudf::test::lists_column_wrapper _col{{{"abc", "def"}, {"xyz"}}}; - std::unique_ptr lhs = cudf::get_element(_col, 0); - std::unique_ptr rhs = cudf::get_element(_col, 0); - - cudf::test::fixed_width_column_wrapper mask{}; - - auto expected = empty_like(_col); - - auto out = cudf::copy_if_else(*lhs, *rhs, mask); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), *expected); - } - - // structs - { - cudf::test::lists_column_wrapper col0{{{"abc", "def"}, {"xyz"}}}; - cudf::test::fixed_width_column_wrapper col1{1}; - - cudf::table_view tbl({col0, col1}); - cudf::struct_scalar lhs(tbl); - cudf::struct_scalar rhs(tbl); - - std::vector> cols; - cols.push_back(col0.release()); - cols.push_back(col1.release()); - cudf::test::structs_column_wrapper struct_col(std::move(cols)); - - cudf::test::fixed_width_column_wrapper mask{}; - - auto expected = cudf::empty_like(struct_col); - - auto out = cudf::copy_if_else(lhs, rhs, mask); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), *expected); - } -} - -template -struct CopyTestNumeric : public cudf::test::BaseFixture { -}; -TYPED_TEST_CASE(CopyTestNumeric, cudf::test::NumericTypes); - -TYPED_TEST(CopyTestNumeric, CopyIfElseTestScalarColumn) -{ - using T = TypeParam; - - int num_els = 4; - - bool mask[] = {1, 0, 0, 1}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); - - cudf::numeric_scalar lhs_w(5); - - const auto rhs = cudf::test::make_type_param_vector({6, 6, 6, 6}); - bool rhs_v[] = {1, 0, 1, 1}; - wrapper rhs_w(rhs.begin(), rhs.end(), rhs_v); - - const auto expected = cudf::test::make_type_param_vector({5, 6, 6, 5}); - wrapper expected_w(expected.begin(), expected.end(), rhs_v); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTestNumeric, CopyIfElseTestColumnScalar) -{ - using T = TypeParam; - - int num_els = 4; - - bool mask[] = {1, 0, 0, 1}; - bool mask_v[] = {1, 1, 1, 0}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els, mask_v); - - const auto lhs = cudf::test::make_type_param_vector({5, 5, 5, 5}); - bool lhs_v[] = {0, 1, 1, 1}; - wrapper lhs_w(lhs.begin(), lhs.end(), lhs_v); - - cudf::numeric_scalar rhs_w(6); - - const auto expected = cudf::test::make_type_param_vector({5, 6, 6, 6}); - wrapper expected_w(expected.begin(), expected.end(), lhs_v); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTestNumeric, CopyIfElseTestScalarScalar) -{ - using T = TypeParam; - - int num_els = 4; - - bool mask[] = {1, 0, 0, 1}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); - - cudf::numeric_scalar lhs_w(5); - cudf::numeric_scalar rhs_w(6, false); - - const auto expected = cudf::test::make_type_param_vector({5, 6, 6, 5}); - wrapper expected_w(expected.begin(), expected.end(), mask); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -template -struct create_chrono_scalar { - template - typename std::enable_if_t< - std::is_same::type, std::true_type>::value, - cudf::timestamp_scalar> - operator()(Args&&... args) const - { - return cudf::timestamp_scalar(std::forward(args)...); - } - - template - typename std::enable_if_t< - std::is_same::type, std::true_type>::value, - cudf::duration_scalar> - operator()(Args&&... args) const - { - return cudf::duration_scalar(std::forward(args)...); - } -}; - -template -struct CopyTestChrono : public cudf::test::BaseFixture { -}; -TYPED_TEST_CASE(CopyTestChrono, cudf::test::ChronoTypes); - -TYPED_TEST(CopyTestChrono, CopyIfElseTestScalarColumn) -{ - using T = TypeParam; - - int num_els = 4; - - bool mask[] = {1, 0, 0, 1}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); - - auto lhs_w = create_chrono_scalar{}(cudf::test::make_type_param_scalar(5), true); - - bool rhs_v[] = {1, 0, 1, 1}; - wrapper rhs_w({6, 6, 6, 6}, rhs_v); - - wrapper expected_w({5, 6, 6, 5}, rhs_v); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTestChrono, CopyIfElseTestColumnScalar) -{ - using T = TypeParam; - - int num_els = 4; - - bool mask[] = {1, 0, 0, 1}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); - - bool lhs_v[] = {0, 1, 1, 1}; - wrapper lhs_w({5, 5, 5, 5}, lhs_v); - - auto rhs_w = create_chrono_scalar{}(cudf::test::make_type_param_scalar(6), true); - - wrapper expected_w({5, 6, 6, 5}, lhs_v); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -TYPED_TEST(CopyTestChrono, CopyIfElseTestScalarScalar) -{ - using T = TypeParam; - - int num_els = 4; - - bool mask[] = {1, 0, 0, 1}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); - - auto lhs_w = create_chrono_scalar{}(cudf::test::make_type_param_scalar(5), true); - auto rhs_w = create_chrono_scalar{}(cudf::test::make_type_param_scalar(6), false); - - wrapper expected_w({5, 6, 6, 5}, mask); - - auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} - -struct CopyTestUntyped : public cudf::test::BaseFixture { -}; - -TEST_F(CopyTestUntyped, CopyIfElseTypeMismatch) -{ - cudf::test::fixed_width_column_wrapper mask_w{1, 1, 1, 1}; - - wrapper lhs_w{5, 5, 5, 5}; - wrapper rhs_w{6, 6, 6, 6}; - - EXPECT_THROW(cudf::copy_if_else(lhs_w, rhs_w, mask_w), cudf::logic_error); -} - -struct StringsCopyIfElseTest : public cudf::test::BaseFixture { -}; - -TEST_F(StringsCopyIfElseTest, CopyIfElse) -{ - auto valids = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return i % 2 == 0 ? true : false; }); - - std::vector h_strings1{"eee", "bb", "", "aa", "bbb", "ééé"}; - cudf::test::strings_column_wrapper strings1(h_strings1.begin(), h_strings1.end(), valids); - std::vector h_strings2{"zz", "", "yyy", "w", "ééé", "ooo"}; - cudf::test::strings_column_wrapper strings2(h_strings2.begin(), h_strings2.end(), valids); - - bool mask[] = {1, 1, 0, 1, 0, 1}; - bool mask_v[] = {1, 1, 1, 1, 1, 0}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + 6, mask_v); - - auto results = cudf::copy_if_else(strings1, strings2, mask_w); - - std::vector h_expected; - for (cudf::size_type idx = 0; idx < static_cast(h_strings1.size()); ++idx) { - if (mask[idx] and mask_v[idx]) - h_expected.push_back(h_strings1[idx]); - else - h_expected.push_back(h_strings2[idx]); - } - cudf::test::strings_column_wrapper expected(h_expected.begin(), h_expected.end(), valids); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); -} - -TEST_F(StringsCopyIfElseTest, CopyIfElseScalarColumn) -{ - auto valids = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return i % 2 == 0 ? true : false; }); - - std::vector h_string1{"eee"}; - cudf::string_scalar strings1{h_string1[0]}; - std::vector h_strings2{"zz", "", "yyy", "w", "ééé", "ooo"}; - cudf::test::strings_column_wrapper strings2(h_strings2.begin(), h_strings2.end(), valids); - - bool mask[] = {1, 0, 1, 0, 1, 0}; - bool mask_v[] = {1, 1, 1, 1, 1, 0}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + 6, mask_v); - - auto results = cudf::copy_if_else(strings1, strings2, mask_w); - - std::vector h_expected; - for (cudf::size_type idx = 0; idx < static_cast(h_strings2.size()); ++idx) { - if (mask[idx] and mask_v[idx]) { - h_expected.push_back(h_string1[0]); - } else { - h_expected.push_back(h_strings2[idx]); - } - } - cudf::test::strings_column_wrapper expected(h_expected.begin(), h_expected.end(), valids); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); -} - -TEST_F(StringsCopyIfElseTest, CopyIfElseColumnScalar) -{ - auto valids = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return i % 2 == 0 ? true : false; }); - - std::vector h_string1{"eee"}; - cudf::string_scalar strings1{h_string1[0]}; - std::vector h_strings2{"zz", "", "yyy", "w", "ééé", "ooo"}; - cudf::test::strings_column_wrapper strings2(h_strings2.begin(), h_strings2.end(), valids); - - bool mask[] = {0, 1, 1, 1, 0, 1}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + 6); - - auto results = cudf::copy_if_else(strings2, strings1, mask_w); - - std::vector h_expected; - for (cudf::size_type idx = 0; idx < static_cast(h_strings2.size()); ++idx) { - if (mask[idx]) { - h_expected.push_back(h_strings2[idx]); - } else { - h_expected.push_back(h_string1[0]); - } - } - cudf::test::strings_column_wrapper expected(h_expected.begin(), h_expected.end(), valids); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); -} - -TEST_F(StringsCopyIfElseTest, CopyIfElseScalarScalar) -{ - auto valids = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return i % 2 == 0 ? true : false; }); - - std::vector h_string1{"eee"}; - cudf::string_scalar string1{h_string1[0]}; - std::vector h_string2{"aaa"}; - cudf::string_scalar string2{h_string2[0], false}; - - constexpr cudf::size_type mask_size = 6; - bool mask[] = {1, 0, 1, 0, 1, 0}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + mask_size); - - auto results = cudf::copy_if_else(string1, string2, mask_w); - - std::vector h_expected; - for (cudf::size_type idx = 0; idx < static_cast(mask_size); ++idx) { - if (mask[idx]) { - h_expected.push_back(h_string1[0]); - } else { - h_expected.push_back(h_string2[0]); - } - } - cudf::test::strings_column_wrapper expected(h_expected.begin(), h_expected.end(), valids); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); -} - -template -struct FixedPointTypes : public cudf::test::BaseFixture { -}; - -TYPED_TEST_CASE(FixedPointTypes, cudf::test::FixedPointTypes); - -TYPED_TEST(FixedPointTypes, FixedPointSimple) -{ - using namespace numeric; - using decimalXX = TypeParam; - using RepType = cudf::device_storage_type_t; - using fp_wrapper = cudf::test::fixed_point_column_wrapper; - - auto const mask = cudf::test::fixed_width_column_wrapper{0, 1, 1, 1, 0, 0}; - auto const a = fp_wrapper{{110, 220, 330, 440, 550, 660}, scale_type{-2}}; - auto const b = fp_wrapper{{0, 0, 0, 0, 0, 0}, scale_type{-2}}; - auto const expected = fp_wrapper{{0, 220, 330, 440, 0, 0}, scale_type{-2}}; - auto const result = cudf::copy_if_else(a, b, mask); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); -} - -TYPED_TEST(FixedPointTypes, FixedPointLarge) -{ - using namespace numeric; - using namespace cudf::test; - using decimalXX = TypeParam; - using RepType = cudf::device_storage_type_t; - using fp_wrapper = cudf::test::fixed_point_column_wrapper; - - auto a = thrust::make_counting_iterator(-1000); - auto b = thrust::make_constant_iterator(0); - auto m = cudf::detail::make_counting_transform_iterator(-1000, [](int i) { return i > 0; }); - auto e = - cudf::detail::make_counting_transform_iterator(-1000, [](int i) { return std::max(0, i); }); - - auto const mask = cudf::test::fixed_width_column_wrapper(m, m + 2000); - auto const A = fp_wrapper{a, a + 2000, scale_type{-3}}; - auto const B = fp_wrapper{b, b + 2000, scale_type{-3}}; - auto const expected = fp_wrapper{e, e + 2000, scale_type{-3}}; - auto const result = cudf::copy_if_else(A, B, mask); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); -} - -TYPED_TEST(FixedPointTypes, FixedPointScaleMismatch) -{ - using namespace numeric; - using decimalXX = TypeParam; - using RepType = cudf::device_storage_type_t; - using fp_wrapper = cudf::test::fixed_point_column_wrapper; - - auto const mask = cudf::test::fixed_width_column_wrapper{0, 1, 1, 1, 0, 0}; - auto const a = fp_wrapper{{110, 220, 330, 440, 550, 660}, scale_type{-2}}; - auto const b = fp_wrapper{{0, 0, 0, 0, 0, 0}, scale_type{-1}}; - - EXPECT_THROW(cudf::copy_if_else(a, b, mask), cudf::logic_error); -} From 9d56d0b10a8c0d64858e9b2dc7768ddec7b9e4fa Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 28 Sep 2021 16:02:56 -0400 Subject: [PATCH 4/9] add missing include --- cpp/src/copying/copy.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu index 72e5582f9aa..6669daf97b5 100644 --- a/cpp/src/copying/copy.cu +++ b/cpp/src/copying/copy.cu @@ -26,6 +26,7 @@ #include #include +#include namespace cudf { namespace detail { From efac52d2e3b549811c36f5fb79bf0060ceef72eb Mon Sep 17 00:00:00 2001 From: David Wendt Date: Wed, 29 Sep 2021 14:31:52 -0400 Subject: [PATCH 5/9] create multi-block test and remove detail_copy_tests.cu --- cpp/tests/copying/copy_tests.cpp | 21 +++++ cpp/tests/copying/detail_copy_tests.cu | 122 ------------------------- 2 files changed, 21 insertions(+), 122 deletions(-) delete mode 100644 cpp/tests/copying/detail_copy_tests.cu diff --git a/cpp/tests/copying/copy_tests.cpp b/cpp/tests/copying/copy_tests.cpp index 9a6a241d5c2..8d4e9295783 100644 --- a/cpp/tests/copying/copy_tests.cpp +++ b/cpp/tests/copying/copy_tests.cpp @@ -101,6 +101,27 @@ TYPED_TEST(CopyTest, CopyIfElseTestLong) CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); } +TYPED_TEST(CopyTest, CopyIfElseTestMultipleBlocks) +{ + using T = TypeParam; + + int num = 1000; // larger than a single block + std::vector h_lhs(num, 5); + std::vector h_rhs(num, 6); + std::vector h_mask(num, false); + std::vector h_validity(num, true); + h_validity[0] = 0; + + cudf::test::fixed_width_column_wrapper lhs_w( + h_lhs.begin(), h_lhs.end(), h_validity.begin()); + cudf::test::fixed_width_column_wrapper rhs_w( + h_rhs.begin(), h_rhs.end(), h_validity.begin()); + cudf::test::fixed_width_column_wrapper mask_w(h_mask.begin(), h_mask.end()); + + auto out = cudf::copy_if_else(lhs_w, rhs_w, mask_w); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), rhs_w); +} + TYPED_TEST(CopyTest, CopyIfElseTestEmptyInputs) { using T = TypeParam; diff --git a/cpp/tests/copying/detail_copy_tests.cu b/cpp/tests/copying/detail_copy_tests.cu deleted file mode 100644 index cd45c6169c7..00000000000 --- a/cpp/tests/copying/detail_copy_tests.cu +++ /dev/null @@ -1,122 +0,0 @@ -/* - * Copyright (c) 2021, 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 -#include -#include -#include -#include - -#include -#include -#include -#include -#include - -#include - -template -struct CopyDetailTest : public cudf::test::BaseFixture { -}; - -TYPED_TEST_CASE(CopyDetailTest, cudf::test::FixedWidthTypesWithoutFixedPoint); - -struct copy_if_else_tiny_grid_functor { - template ())> - std::unique_ptr operator()(cudf::column_view const& lhs, - cudf::column_view const& rhs, - Filter filter, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) - { - // output - std::unique_ptr out = - cudf::allocate_like(lhs, lhs.size(), cudf::mask_allocation_policy::RETAIN, mr); - - // device views - auto lhs_view = cudf::column_device_view::create(lhs); - auto rhs_view = cudf::column_device_view::create(rhs); - auto lhs_iter = cudf::detail::make_optional_iterator(*lhs_view, cudf::contains_nulls::NO{}); - auto rhs_iter = cudf::detail::make_optional_iterator(*rhs_view, cudf::contains_nulls::NO{}); - auto out_dv = cudf::mutable_column_device_view::create(*out); - - // call the kernel with an artificially small grid - cudf::detail::copy_if_else_kernel<32, T, decltype(lhs_iter), decltype(rhs_iter), Filter, false> - <<<1, 32, 0, stream.value()>>>(lhs_iter, rhs_iter, filter, *out_dv, nullptr); - - return out; - } - - template ())> - std::unique_ptr operator()(cudf::column_view const& lhs, - cudf::column_view const& rhs, - Filter filter, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) - { - CUDF_FAIL("Unexpected test execution"); - } -}; - -std::unique_ptr tiny_grid_launch(cudf::column_view const& lhs, - cudf::column_view const& rhs, - cudf::column_view const& boolean_mask) -{ - auto bool_mask_device_p = cudf::column_device_view::create(boolean_mask); - cudf::column_device_view bool_mask_device = *bool_mask_device_p; - auto filter = [bool_mask_device] __device__(cudf::size_type i) { - return bool_mask_device.element(i); - }; - return cudf::type_dispatcher(lhs.type(), - copy_if_else_tiny_grid_functor{}, - lhs, - rhs, - filter, - rmm::cuda_stream_default, - rmm::mr::get_current_device_resource()); -} - -TYPED_TEST(CopyDetailTest, CopyIfElseTestTinyGrid) -{ - using T = TypeParam; - - // make sure we span at least 2 warps - int num_els = 64; - - bool mask[] = {1, 0, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, - 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; - cudf::test::fixed_width_column_wrapper mask_w(mask, mask + num_els); - - cudf::test::fixed_width_column_wrapper lhs_w( - {5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, - 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, - 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5}); - - cudf::test::fixed_width_column_wrapper rhs_w( - {6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, - 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, - 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6}); - - cudf::test::fixed_width_column_wrapper expected_w( - {5, 6, 5, 6, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 6, 6, 6, - 6, 5, 5, 5, 5, 5, 5, 5, 5, 5, 6, 6, 6, 6, 5, 5, 5, 5, 5, 5, 5, 5, - 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5, 5}); - - auto out = tiny_grid_launch(lhs_w, rhs_w, mask_w); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), expected_w); -} From 3b9f51f65141df126c3552e6f6de60915cf8a3c0 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Thu, 30 Sep 2021 10:19:16 -0400 Subject: [PATCH 6/9] remove valid var --- cpp/include/cudf/detail/copy_if_else.cuh | 17 +++++++---------- cpp/tests/CMakeLists.txt | 1 - 2 files changed, 7 insertions(+), 11 deletions(-) diff --git a/cpp/include/cudf/detail/copy_if_else.cuh b/cpp/include/cudf/detail/copy_if_else.cuh index 596e119fd2f..aead26e0958 100644 --- a/cpp/include/cudf/detail/copy_if_else.cuh +++ b/cpp/include/cudf/detail/copy_if_else.cuh @@ -33,7 +33,7 @@ template + bool has_nulls> __launch_bounds__(block_size) __global__ void copy_if_else_kernel(LeftIter lhs, RightIter rhs, @@ -64,17 +64,14 @@ __launch_bounds__(block_size) __global__ size_type warp_cur = warp_begin + warp_id; size_type index = tid; while (warp_cur <= warp_end) { - bool valid = false; - if (index >= begin && index < end) { - auto value = filter(index) ? lhs[index] : rhs[index]; - valid = !has_validity || value.has_value(); - if (valid) { out.element(index) = static_cast(value.value()); } - } + auto const opt_value = + (index < end) ? (filter(index) ? lhs[index] : rhs[index]) : thrust::nullopt; + if (not has_nulls or opt_value) { out.element(index) = static_cast(opt_value.value()); } // update validity - if (has_validity) { + if (has_nulls) { // the final validity mask for this warp - int warp_mask = __ballot_sync(0xFFFF'FFFF, valid); + int warp_mask = __ballot_sync(0xFFFF'FFFF, opt_value.has_value()); // only one guy in the warp needs to update the mask and count if (lane_id == 0) { out.set_mask_word(warp_cur, warp_mask); @@ -87,7 +84,7 @@ __launch_bounds__(block_size) __global__ index += block_size * gridDim.x; } - if (has_validity) { + if (has_nulls) { // sum all null counts across all warps size_type block_valid_count = single_lane_block_sum_reduce(warp_valid_count); diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 09d0c07be1c..f7bde31e6ea 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -219,7 +219,6 @@ ConfigureTest(COPYING_TEST copying/copy_if_else_nested_tests.cpp copying/copy_range_tests.cpp copying/copy_tests.cpp - copying/detail_copy_tests.cu copying/detail_gather_tests.cu copying/gather_list_tests.cpp copying/gather_str_tests.cpp From f0fed625010746092242ca849b3ee2f595217585 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Thu, 30 Sep 2021 11:41:42 -0400 Subject: [PATCH 7/9] use operator* instead of value() --- cpp/include/cudf/detail/copy_if_else.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cudf/detail/copy_if_else.cuh b/cpp/include/cudf/detail/copy_if_else.cuh index aead26e0958..c2455e8502a 100644 --- a/cpp/include/cudf/detail/copy_if_else.cuh +++ b/cpp/include/cudf/detail/copy_if_else.cuh @@ -66,7 +66,7 @@ __launch_bounds__(block_size) __global__ while (warp_cur <= warp_end) { auto const opt_value = (index < end) ? (filter(index) ? lhs[index] : rhs[index]) : thrust::nullopt; - if (not has_nulls or opt_value) { out.element(index) = static_cast(opt_value.value()); } + if (not has_nulls or opt_value) { out.element(index) = static_cast(*opt_value); } // update validity if (has_nulls) { From a5f17048764d7cfd665e330a3ecf834945f8f9e0 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Fri, 1 Oct 2021 09:55:31 -0400 Subject: [PATCH 8/9] undo indexalator commit --- cpp/include/cudf/detail/indexalator.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cudf/detail/indexalator.cuh b/cpp/include/cudf/detail/indexalator.cuh index d6a07d5ec4b..d0fa4e02440 100644 --- a/cpp/include/cudf/detail/indexalator.cuh +++ b/cpp/include/cudf/detail/indexalator.cuh @@ -484,7 +484,7 @@ struct indexalator_factory { /** * @brief An index accessor that returns a validity flag along with the index value. * - * This is suitable as a `pair_iterator`. + * This is suitable as a `pair_iterator` for calling functions like `copy_if_else`. */ struct nullable_index_accessor { input_indexalator iter; From 37c513680a32a5e52d7e325c29f6f28980ebdc20 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 11 Oct 2021 08:13:11 -0400 Subject: [PATCH 9/9] remove unneeded include --- cpp/benchmarks/copying/copy_if_else_benchmark.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/cpp/benchmarks/copying/copy_if_else_benchmark.cpp b/cpp/benchmarks/copying/copy_if_else_benchmark.cpp index 0a5c425e601..513e4f4c179 100644 --- a/cpp/benchmarks/copying/copy_if_else_benchmark.cpp +++ b/cpp/benchmarks/copying/copy_if_else_benchmark.cpp @@ -23,8 +23,6 @@ #include -// #include - class CopyIfElse : public cudf::benchmark { };