diff --git a/cpp/include/cudf/detail/utilities/vector_factories.hpp b/cpp/include/cudf/detail/utilities/vector_factories.hpp index db8e6e4a156..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( @@ -281,6 +281,16 @@ rmm::device_uvector make_device_uvector_sync( return make_device_uvector_sync(device_span{c}, stream, mr); } +// 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) +{ + OutContainer result(v.size()); + CUDA_TRY(cudaMemcpyAsync( + result.data(), v.data(), v.size() * sizeof(T), cudaMemcpyDeviceToHost, stream.value())); + return result; +} + /** * @brief Asynchronously construct a `std::vector` containing a copy of data from a * `device_span` @@ -293,13 +303,9 @@ rmm::device_uvector make_device_uvector_sync( * @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) { - 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); } /** @@ -319,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); } @@ -337,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(); @@ -368,6 +373,89 @@ 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) +{ + 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) +{ + 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 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>; 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/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( diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 3246f565443..2766cbb86fc 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -342,7 +342,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 @@ -375,7 +375,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 bitmask/is_element_valid_tests.cpp) diff --git a/cpp/tests/bitmask/bitmask_tests.cu b/cpp/tests/bitmask/bitmask_tests.cpp similarity index 73% rename from cpp/tests/bitmask/bitmask_tests.cu rename to cpp/tests/bitmask/bitmask_tests.cpp index 2f820da687e..3fb12efcc93 100644 --- a/cpp/tests/bitmask/bitmask_tests.cu +++ b/cpp/tests/bitmask/bitmask_tests.cpp @@ -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}; diff --git a/cpp/tests/bitmask/set_nullmask_tests.cu b/cpp/tests/bitmask/set_nullmask_tests.cu index ae4896827fd..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. @@ -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)); } diff --git a/cpp/tests/column/compound_test.cu b/cpp/tests/column/compound_test.cu index 97a6dbb0c22..0df1cfaeccc 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. @@ -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(), 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(), 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(), 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(), 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(), 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(), thrust::make_counting_iterator(0), thrust::make_counting_iterator(100), checker_for_level2{*column})); 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] << "}" 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); } diff --git a/cpp/tests/hash_map/map_test.cu b/cpp/tests/hash_map/map_test.cu index 49cfda078b1..a747646d894 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,24 +14,25 @@ * limitations under the License. */ +#include +#include + #include #include -#include #include -#include +#include #include -#include - #include #include #include #include #include #include +#include "rmm/exec_policy.hpp" template struct key_value_types { @@ -53,13 +54,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; }; @@ -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() 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 diff --git a/cpp/tests/iterator/iterator_tests.cuh b/cpp/tests/iterator/iterator_tests.cuh index 69a905386e2..06ac472d6d5 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. @@ -14,27 +14,30 @@ */ #pragma once -#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 struct IteratorTest : public cudf::test::BaseFixture { @@ -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,57 +75,41 @@ 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) { - 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()); -#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_numeric.cu b/cpp/tests/iterator/optional_iterator_test_numeric.cu index 90dc33ba628..313fd1358f6 100644 --- a/cpp/tests/iterator/optional_iterator_test_numeric.cu +++ b/cpp/tests/iterator/optional_iterator_test_numeric.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/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()); 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/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 32ddf11f16d..58ef08f6052 100644 --- a/cpp/tests/replace/replace_tests.cpp +++ b/cpp/tests/replace/replace_tests.cpp @@ -25,17 +25,17 @@ #include #include +#include #include #include -#include "cudf/fixed_point/fixed_point.hpp" -#include #include #include #include #include #include +#include "cudf/types.hpp" struct ReplaceErrorTest : public cudf::test::BaseFixture { }; @@ -315,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) { @@ -346,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); } @@ -396,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); } @@ -407,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); } @@ -418,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); } @@ -430,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); } @@ -440,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); } @@ -451,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); @@ -464,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, @@ -480,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/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()); 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; 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())); 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.cpp similarity index 96% rename from cpp/tests/strings/integers_tests.cu rename to cpp/tests/strings/integers_tests.cpp index f15116ae4c2..d5f17954c50 100644 --- a/cpp/tests/strings/integers_tests.cu +++ b/cpp/tests/strings/integers_tests.cpp @@ -18,11 +18,15 @@ #include #include +#include #include #include #include #include +#include +#include + #include #include @@ -287,16 +291,16 @@ TYPED_TEST_CASE(StringsIntegerConvertTest, cudf::test::IntegralTypesNotBool); TYPED_TEST(StringsIntegerConvertTest, FromToInteger) { - thrust::device_vector d_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::host_vector h_integers(255); + 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); 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 +308,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)); 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), 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 +} 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, 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]; 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,