From 9fa89048f154df78bccfa8d3048c03e6c02cd4bc Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Sun, 24 Nov 2024 16:44:36 -0800 Subject: [PATCH] Move kernels to their own namespaces to avoid build conflicts --- .../detail/bloom_filter/bloom_filter_impl.cuh | 9 ++++--- include/cuco/detail/bloom_filter/kernels.cuh | 4 +-- .../cuco/detail/open_addressing/kernels.cuh | 4 +-- .../open_addressing/open_addressing_impl.cuh | 25 ++++++++++--------- 4 files changed, 22 insertions(+), 20 deletions(-) diff --git a/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh index a67db3475..d1a81affd 100644 --- a/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh +++ b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh @@ -198,7 +198,7 @@ class bloom_filter_impl { auto const grid_size = cuco::detail::grid_size(num_keys, cg_size, cuco::detail::default_stride(), block_size); - detail::add_if_n + bloom_filter_ns::detail::add_if_n <<>>(first, num_keys, stencil, pred, *this); } @@ -303,8 +303,9 @@ class bloom_filter_impl { auto const grid_size = cuco::detail::grid_size(num_keys, cg_size, cuco::detail::default_stride(), block_size); - detail::contains_if_n<<>>( - first, num_keys, stencil, pred, output_begin, *this); + bloom_filter_ns::detail::contains_if_n + <<>>( + first, num_keys, stencil, pred, output_begin, *this); } [[nodiscard]] __host__ __device__ constexpr word_type* data() noexcept { return words_; } @@ -365,4 +366,4 @@ class bloom_filter_impl { policy_type policy_; }; -} // namespace cuco::detail \ No newline at end of file +} // namespace cuco::detail diff --git a/include/cuco/detail/bloom_filter/kernels.cuh b/include/cuco/detail/bloom_filter/kernels.cuh index 1d514aa6b..6b02af322 100644 --- a/include/cuco/detail/bloom_filter/kernels.cuh +++ b/include/cuco/detail/bloom_filter/kernels.cuh @@ -22,7 +22,7 @@ #include #include -namespace cuco::detail { +namespace cuco::bloom_filter_ns::detail { CUCO_SUPPRESS_KERNEL_WARNINGS @@ -89,4 +89,4 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void contains_if_n(InputIt first, } } -} // namespace cuco::detail \ No newline at end of file +} // namespace cuco::bloom_filter_ns::detail diff --git a/include/cuco/detail/open_addressing/kernels.cuh b/include/cuco/detail/open_addressing/kernels.cuh index 56f329485..c93edc5f4 100644 --- a/include/cuco/detail/open_addressing/kernels.cuh +++ b/include/cuco/detail/open_addressing/kernels.cuh @@ -25,7 +25,7 @@ #include -namespace cuco::detail { +namespace cuco::open_addressing_ns::detail { CUCO_SUPPRESS_KERNEL_WARNINGS /** @@ -729,4 +729,4 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void rehash( } } -} // namespace cuco::detail +} // namespace cuco::open_addressing_ns::detail diff --git a/include/cuco/detail/open_addressing/open_addressing_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_impl.cuh index f8d36b556..cdec2c019 100644 --- a/include/cuco/detail/open_addressing/open_addressing_impl.cuh +++ b/include/cuco/detail/open_addressing/open_addressing_impl.cuh @@ -342,7 +342,7 @@ class open_addressing_impl { auto const grid_size = cuco::detail::grid_size(num_keys, cg_size); - detail::insert_if_n + open_addressing_ns::detail::insert_if_n <<>>( first, num_keys, stencil, pred, counter.data(), container_ref); @@ -384,7 +384,7 @@ class open_addressing_impl { auto const grid_size = cuco::detail::grid_size(num_keys, cg_size); - detail::insert_if_n + open_addressing_ns::detail::insert_if_n <<>>( first, num_keys, stencil, pred, container_ref); } @@ -426,7 +426,7 @@ class open_addressing_impl { auto const grid_size = cuco::detail::grid_size(num_keys, cg_size); - detail::insert_and_find + open_addressing_ns::detail::insert_and_find <<>>( first, num_keys, found_begin, inserted_begin, container_ref); } @@ -466,7 +466,7 @@ class open_addressing_impl { auto const grid_size = cuco::detail::grid_size(num_keys, cg_size); - detail::erase + open_addressing_ns::detail::erase <<>>( first, num_keys, container_ref); } @@ -540,7 +540,7 @@ class open_addressing_impl { auto const grid_size = cuco::detail::grid_size(num_keys, cg_size); - detail::contains_if_n + open_addressing_ns::detail::contains_if_n <<>>( first, num_keys, stencil, pred, output_begin, container_ref); } @@ -615,7 +615,7 @@ class open_addressing_impl { auto const grid_size = cuco::detail::grid_size(num_keys, cg_size); - detail::find_if_n + open_addressing_ns::detail::find_if_n <<>>( first, num_keys, stencil, pred, output_begin, container_ref); } @@ -886,7 +886,7 @@ class open_addressing_impl { auto const grid_size = cuco::detail::grid_size(num_keys, cg_size); - detail::for_each_n + open_addressing_ns::detail::for_each_n <<>>( first, num_keys, std::forward(callback_op), container_ref); } @@ -912,7 +912,7 @@ class open_addressing_impl { // TODO: custom kernel to be replaced by cub::DeviceReduce::Sum when cub version is bumped to // v2.1.0 - detail::size + open_addressing_ns::detail::size <<>>( storage_.ref(), is_filled, counter.data()); @@ -1017,7 +1017,7 @@ class open_addressing_impl { auto const is_filled = open_addressing_ns::detail::slot_is_filled{ this->empty_key_sentinel(), this->erased_key_sentinel()}; - detail::rehash<<>>( + open_addressing_ns::detail::rehash<<>>( old_storage.ref(), container.ref(op::insert), is_filled); } @@ -1120,7 +1120,7 @@ class open_addressing_impl { auto const grid_size = cuco::detail::grid_size(num_keys, cg_size); - detail::count + open_addressing_ns::detail::count <<>>( first, num_keys, counter.data(), container_ref); @@ -1180,8 +1180,9 @@ class open_addressing_impl { auto constexpr grid_stride = 1; auto const grid_size = cuco::detail::grid_size(n, cg_size, grid_stride, block_size); - detail::retrieve<<>>( - first, n, output_probe, output_match, counter.data(), container_ref); + open_addressing_ns::detail::retrieve + <<>>( + first, n, output_probe, output_match, counter.data(), container_ref); auto const num_retrieved = counter.load_to_host(stream.get());