From 1f6ec29747ac6c9f8b924ba0880b1c9a33168bf9 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 4 May 2021 15:28:43 +1000 Subject: [PATCH 01/27] Convert to uvector in bitmask_tests.cu --- cpp/tests/bitmask/bitmask_tests.cu | 205 ++++++++++++++++------------- 1 file changed, 111 insertions(+), 94 deletions(-) diff --git a/cpp/tests/bitmask/bitmask_tests.cu b/cpp/tests/bitmask/bitmask_tests.cu index 2f820da687e..3fb12efcc93 100644 --- a/cpp/tests/bitmask/bitmask_tests.cu +++ b/cpp/tests/bitmask/bitmask_tests.cu @@ -15,6 +15,7 @@ */ #include #include +#include #include #include #include @@ -23,9 +24,9 @@ #include #include -#include -#include #include +#include "rmm/cuda_stream_view.hpp" +#include "rmm/device_uvector.hpp" struct BitmaskUtilitiesTest : public cudf::test::BaseFixture { }; @@ -76,161 +77,177 @@ TEST_F(CountBitmaskTest, NullMask) } } +// Utility to construct a mask vector. If fill_valid is false (default), it is initialized to all +// null. Otherwise it is initialized to all valid. +rmm::device_uvector make_mask(cudf::size_type size, bool fill_valid = false) +{ + if (!fill_valid) { + return cudf::detail::make_zeroed_device_uvector_sync(size); + } else { + auto ret = rmm::device_uvector(size, rmm::cuda_stream_default); + CUDA_TRY(cudaMemsetAsync(ret.data(), + ~cudf::bitmask_type{0}, + size * sizeof(cudf::bitmask_type), + rmm::cuda_stream_default.value())); + return ret; + } +} + TEST_F(CountBitmaskTest, NegativeStart) { - thrust::device_vector mask(1, 0); - EXPECT_THROW(cudf::count_set_bits(mask.data().get(), -1, 32), cudf::logic_error); + auto mask = make_mask(1); + EXPECT_THROW(cudf::count_set_bits(mask.data(), -1, 32), cudf::logic_error); std::vector indices = {0, 16, -1, 32}; - EXPECT_THROW(cudf::segmented_count_set_bits(mask.data().get(), indices), cudf::logic_error); + EXPECT_THROW(cudf::segmented_count_set_bits(mask.data(), indices), cudf::logic_error); } TEST_F(CountBitmaskTest, StartLargerThanStop) { - thrust::device_vector mask(1, 0); - EXPECT_THROW(cudf::count_set_bits(mask.data().get(), 32, 31), cudf::logic_error); + auto mask = make_mask(1); + EXPECT_THROW(cudf::count_set_bits(mask.data(), 32, 31), cudf::logic_error); std::vector indices = {0, 16, 31, 30}; - EXPECT_THROW(cudf::segmented_count_set_bits(mask.data().get(), indices), cudf::logic_error); + EXPECT_THROW(cudf::segmented_count_set_bits(mask.data(), indices), cudf::logic_error); } TEST_F(CountBitmaskTest, EmptyRange) { - thrust::device_vector mask(1, 0); - EXPECT_EQ(0, cudf::count_set_bits(mask.data().get(), 17, 17)); + auto mask = make_mask(1); + EXPECT_EQ(0, cudf::count_set_bits(mask.data(), 17, 17)); std::vector indices = {0, 0, 17, 17}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{0, 0})); } TEST_F(CountBitmaskTest, SingleWordAllZero) { - thrust::device_vector mask(1, 0); - EXPECT_EQ(0, cudf::count_set_bits(mask.data().get(), 0, 32)); + auto mask = make_mask(1); + EXPECT_EQ(0, cudf::count_set_bits(mask.data(), 0, 32)); std::vector indices = {0, 32, 0, 32}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{0, 0})); } TEST_F(CountBitmaskTest, SingleBitAllZero) { - thrust::device_vector mask(1, 0); - EXPECT_EQ(0, cudf::count_set_bits(mask.data().get(), 17, 18)); + auto mask = make_mask(1); + EXPECT_EQ(0, cudf::count_set_bits(mask.data(), 17, 18)); std::vector indices = {17, 18, 7, 8}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{0, 0})); } TEST_F(CountBitmaskTest, SingleBitAllSet) { - thrust::device_vector mask(1, ~cudf::bitmask_type{0}); - EXPECT_EQ(1, cudf::count_set_bits(mask.data().get(), 13, 14)); + auto mask = make_mask(1, true); + EXPECT_EQ(1, cudf::count_set_bits(mask.data(), 13, 14)); std::vector indices = {13, 14, 0, 1}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{1, 1})); } TEST_F(CountBitmaskTest, SingleWordAllBitsSet) { - thrust::device_vector mask(1, ~cudf::bitmask_type{0}); - EXPECT_EQ(32, cudf::count_set_bits(mask.data().get(), 0, 32)); + auto mask = make_mask(1, true); + EXPECT_EQ(32, cudf::count_set_bits(mask.data(), 0, 32)); std::vector indices = {0, 32, 0, 32}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{32, 32})); } TEST_F(CountBitmaskTest, SingleWordPreSlack) { - thrust::device_vector mask(1, ~cudf::bitmask_type{0}); - EXPECT_EQ(25, cudf::count_set_bits(mask.data().get(), 7, 32)); + auto mask = make_mask(1, true); + EXPECT_EQ(25, cudf::count_set_bits(mask.data(), 7, 32)); std::vector indices = {7, 32, 8, 32}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{25, 24})); } TEST_F(CountBitmaskTest, SingleWordPostSlack) { - thrust::device_vector mask(1, ~cudf::bitmask_type{0}); - EXPECT_EQ(17, cudf::count_set_bits(mask.data().get(), 0, 17)); + auto mask = make_mask(1, true); + EXPECT_EQ(17, cudf::count_set_bits(mask.data(), 0, 17)); std::vector indices = {0, 17, 0, 18}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{17, 18})); } TEST_F(CountBitmaskTest, SingleWordSubset) { - thrust::device_vector mask(1, ~cudf::bitmask_type{0}); - EXPECT_EQ(30, cudf::count_set_bits(mask.data().get(), 1, 31)); + auto mask = make_mask(1, true); + EXPECT_EQ(30, cudf::count_set_bits(mask.data(), 1, 31)); std::vector indices = {1, 31, 7, 17}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{30, 10})); } TEST_F(CountBitmaskTest, SingleWordSubset2) { - thrust::device_vector mask(1, ~cudf::bitmask_type{0}); - EXPECT_EQ(28, cudf::count_set_bits(mask.data().get(), 2, 30)); + auto mask = make_mask(1, true); + EXPECT_EQ(28, cudf::count_set_bits(mask.data(), 2, 30)); std::vector indices = {4, 16, 2, 30}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{12, 28})); } TEST_F(CountBitmaskTest, MultipleWordsAllBits) { - thrust::device_vector mask(10, ~cudf::bitmask_type{0}); - EXPECT_EQ(320, cudf::count_set_bits(mask.data().get(), 0, 320)); + auto mask = make_mask(10, true); + EXPECT_EQ(320, cudf::count_set_bits(mask.data(), 0, 320)); std::vector indices = {0, 320, 0, 320}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{320, 320})); } TEST_F(CountBitmaskTest, MultipleWordsSubsetWordBoundary) { - thrust::device_vector mask(10, ~cudf::bitmask_type{0}); - EXPECT_EQ(256, cudf::count_set_bits(mask.data().get(), 32, 288)); + auto mask = make_mask(10, true); + EXPECT_EQ(256, cudf::count_set_bits(mask.data(), 32, 288)); std::vector indices = {32, 192, 32, 288}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{160, 256})); } TEST_F(CountBitmaskTest, MultipleWordsSplitWordBoundary) { - thrust::device_vector mask(10, ~cudf::bitmask_type{0}); - EXPECT_EQ(2, cudf::count_set_bits(mask.data().get(), 31, 33)); + auto mask = make_mask(10, true); + EXPECT_EQ(2, cudf::count_set_bits(mask.data(), 31, 33)); std::vector indices = {31, 33, 60, 67}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{2, 7})); } TEST_F(CountBitmaskTest, MultipleWordsSubset) { - thrust::device_vector mask(10, ~cudf::bitmask_type{0}); - EXPECT_EQ(226, cudf::count_set_bits(mask.data().get(), 67, 293)); + auto mask = make_mask(10, true); + EXPECT_EQ(226, cudf::count_set_bits(mask.data(), 67, 293)); std::vector indices = {67, 293, 37, 319}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{226, 282})); } TEST_F(CountBitmaskTest, MultipleWordsSingleBit) { - thrust::device_vector mask(10, ~cudf::bitmask_type{0}); - EXPECT_EQ(1, cudf::count_set_bits(mask.data().get(), 67, 68)); + auto mask = make_mask(10, true); + EXPECT_EQ(1, cudf::count_set_bits(mask.data(), 67, 68)); std::vector indices = {67, 68, 31, 32, 192, 193}; - auto counts = cudf::segmented_count_set_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_set_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{1, 1, 1})); } @@ -238,11 +255,11 @@ using CountUnsetBitsTest = CountBitmaskTest; TEST_F(CountUnsetBitsTest, SingleBitAllSet) { - thrust::device_vector mask(1, ~cudf::bitmask_type{0}); - EXPECT_EQ(0, cudf::count_unset_bits(mask.data().get(), 13, 14)); + auto mask = make_mask(1, true); + EXPECT_EQ(0, cudf::count_unset_bits(mask.data(), 13, 14)); std::vector indices = {13, 14, 31, 32}; - auto counts = cudf::segmented_count_unset_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_unset_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{0, 0})); } @@ -258,101 +275,101 @@ TEST_F(CountUnsetBitsTest, NullMask) TEST_F(CountUnsetBitsTest, SingleWordAllBits) { - thrust::device_vector mask(1, cudf::bitmask_type{0}); - EXPECT_EQ(32, cudf::count_unset_bits(mask.data().get(), 0, 32)); + auto mask = make_mask(1); + EXPECT_EQ(32, cudf::count_unset_bits(mask.data(), 0, 32)); std::vector indices = {0, 32, 0, 32}; - auto counts = cudf::segmented_count_unset_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_unset_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{32, 32})); } TEST_F(CountUnsetBitsTest, SingleWordPreSlack) { - thrust::device_vector mask(1, cudf::bitmask_type{0}); - EXPECT_EQ(25, cudf::count_unset_bits(mask.data().get(), 7, 32)); + auto mask = make_mask(1); + EXPECT_EQ(25, cudf::count_unset_bits(mask.data(), 7, 32)); std::vector indices = {7, 32, 8, 32}; - auto counts = cudf::segmented_count_unset_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_unset_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{25, 24})); } TEST_F(CountUnsetBitsTest, SingleWordPostSlack) { - thrust::device_vector mask(1, cudf::bitmask_type{0}); - EXPECT_EQ(17, cudf::count_unset_bits(mask.data().get(), 0, 17)); + auto mask = make_mask(1); + EXPECT_EQ(17, cudf::count_unset_bits(mask.data(), 0, 17)); std::vector indices = {0, 17, 0, 18}; - auto counts = cudf::segmented_count_unset_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_unset_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{17, 18})); } TEST_F(CountUnsetBitsTest, SingleWordSubset) { - thrust::device_vector mask(1, cudf::bitmask_type{0}); - EXPECT_EQ(30, cudf::count_unset_bits(mask.data().get(), 1, 31)); + auto mask = make_mask(1); + EXPECT_EQ(30, cudf::count_unset_bits(mask.data(), 1, 31)); std::vector indices = {1, 31, 7, 17}; - auto counts = cudf::segmented_count_unset_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_unset_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{30, 10})); } TEST_F(CountUnsetBitsTest, SingleWordSubset2) { - thrust::device_vector mask(1, cudf::bitmask_type{0}); - EXPECT_EQ(28, cudf::count_unset_bits(mask.data().get(), 2, 30)); + auto mask = make_mask(1); + EXPECT_EQ(28, cudf::count_unset_bits(mask.data(), 2, 30)); std::vector indices = {4, 16, 2, 30}; - auto counts = cudf::segmented_count_unset_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_unset_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{12, 28})); } TEST_F(CountUnsetBitsTest, MultipleWordsAllBits) { - thrust::device_vector mask(10, cudf::bitmask_type{0}); - EXPECT_EQ(320, cudf::count_unset_bits(mask.data().get(), 0, 320)); + auto mask = make_mask(10); + EXPECT_EQ(320, cudf::count_unset_bits(mask.data(), 0, 320)); std::vector indices = {0, 320, 0, 320}; - auto counts = cudf::segmented_count_unset_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_unset_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{320, 320})); } TEST_F(CountUnsetBitsTest, MultipleWordsSubsetWordBoundary) { - thrust::device_vector mask(10, cudf::bitmask_type{0}); - EXPECT_EQ(256, cudf::count_unset_bits(mask.data().get(), 32, 288)); + auto mask = make_mask(10); + EXPECT_EQ(256, cudf::count_unset_bits(mask.data(), 32, 288)); std::vector indices = {32, 192, 32, 288}; - auto counts = cudf::segmented_count_unset_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_unset_bits(mask.data(), indices); EXPECT_THAT(counts, testing::ContainerEq(std::vector{160, 256})); } TEST_F(CountUnsetBitsTest, MultipleWordsSplitWordBoundary) { - thrust::device_vector mask(10, cudf::bitmask_type{0}); - EXPECT_EQ(2, cudf::count_unset_bits(mask.data().get(), 31, 33)); + auto mask = make_mask(10); + EXPECT_EQ(2, cudf::count_unset_bits(mask.data(), 31, 33)); std::vector indices = {31, 33, 60, 67}; - auto counts = cudf::segmented_count_unset_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_unset_bits(mask.data(), indices); EXPECT_THAT(counts, ::testing::ContainerEq(std::vector{2, 7})); } TEST_F(CountUnsetBitsTest, MultipleWordsSubset) { - thrust::device_vector mask(10, cudf::bitmask_type{0}); - EXPECT_EQ(226, cudf::count_unset_bits(mask.data().get(), 67, 293)); + auto mask = make_mask(10); + EXPECT_EQ(226, cudf::count_unset_bits(mask.data(), 67, 293)); std::vector indices = {67, 293, 37, 319}; - auto counts = cudf::segmented_count_unset_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_unset_bits(mask.data(), indices); EXPECT_THAT(counts, ::testing::ContainerEq(std::vector{226, 282})); } TEST_F(CountUnsetBitsTest, MultipleWordsSingleBit) { - thrust::device_vector mask(10, cudf::bitmask_type{0}); - EXPECT_EQ(1, cudf::count_unset_bits(mask.data().get(), 67, 68)); + auto mask = make_mask(10); + EXPECT_EQ(1, cudf::count_unset_bits(mask.data(), 67, 68)); std::vector indices = {67, 68, 31, 32, 192, 193}; - auto counts = cudf::segmented_count_unset_bits(mask.data().get(), indices); + auto counts = cudf::segmented_count_unset_bits(mask.data(), indices); EXPECT_THAT(counts, ::testing::ContainerEq(std::vector{1, 1, 1})); } @@ -362,7 +379,7 @@ struct CopyBitmaskTest : public cudf::test::BaseFixture, cudf::test::UniformRand void cleanEndWord(rmm::device_buffer &mask, int begin_bit, int end_bit) { - thrust::device_ptr ptr(static_cast(mask.data())); + auto ptr = static_cast(mask.data()); auto number_of_mask_words = cudf::num_bitmask_words(static_cast(end_bit - begin_bit)); auto number_of_bits = end_bit - begin_bit; @@ -374,20 +391,20 @@ void cleanEndWord(rmm::device_buffer &mask, int begin_bit, int end_bit) TEST_F(CopyBitmaskTest, NegativeStart) { - thrust::device_vector mask(1, 0); - EXPECT_THROW(cudf::copy_bitmask(mask.data().get(), -1, 32), cudf::logic_error); + auto mask = make_mask(1); + EXPECT_THROW(cudf::copy_bitmask(mask.data(), -1, 32), cudf::logic_error); } TEST_F(CopyBitmaskTest, StartLargerThanStop) { - thrust::device_vector mask(1, 0); - EXPECT_THROW(cudf::copy_bitmask(mask.data().get(), 32, 31), cudf::logic_error); + auto mask = make_mask(1); + EXPECT_THROW(cudf::copy_bitmask(mask.data(), 32, 31), cudf::logic_error); } TEST_F(CopyBitmaskTest, EmptyRange) { - thrust::device_vector mask(1, 0); - auto buff = cudf::copy_bitmask(mask.data().get(), 17, 17); + auto mask = make_mask(1); + auto buff = cudf::copy_bitmask(mask.data(), 17, 17); EXPECT_EQ(0, static_cast(buff.size())); } @@ -399,7 +416,7 @@ TEST_F(CopyBitmaskTest, NullPtr) TEST_F(CopyBitmaskTest, TestZeroOffset) { - thrust::host_vector validity_bit(1000); + std::vector validity_bit(1000); for (auto &m : validity_bit) { m = this->generate(); } auto input_mask = cudf::test::detail::make_null_mask(validity_bit.begin(), validity_bit.end()); @@ -419,7 +436,7 @@ TEST_F(CopyBitmaskTest, TestZeroOffset) TEST_F(CopyBitmaskTest, TestNonZeroOffset) { - thrust::host_vector validity_bit(1000); + std::vector validity_bit(1000); for (auto &m : validity_bit) { m = this->generate(); } auto input_mask = cudf::test::detail::make_null_mask(validity_bit.begin(), validity_bit.end()); @@ -441,7 +458,7 @@ TEST_F(CopyBitmaskTest, TestCopyColumnViewVectorContiguous) { cudf::data_type t{cudf::type_id::INT32}; cudf::size_type num_elements = 1001; - thrust::host_vector validity_bit(num_elements); + std::vector validity_bit(num_elements); for (auto &m : validity_bit) { m = this->generate(); } auto gold_mask = cudf::test::detail::make_null_mask(validity_bit.begin(), validity_bit.end()); @@ -476,7 +493,7 @@ TEST_F(CopyBitmaskTest, TestCopyColumnViewVectorDiscontiguous) { cudf::data_type t{cudf::type_id::INT32}; cudf::size_type num_elements = 1001; - thrust::host_vector validity_bit(num_elements); + std::vector validity_bit(num_elements); for (auto &m : validity_bit) { m = this->generate(); } auto gold_mask = cudf::test::detail::make_null_mask(validity_bit.begin(), validity_bit.end()); std::vector split{0, 104, 128, 152, 311, 491, 583, 734, 760, num_elements}; From 2708c0160a93f71e3bdd7342ceed94d0fe01c57d Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 4 May 2021 15:30:43 +1000 Subject: [PATCH 02/27] Rename bitmask_tests .cu -> .cpp --- cpp/tests/CMakeLists.txt | 3 +-- cpp/tests/bitmask/{bitmask_tests.cu => bitmask_tests.cpp} | 0 2 files changed, 1 insertion(+), 2 deletions(-) rename cpp/tests/bitmask/{bitmask_tests.cu => bitmask_tests.cpp} (100%) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 6dd50592274..8fb783e2281 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -368,8 +368,7 @@ ConfigureTest(TEXT_TEST ConfigureTest(BITMASK_TEST bitmask/valid_if_tests.cu bitmask/set_nullmask_tests.cu - bitmask/bitmask_tests.cu) - + bitmask/bitmask_tests.cpp) ################################################################################################### # - dictionary tests ------------------------------------------------------------------------------ diff --git a/cpp/tests/bitmask/bitmask_tests.cu b/cpp/tests/bitmask/bitmask_tests.cpp similarity index 100% rename from cpp/tests/bitmask/bitmask_tests.cu rename to cpp/tests/bitmask/bitmask_tests.cpp From e95078bec29004bed11c8c7bda93eea50ec6f36b Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 4 May 2021 15:37:13 +1000 Subject: [PATCH 03/27] Convert compound_test to uvector --- cpp/tests/column/compound_test.cu | 41 +++++++++++++++++-------------- 1 file changed, 22 insertions(+), 19 deletions(-) diff --git a/cpp/tests/column/compound_test.cu b/cpp/tests/column/compound_test.cu index 97a6dbb0c22..57ec93be55a 100644 --- a/cpp/tests/column/compound_test.cu +++ b/cpp/tests/column/compound_test.cu @@ -21,10 +21,13 @@ #include #include -#include -#include +#include +#include +#include + #include #include + #include struct CompoundColumnTest : public cudf::test::BaseFixture { @@ -61,13 +64,13 @@ struct checker_for_level2 { TEST_F(CompoundColumnTest, ChildrenLevel1) { - thrust::device_vector data(1000); - thrust::sequence(thrust::device, data.begin(), data.end(), 1); + rmm::device_uvector data(1000, rmm::cuda_stream_default); + thrust::sequence(rmm::exec_policy(rmm::cuda_stream_default), data.begin(), data.end(), 1); auto null_mask = cudf::create_null_mask(100, cudf::mask_state::UNALLOCATED); - rmm::device_buffer data1(data.data().get() + 100, 100 * sizeof(int32_t)); - rmm::device_buffer data2(data.data().get() + 200, 100 * sizeof(int32_t)); - rmm::device_buffer data3(data.data().get() + 300, 100 * sizeof(int32_t)); + rmm::device_buffer data1(data.data() + 100, 100 * sizeof(int32_t)); + rmm::device_buffer data2(data.data() + 200, 100 * sizeof(int32_t)); + rmm::device_buffer data3(data.data() + 300, 100 * sizeof(int32_t)); auto child1 = std::make_unique(cudf::data_type{cudf::type_id::INT32}, 100, data1, null_mask, 0); auto child2 = @@ -89,14 +92,14 @@ TEST_F(CompoundColumnTest, ChildrenLevel1) { auto column = cudf::column_device_view::create(parent->view()); - EXPECT_TRUE(thrust::any_of(thrust::device, + EXPECT_TRUE(thrust::any_of(rmm::exec_policy(rmm::cuda_stream_default), thrust::make_counting_iterator(0), thrust::make_counting_iterator(100), checker_for_level1{*column})); } { auto column = cudf::mutable_column_device_view::create(parent->mutable_view()); - EXPECT_TRUE(thrust::any_of(thrust::device, + EXPECT_TRUE(thrust::any_of(rmm::exec_policy(rmm::cuda_stream_default), thrust::make_counting_iterator(0), thrust::make_counting_iterator(100), checker_for_level1{*column})); @@ -105,16 +108,16 @@ TEST_F(CompoundColumnTest, ChildrenLevel1) TEST_F(CompoundColumnTest, ChildrenLevel2) { - thrust::device_vector data(1000); - thrust::sequence(thrust::device, data.begin(), data.end(), 1); + rmm::device_uvector data(1000, rmm::cuda_stream_default); + thrust::sequence(rmm::exec_policy(rmm::cuda_stream_default), data.begin(), data.end(), 1); auto null_mask = cudf::create_null_mask(100, cudf::mask_state::UNALLOCATED); - rmm::device_buffer data11(data.data().get() + 100, 100 * sizeof(int32_t)); - rmm::device_buffer data12(data.data().get() + 200, 100 * sizeof(int32_t)); - rmm::device_buffer data13(data.data().get() + 300, 100 * sizeof(int32_t)); - rmm::device_buffer data21(data.data().get() + 400, 100 * sizeof(int32_t)); - rmm::device_buffer data22(data.data().get() + 500, 100 * sizeof(int32_t)); - rmm::device_buffer data23(data.data().get() + 600, 100 * sizeof(int32_t)); + rmm::device_buffer data11(data.data() + 100, 100 * sizeof(int32_t)); + rmm::device_buffer data12(data.data() + 200, 100 * sizeof(int32_t)); + rmm::device_buffer data13(data.data() + 300, 100 * sizeof(int32_t)); + rmm::device_buffer data21(data.data() + 400, 100 * sizeof(int32_t)); + rmm::device_buffer data22(data.data() + 500, 100 * sizeof(int32_t)); + rmm::device_buffer data23(data.data() + 600, 100 * sizeof(int32_t)); auto gchild11 = std::make_unique( cudf::data_type{cudf::type_id::INT32}, 100, data11, null_mask, 0); auto gchild12 = std::make_unique( @@ -162,14 +165,14 @@ TEST_F(CompoundColumnTest, ChildrenLevel2) { auto column = cudf::column_device_view::create(parent->view()); - EXPECT_TRUE(thrust::any_of(thrust::device, + EXPECT_TRUE(thrust::any_of(rmm::exec_policy(rmm::cuda_stream_default), thrust::make_counting_iterator(0), thrust::make_counting_iterator(100), checker_for_level2{*column})); } { auto column = cudf::mutable_column_device_view::create(parent->mutable_view()); - EXPECT_TRUE(thrust::any_of(thrust::device, + EXPECT_TRUE(thrust::any_of(rmm::exec_policy(rmm::cuda_stream_default), thrust::make_counting_iterator(0), thrust::make_counting_iterator(100), checker_for_level2{*column})); From b6a212effd944aa2894c7d8958baca04909e336f Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 May 2021 07:47:30 +1000 Subject: [PATCH 04/27] copyright --- cpp/tests/column/compound_test.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/column/compound_test.cu b/cpp/tests/column/compound_test.cu index 57ec93be55a..b5ed2ab7a49 100644 --- a/cpp/tests/column/compound_test.cu +++ b/cpp/tests/column/compound_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. From afe19e7e2228a0be357664ee2a2b32bcdfd2837d Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 May 2021 09:25:18 +1000 Subject: [PATCH 05/27] Add `make_host_vector` factories --- .../detail/utilities/vector_factories.hpp | 103 +++++++++++++++++- 1 file changed, 99 insertions(+), 4 deletions(-) diff --git a/cpp/include/cudf/detail/utilities/vector_factories.hpp b/cpp/include/cudf/detail/utilities/vector_factories.hpp index db8e6e4a156..8bcf75a9769 100644 --- a/cpp/include/cudf/detail/utilities/vector_factories.hpp +++ b/cpp/include/cudf/detail/utilities/vector_factories.hpp @@ -281,6 +281,20 @@ rmm::device_uvector make_device_uvector_sync( return make_device_uvector_sync(device_span{c}, stream, mr); } +namespace { + +// Utility function template to allow copying to either a thrust::host_vector or std::vector +template +OutContainer make_vector_async(device_span v, + rmm::cuda_stream_view stream = rmm::cuda_stream_default) +{ + OutContainer result(v.size()); + CUDA_TRY(cudaMemcpyAsync( + result.data(), v.data(), v.size() * sizeof(T), cudaMemcpyDeviceToHost, stream.value())); + return result; +} +} // namespace + /** * @brief Asynchronously construct a `std::vector` containing a copy of data from a * `device_span` @@ -296,10 +310,7 @@ template std::vector make_std_vector_async(device_span v, rmm::cuda_stream_view stream = rmm::cuda_stream_default) { - std::vector result(v.size()); - CUDA_TRY(cudaMemcpyAsync( - result.data(), v.data(), v.size() * sizeof(T), cudaMemcpyDeviceToHost, stream.value())); - return result; + return make_vector_async>(v, stream); } /** @@ -368,6 +379,90 @@ std::vector make_std_vector_sync( return make_std_vector_sync(device_span{c}, stream); } +/** + * @brief Asynchronously construct a `thrust::host_vector` containing a copy of data from a + * `device_span` + * + * @note This function does not synchronize `stream`. + * + * @tparam T The type of the data to copy + * @param source_data The device data to copy + * @param stream The stream on which to perform the copy + * @return The data copied to the host + */ +template +thrust::host_vector make_host_vector_async( + device_span v, rmm::cuda_stream_view stream = rmm::cuda_stream_default) +{ + return make_vector_async>(v, stream); +} + +/** + * @brief Asynchronously construct a `std::vector` containing a copy of data from a device + * container + * + * @note This function synchronizes `stream`. + * + * @tparam Container The type of the container to copy from + * @tparam T The type of the data to copy + * @param c The input device container from which to copy + * @param stream The stream on which to perform the copy + * @return The data copied to the host + */ +template < + typename Container, + std::enable_if_t< + std::is_convertible>::value>* = + nullptr> +thrust::host_vector make_host_vector_async( + Container const& c, rmm::cuda_stream_view stream = rmm::cuda_stream_default) +{ + return make_host_vector_async(device_span{c}, stream); +} + +/** + * @brief Synchronously construct a `std::vector` containing a copy of data from a + * `device_span` + * + * @note This function does a synchronize on `stream`. + * + * @tparam T The type of the data to copy + * @param source_data The device data to copy + * @param stream The stream on which to perform the copy + * @return The data copied to the host + */ +template +thrust::host_vector make_host_vector_sync( + device_span v, rmm::cuda_stream_view stream = rmm::cuda_stream_default) +{ + auto result = make_host_vector_async(v, stream); + stream.synchronize(); + return result; +} + +/** + * @brief Synchronously construct a `std::vector` containing a copy of data from a device + * container + * + * @note This function synchronizes `stream`. + * + * @tparam Container The type of the container to copy from + * @tparam T The type of the data to copy + * @param c The input device container from which to copy + * @param stream The stream on which to perform the copy + * @return The data copied to the host + */ +template < + typename Container, + std::enable_if_t< + std::is_convertible>::value>* = + nullptr> +thrust::host_vector make_host_vector_sync( + Container const& c, rmm::cuda_stream_view stream = rmm::cuda_stream_default) +{ + return make_host_vector_sync(device_span{c}, stream); +} + } // namespace detail } // namespace cudf From 3457ef0351414bf885e729e3c231a26b673fa458 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 May 2021 09:25:30 +1000 Subject: [PATCH 06/27] Convert device_atomics_test to device_uvector --- .../device_atomics/device_atomics_test.cu | 72 +++++++++---------- 1 file changed, 33 insertions(+), 39 deletions(-) diff --git a/cpp/tests/device_atomics/device_atomics_test.cu b/cpp/tests/device_atomics/device_atomics_test.cu index 1cdfd6ad8ef..aa53877f27d 100644 --- a/cpp/tests/device_atomics/device_atomics_test.cu +++ b/cpp/tests/device_atomics/device_atomics_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019, NVIDIA CORPORATION. + * 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. @@ -14,16 +14,19 @@ * limitations under the License. */ -#include +#include +#include #include #include -#include - #include #include #include +#include + +#include + template __global__ void gpu_atomic_test(T* result, T* data, size_t size) { @@ -89,13 +92,13 @@ __global__ void gpu_atomicCAS_test(T* result, T* data, size_t size) } template -typename std::enable_if_t(), T> accumulate(std::vector const& xs) +typename std::enable_if_t(), T> accumulate(cudf::host_span xs) { return std::accumulate(xs.begin(), xs.end(), T{0}); } template -typename std::enable_if_t(), T> accumulate(std::vector const& xs) +typename std::enable_if_t(), T> accumulate(cudf::host_span xs) { auto ys = std::vector(xs.size()); std::transform( @@ -112,8 +115,8 @@ struct AtomicsTest : public cudf::test::BaseFixture { { size_t vec_size = v_input.size(); - // use transform from std::vector instead. - std::vector v(vec_size); + // use transform from thrust::host_vector instead. + thrust::host_vector v(vec_size); std::transform(v_input.begin(), v_input.end(), v.begin(), [](int x) { T t = cudf::test::make_type_param_scalar(x); return t; @@ -124,7 +127,7 @@ struct AtomicsTest : public cudf::test::BaseFixture { exact[1] = *(std::min_element(v.begin(), v.end())); exact[2] = *(std::max_element(v.begin(), v.end())); - std::vector result_init(9); // +3 padding for int8 tests + thrust::host_vector result_init(9); // +3 padding for int8 tests result_init[0] = cudf::test::make_type_param_scalar(0); result_init[1] = std::numeric_limits::max(); result_init[2] = std::numeric_limits::min(); @@ -132,22 +135,20 @@ struct AtomicsTest : public cudf::test::BaseFixture { result_init[4] = result_init[1]; result_init[5] = result_init[2]; - thrust::device_vector dev_data(v); - thrust::device_vector dev_result(result_init); + auto dev_data = cudf::detail::make_device_uvector_sync(v); + auto dev_result = cudf::detail::make_device_uvector_sync(result_init); if (block_size == 0) { block_size = vec_size; } if (is_cas_test) { - gpu_atomicCAS_test<<>>( - dev_result.data().get(), dev_data.data().get(), vec_size); + gpu_atomicCAS_test<<>>(dev_result.data(), dev_data.data(), vec_size); } else { - gpu_atomic_test<<>>( - dev_result.data().get(), dev_data.data().get(), vec_size); + gpu_atomic_test<<>>(dev_result.data(), dev_data.data(), vec_size); } - thrust::host_vector host_result(dev_result); - CUDA_TRY(cudaDeviceSynchronize()); - CHECK_CUDA(0); + auto host_result = cudf::detail::make_host_vector_sync(dev_result); + + CHECK_CUDA(rmm::cuda_stream_default.value()); if (!is_timestamp_sum()) { EXPECT_EQ(host_result[0], exact[0]) << "atomicAdd test failed"; @@ -272,15 +273,10 @@ struct AtomicsBitwiseOpTest : public cudf::test::BaseFixture { return t; }); - std::vector identity = {T(~0ull), - T(0), - T(0), - T(~0ull), - T(0), - T(0), - T(0), - T(0), - T(0)}; // +3 elements padding for int8 tests + thrust::host_vector identity(9, T{0}); // +3 elements padding for int8 tests + identity[0] = T(~0ull); + identity[3] = T(~0ull); + T exact[3]; exact[0] = std::accumulate( v.begin(), v.end(), identity[0], [](T acc, uint64_t i) { return acc & T(i); }); @@ -289,22 +285,20 @@ struct AtomicsBitwiseOpTest : public cudf::test::BaseFixture { exact[2] = std::accumulate( v.begin(), v.end(), identity[2], [](T acc, uint64_t i) { return acc ^ T(i); }); - thrust::device_vector dev_result(identity); - thrust::device_vector dev_data(v); + auto dev_result = cudf::detail::make_device_uvector_sync(identity); + auto dev_data = cudf::detail::make_device_uvector_sync(v); if (block_size == 0) { block_size = vec_size; } - gpu_atomic_bitwiseOp_test - <<>>(reinterpret_cast(dev_result.data().get()), - reinterpret_cast(dev_data.data().get()), - vec_size); + gpu_atomic_bitwiseOp_test<<>>( + reinterpret_cast(dev_result.data()), reinterpret_cast(dev_data.data()), vec_size); + + auto host_result = cudf::detail::make_host_vector_sync(dev_result); - thrust::host_vector host_result(dev_result); - CUDA_TRY(cudaDeviceSynchronize()); - CHECK_CUDA(0); + CHECK_CUDA(rmm::cuda_stream_default.value()); - print_exact(exact, "exact"); - print_exact(host_result.data(), "result"); + // print_exact(exact, "exact"); + // print_exact(host_result.data(), "result"); EXPECT_EQ(host_result[0], exact[0]) << "atomicAnd test failed"; EXPECT_EQ(host_result[1], exact[1]) << "atomicOr test failed"; @@ -314,7 +308,7 @@ struct AtomicsBitwiseOpTest : public cudf::test::BaseFixture { EXPECT_EQ(host_result[5], exact[2]) << "atomicXor test(2) failed"; } - void print_exact(const T* v, const char* msg) + [[maybe_unused]] void print_exact(const T* v, const char* msg) { std::cout << std::hex << std::showbase; std::cout << "The " << msg << " = {" << +v[0] << ", " << +v[1] << ", " << +v[2] << "}" From d307b5594793b6fbf5d1cd73ac904267dafcc918 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 10:54:57 +1000 Subject: [PATCH 07/27] Convert fixed_point_tests --- cpp/tests/fixed_point/fixed_point_tests.cu | 26 +++++++++++++--------- 1 file changed, 15 insertions(+), 11 deletions(-) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 5f74e459bb1..124d9339ebf 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -35,6 +35,8 @@ #include #include #include +#include "cudf/detail/utilities/vector_factories.hpp" +#include "rmm/cuda_stream_view.hpp" using namespace numeric; @@ -507,37 +509,39 @@ TEST_F(FixedPointTest, DecimalXXThrustOnDevice) { using decimal32 = fixed_point; - thrust::device_vector vec1(1000, decimal32{1, scale_type{-2}}); + std::vector vec1(1000, decimal32{1, scale_type{-2}}); + auto d_vec1 = cudf::detail::make_device_uvector_sync(vec1); auto const sum = thrust::reduce( - rmm::exec_policy(), std::cbegin(vec1), std::cend(vec1), decimal32{0, scale_type{-2}}); + rmm::exec_policy(), std::cbegin(d_vec1), std::cend(d_vec1), decimal32{0, scale_type{-2}}); EXPECT_EQ(static_cast(sum), 1000); // TODO: Once nvbugs/1990211 is fixed (ExclusiveSum initial_value = 0 bug) // change inclusive scan to run on device (avoid copying to host) - thrust::host_vector vec1_host = vec1; + thrust::inclusive_scan(std::cbegin(vec1), std::cend(vec1), std::begin(vec1)); - thrust::inclusive_scan(std::cbegin(vec1_host), std::cend(vec1_host), std::begin(vec1_host)); - - vec1 = vec1_host; + d_vec1 = cudf::detail::make_device_uvector_sync(vec1); std::vector vec2(1000); std::iota(std::begin(vec2), std::end(vec2), 1); auto const res1 = thrust::reduce( - rmm::exec_policy(), std::cbegin(vec1), std::cend(vec1), decimal32{0, scale_type{-2}}); + rmm::exec_policy(), std::cbegin(d_vec1), std::cend(d_vec1), decimal32{0, scale_type{-2}}); auto const res2 = std::accumulate(std::cbegin(vec2), std::cend(vec2), 0); EXPECT_EQ(static_cast(res1), res2); - thrust::device_vector vec3(1000); + rmm::device_uvector d_vec3(1000, rmm::cuda_stream_default); - thrust::transform( - rmm::exec_policy(), std::cbegin(vec1), std::cend(vec1), std::begin(vec3), cast_to_int32_fn{}); + thrust::transform(rmm::exec_policy(), + std::cbegin(d_vec1), + std::cend(d_vec1), + std::begin(d_vec3), + cast_to_int32_fn{}); - thrust::host_vector vec3_host = vec3; + auto vec3 = cudf::detail::make_std_vector_sync(d_vec3); EXPECT_EQ(vec2, vec3); } From 944cb3abbc3dec5c9dd24734f144922b2d371000 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 11:15:48 +1000 Subject: [PATCH 08/27] Convert map tests --- cpp/tests/hash_map/map_test.cu | 14 +++++++------- cpp/tests/hash_map/multimap_test.cu | 12 ++++-------- 2 files changed, 11 insertions(+), 15 deletions(-) diff --git a/cpp/tests/hash_map/map_test.cu b/cpp/tests/hash_map/map_test.cu index 49cfda078b1..a25c35a63e0 100644 --- a/cpp/tests/hash_map/map_test.cu +++ b/cpp/tests/hash_map/map_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2020, NVIDIA CORPORATION. + * Copyright (c) 2018-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. @@ -14,18 +14,18 @@ * limitations under the License. */ +#include +#include + #include #include -#include #include -#include +#include #include -#include - #include #include #include @@ -53,13 +53,13 @@ struct InsertTest : public cudf::test::BaseFixture { // prevent overflow of small types const size_t input_size = std::min(static_cast(size), std::numeric_limits::max()); - pairs.resize(input_size); + pairs.resize(input_size, rmm::cuda_stream_default); map = std::move(map_type::create(compute_hash_table_size(size))); rmm::cuda_stream_default.synchronize(); } const cudf::size_type size{10000}; - rmm::device_vector pairs; + rmm::device_uvector pairs{static_cast(size), rmm::cuda_stream_default}; std::unique_ptr> map; }; diff --git a/cpp/tests/hash_map/multimap_test.cu b/cpp/tests/hash_map/multimap_test.cu index 7fd69e90efd..21135746227 100644 --- a/cpp/tests/hash_map/multimap_test.cu +++ b/cpp/tests/hash_map/multimap_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2020, NVIDIA CORPORATION. + * Copyright (c) 2018-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. @@ -14,21 +14,17 @@ * limitations under the License. */ +#include +#include + #include #include -#include - #include -#include - #include -#include -#include #include -#include // This is necessary to do a parametrized typed-test over multiple template // arguments From 884b29bb4d1dd9e1379af7bce6f5d3b7fb1b010d Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 12:02:49 +1000 Subject: [PATCH 09/27] Convert iterator tests to device_uvector --- cpp/tests/iterator/iterator_tests.cuh | 58 ++++++++------------ cpp/tests/iterator/optional_iterator_test.cu | 30 ++++++---- 2 files changed, 41 insertions(+), 47 deletions(-) diff --git a/cpp/tests/iterator/iterator_tests.cuh b/cpp/tests/iterator/iterator_tests.cuh index 34ce93c955e..c15e0542afe 100644 --- a/cpp/tests/iterator/iterator_tests.cuh +++ b/cpp/tests/iterator/iterator_tests.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * 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. @@ -13,27 +13,30 @@ * the License. */ -#include // include iterator header -#include //for meanvar - -#include -#include -#include -#include -#include - #include #include #include +#include // include iterator header +#include //for meanvar +#include + +#include +#include +#include + #include #include #include -// for reduction tests -#include #include +#include +#include +#include +#include +#include + // Base Typed test fixture for iterator test template @@ -43,7 +46,7 @@ struct IteratorTest : public cudf::test::BaseFixture { void iterator_test_cub(T_output expected, InputIterator d_in, int num_items) { T_output init = cudf::test::make_type_param_scalar(0); - thrust::device_vector dev_result(1, init); + rmm::device_uvector dev_result(1, rmm::cuda_stream_default); // Get temporary storage size size_t temp_storage_bytes = 0; @@ -72,18 +75,18 @@ struct IteratorTest : public cudf::test::BaseFixture { // iterator test case which uses thrust template - void iterator_test_thrust(thrust::host_vector& expected, + void iterator_test_thrust(thrust::host_vector const& expected, InputIterator d_in, int num_items) { InputIterator d_in_last = d_in + num_items; EXPECT_EQ(thrust::distance(d_in, d_in_last), num_items); - thrust::device_vector dev_expected(expected); + auto dev_expected = cudf::detail::make_device_uvector_sync(expected); // Can't use this because time_point make_pair bug in libcudacxx // bool result = thrust::equal(thrust::device, d_in, d_in_last, dev_expected.begin()); bool result = thrust::transform_reduce( - thrust::device, + rmm::exec_policy(), thrust::make_zip_iterator(thrust::make_tuple(d_in, dev_expected.begin())), thrust::make_zip_iterator(thrust::make_tuple(d_in_last, dev_expected.end())), [] __device__(auto it) { @@ -92,37 +95,22 @@ struct IteratorTest : public cudf::test::BaseFixture { }, true, thrust::logical_and()); -#ifndef NDEBUG - thrust::device_vector vec(expected.size(), false); - thrust::transform( - thrust::device, - thrust::make_zip_iterator(thrust::make_tuple(d_in, dev_expected.begin())), - thrust::make_zip_iterator(thrust::make_tuple(d_in_last, dev_expected.end())), - vec.begin(), - [] __device__(auto it) { return (thrust::get<0>(it)) == T_output(thrust::get<1>(it)); }); - thrust::copy(vec.begin(), vec.end(), std::ostream_iterator(std::cout, " ")); - std::cout << std::endl; -#endif EXPECT_TRUE(result) << "thrust test"; } template void evaluate(T_output expected, - thrust::device_vector& dev_result, + rmm::device_uvector const& dev_result, const char* msg = nullptr) { - thrust::host_vector hos_result(dev_result); + auto host_result = cudf::detail::make_host_vector_sync(dev_result); - EXPECT_EQ(expected, hos_result[0]) << msg; - std::cout << "Done: expected <" << msg - << "> = " - //<< hos_result[0] //TODO uncomment after time_point ostream operator<< - << std::endl; + EXPECT_EQ(expected, host_result[0]) << msg; } template - void values_equal_test(thrust::host_vector& expected, + void values_equal_test(thrust::host_vector const& expected, const cudf::column_device_view& col) { if (col.nullable()) { diff --git a/cpp/tests/iterator/optional_iterator_test.cu b/cpp/tests/iterator/optional_iterator_test.cu index 460f692ee14..020a63bc00a 100644 --- a/cpp/tests/iterator/optional_iterator_test.cu +++ b/cpp/tests/iterator/optional_iterator_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -177,22 +177,28 @@ TYPED_TEST(IteratorTest, null_optional_iterator) [](auto s, bool b) { return thrust::optional{s}; }); // GPU test for correct null mapping + this->iterator_test_thrust( + optional_values, + d_col->template optional_begin(cudf::contains_nulls::DYNAMIC{}, true), + host_values.size()); + this->iterator_test_thrust(optional_values, - d_col->optional_begin(cudf::contains_nulls::DYNAMIC{}, true), + d_col->template optional_begin(cudf::contains_nulls::YES{}), + host_values.size()); + this->iterator_test_thrust(optional_values, + d_col->template optional_begin(cudf::contains_nulls::YES{}), host_values.size()); + // GPU test for ignoring null mapping this->iterator_test_thrust( - optional_values, d_col->optional_begin(cudf::contains_nulls::YES{}), host_values.size()); - this->iterator_test_thrust( - optional_values, d_col->optional_begin(cudf::contains_nulls::YES{}), host_values.size()); + value_all_valid, + d_col->template optional_begin(cudf::contains_nulls::DYNAMIC{}, false), + host_values.size()); - // GPU test for ignoring null mapping this->iterator_test_thrust(value_all_valid, - d_col->optional_begin(cudf::contains_nulls::DYNAMIC{}, false), + d_col->template optional_begin(cudf::contains_nulls::NO{}), + host_values.size()); + this->iterator_test_thrust(value_all_valid, + d_col->template optional_begin(cudf::contains_nulls::NO{}), host_values.size()); - - this->iterator_test_thrust( - value_all_valid, d_col->optional_begin(cudf::contains_nulls::NO{}), host_values.size()); - this->iterator_test_thrust( - value_all_valid, d_col->optional_begin(cudf::contains_nulls::NO{}), host_values.size()); } From 531f591132d2222b0e9bc6e1ebb0af0735df2288 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 12:22:56 +1000 Subject: [PATCH 10/27] Convert set_nullmask_tests to uvector --- cpp/tests/bitmask/set_nullmask_tests.cu | 24 +++++++++++++++++------- 1 file changed, 17 insertions(+), 7 deletions(-) diff --git a/cpp/tests/bitmask/set_nullmask_tests.cu b/cpp/tests/bitmask/set_nullmask_tests.cu index ae4896827fd..d1560acea23 100644 --- a/cpp/tests/bitmask/set_nullmask_tests.cu +++ b/cpp/tests/bitmask/set_nullmask_tests.cu @@ -18,11 +18,15 @@ #include +#include #include +#include + +#include +#include +#include -#include #include -#include struct valid_bit_functor { cudf::bitmask_type const* _null_mask; @@ -41,12 +45,18 @@ std::ostream& operator<<(std::ostream& stream, thrust::host_vector const& struct SetBitmaskTest : public cudf::test::BaseFixture { void expect_bitmask_equal(cudf::bitmask_type const* bitmask, // Device Ptr cudf::size_type start_bit, - thrust::host_vector const& expect) + thrust::host_vector const& expect, + rmm::cuda_stream_view stream = rmm::cuda_stream_default) { - auto itb_dev = thrust::make_transform_iterator(thrust::counting_iterator{0}, - valid_bit_functor{bitmask}); - thrust::device_vector result(itb_dev + start_bit, itb_dev + start_bit + expect.size()); - thrust::host_vector host_result(result); + rmm::device_uvector result(expect.size(), stream); + auto counting_iter = thrust::counting_iterator{0}; + thrust::transform(rmm::exec_policy(stream), + counting_iter + start_bit, + counting_iter + start_bit + expect.size(), + result.begin(), + valid_bit_functor{bitmask}); + + auto host_result = cudf::detail::make_host_vector_sync(result, stream); EXPECT_THAT(host_result, testing::ElementsAreArray(expect)); } From 8410543a8bf84fc331b763168c53eaced1541fbd Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 13:34:05 +1000 Subject: [PATCH 11/27] Allow constructing a span from a `std::string` --- cpp/include/cudf/utilities/span.hpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/cpp/include/cudf/utilities/span.hpp b/cpp/include/cudf/utilities/span.hpp index 2f3577623a9..52ad0648e23 100644 --- a/cpp/include/cudf/utilities/span.hpp +++ b/cpp/include/cudf/utilities/span.hpp @@ -120,6 +120,11 @@ struct is_host_span_supported_container< // thrust::host_vector> : std::true_type { }; +template +struct is_host_span_supported_container< // + std::basic_string, Alloc>> : std::true_type { +}; + template struct host_span : public cudf::detail::span_base> { using base = cudf::detail::span_base>; From 3620171ba17eacbd2bbe9e3678eed462f43ebb64 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 13:34:33 +1000 Subject: [PATCH 12/27] Convert value_iterator_tests.cu --- cpp/tests/bitmask/set_nullmask_tests.cu | 2 +- cpp/tests/iterator/iterator_tests.cuh | 3 +- cpp/tests/iterator/value_iterator_test.cu | 56 +++++++++++------------ 3 files changed, 28 insertions(+), 33 deletions(-) diff --git a/cpp/tests/bitmask/set_nullmask_tests.cu b/cpp/tests/bitmask/set_nullmask_tests.cu index d1560acea23..235aec7ddf8 100644 --- a/cpp/tests/bitmask/set_nullmask_tests.cu +++ b/cpp/tests/bitmask/set_nullmask_tests.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. diff --git a/cpp/tests/iterator/iterator_tests.cuh b/cpp/tests/iterator/iterator_tests.cuh index c15e0542afe..68ff29455f9 100644 --- a/cpp/tests/iterator/iterator_tests.cuh +++ b/cpp/tests/iterator/iterator_tests.cuh @@ -90,8 +90,7 @@ struct IteratorTest : public cudf::test::BaseFixture { thrust::make_zip_iterator(thrust::make_tuple(d_in, dev_expected.begin())), thrust::make_zip_iterator(thrust::make_tuple(d_in_last, dev_expected.end())), [] __device__(auto it) { - return static_cast(thrust::get<0>(it)) == - T_output(thrust::get<1>(it)); + return static_cast(thrust::get<0>(it)) == T_output(thrust::get<1>(it)); }, true, thrust::logical_and()); diff --git a/cpp/tests/iterator/value_iterator_test.cu b/cpp/tests/iterator/value_iterator_test.cu index 542123ffd25..15a2ad0459b 100644 --- a/cpp/tests/iterator/value_iterator_test.cu +++ b/cpp/tests/iterator/value_iterator_test.cu @@ -14,22 +14,22 @@ */ #include +#include + auto strings_to_string_views(std::vector& input_strings) { auto all_valid = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return true; }); - std::vector chars; - std::vector offsets; - std::tie(chars, offsets) = cudf::test::detail::make_chars_and_offsets( + auto [chars, offsets] = cudf::test::detail::make_chars_and_offsets( input_strings.begin(), input_strings.end(), all_valid); - thrust::device_vector dev_chars(chars); - char* c_start = thrust::raw_pointer_cast(dev_chars.data()); + auto dev_chars = cudf::detail::make_device_uvector_sync(chars); + char* c_start = thrust::raw_pointer_cast(dev_chars.data()); // calculate the expected value by CPU. (but contains device pointers) - std::vector replaced_array(input_strings.size()); + thrust::host_vector replaced_array(input_strings.size()); std::transform(thrust::counting_iterator(0), thrust::counting_iterator(replaced_array.size()), replaced_array.begin(), - [c_start, offsets](auto i) { + [c_start, offsets = offsets](auto i) { return cudf::string_view(c_start + offsets[i], offsets[i + 1] - offsets[i]); }); return std::make_tuple(std::move(dev_chars), replaced_array); @@ -42,9 +42,11 @@ TYPED_TEST_CASE(IteratorTest, TestingTypes); // tests for non-null iterator (pointer of device array) TYPED_TEST(IteratorTest, non_null_iterator) { - using T = TypeParam; - auto host_array = cudf::test::make_type_param_vector({0, 6, 0, -14, 13, 64, -13, -20, 45}); - thrust::device_vector dev_array(host_array); + using T = TypeParam; + // wrap in host_vector to avoid vector issues + auto host_array = thrust::host_vector( + cudf::test::make_type_param_vector({0, 6, 0, -14, 13, 64, -13, -20, 45})); + auto dev_array = cudf::detail::make_device_uvector_sync(host_array); // calculate the expected value by CPU. thrust::host_vector replaced_array(host_array); @@ -225,11 +227,11 @@ struct StringIteratorTest : public IteratorTest { TEST_F(StringIteratorTest, string_view_null_iterator) { using T = cudf::string_view; - // T init = T{"", 0}; + std::string zero("zero"); // the char data has to be in GPU - thrust::device_vector initmsg(zero.begin(), zero.end()); - T init = T{initmsg.data().get(), int(initmsg.size())}; + auto initmsg = cudf::detail::make_device_uvector_sync(zero); + T init = T{initmsg.data(), int(initmsg.size())}; // data and valid arrays std::vector host_values( @@ -244,9 +246,7 @@ TEST_F(StringIteratorTest, string_view_null_iterator) replaced_strings.begin(), [zero](auto s, auto b) { return b ? s : zero; }); - thrust::device_vector dev_chars; - thrust::host_vector replaced_array(host_values.size()); - std::tie(dev_chars, replaced_array) = strings_to_string_views(replaced_strings); + auto [dev_chars, replaced_array] = strings_to_string_views(replaced_strings); // create a column with bool vector cudf::test::strings_column_wrapper w_col( @@ -262,19 +262,16 @@ TEST_F(StringIteratorTest, string_view_null_iterator) TEST_F(StringIteratorTest, string_view_no_null_iterator) { using T = cudf::string_view; - // T init = T{"", 0}; - std::string zero("zero"); + // the char data has to be in GPU - thrust::device_vector initmsg(zero.begin(), zero.end()); - T init = T{initmsg.data().get(), int(initmsg.size())}; + auto initmsg = cudf::detail::make_device_uvector_sync(std::string{"zero"}); + T init = T{initmsg.data(), int(initmsg.size())}; // data array std::vector host_values( {"one", "two", "three", "four", "five", "six", "eight", "nine"}); - thrust::device_vector dev_chars; - thrust::host_vector all_array(host_values.size()); - std::tie(dev_chars, all_array) = strings_to_string_views(host_values); + auto [dev_chars, all_array] = strings_to_string_views(host_values); // create a column with bool vector cudf::test::strings_column_wrapper w_col(host_values.begin(), host_values.end()); @@ -288,18 +285,17 @@ TEST_F(StringIteratorTest, string_view_no_null_iterator) TEST_F(StringIteratorTest, string_scalar_iterator) { using T = cudf::string_view; - // T init = T{"", 0}; - std::string zero("zero"); + + std::string zero{"zero"}; + // the char data has to be in GPU - thrust::device_vector initmsg(zero.begin(), zero.end()); - T init = T{initmsg.data().get(), int(initmsg.size())}; + auto initmsg = cudf::detail::make_device_uvector_sync(zero); + T init = T{initmsg.data(), int(initmsg.size())}; // data array std::vector host_values(100, zero); - thrust::device_vector dev_chars; - thrust::host_vector all_array(host_values.size()); - std::tie(dev_chars, all_array) = strings_to_string_views(host_values); + auto [dev_chars, all_array] = strings_to_string_views(host_values); // calculate the expected value by CPU. thrust::host_vector> value_and_validity(host_values.size()); From 969b69077ddf7c5a8cf95393ca51e7949724637a Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 13:36:19 +1000 Subject: [PATCH 13/27] Unused header --- cpp/tests/reductions/reduction_tests.cpp | 1 - cpp/tests/replace/replace_tests.cpp | 3 +-- 2 files changed, 1 insertion(+), 3 deletions(-) diff --git a/cpp/tests/reductions/reduction_tests.cpp b/cpp/tests/reductions/reduction_tests.cpp index fce9e77dc55..9c66ccd4623 100644 --- a/cpp/tests/reductions/reduction_tests.cpp +++ b/cpp/tests/reductions/reduction_tests.cpp @@ -27,7 +27,6 @@ #include #include -#include #include #include diff --git a/cpp/tests/replace/replace_tests.cpp b/cpp/tests/replace/replace_tests.cpp index 32ddf11f16d..fede13d9aa5 100644 --- a/cpp/tests/replace/replace_tests.cpp +++ b/cpp/tests/replace/replace_tests.cpp @@ -25,11 +25,10 @@ #include #include +#include #include #include -#include "cudf/fixed_point/fixed_point.hpp" -#include #include #include From c0c46d4c147af111fd639470f5a8074fa4ec6426 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 13:41:48 +1000 Subject: [PATCH 14/27] Convert scalar test --- cpp/tests/scalar/scalar_device_view_test.cu | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/cpp/tests/scalar/scalar_device_view_test.cu b/cpp/tests/scalar/scalar_device_view_test.cu index c501071ccbe..d0b6b0db44a 100644 --- a/cpp/tests/scalar/scalar_device_view_test.cu +++ b/cpp/tests/scalar/scalar_device_view_test.cu @@ -14,6 +14,7 @@ * limitations under the License. */ +#include #include #include #include @@ -24,8 +25,6 @@ #include #include -#include - #include #include @@ -126,10 +125,9 @@ TEST_F(StringScalarDeviceViewTest, Value) auto scalar_device_view = cudf::get_scalar_device_view(s); rmm::device_scalar result; - rmm::device_vector value_v(value.begin(), value.end()); + auto value_v = cudf::detail::make_device_uvector_sync(value); - test_string_value<<<1, 1>>>( - scalar_device_view, value_v.data().get(), value.size(), result.data()); + test_string_value<<<1, 1>>>(scalar_device_view, value_v.data(), value.size(), result.data()); CHECK_CUDA(0); EXPECT_TRUE(result.value()); From 96b37a41b72c707b3b5cd473dfd9eea6ce800038 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 13:42:35 +1000 Subject: [PATCH 15/27] Convert strings array_tests.cu --- cpp/tests/strings/array_tests.cu | 23 ++++++++++++----------- 1 file changed, 12 insertions(+), 11 deletions(-) diff --git a/cpp/tests/strings/array_tests.cu b/cpp/tests/strings/array_tests.cu index 2d1ae1a862d..a4d8ecb2bec 100644 --- a/cpp/tests/strings/array_tests.cu +++ b/cpp/tests/strings/array_tests.cu @@ -14,10 +14,16 @@ * limitations under the License. */ +#include +#include +#include +#include + #include #include #include #include +#include #include #include #include @@ -27,10 +33,7 @@ #include #include -#include -#include -#include -#include +#include #include @@ -192,9 +195,8 @@ TEST_F(StringsColumnTest, Scatter) thrust::make_transform_iterator(h_strings2.begin(), [](auto str) { return str != nullptr; })); auto source = cudf::strings_column_view(strings2); - rmm::device_vector scatter_map; - scatter_map.push_back(4); - scatter_map.push_back(1); + std::vector h_scatter_map({4, 1}); + auto scatter_map = cudf::detail::make_device_uvector_sync(h_scatter_map); auto source_column = cudf::column_device_view::create(source.parent()); auto begin = @@ -220,9 +222,8 @@ TEST_F(StringsColumnTest, ScatterScalar) thrust::make_transform_iterator(h_strings1.begin(), [](auto str) { return str != nullptr; })); auto target = cudf::strings_column_view(strings1); - rmm::device_vector scatter_map; - scatter_map.push_back(0); - scatter_map.push_back(5); + std::vector h_scatter_map({0, 5}); + auto scatter_map = cudf::detail::make_device_uvector_sync(h_scatter_map); cudf::string_scalar scalar("__"); auto begin = thrust::make_constant_iterator(cudf::string_view(scalar.data(), scalar.size())); @@ -246,7 +247,7 @@ TEST_F(StringsColumnTest, ScatterZeroSizeStringsColumn) cudf::column_view values(cudf::data_type{cudf::type_id::STRING}, 0, nullptr, nullptr, 0); auto target = cudf::strings_column_view(values); - rmm::device_vector scatter_map; + rmm::device_uvector scatter_map(0, rmm::cuda_stream_default); cudf::string_scalar scalar(""); auto begin = thrust::make_constant_iterator(cudf::string_view(scalar.data(), scalar.size())); From fd1e5081e426dbcf07f9acc19ae451ef9e2919de Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 14:20:35 +1000 Subject: [PATCH 16/27] Convert strings_tests --- cpp/tests/strings/factories_test.cu | 30 +++++++++++++++-------------- cpp/tests/strings/hash_string.cu | 18 +++++++++++------ cpp/tests/strings/integers_tests.cu | 18 +++++++++++------ 3 files changed, 40 insertions(+), 26 deletions(-) diff --git a/cpp/tests/strings/factories_test.cu b/cpp/tests/strings/factories_test.cu index be592478b13..854194d13c8 100644 --- a/cpp/tests/strings/factories_test.cu +++ b/cpp/tests/strings/factories_test.cu @@ -28,6 +28,7 @@ #include #include +#include #include #include @@ -55,7 +56,7 @@ TEST_F(StringsFactoriesTest, CreateColumnFromPair) memsize += *itr ? (cudf::size_type)strlen(*itr) : 0; cudf::size_type count = (cudf::size_type)h_test_strings.size(); thrust::host_vector h_buffer(memsize); - thrust::device_vector d_buffer(memsize); + rmm::device_uvector d_buffer(memsize, rmm::cuda_stream_default); thrust::host_vector> strings(count); thrust::host_vector h_offsets(count + 1); cudf::size_type offset = 0; @@ -69,14 +70,13 @@ TEST_F(StringsFactoriesTest, CreateColumnFromPair) } else { cudf::size_type length = (cudf::size_type)strlen(str); memcpy(h_buffer.data() + offset, str, length); - strings[idx] = - thrust::pair{d_buffer.data().get() + offset, length}; + strings[idx] = thrust::pair{d_buffer.data() + offset, length}; offset += length; } h_offsets[idx + 1] = offset; } - rmm::device_vector> d_strings(strings); - CUDA_TRY(cudaMemcpy(d_buffer.data().get(), h_buffer.data(), memsize, cudaMemcpyHostToDevice)); + auto d_strings = cudf::detail::make_device_uvector_sync(strings); + CUDA_TRY(cudaMemcpy(d_buffer.data(), h_buffer.data(), memsize, cudaMemcpyHostToDevice)); auto column = cudf::make_strings_column(d_strings); EXPECT_EQ(column->type(), cudf::data_type{cudf::type_id::STRING}); EXPECT_EQ(column->null_count(), nulls); @@ -133,11 +133,12 @@ TEST_F(StringsFactoriesTest, CreateColumnFromOffsets) null_count++; h_offsets[idx + 1] = offset; } + std::vector h_nulls{h_null_mask}; - rmm::device_vector d_buffer(h_buffer); - rmm::device_vector d_offsets(h_offsets); - rmm::device_vector d_nulls(h_nulls); - auto column = cudf::make_strings_column(d_buffer, d_offsets, d_nulls, null_count); + auto d_buffer = cudf::detail::make_device_uvector_sync(h_buffer); + auto d_offsets = cudf::detail::make_device_uvector_sync(h_offsets); + auto d_nulls = cudf::detail::make_device_uvector_sync(h_nulls); + auto column = cudf::make_strings_column(d_buffer, d_offsets, d_nulls, null_count); EXPECT_EQ(column->type(), cudf::data_type{cudf::type_id::STRING}); EXPECT_EQ(column->null_count(), null_count); EXPECT_EQ(2, column->num_children()); @@ -169,14 +170,15 @@ TEST_F(StringsFactoriesTest, CreateScalar) TEST_F(StringsFactoriesTest, EmptyStringsColumn) { - rmm::device_vector d_chars; - rmm::device_vector d_offsets(1, 0); - rmm::device_vector d_nulls; + rmm::device_uvector d_chars{0, rmm::cuda_stream_default}; + auto d_offsets = cudf::detail::make_zeroed_device_uvector_sync(1); + rmm::device_uvector d_nulls{0, rmm::cuda_stream_default}; auto results = cudf::make_strings_column(d_chars, d_offsets, d_nulls, 0); cudf::test::expect_strings_empty(results->view()); - rmm::device_vector> d_strings; + rmm::device_uvector> d_strings{ + 0, rmm::cuda_stream_default}; results = cudf::make_strings_column(d_strings); cudf::test::expect_strings_empty(results->view()); } @@ -224,7 +226,7 @@ TEST_F(StringsFactoriesTest, StringPairWithNullsAndEmpty) {0, 1, 1, 1, 1, 0, 1, 1, 1, 0, 1}); auto d_column = cudf::column_device_view::create(data); - rmm::device_vector pairs(d_column->size()); + rmm::device_uvector pairs(d_column->size(), rmm::cuda_stream_default); thrust::transform(thrust::device, d_column->pair_begin(), d_column->pair_end(), diff --git a/cpp/tests/strings/hash_string.cu b/cpp/tests/strings/hash_string.cu index 629c02a989e..023d648cfdf 100644 --- a/cpp/tests/strings/hash_string.cu +++ b/cpp/tests/strings/hash_string.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019, NVIDIA CORPORATION. + * 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. @@ -14,15 +14,21 @@ * limitations under the License. */ +#include "./utilities.h" +#include "rmm/exec_policy.hpp" + #include #include #include +#include #include #include #include -#include "./utilities.h" -#include +#include +#include +#include + #include #include @@ -58,8 +64,8 @@ TEST_F(StringsHashTest, HashTest) auto strings_column = cudf::column_device_view::create(strings_view.parent()); auto d_view = *strings_column; - thrust::device_vector d_values(strings_view.size()); - thrust::transform(thrust::device, + rmm::device_uvector d_values(strings_view.size(), rmm::cuda_stream_default); + thrust::transform(rmm::exec_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(strings_view.size()), d_values.begin(), @@ -67,6 +73,6 @@ TEST_F(StringsHashTest, HashTest) uint32_t h_expected[] = { 2739798893, 2739798893, 3506676360, 1891213601, 3778137224, 0, 0, 1551088011}; - thrust::host_vector h_values(d_values); + auto h_values = cudf::detail::make_host_vector_sync(d_values); for (uint32_t idx = 0; idx < h_values.size(); ++idx) EXPECT_EQ(h_values[idx], h_expected[idx]); } diff --git a/cpp/tests/strings/integers_tests.cu b/cpp/tests/strings/integers_tests.cu index f15116ae4c2..351807dbdf2 100644 --- a/cpp/tests/strings/integers_tests.cu +++ b/cpp/tests/strings/integers_tests.cu @@ -18,11 +18,15 @@ #include #include +#include #include #include #include #include +#include +#include + #include #include @@ -287,16 +291,17 @@ TYPED_TEST_CASE(StringsIntegerConvertTest, cudf::test::IntegralTypesNotBool); TYPED_TEST(StringsIntegerConvertTest, FromToInteger) { - thrust::device_vector d_integers(255); + thrust::host_vector h_integers(255); thrust::sequence( - thrust::device, d_integers.begin(), d_integers.end(), -(TypeParam)(d_integers.size() / 2)); - d_integers.push_back(std::numeric_limits::min()); - d_integers.push_back(std::numeric_limits::max()); + thrust::seq, h_integers.begin(), h_integers.end(), -(TypeParam)(h_integers.size() / 2)); + h_integers.push_back(std::numeric_limits::min()); + h_integers.push_back(std::numeric_limits::max()); + auto d_integers = cudf::detail::make_device_uvector_sync(h_integers); auto integers = cudf::make_numeric_column(cudf::data_type{cudf::type_to_id()}, (cudf::size_type)d_integers.size()); auto integers_view = integers->mutable_view(); CUDA_TRY(cudaMemcpy(integers_view.data(), - d_integers.data().get(), + d_integers.data(), d_integers.size() * sizeof(TypeParam), cudaMemcpyDeviceToDevice)); integers_view.set_null_count(0); @@ -304,7 +309,8 @@ TYPED_TEST(StringsIntegerConvertTest, FromToInteger) // convert to strings auto results_strings = cudf::strings::from_integers(integers->view()); - thrust::host_vector h_integers(d_integers); + // copy back to host + h_integers = cudf::detail::make_host_vector_sync(d_integers); std::vector h_strings; for (auto itr = h_integers.begin(); itr != h_integers.end(); ++itr) h_strings.push_back(std::to_string(*itr)); From af34140e38d380ff0b9cc1f68b9ef0f20332d737 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 14:22:15 +1000 Subject: [PATCH 17/27] convert table_view_tests --- cpp/tests/table/table_view_tests.cu | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/cpp/tests/table/table_view_tests.cu b/cpp/tests/table/table_view_tests.cu index d700892de78..1fb4b88c79e 100644 --- a/cpp/tests/table/table_view_tests.cu +++ b/cpp/tests/table/table_view_tests.cu @@ -20,6 +20,7 @@ #include #include +#include #include #include #include @@ -42,10 +43,10 @@ void row_comparison(cudf::table_view input1, auto device_table_1 = cudf::table_device_view::create(input1, stream); auto device_table_2 = cudf::table_device_view::create(input2, stream); - rmm::device_vector d_column_order(column_order); + auto d_column_order = cudf::detail::make_device_uvector_sync(column_order); auto comparator = cudf::row_lexicographic_comparator( - *device_table_1, *device_table_2, d_column_order.data().get()); + *device_table_1, *device_table_2, d_column_order.data()); thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), From 77841dfca90f81c111220e30e0ff7a62ad6642aa Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 6 May 2021 14:24:33 +1000 Subject: [PATCH 18/27] integers_tests.cu --> .cpp --- cpp/tests/CMakeLists.txt | 2 +- cpp/tests/strings/{integers_tests.cu => integers_tests.cpp} | 3 +-- 2 files changed, 2 insertions(+), 3 deletions(-) rename cpp/tests/strings/{integers_tests.cu => integers_tests.cpp} (99%) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index c8dd148175e..968e41ecebd 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -335,7 +335,7 @@ ConfigureTest(STRINGS_TEST strings/fixed_point_tests.cpp strings/floats_tests.cpp strings/hash_string.cu - strings/integers_tests.cu + strings/integers_tests.cpp strings/ipv4_tests.cpp strings/json_tests.cpp strings/pad_tests.cpp diff --git a/cpp/tests/strings/integers_tests.cu b/cpp/tests/strings/integers_tests.cpp similarity index 99% rename from cpp/tests/strings/integers_tests.cu rename to cpp/tests/strings/integers_tests.cpp index 351807dbdf2..d5f17954c50 100644 --- a/cpp/tests/strings/integers_tests.cu +++ b/cpp/tests/strings/integers_tests.cpp @@ -292,8 +292,7 @@ TYPED_TEST_CASE(StringsIntegerConvertTest, cudf::test::IntegralTypesNotBool); TYPED_TEST(StringsIntegerConvertTest, FromToInteger) { thrust::host_vector h_integers(255); - thrust::sequence( - thrust::seq, h_integers.begin(), h_integers.end(), -(TypeParam)(h_integers.size() / 2)); + std::iota(h_integers.begin(), h_integers.end(), -(TypeParam)(h_integers.size() / 2)); h_integers.push_back(std::numeric_limits::min()); h_integers.push_back(std::numeric_limits::max()); auto d_integers = cudf::detail::make_device_uvector_sync(h_integers); From 6acd40485885d74655082e1801289b61aaaf8bfb Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 11 May 2021 10:03:27 +1000 Subject: [PATCH 19/27] Convert timestamps_test --- cpp/tests/wrappers/timestamps_test.cu | 17 ++++++++++++----- 1 file changed, 12 insertions(+), 5 deletions(-) diff --git a/cpp/tests/wrappers/timestamps_test.cu b/cpp/tests/wrappers/timestamps_test.cu index d1c0ad5d840..64d9ad6fc3f 100644 --- a/cpp/tests/wrappers/timestamps_test.cu +++ b/cpp/tests/wrappers/timestamps_test.cu @@ -29,6 +29,8 @@ #include #include +#include +#include template struct ChronoColumnTest : public cudf::test::BaseFixture { @@ -88,9 +90,10 @@ TYPED_TEST(ChronoColumnTest, ChronoDurationsMatchPrimitiveRepresentation) auto primitive_col = fixed_width_column_wrapper(chrono_col_data.begin(), chrono_col_data.end()); - thrust::device_vector indices(this->size()); - thrust::sequence(indices.begin(), indices.end()); - EXPECT_TRUE(thrust::all_of(indices.begin(), + rmm::device_uvector indices(this->size(), rmm::cuda_stream_default); + thrust::sequence(rmm::exec_policy(), indices.begin(), indices.end()); + EXPECT_TRUE(thrust::all_of(rmm::exec_policy(), + indices.begin(), indices.end(), compare_chrono_elements_to_primitive_representation{ *cudf::column_device_view::create(primitive_col), @@ -141,10 +144,11 @@ TYPED_TEST(ChronoColumnTest, ChronosCanBeComparedInDeviceCode) auto chrono_rhs_col = generate_timestamps(this->size(), time_point_ms(start_rhs), time_point_ms(stop_rhs_)); - thrust::device_vector indices(this->size()); - thrust::sequence(indices.begin(), indices.end()); + rmm::device_uvector indices(this->size(), rmm::cuda_stream_default); + thrust::sequence(rmm::exec_policy(), indices.begin(), indices.end()); EXPECT_TRUE(thrust::all_of( + rmm::exec_policy(), indices.begin(), indices.end(), compare_chrono_elements{cudf::binary_operator::LESS, @@ -152,6 +156,7 @@ TYPED_TEST(ChronoColumnTest, ChronosCanBeComparedInDeviceCode) *cudf::column_device_view::create(chrono_rhs_col)})); EXPECT_TRUE(thrust::all_of( + rmm::exec_policy(), indices.begin(), indices.end(), compare_chrono_elements{cudf::binary_operator::GREATER, @@ -159,6 +164,7 @@ TYPED_TEST(ChronoColumnTest, ChronosCanBeComparedInDeviceCode) *cudf::column_device_view::create(chrono_lhs_col)})); EXPECT_TRUE(thrust::all_of( + rmm::exec_policy(), indices.begin(), indices.end(), compare_chrono_elements{cudf::binary_operator::LESS_EQUAL, @@ -166,6 +172,7 @@ TYPED_TEST(ChronoColumnTest, ChronosCanBeComparedInDeviceCode) *cudf::column_device_view::create(chrono_lhs_col)})); EXPECT_TRUE(thrust::all_of( + rmm::exec_policy(), indices.begin(), indices.end(), compare_chrono_elements{cudf::binary_operator::GREATER_EQUAL, From 05e4ec276acb6c9300eb30085bbf3801a5665948 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 11 May 2021 11:25:28 +1000 Subject: [PATCH 20/27] Change `make_type_param_vector` to return a thrust::host_vector rather than std::vector --- cpp/include/cudf_test/type_lists.hpp | 11 ++- cpp/tests/replace/clamp_test.cpp | 4 +- cpp/tests/replace/replace_nulls_tests.cpp | 67 ++++++------- cpp/tests/replace/replace_tests.cpp | 95 ++++++++++--------- cpp/tests/rolling/rolling_test.cpp | 9 +- .../drop_duplicates_tests.cpp | 8 +- 6 files changed, 97 insertions(+), 97 deletions(-) diff --git a/cpp/include/cudf_test/type_lists.hpp b/cpp/include/cudf_test/type_lists.hpp index 71c2b74b37b..a344173144d 100644 --- a/cpp/include/cudf_test/type_lists.hpp +++ b/cpp/include/cudf_test/type_lists.hpp @@ -25,6 +25,8 @@ #include #include +#include + #include #include @@ -79,10 +81,10 @@ constexpr auto types_to_ids() template typename std::enable_if() && !cudf::is_timestamp_t::value, - std::vector>::type + thrust::host_vector>::type make_type_param_vector(std::initializer_list const& init_list) { - std::vector vec(init_list.size()); + thrust::host_vector vec(init_list.size()); std::transform(std::cbegin(init_list), std::cend(init_list), std::begin(vec), [](auto const& e) { if (std::is_unsigned::value) return static_cast(std::abs(e)); @@ -93,10 +95,11 @@ make_type_param_vector(std::initializer_list const& init_list) } template -typename std::enable_if::value, std::vector>::type +typename std::enable_if::value, + thrust::host_vector>::type make_type_param_vector(std::initializer_list const& init_list) { - std::vector vec(init_list.size()); + thrust::host_vector vec(init_list.size()); std::transform(std::cbegin(init_list), std::cend(init_list), std::begin(vec), [](auto const& e) { return TypeParam{typename TypeParam::duration{e}}; }); diff --git a/cpp/tests/replace/clamp_test.cpp b/cpp/tests/replace/clamp_test.cpp index 47599035709..499745c7dc4 100644 --- a/cpp/tests/replace/clamp_test.cpp +++ b/cpp/tests/replace/clamp_test.cpp @@ -135,8 +135,8 @@ TEST_F(ClampEmptyCaseTest, EmptyInput) template struct ClampTestNumeric : public cudf::test::BaseFixture { - std::unique_ptr run_clamp(std::vector input, - std::vector input_validity, + std::unique_ptr run_clamp(cudf::host_span input, + cudf::host_span input_validity, T lo, bool lo_validity, T hi, diff --git a/cpp/tests/replace/replace_nulls_tests.cpp b/cpp/tests/replace/replace_nulls_tests.cpp index f6937c29d04..cd19b0a70f3 100644 --- a/cpp/tests/replace/replace_nulls_tests.cpp +++ b/cpp/tests/replace/replace_nulls_tests.cpp @@ -267,10 +267,11 @@ void ReplaceNullsScalar(cudf::test::fixed_width_column_wrapper input, TYPED_TEST(ReplaceNullsTest, ReplaceColumn) { - std::vector inputColumn = + auto const inputColumn = cudf::test::make_type_param_vector({0, 1, 2, 3, 4, 5, 6, 7, 8, 9}); - std::vector inputValid{0, 0, 0, 0, 0, 1, 1, 1, 1, 1}; - std::vector replacementColumn = + auto const inputValid = + cudf::test::make_type_param_vector({0, 0, 0, 0, 0, 1, 1, 1, 1, 1}); + auto const replacementColumn = cudf::test::make_type_param_vector({0, 1, 2, 3, 4, 5, 6, 7, 8, 9}); ReplaceNullsColumn(cudf::test::fixed_width_column_wrapper( @@ -290,10 +291,11 @@ TYPED_TEST(ReplaceNullsTest, ReplaceColumn_Empty) TYPED_TEST(ReplaceNullsTest, ReplaceScalar) { - std::vector inputColumn = + auto const inputColumn = cudf::test::make_type_param_vector({0, 1, 2, 3, 4, 5, 6, 7, 8, 9}); - std::vector inputValid{0, 0, 0, 0, 0, 1, 1, 1, 1, 1}; - std::vector expectedColumn = + auto const inputValid = + cudf::test::make_type_param_vector({0, 0, 0, 0, 0, 1, 1, 1, 1, 1}); + auto const expectedColumn = cudf::test::make_type_param_vector({1, 1, 1, 1, 1, 5, 6, 7, 8, 9}); cudf::numeric_scalar replacement(1); @@ -308,13 +310,16 @@ TYPED_TEST(ReplaceNullsTest, ReplacementHasNulls) { using T = TypeParam; - std::vector input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); - std::vector replace_column = cudf::test::make_type_param_vector({4, 5, 6, 7, 8, 9, 0, 1}); - std::vector result_column = cudf::test::make_type_param_vector({4, 5, 6, 3, 1, 2, 8, 4}); + auto const input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); + auto const replace_column = cudf::test::make_type_param_vector({4, 5, 6, 7, 8, 9, 0, 1}); + auto const result_column = cudf::test::make_type_param_vector({4, 5, 6, 3, 1, 2, 8, 4}); - std::vector input_valid{0, 0, 1, 1, 1, 1, 1, 1}; - std::vector replace_valid{1, 0, 1, 1, 1, 1, 1, 1}; - std::vector result_valid{1, 0, 1, 1, 1, 1, 1, 1}; + auto const input_valid = + cudf::test::make_type_param_vector({0, 0, 1, 1, 1, 1, 1, 1}); + auto const replace_valid = + cudf::test::make_type_param_vector({1, 0, 1, 1, 1, 1, 1, 1}); + auto const result_valid = + cudf::test::make_type_param_vector({1, 0, 1, 1, 1, 1, 1, 1}); ReplaceNullsColumn(cudf::test::fixed_width_column_wrapper( input_column.begin(), input_column.end(), input_valid.begin()), @@ -375,11 +380,9 @@ void TestReplaceNullsWithPolicy(cudf::test::fixed_width_column_wrapper input, TYPED_TEST(ReplaceNullsPolicyTest, PrecedingFill) { - std::vector col = - cudf::test::make_type_param_vector({42, 2, 1, -10, 20, -30}); - std::vector mask = - cudf::test::make_type_param_vector({1, 0, 0, 1, 0, 1}); - std::vector expect_col = + auto const col = cudf::test::make_type_param_vector({42, 2, 1, -10, 20, -30}); + auto const mask = cudf::test::make_type_param_vector({1, 0, 0, 1, 0, 1}); + auto const expect_col = cudf::test::make_type_param_vector({42, 42, 42, -10, -10, -30}); TestReplaceNullsWithPolicy( @@ -391,11 +394,9 @@ TYPED_TEST(ReplaceNullsPolicyTest, PrecedingFill) TYPED_TEST(ReplaceNullsPolicyTest, FollowingFill) { - std::vector col = - cudf::test::make_type_param_vector({42, 2, 1, -10, 20, -30}); - std::vector mask = - cudf::test::make_type_param_vector({1, 0, 0, 1, 0, 1}); - std::vector expect_col = + auto const col = cudf::test::make_type_param_vector({42, 2, 1, -10, 20, -30}); + auto const mask = cudf::test::make_type_param_vector({1, 0, 0, 1, 0, 1}); + auto const expect_col = cudf::test::make_type_param_vector({42, -10, -10, -10, -30, -30}); TestReplaceNullsWithPolicy( @@ -407,13 +408,10 @@ TYPED_TEST(ReplaceNullsPolicyTest, FollowingFill) TYPED_TEST(ReplaceNullsPolicyTest, PrecedingFillLeadingNulls) { - std::vector col = cudf::test::make_type_param_vector({1, 2, 3, 4, 5}); - std::vector mask = - cudf::test::make_type_param_vector({0, 0, 1, 0, 1}); - std::vector expect_col = - cudf::test::make_type_param_vector({1, 2, 3, 3, 5}); - std::vector expect_mask = - cudf::test::make_type_param_vector({0, 0, 1, 1, 1}); + auto const col = cudf::test::make_type_param_vector({1, 2, 3, 4, 5}); + auto const mask = cudf::test::make_type_param_vector({0, 0, 1, 0, 1}); + auto const expect_col = cudf::test::make_type_param_vector({1, 2, 3, 3, 5}); + auto const expect_mask = cudf::test::make_type_param_vector({0, 0, 1, 1, 1}); TestReplaceNullsWithPolicy( cudf::test::fixed_width_column_wrapper(col.begin(), col.end(), mask.begin()), @@ -424,13 +422,10 @@ TYPED_TEST(ReplaceNullsPolicyTest, PrecedingFillLeadingNulls) TYPED_TEST(ReplaceNullsPolicyTest, FollowingFillTrailingNulls) { - std::vector col = cudf::test::make_type_param_vector({1, 2, 3, 4, 5}); - std::vector mask = - cudf::test::make_type_param_vector({1, 0, 1, 0, 0}); - std::vector expect_col = - cudf::test::make_type_param_vector({1, 3, 3, 4, 5}); - std::vector expect_mask = - cudf::test::make_type_param_vector({1, 1, 1, 0, 0}); + auto const col = cudf::test::make_type_param_vector({1, 2, 3, 4, 5}); + auto const mask = cudf::test::make_type_param_vector({1, 0, 1, 0, 0}); + auto const expect_col = cudf::test::make_type_param_vector({1, 3, 3, 4, 5}); + auto const expect_mask = cudf::test::make_type_param_vector({1, 1, 1, 0, 0}); TestReplaceNullsWithPolicy( cudf::test::fixed_width_column_wrapper(col.begin(), col.end(), mask.begin()), diff --git a/cpp/tests/replace/replace_tests.cpp b/cpp/tests/replace/replace_tests.cpp index fede13d9aa5..58ef08f6052 100644 --- a/cpp/tests/replace/replace_tests.cpp +++ b/cpp/tests/replace/replace_tests.cpp @@ -35,6 +35,7 @@ #include #include #include +#include "cudf/types.hpp" struct ReplaceErrorTest : public cudf::test::BaseFixture { }; @@ -314,13 +315,12 @@ struct ReplaceTest : cudf::test::BaseFixture { * @param print Optionally print the set of columns for debug */ template -void test_replace( - std::vector const& input_column, - std::vector const& values_to_replace_column, - std::vector const& replacement_values_column, - std::vector const& input_column_valid = std::vector{}, - std::vector const& replacement_values_valid = std::vector{}, - bool print = false) +void test_replace(cudf::host_span input_column, + cudf::host_span values_to_replace_column, + cudf::host_span replacement_values_column, + cudf::host_span input_column_valid = {}, + cudf::host_span replacement_values_valid = {}, + bool print = false) { cudf::test::fixed_width_column_wrapper _input_column(input_column.begin(), input_column.end()); if (input_column_valid.size() > 0) { @@ -345,9 +345,10 @@ void test_replace( _input_column, _values_to_replace_column, _replacement_values_column)); /* computing the expected result */ - std::vector reference_result(input_column); - std::vector isReplaced(reference_result.size(), false); - std::vector expected_valid(input_column_valid); + thrust::host_vector reference_result(input_column.begin(), input_column.end()); + thrust::host_vector isReplaced(reference_result.size(), false); + thrust::host_vector expected_valid(input_column_valid.begin(), + input_column_valid.end()); if (replacement_values_valid.size() > 0 && 0 == input_column_valid.size()) { expected_valid.assign(input_column.size(), true); } @@ -395,10 +396,10 @@ TYPED_TEST_CASE(ReplaceTest, Types); // Simple test, replacing all even replacement_values_column TYPED_TEST(ReplaceTest, ReplaceEvenPosition) { - using T = TypeParam; - std::vector input_column = cudf::test::make_type_param_vector({1, 2, 3, 4, 5, 6, 7, 8}); - std::vector values_to_replace_column = cudf::test::make_type_param_vector({2, 6, 4, 8}); - std::vector replacement_values_column = cudf::test::make_type_param_vector({0, 4, 2, 6}); + using T = TypeParam; + auto const input_column = cudf::test::make_type_param_vector({1, 2, 3, 4, 5, 6, 7, 8}); + auto const values_to_replace_column = cudf::test::make_type_param_vector({2, 6, 4, 8}); + auto const replacement_values_column = cudf::test::make_type_param_vector({0, 4, 2, 6}); test_replace(input_column, values_to_replace_column, replacement_values_column); } @@ -406,10 +407,10 @@ TYPED_TEST(ReplaceTest, ReplaceEvenPosition) // Similar test as ReplaceEvenPosition, but with unordered data TYPED_TEST(ReplaceTest, Unordered) { - using T = TypeParam; - std::vector input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); - std::vector values_to_replace_column = cudf::test::make_type_param_vector({2, 6, 4, 8}); - std::vector replacement_values_column = cudf::test::make_type_param_vector({0, 4, 2, 6}); + using T = TypeParam; + auto const input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); + auto const values_to_replace_column = cudf::test::make_type_param_vector({2, 6, 4, 8}); + auto const replacement_values_column = cudf::test::make_type_param_vector({0, 4, 2, 6}); test_replace(input_column, values_to_replace_column, replacement_values_column); } @@ -417,10 +418,10 @@ TYPED_TEST(ReplaceTest, Unordered) // Testing with Nothing To Replace TYPED_TEST(ReplaceTest, NothingToReplace) { - using T = TypeParam; - std::vector input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); - std::vector values_to_replace_column = cudf::test::make_type_param_vector({10, 11, 12}); - std::vector replacement_values_column = cudf::test::make_type_param_vector({15, 16, 17}); + using T = TypeParam; + auto const input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); + auto const values_to_replace_column = cudf::test::make_type_param_vector({10, 11, 12}); + auto const replacement_values_column = cudf::test::make_type_param_vector({15, 16, 17}); test_replace(input_column, values_to_replace_column, replacement_values_column); } @@ -429,9 +430,9 @@ TYPED_TEST(ReplaceTest, NothingToReplace) TYPED_TEST(ReplaceTest, EmptyData) { using T = TypeParam; - std::vector input_column{{}}; - std::vector values_to_replace_column = cudf::test::make_type_param_vector({10, 11, 12}); - std::vector replacement_values_column = cudf::test::make_type_param_vector({15, 16, 17}); + thrust::host_vector input_column{{}}; + auto const values_to_replace_column = cudf::test::make_type_param_vector({10, 11, 12}); + auto const replacement_values_column = cudf::test::make_type_param_vector({15, 16, 17}); test_replace(input_column, values_to_replace_column, replacement_values_column); } @@ -439,10 +440,10 @@ TYPED_TEST(ReplaceTest, EmptyData) // Testing with empty Replace TYPED_TEST(ReplaceTest, EmptyReplace) { - using T = TypeParam; - std::vector input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); - std::vector values_to_replace_column{}; - std::vector replacement_values_column{}; + using T = TypeParam; + auto const input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); + thrust::host_vector values_to_replace_column{}; + thrust::host_vector replacement_values_column{}; test_replace(input_column, values_to_replace_column, replacement_values_column); } @@ -450,11 +451,12 @@ TYPED_TEST(ReplaceTest, EmptyReplace) // Testing with input column containing nulls TYPED_TEST(ReplaceTest, NullsInData) { - using T = TypeParam; - std::vector input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); - std::vector input_column_valid{1, 1, 1, 0, 0, 1, 1, 1}; - std::vector values_to_replace_column = cudf::test::make_type_param_vector({2, 6, 4, 8}); - std::vector replacement_values_column = cudf::test::make_type_param_vector({0, 4, 2, 6}); + using T = TypeParam; + auto const input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); + auto const input_column_valid = + cudf::test::make_type_param_vector({1, 1, 1, 0, 0, 1, 1, 1}); + auto const values_to_replace_column = cudf::test::make_type_param_vector({2, 6, 4, 8}); + auto const replacement_values_column = cudf::test::make_type_param_vector({0, 4, 2, 6}); test_replace( input_column, values_to_replace_column, replacement_values_column, input_column_valid); @@ -463,11 +465,12 @@ TYPED_TEST(ReplaceTest, NullsInData) // Testing with replacement column containing nulls TYPED_TEST(ReplaceTest, NullsInNewValues) { - using T = TypeParam; - std::vector input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); - std::vector values_to_replace_column = cudf::test::make_type_param_vector({2, 6, 4, 8}); - std::vector replacement_values_column = cudf::test::make_type_param_vector({0, 4, 2, 6}); - std::vector replacement_values_valid{0, 1, 1, 1}; + using T = TypeParam; + auto const input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); + auto const values_to_replace_column = cudf::test::make_type_param_vector({2, 6, 4, 8}); + auto const replacement_values_column = cudf::test::make_type_param_vector({0, 4, 2, 6}); + auto const replacement_values_valid = + cudf::test::make_type_param_vector({0, 1, 1, 1}); test_replace(input_column, values_to_replace_column, @@ -479,12 +482,14 @@ TYPED_TEST(ReplaceTest, NullsInNewValues) // Testing with both replacement and input column containing nulls TYPED_TEST(ReplaceTest, NullsInBoth) { - using T = TypeParam; - std::vector input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); - std::vector input_column_valid{1, 1, 1, 0, 0, 1, 1, 1}; - std::vector values_to_replace_column = cudf::test::make_type_param_vector({2, 6, 4, 8}); - std::vector replacement_values_column = cudf::test::make_type_param_vector({0, 4, 2, 6}); - std::vector replacement_values_valid{1, 1, 0, 1}; + using T = TypeParam; + auto const input_column = cudf::test::make_type_param_vector({7, 5, 6, 3, 1, 2, 8, 4}); + auto const input_column_valid = + cudf::test::make_type_param_vector({1, 1, 1, 0, 0, 1, 1, 1}); + auto const values_to_replace_column = cudf::test::make_type_param_vector({2, 6, 4, 8}); + auto const replacement_values_column = cudf::test::make_type_param_vector({0, 4, 2, 6}); + auto const replacement_values_valid = + cudf::test::make_type_param_vector({1, 1, 0, 1}); test_replace(input_column, values_to_replace_column, diff --git a/cpp/tests/rolling/rolling_test.cpp b/cpp/tests/rolling/rolling_test.cpp index 33171b269ce..c22acf6b022 100644 --- a/cpp/tests/rolling/rolling_test.cpp +++ b/cpp/tests/rolling/rolling_test.cpp @@ -603,8 +603,7 @@ TYPED_TEST_CASE(RollingTest, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(RollingTest, SimpleStatic) { // https://pandas.pydata.org/pandas-docs/stable/reference/api/pandas.DataFrame.rolling.html - const std::vector col_data = - cudf::test::make_type_param_vector({0, 1, 2, 0, 4}); + auto const col_data = cudf::test::make_type_param_vector({0, 1, 2, 0, 4}); const std::vector col_mask = {1, 1, 1, 0, 1}; fixed_width_column_wrapper input(col_data.begin(), col_data.end(), col_mask.begin()); @@ -632,8 +631,7 @@ TYPED_TEST(RollingTest, NegativeWindowSizes) TYPED_TEST(RollingTest, SimpleDynamic) { // https://pandas.pydata.org/pandas-docs/stable/reference/api/pandas.DataFrame.rolling.html - const std::vector col_data = - cudf::test::make_type_param_vector({0, 1, 2, 0, 4}); + auto const col_data = cudf::test::make_type_param_vector({0, 1, 2, 0, 4}); const std::vector col_mask = {1, 1, 1, 0, 1}; fixed_width_column_wrapper input(col_data.begin(), col_data.end(), col_mask.begin()); @@ -647,8 +645,7 @@ TYPED_TEST(RollingTest, SimpleDynamic) // this is a special test to check the volatile count variable issue (see rolling.cu for detail) TYPED_TEST(RollingTest, VolatileCount) { - const std::vector col_data = - cudf::test::make_type_param_vector({8, 70, 45, 20, 59, 80}); + auto const col_data = cudf::test::make_type_param_vector({8, 70, 45, 20, 59, 80}); const std::vector col_mask = {1, 1, 0, 0, 1, 0}; fixed_width_column_wrapper input(col_data.begin(), col_data.end(), col_mask.begin()); diff --git a/cpp/tests/stream_compaction/drop_duplicates_tests.cpp b/cpp/tests/stream_compaction/drop_duplicates_tests.cpp index ee4a4df38e8..373cd50fb1f 100644 --- a/cpp/tests/stream_compaction/drop_duplicates_tests.cpp +++ b/cpp/tests/stream_compaction/drop_duplicates_tests.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019, NVIDIA CORPORATION. + * 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. @@ -41,7 +41,7 @@ TYPED_TEST(DistinctCountCommon, NoNull) { using T = TypeParam; - std::vector input = cudf::test::make_type_param_vector( + auto const input = cudf::test::make_type_param_vector( {1, 3, 3, 4, 31, 1, 8, 2, 0, 4, 1, 4, 10, 40, 31, 42, 0, 42, 8, 5, 4}); cudf::test::fixed_width_column_wrapper input_col(input.begin(), input.end()); @@ -55,9 +55,9 @@ TYPED_TEST(DistinctCountCommon, TableNoNull) { using T = TypeParam; - std::vector input1 = cudf::test::make_type_param_vector( + auto const input1 = cudf::test::make_type_param_vector( {1, 3, 3, 4, 31, 1, 8, 2, 0, 4, 1, 4, 10, 40, 31, 42, 0, 42, 8, 5, 4}); - std::vector input2 = cudf::test::make_type_param_vector( + auto const input2 = cudf::test::make_type_param_vector( {3, 3, 4, 31, 1, 8, 5, 0, 4, 1, 4, 10, 40, 31, 42, 0, 42, 8, 5, 4, 1}); std::vector> pair_input; From 4e5af1c77f4adb4263b181c620012f3b2c740e0f Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 11 May 2021 11:42:39 +1000 Subject: [PATCH 21/27] Fix map_test race condition --- cpp/tests/hash_map/map_test.cu | 70 +++++++++++++++++++++++----------- 1 file changed, 48 insertions(+), 22 deletions(-) diff --git a/cpp/tests/hash_map/map_test.cu b/cpp/tests/hash_map/map_test.cu index a25c35a63e0..a747646d894 100644 --- a/cpp/tests/hash_map/map_test.cu +++ b/cpp/tests/hash_map/map_test.cu @@ -32,6 +32,7 @@ #include #include #include +#include "rmm/exec_policy.hpp" template struct key_value_types { @@ -137,53 +138,78 @@ TYPED_TEST(InsertTest, UniqueKeysUniqueValues) { using map_type = typename TypeParam::map_type; using pair_type = typename TypeParam::pair_type; - thrust::tabulate(this->pairs.begin(), this->pairs.end(), unique_pair_generator{}); + thrust::tabulate( + rmm::exec_policy(), this->pairs.begin(), this->pairs.end(), unique_pair_generator{}); // All pairs should be new inserts - EXPECT_TRUE(thrust::all_of( - this->pairs.begin(), this->pairs.end(), insert_pair{*this->map})); + EXPECT_TRUE(thrust::all_of(rmm::exec_policy(), + this->pairs.begin(), + this->pairs.end(), + insert_pair{*this->map})); // All pairs should be present in the map - EXPECT_TRUE(thrust::all_of( - this->pairs.begin(), this->pairs.end(), find_pair{*this->map})); + EXPECT_TRUE(thrust::all_of(rmm::exec_policy(), + this->pairs.begin(), + this->pairs.end(), + find_pair{*this->map})); } TYPED_TEST(InsertTest, IdenticalKeysIdenticalValues) { using map_type = typename TypeParam::map_type; using pair_type = typename TypeParam::pair_type; - thrust::tabulate(this->pairs.begin(), this->pairs.end(), identical_pair_generator{}); + thrust::tabulate(rmm::exec_policy(), + this->pairs.begin(), + this->pairs.end(), + identical_pair_generator{}); // Insert a single pair - EXPECT_TRUE(thrust::all_of( - this->pairs.begin(), this->pairs.begin() + 1, insert_pair{*this->map})); + EXPECT_TRUE(thrust::all_of(rmm::exec_policy(), + this->pairs.begin(), + this->pairs.begin() + 1, + insert_pair{*this->map})); // Identical inserts should all return false (no new insert) - EXPECT_FALSE(thrust::all_of( - this->pairs.begin(), this->pairs.end(), insert_pair{*this->map})); + EXPECT_FALSE(thrust::all_of(rmm::exec_policy(), + this->pairs.begin(), + this->pairs.end(), + insert_pair{*this->map})); // All pairs should be present in the map - EXPECT_TRUE(thrust::all_of( - this->pairs.begin(), this->pairs.end(), find_pair{*this->map})); + EXPECT_TRUE(thrust::all_of(rmm::exec_policy(), + this->pairs.begin(), + this->pairs.end(), + find_pair{*this->map})); } TYPED_TEST(InsertTest, IdenticalKeysUniqueValues) { using map_type = typename TypeParam::map_type; using pair_type = typename TypeParam::pair_type; - thrust::tabulate(this->pairs.begin(), this->pairs.end(), identical_key_generator{}); + thrust::tabulate(rmm::exec_policy(), + this->pairs.begin(), + this->pairs.end(), + identical_key_generator{}); // Insert a single pair - EXPECT_TRUE(thrust::all_of( - this->pairs.begin(), this->pairs.begin() + 1, insert_pair{*this->map})); + EXPECT_TRUE(thrust::all_of(rmm::exec_policy(), + this->pairs.begin(), + this->pairs.begin() + 1, + insert_pair{*this->map})); // Identical key inserts should all return false (no new insert) - EXPECT_FALSE(thrust::all_of( - this->pairs.begin() + 1, this->pairs.end(), insert_pair{*this->map})); + EXPECT_FALSE(thrust::all_of(rmm::exec_policy(), + this->pairs.begin() + 1, + this->pairs.end(), + insert_pair{*this->map})); // Only first pair is present in map - EXPECT_TRUE(thrust::all_of( - this->pairs.begin(), this->pairs.begin() + 1, find_pair{*this->map})); - - EXPECT_FALSE(thrust::all_of( - this->pairs.begin() + 1, this->pairs.end(), find_pair{*this->map})); + EXPECT_TRUE(thrust::all_of(rmm::exec_policy(), + this->pairs.begin(), + this->pairs.begin() + 1, + find_pair{*this->map})); + + EXPECT_FALSE(thrust::all_of(rmm::exec_policy(), + this->pairs.begin() + 1, + this->pairs.end(), + find_pair{*this->map})); } CUDF_TEST_PROGRAM_MAIN() From e6e622f3d59a23158d95088651a9891491ad7ed3 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 11 May 2021 12:04:13 +1000 Subject: [PATCH 22/27] Convert type_dispatcher_test to uvector --- cpp/tests/types/type_dispatcher_test.cu | 16 +++++++++------- 1 file changed, 9 insertions(+), 7 deletions(-) diff --git a/cpp/tests/types/type_dispatcher_test.cu b/cpp/tests/types/type_dispatcher_test.cu index 3924fa1ac19..bc690e04f21 100644 --- a/cpp/tests/types/type_dispatcher_test.cu +++ b/cpp/tests/types/type_dispatcher_test.cu @@ -14,6 +14,7 @@ * limitations under the License. */ +#include #include #include #include @@ -21,7 +22,8 @@ #include #include -#include +#include +#include struct DispatcherTest : public cudf::test::BaseFixture { }; @@ -67,10 +69,10 @@ __global__ void dispatch_test_kernel(cudf::type_id id, bool* d_result) TYPED_TEST(TypedDispatcherTest, DeviceDispatch) { - thrust::device_vector result(1, false); - dispatch_test_kernel<<<1, 1>>>(cudf::type_to_id(), result.data().get()); + auto result = cudf::detail::make_zeroed_device_uvector_sync(1); + dispatch_test_kernel<<<1, 1>>>(cudf::type_to_id(), result.data()); CUDA_TRY(cudaDeviceSynchronize()); - EXPECT_EQ(true, result[0]); + EXPECT_EQ(true, result.front_element(rmm::cuda_stream_default)); } struct IdDispatcherTest : public DispatcherTest, public testing::WithParamInterface { @@ -128,11 +130,11 @@ __global__ void double_dispatch_test_kernel(cudf::type_id id1, cudf::type_id id2 TYPED_TEST(TypedDoubleDispatcherTest, DeviceDoubleDispatch) { - thrust::device_vector result(1, false); + auto result = cudf::detail::make_zeroed_device_uvector_sync(1); double_dispatch_test_kernel<<<1, 1>>>( - cudf::type_to_id(), cudf::type_to_id(), result.data().get()); + cudf::type_to_id(), cudf::type_to_id(), result.data()); CUDA_TRY(cudaDeviceSynchronize()); - EXPECT_EQ(true, result[0]); + EXPECT_EQ(true, result.front_element(rmm::cuda_stream_default)); } struct IdDoubleDispatcherTest : public DispatcherTest, From fa9856b4013f97334e3eb2ecba49ac0bfe4e4c99 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 11 May 2021 12:36:43 +1000 Subject: [PATCH 23/27] Convert column_utilities.cu to uvector --- cpp/tests/utilities/column_utilities.cu | 33 ++++++++++++++----------- 1 file changed, 19 insertions(+), 14 deletions(-) diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index f17446ca1dc..8aac7370b13 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -42,6 +42,8 @@ #include #include +#include "cudf/detail/utilities/vector_factories.hpp" +#include "rmm/cuda_stream_view.hpp" namespace cudf { namespace test { @@ -170,7 +172,7 @@ class corresponding_rows_not_equivalent { }; // Stringify the inconsistent values resulted from the comparison of two columns element-wise -std::string stringify_column_differences(thrust::device_vector const& differences, +std::string stringify_column_differences(cudf::device_span differences, column_view const& lhs, column_view const& rhs, bool print_all_differences, @@ -182,10 +184,10 @@ std::string stringify_column_differences(thrust::device_vector const& diffe std::ostringstream buffer; buffer << depth_str << "differences:" << std::endl; - // thrust may crash if a device_vector is passed to fixed_width_column_wrapper, + // thrust may crash if a device vector is passed to fixed_width_column_wrapper, // thus we construct fixed_width_column_wrapper from a host_vector instead - thrust::host_vector h_differences(differences); - auto source_table = cudf::table_view({lhs, rhs}); + auto h_differences = cudf::detail::make_host_vector_sync(differences); + auto source_table = cudf::table_view({lhs, rhs}); auto diff_column = fixed_width_column_wrapper(h_differences.begin(), h_differences.end()); auto diff_table = cudf::gather(source_table, diff_column); @@ -222,16 +224,18 @@ struct column_comparator_impl { corresponding_rows_unequal, corresponding_rows_not_equivalent>; - auto differences = thrust::device_vector(lhs.size()); // worst case: everything different - auto diff_iter = thrust::copy_if(thrust::device, + auto differences = rmm::device_uvector( + lhs.size(), rmm::cuda_stream_default); // worst case: everything different + auto diff_iter = thrust::copy_if(rmm::exec_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(lhs.size()), differences.begin(), ComparatorType(*d_lhs, *d_rhs)); - differences.resize(thrust::distance(differences.begin(), diff_iter)); // shrink back down + differences.resize(thrust::distance(differences.begin(), diff_iter), + rmm::cuda_stream_default); // shrink back down - if (not differences.empty()) + if (not differences.is_empty()) GTEST_FAIL() << stringify_column_differences( differences, lhs, rhs, print_all_differences, depth); } @@ -256,7 +260,7 @@ struct column_comparator_impl { if (lhs_l.is_empty()) { return; } // worst case - everything is different - thrust::device_vector differences(lhs.size()); + rmm::device_uvector differences(lhs.size(), rmm::cuda_stream_default); // TODO : determine how equals/equivalency should work for columns with divergent underlying // data, but equivalent null masks. Example: @@ -307,7 +311,7 @@ struct column_comparator_impl { }); auto diff_iter = thrust::copy_if( - thrust::device, + rmm::exec_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(lhs_l.size() + 1), differences.begin(), @@ -323,9 +327,10 @@ struct column_comparator_impl { return lhs_offsets[index] == rhs_offsets[index] ? false : true; }); - differences.resize(thrust::distance(differences.begin(), diff_iter)); // shrink back down + differences.resize(thrust::distance(differences.begin(), diff_iter), + rmm::cuda_stream_default); // shrink back down - if (not differences.empty()) + if (not differences.is_empty()) GTEST_FAIL() << stringify_column_differences( differences, lhs, rhs, print_all_differences, depth); @@ -522,7 +527,7 @@ std::string nested_offsets_to_string(NestedColumnView const& c, std::string cons // the first offset value to normalize everything against size_type first = cudf::detail::get_value(offsets, c.offset(), rmm::cuda_stream_default); - rmm::device_vector shifted_offsets(output_size); + rmm::device_uvector shifted_offsets(output_size, rmm::cuda_stream_default); // normalize the offset values for the column offset size_type const* d_offsets = offsets.head() + c.offset(); @@ -533,7 +538,7 @@ std::string nested_offsets_to_string(NestedColumnView const& c, std::string cons shifted_offsets.begin(), [first] __device__(int32_t offset) { return static_cast(offset - first); }); - thrust::host_vector h_shifted_offsets(shifted_offsets); + auto const h_shifted_offsets = cudf::detail::make_host_vector_sync(shifted_offsets); std::ostringstream buffer; for (size_t idx = 0; idx < h_shifted_offsets.size(); idx++) { buffer << h_shifted_offsets[idx]; From c48ab134a92057c2c858edf5ecfb972924aabc3c Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 11 May 2021 12:55:55 +1000 Subject: [PATCH 24/27] Convert value_iterator_test* to uvector --- cpp/tests/iterator/value_iterator_test.cuh | 3 +- .../iterator/value_iterator_test_strings.cu | 35 ++++++++----------- 2 files changed, 17 insertions(+), 21 deletions(-) diff --git a/cpp/tests/iterator/value_iterator_test.cuh b/cpp/tests/iterator/value_iterator_test.cuh index f8dab90b2b5..3a7ef075a41 100644 --- a/cpp/tests/iterator/value_iterator_test.cuh +++ b/cpp/tests/iterator/value_iterator_test.cuh @@ -13,13 +13,14 @@ * the License. */ #include +#include "cudf/detail/utilities/vector_factories.hpp" // tests for non-null iterator (pointer of device array) template void non_null_iterator(IteratorTest& testFixture) { auto host_array = cudf::test::make_type_param_vector({0, 6, 0, -14, 13, 64, -13, -20, 45}); - thrust::device_vector dev_array(host_array); + auto dev_array = cudf::detail::make_device_uvector_sync(host_array); // calculate the expected value by CPU. thrust::host_vector replaced_array(host_array); diff --git a/cpp/tests/iterator/value_iterator_test_strings.cu b/cpp/tests/iterator/value_iterator_test_strings.cu index 2d343bf7cdd..f28067649fd 100644 --- a/cpp/tests/iterator/value_iterator_test_strings.cu +++ b/cpp/tests/iterator/value_iterator_test_strings.cu @@ -13,6 +13,9 @@ * the License. */ #include +#include "cudf/detail/utilities/vector_factories.hpp" +#include "rmm/cuda_stream_view.hpp" +#include "rmm/device_uvector.hpp" auto strings_to_string_views(std::vector& input_strings) { @@ -21,15 +24,14 @@ auto strings_to_string_views(std::vector& input_strings) std::vector offsets; std::tie(chars, offsets) = cudf::test::detail::make_chars_and_offsets( input_strings.begin(), input_strings.end(), all_valid); - thrust::device_vector dev_chars(chars); - char* c_start = thrust::raw_pointer_cast(dev_chars.data()); + auto dev_chars = cudf::detail::make_device_uvector_sync(chars); // calculate the expected value by CPU. (but contains device pointers) - std::vector replaced_array(input_strings.size()); + thrust::host_vector replaced_array(input_strings.size()); std::transform(thrust::counting_iterator(0), thrust::counting_iterator(replaced_array.size()), replaced_array.begin(), - [c_start, offsets](auto i) { + [c_start = dev_chars.begin(), offsets](auto i) { return cudf::string_view(c_start + offsets[i], offsets[i + 1] - offsets[i]); }); return std::make_tuple(std::move(dev_chars), replaced_array); @@ -41,11 +43,10 @@ struct StringIteratorTest : public IteratorTest { TEST_F(StringIteratorTest, string_view_null_iterator) { using T = cudf::string_view; - // T init = T{"", 0}; std::string zero("zero"); // the char data has to be in GPU - thrust::device_vector initmsg(zero.begin(), zero.end()); - T init = T{initmsg.data().get(), int(initmsg.size())}; + auto initmsg = cudf::detail::make_device_uvector_sync(zero); + T init = T{initmsg.data(), int(initmsg.size())}; // data and valid arrays std::vector host_values( @@ -60,9 +61,7 @@ TEST_F(StringIteratorTest, string_view_null_iterator) replaced_strings.begin(), [zero](auto s, auto b) { return b ? s : zero; }); - thrust::device_vector dev_chars; - thrust::host_vector replaced_array(host_values.size()); - std::tie(dev_chars, replaced_array) = strings_to_string_views(replaced_strings); + auto [dev_chars, replaced_array] = strings_to_string_views(replaced_strings); // create a column with bool vector cudf::test::strings_column_wrapper w_col( @@ -81,16 +80,14 @@ TEST_F(StringIteratorTest, string_view_no_null_iterator) // T init = T{"", 0}; std::string zero("zero"); // the char data has to be in GPU - thrust::device_vector initmsg(zero.begin(), zero.end()); - T init = T{initmsg.data().get(), int(initmsg.size())}; + auto initmsg = cudf::detail::make_device_uvector_sync(zero); + T init = T{initmsg.data(), int(initmsg.size())}; // data array std::vector host_values( {"one", "two", "three", "four", "five", "six", "eight", "nine"}); - thrust::device_vector dev_chars; - thrust::host_vector all_array(host_values.size()); - std::tie(dev_chars, all_array) = strings_to_string_views(host_values); + auto [dev_chars, all_array] = strings_to_string_views(host_values); // create a column with bool vector cudf::test::strings_column_wrapper w_col(host_values.begin(), host_values.end()); @@ -107,15 +104,13 @@ TEST_F(StringIteratorTest, string_scalar_iterator) // T init = T{"", 0}; std::string zero("zero"); // the char data has to be in GPU - thrust::device_vector initmsg(zero.begin(), zero.end()); - T init = T{initmsg.data().get(), int(initmsg.size())}; + auto initmsg = cudf::detail::make_device_uvector_sync(zero); + T init = T{initmsg.data(), int(initmsg.size())}; // data array std::vector host_values(100, zero); - thrust::device_vector dev_chars; - thrust::host_vector all_array(host_values.size()); - std::tie(dev_chars, all_array) = strings_to_string_views(host_values); + auto [dev_chars, all_array] = strings_to_string_views(host_values); // calculate the expected value by CPU. thrust::host_vector> value_and_validity(host_values.size()); From 3536c8f39b273da5c22b2bbc3b37c1a39c0e0595 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 12 May 2021 07:57:58 +1000 Subject: [PATCH 25/27] Remove anonymous namespace --- cpp/include/cudf/detail/utilities/vector_factories.hpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/cpp/include/cudf/detail/utilities/vector_factories.hpp b/cpp/include/cudf/detail/utilities/vector_factories.hpp index 8bcf75a9769..ebee4439c40 100644 --- a/cpp/include/cudf/detail/utilities/vector_factories.hpp +++ b/cpp/include/cudf/detail/utilities/vector_factories.hpp @@ -281,8 +281,6 @@ rmm::device_uvector make_device_uvector_sync( return make_device_uvector_sync(device_span{c}, stream, mr); } -namespace { - // Utility function template to allow copying to either a thrust::host_vector or std::vector template OutContainer make_vector_async(device_span v, @@ -293,7 +291,6 @@ OutContainer make_vector_async(device_span v, result.data(), v.data(), v.size() * sizeof(T), cudaMemcpyDeviceToHost, stream.value())); return result; } -} // namespace /** * @brief Asynchronously construct a `std::vector` containing a copy of data from a From d26371d46a46b5f574209f4dc102036ac0a0c513 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 12 May 2021 08:24:48 +1000 Subject: [PATCH 26/27] Require explicit stream parameter to async vector factories --- .../detail/utilities/vector_factories.hpp | 26 ++++++++----------- cpp/src/copying/concatenate.cu | 11 +++----- 2 files changed, 14 insertions(+), 23 deletions(-) diff --git a/cpp/include/cudf/detail/utilities/vector_factories.hpp b/cpp/include/cudf/detail/utilities/vector_factories.hpp index ebee4439c40..1e735719400 100644 --- a/cpp/include/cudf/detail/utilities/vector_factories.hpp +++ b/cpp/include/cudf/detail/utilities/vector_factories.hpp @@ -47,7 +47,7 @@ namespace detail { template rmm::device_uvector make_zeroed_device_uvector_async( std::size_t size, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { rmm::device_uvector ret(size, stream, mr); @@ -93,7 +93,7 @@ rmm::device_uvector make_zeroed_device_uvector_sync( template rmm::device_uvector make_device_uvector_async( host_span source_data, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { rmm::device_uvector ret(source_data.size(), stream, mr); @@ -124,7 +124,7 @@ template >::value>* = nullptr> rmm::device_uvector make_device_uvector_async( Container const& c, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { return make_device_uvector_async(host_span{c}, stream, mr); @@ -177,7 +177,7 @@ template < nullptr> rmm::device_uvector make_device_uvector_async( Container const& c, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { return make_device_uvector_async( @@ -283,8 +283,7 @@ rmm::device_uvector make_device_uvector_sync( // Utility function template to allow copying to either a thrust::host_vector or std::vector template -OutContainer make_vector_async(device_span v, - rmm::cuda_stream_view stream = rmm::cuda_stream_default) +OutContainer make_vector_async(device_span v, rmm::cuda_stream_view stream) { OutContainer result(v.size()); CUDA_TRY(cudaMemcpyAsync( @@ -304,8 +303,7 @@ OutContainer make_vector_async(device_span v, * @return The data copied to the host */ template -std::vector make_std_vector_async(device_span v, - rmm::cuda_stream_view stream = rmm::cuda_stream_default) +std::vector make_std_vector_async(device_span v, rmm::cuda_stream_view stream) { return make_vector_async>(v, stream); } @@ -327,8 +325,8 @@ template < std::enable_if_t< std::is_convertible>::value>* = nullptr> -std::vector make_std_vector_async( - Container const& c, rmm::cuda_stream_view stream = rmm::cuda_stream_default) +std::vector make_std_vector_async(Container const& c, + rmm::cuda_stream_view stream) { return make_std_vector_async(device_span{c}, stream); } @@ -345,8 +343,7 @@ std::vector make_std_vector_async( * @return The data copied to the host */ template -std::vector make_std_vector_sync(device_span v, - rmm::cuda_stream_view stream = rmm::cuda_stream_default) +std::vector make_std_vector_sync(device_span v, rmm::cuda_stream_view stream) { auto result = make_std_vector_async(v, stream); stream.synchronize(); @@ -388,8 +385,7 @@ std::vector make_std_vector_sync( * @return The data copied to the host */ template -thrust::host_vector make_host_vector_async( - device_span v, rmm::cuda_stream_view stream = rmm::cuda_stream_default) +thrust::host_vector make_host_vector_async(device_span v, rmm::cuda_stream_view stream) { return make_vector_async>(v, stream); } @@ -412,7 +408,7 @@ template < std::is_convertible>::value>* = nullptr> thrust::host_vector make_host_vector_async( - Container const& c, rmm::cuda_stream_view stream = rmm::cuda_stream_default) + Container const& c, rmm::cuda_stream_view stream) { return make_host_vector_async(device_span{c}, stream); } diff --git a/cpp/src/copying/concatenate.cu b/cpp/src/copying/concatenate.cu index e87cadbffe8..6ba10bef396 100644 --- a/cpp/src/copying/concatenate.cu +++ b/cpp/src/copying/concatenate.cu @@ -57,9 +57,6 @@ auto create_device_views(host_span views, rmm::cuda_stream_vi column_device_view::create(std::declval(), std::declval())); auto device_view_owners = std::vector(views.size()); std::transform(views.begin(), views.end(), device_view_owners.begin(), [stream](auto const& col) { - // TODO creating this device view can invoke null count computation - // even though it isn't used. See this issue: - // https://github.com/rapidsai/cudf/issues/4368 return column_device_view::create(col, stream); }); @@ -70,10 +67,8 @@ auto create_device_views(host_span views, rmm::cuda_stream_vi device_view_owners.cend(), std::back_inserter(device_views), [](auto const& col) { return *col; }); - // TODO each of these device vector copies invoke stream synchronization - // which appears to add unnecessary overhead. See this issue: - // https://github.com/rapidsai/rmm/issues/120 - auto d_views = make_device_uvector_async(device_views); + + auto d_views = make_device_uvector_async(device_views, stream); // Compute the partition offsets auto offsets = thrust::host_vector(views.size() + 1); @@ -84,7 +79,7 @@ auto create_device_views(host_span views, rmm::cuda_stream_vi std::next(offsets.begin()), [](auto const& col) { return col.size(); }, thrust::plus{}); - auto d_offsets = make_device_uvector_async(offsets); + auto d_offsets = make_device_uvector_async(offsets, stream); auto const output_size = offsets.back(); return std::make_tuple( From d1866874f2f6759f80a4abe4a5712c7e33f471ea Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 12 May 2021 08:36:36 +1000 Subject: [PATCH 27/27] Remove optional default stream from exec_policy calls --- cpp/tests/column/compound_test.cu | 12 ++++++------ cpp/tests/transform/row_bit_count_test.cu | 8 ++++---- 2 files changed, 10 insertions(+), 10 deletions(-) diff --git a/cpp/tests/column/compound_test.cu b/cpp/tests/column/compound_test.cu index b5ed2ab7a49..0df1cfaeccc 100644 --- a/cpp/tests/column/compound_test.cu +++ b/cpp/tests/column/compound_test.cu @@ -65,7 +65,7 @@ struct checker_for_level2 { TEST_F(CompoundColumnTest, ChildrenLevel1) { rmm::device_uvector data(1000, rmm::cuda_stream_default); - thrust::sequence(rmm::exec_policy(rmm::cuda_stream_default), data.begin(), data.end(), 1); + thrust::sequence(rmm::exec_policy(), data.begin(), data.end(), 1); auto null_mask = cudf::create_null_mask(100, cudf::mask_state::UNALLOCATED); rmm::device_buffer data1(data.data() + 100, 100 * sizeof(int32_t)); @@ -92,14 +92,14 @@ TEST_F(CompoundColumnTest, ChildrenLevel1) { auto column = cudf::column_device_view::create(parent->view()); - EXPECT_TRUE(thrust::any_of(rmm::exec_policy(rmm::cuda_stream_default), + EXPECT_TRUE(thrust::any_of(rmm::exec_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(100), checker_for_level1{*column})); } { auto column = cudf::mutable_column_device_view::create(parent->mutable_view()); - EXPECT_TRUE(thrust::any_of(rmm::exec_policy(rmm::cuda_stream_default), + EXPECT_TRUE(thrust::any_of(rmm::exec_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(100), checker_for_level1{*column})); @@ -109,7 +109,7 @@ TEST_F(CompoundColumnTest, ChildrenLevel1) TEST_F(CompoundColumnTest, ChildrenLevel2) { rmm::device_uvector data(1000, rmm::cuda_stream_default); - thrust::sequence(rmm::exec_policy(rmm::cuda_stream_default), data.begin(), data.end(), 1); + thrust::sequence(rmm::exec_policy(), data.begin(), data.end(), 1); auto null_mask = cudf::create_null_mask(100, cudf::mask_state::UNALLOCATED); rmm::device_buffer data11(data.data() + 100, 100 * sizeof(int32_t)); @@ -165,14 +165,14 @@ TEST_F(CompoundColumnTest, ChildrenLevel2) { auto column = cudf::column_device_view::create(parent->view()); - EXPECT_TRUE(thrust::any_of(rmm::exec_policy(rmm::cuda_stream_default), + EXPECT_TRUE(thrust::any_of(rmm::exec_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(100), checker_for_level2{*column})); } { auto column = cudf::mutable_column_device_view::create(parent->mutable_view()); - EXPECT_TRUE(thrust::any_of(rmm::exec_policy(rmm::cuda_stream_default), + EXPECT_TRUE(thrust::any_of(rmm::exec_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(100), checker_for_level2{*column})); diff --git a/cpp/tests/transform/row_bit_count_test.cu b/cpp/tests/transform/row_bit_count_test.cu index 313113a58e0..1431710f3ca 100644 --- a/cpp/tests/transform/row_bit_count_test.cu +++ b/cpp/tests/transform/row_bit_count_test.cu @@ -45,7 +45,7 @@ TYPED_TEST(RowBitCountTyped, SimpleTypes) // expect size of the type per row auto expected = make_fixed_width_column(data_type{type_id::INT32}, 16); cudf::mutable_column_view mcv(*expected); - thrust::fill(rmm::exec_policy(rmm::cuda_stream_default), + thrust::fill(rmm::exec_policy(), mcv.begin(), mcv.end(), sizeof(device_storage_type_t) * CHAR_BIT); @@ -68,7 +68,7 @@ TYPED_TEST(RowBitCountTyped, SimpleTypesWithNulls) // expect size of the type + 1 bit per row auto expected = make_fixed_width_column(data_type{type_id::INT32}, 16); cudf::mutable_column_view mcv(*expected); - thrust::fill(rmm::exec_policy(rmm::cuda_stream_default), + thrust::fill(rmm::exec_policy(), mcv.begin(), mcv.end(), (sizeof(device_storage_type_t) * CHAR_BIT) + 1); @@ -488,7 +488,7 @@ TEST_F(RowBitCount, Table) auto expected = cudf::make_fixed_width_column(data_type{type_id::INT32}, t.num_rows()); cudf::mutable_column_view mcv(*expected); thrust::transform( - rmm::exec_policy(rmm::cuda_stream_default), + rmm::exec_policy(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + t.num_rows(), mcv.begin(), @@ -591,4 +591,4 @@ TEST_F(RowBitCount, EmptyTable) auto result = cudf::row_bit_count(empty); CUDF_EXPECTS(result != nullptr && result->size() == 0, "Expected an empty column"); } -} \ No newline at end of file +}