Skip to content

Commit

Permalink
Move kernels to their own namespaces to avoid build conflicts
Browse files Browse the repository at this point in the history
  • Loading branch information
@PointKernel
PointKernel committed Nov 25, 2024
1 parent 1bb0ae6 commit 9fa8904
Show file tree
Hide file tree
Showing 4 changed files with 22 additions and 20 deletions.
9 changes: 5 additions & 4 deletions include/cuco/detail/bloom_filter/bloom_filter_impl.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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<cg_size, block_size>
bloom_filter_ns::detail::add_if_n<cg_size, block_size>
<<<grid_size, block_size, 0, stream.get()>>>(first, num_keys, stencil, pred, *this);
}

Expand Down Expand Up @@ -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<cg_size, block_size><<<grid_size, block_size, 0, stream.get()>>>(
first, num_keys, stencil, pred, output_begin, *this);
bloom_filter_ns::detail::contains_if_n<cg_size, block_size>
<<<grid_size, block_size, 0, stream.get()>>>(
first, num_keys, stencil, pred, output_begin, *this);
}

[[nodiscard]] __host__ __device__ constexpr word_type* data() noexcept { return words_; }
Expand Down Expand Up @@ -365,4 +366,4 @@ class bloom_filter_impl {
policy_type policy_;
};

} // namespace cuco::detail
} // namespace cuco::detail
4 changes: 2 additions & 2 deletions include/cuco/detail/bloom_filter/kernels.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cstdint>
#include <iterator>

namespace cuco::detail {
namespace cuco::bloom_filter_ns::detail {

CUCO_SUPPRESS_KERNEL_WARNINGS

Expand Down Expand Up @@ -89,4 +89,4 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void contains_if_n(InputIt first,
}
}

} // namespace cuco::detail
} // namespace cuco::bloom_filter_ns::detail
4 changes: 2 additions & 2 deletions include/cuco/detail/open_addressing/kernels.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@

#include <iterator>

namespace cuco::detail {
namespace cuco::open_addressing_ns::detail {
CUCO_SUPPRESS_KERNEL_WARNINGS

/**
Expand Down Expand Up @@ -729,4 +729,4 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void rehash(
}
}

} // namespace cuco::detail
} // namespace cuco::open_addressing_ns::detail
25 changes: 13 additions & 12 deletions include/cuco/detail/open_addressing/open_addressing_impl.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -342,7 +342,7 @@ class open_addressing_impl {

auto const grid_size = cuco::detail::grid_size(num_keys, cg_size);

detail::insert_if_n<cg_size, cuco::detail::default_block_size()>
open_addressing_ns::detail::insert_if_n<cg_size, cuco::detail::default_block_size()>
<<<grid_size, cuco::detail::default_block_size(), 0, stream.get()>>>(
first, num_keys, stencil, pred, counter.data(), container_ref);

Expand Down Expand Up @@ -384,7 +384,7 @@ class open_addressing_impl {

auto const grid_size = cuco::detail::grid_size(num_keys, cg_size);

detail::insert_if_n<cg_size, cuco::detail::default_block_size()>
open_addressing_ns::detail::insert_if_n<cg_size, cuco::detail::default_block_size()>
<<<grid_size, cuco::detail::default_block_size(), 0, stream.get()>>>(
first, num_keys, stencil, pred, container_ref);
}
Expand Down Expand Up @@ -426,7 +426,7 @@ class open_addressing_impl {

auto const grid_size = cuco::detail::grid_size(num_keys, cg_size);

detail::insert_and_find<cg_size, cuco::detail::default_block_size()>
open_addressing_ns::detail::insert_and_find<cg_size, cuco::detail::default_block_size()>
<<<grid_size, cuco::detail::default_block_size(), 0, stream.get()>>>(
first, num_keys, found_begin, inserted_begin, container_ref);
}
Expand Down Expand Up @@ -466,7 +466,7 @@ class open_addressing_impl {

auto const grid_size = cuco::detail::grid_size(num_keys, cg_size);

detail::erase<cg_size, cuco::detail::default_block_size()>
open_addressing_ns::detail::erase<cg_size, cuco::detail::default_block_size()>
<<<grid_size, cuco::detail::default_block_size(), 0, stream.get()>>>(
first, num_keys, container_ref);
}
Expand Down Expand Up @@ -540,7 +540,7 @@ class open_addressing_impl {

auto const grid_size = cuco::detail::grid_size(num_keys, cg_size);

detail::contains_if_n<cg_size, cuco::detail::default_block_size()>
open_addressing_ns::detail::contains_if_n<cg_size, cuco::detail::default_block_size()>
<<<grid_size, cuco::detail::default_block_size(), 0, stream.get()>>>(
first, num_keys, stencil, pred, output_begin, container_ref);
}
Expand Down Expand Up @@ -615,7 +615,7 @@ class open_addressing_impl {

auto const grid_size = cuco::detail::grid_size(num_keys, cg_size);

detail::find_if_n<cg_size, cuco::detail::default_block_size()>
open_addressing_ns::detail::find_if_n<cg_size, cuco::detail::default_block_size()>
<<<grid_size, cuco::detail::default_block_size(), 0, stream.get()>>>(
first, num_keys, stencil, pred, output_begin, container_ref);
}
Expand Down Expand Up @@ -886,7 +886,7 @@ class open_addressing_impl {

auto const grid_size = cuco::detail::grid_size(num_keys, cg_size);

detail::for_each_n<cg_size, cuco::detail::default_block_size()>
open_addressing_ns::detail::for_each_n<cg_size, cuco::detail::default_block_size()>
<<<grid_size, cuco::detail::default_block_size(), 0, stream.get()>>>(
first, num_keys, std::forward<CallbackOp>(callback_op), container_ref);
}
Expand All @@ -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<cuco::detail::default_block_size()>
open_addressing_ns::detail::size<cuco::detail::default_block_size()>
<<<grid_size, cuco::detail::default_block_size(), 0, stream.get()>>>(
storage_.ref(), is_filled, counter.data());

Expand Down Expand Up @@ -1017,7 +1017,7 @@ class open_addressing_impl {
auto const is_filled = open_addressing_ns::detail::slot_is_filled<has_payload, key_type>{
this->empty_key_sentinel(), this->erased_key_sentinel()};

detail::rehash<block_size><<<grid_size, block_size, 0, stream.get()>>>(
open_addressing_ns::detail::rehash<block_size><<<grid_size, block_size, 0, stream.get()>>>(
old_storage.ref(), container.ref(op::insert), is_filled);
}

Expand Down Expand Up @@ -1120,7 +1120,7 @@ class open_addressing_impl {

auto const grid_size = cuco::detail::grid_size(num_keys, cg_size);

detail::count<IsOuter, cg_size, cuco::detail::default_block_size()>
open_addressing_ns::detail::count<IsOuter, cg_size, cuco::detail::default_block_size()>
<<<grid_size, cuco::detail::default_block_size(), 0, stream.get()>>>(
first, num_keys, counter.data(), container_ref);

Expand Down Expand Up @@ -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<IsOuter, block_size><<<grid_size, block_size, 0, stream.get()>>>(
first, n, output_probe, output_match, counter.data(), container_ref);
open_addressing_ns::detail::retrieve<IsOuter, block_size>
<<<grid_size, block_size, 0, stream.get()>>>(
first, n, output_probe, output_match, counter.data(), container_ref);

auto const num_retrieved = counter.load_to_host(stream.get());

Expand Down

0 comments on commit 9fa8904

Please sign in to comment.