diff --git a/benchmarks/hash_table/static_set/insert_bench.cu b/benchmarks/hash_table/static_set/insert_bench.cu index 48bc37fa4..c2443c107 100644 --- a/benchmarks/hash_table/static_set/insert_bench.cu +++ b/benchmarks/hash_table/static_set/insert_bench.cu @@ -47,8 +47,10 @@ void static_set_insert(nvbench::state& state, nvbench::type_list) state.exec(nvbench::exec_tag::sync | nvbench::exec_tag::timer, [&](nvbench::launch& launch, auto& timer) { - cuco::experimental::static_set set{ - size, cuco::empty_key{-1}, {}, {}, {}, {launch.get_stream()}}; + auto set = cuco::experimental::make_static_set( + cuco::experimental::extent{size}, + cuco::empty_key{-1}, + cuco::experimental::cuda_stream_ref{launch.get_stream()}); timer.start(); set.insert(keys.begin(), keys.end(), {launch.get_stream()}); diff --git a/include/cuco/detail/extent_base.cuh b/include/cuco/detail/extent_base.cuh new file mode 100644 index 000000000..063fcc741 --- /dev/null +++ b/include/cuco/detail/extent_base.cuh @@ -0,0 +1,37 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include + +namespace cuco::experimental::detail { + +/** + * @brief Base class of public extent class. + * + * This class should not be used directly. + * + * @tparam SizeType Size type + */ +template +class extent_base { + static_assert(std::is_integral_v, "SizeType bust be integral."); + + public: + using value_type = SizeType; ///< Extent value type +}; +} // namespace cuco::experimental::detail \ No newline at end of file diff --git a/include/cuco/detail/static_set/static_set.inl b/include/cuco/detail/static_set/static_set.inl index 1800c8910..c296dd46d 100644 --- a/include/cuco/detail/static_set/static_set.inl +++ b/include/cuco/detail/static_set/static_set.inl @@ -20,6 +20,7 @@ #include #include #include +#include #include #include #include @@ -39,18 +40,19 @@ namespace cuco { namespace experimental { template -constexpr static_set::static_set( +constexpr static_set::static_set( Extent capacity, empty_key empty_key_sentinel, KeyEqual pred, ProbingScheme const& probing_scheme, Allocator const& alloc, + [[maybe_unused]] Storage const& storage, cuda_stream_ref stream) : empty_key_sentinel_{empty_key_sentinel}, predicate_{pred}, @@ -62,15 +64,15 @@ constexpr static_set template -static_set::size_type -static_set::insert( +static_set::size_type +static_set::insert( InputIt first, InputIt last, cuda_stream_ref stream) { auto const num_keys = cuco::detail::distance(first, last); @@ -92,14 +94,14 @@ static_set::ins } template template -void static_set::insert_async( +void static_set::insert_async( InputIt first, InputIt last, cuda_stream_ref stream) { auto const num_keys = cuco::detail::distance(first, last); @@ -116,15 +118,15 @@ void static_set } template template -static_set::size_type -static_set::insert_if( +static_set::size_type +static_set::insert_if( InputIt first, InputIt last, StencilIt stencil, Predicate pred, cuda_stream_ref stream) { auto const num_keys = cuco::detail::distance(first, last); @@ -145,14 +147,14 @@ static_set::ins } template template -void static_set::insert_if_async( +void static_set::insert_if_async( InputIt first, InputIt last, StencilIt stencil, Predicate pred, cuda_stream_ref stream) { auto const num_keys = cuco::detail::distance(first, last); @@ -168,14 +170,14 @@ void static_set } template template -void static_set::contains( +void static_set::contains( InputIt first, InputIt last, OutputIt output_begin, cuda_stream_ref stream) const { contains_async(first, last, output_begin, stream); @@ -183,14 +185,14 @@ void static_set } template template -void static_set::contains_async( +void static_set::contains_async( InputIt first, InputIt last, OutputIt output_begin, cuda_stream_ref stream) const { auto const num_keys = cuco::detail::distance(first, last); @@ -212,14 +214,14 @@ void static_set } template template -void static_set::find( +void static_set::find( InputIt first, InputIt last, OutputIt output_begin, cuda_stream_ref stream) const { find_async(first, last, output_begin, stream); @@ -227,14 +229,14 @@ void static_set } template template -void static_set::find_async( +void static_set::find_async( InputIt first, InputIt last, OutputIt output_begin, cuda_stream_ref stream) const { auto const num_keys = cuco::detail::distance(first, last); @@ -250,14 +252,14 @@ void static_set } template template -OutputIt static_set::retrieve_all( +OutputIt static_set::retrieve_all( OutputIt output_begin, cuda_stream_ref stream) const { auto begin = thrust::make_transform_iterator(thrust::counting_iterator(0), @@ -296,14 +298,14 @@ OutputIt static_set -static_set::size_type -static_set::size( +static_set::size_type +static_set::size( cuda_stream_ref stream) const { auto counter = detail::counter_storage{allocator_}; @@ -323,47 +325,68 @@ static_set::siz } template constexpr auto -static_set::capacity() +static_set::capacity() const noexcept { return storage_.capacity(); } template -constexpr static_set::key_type -static_set::empty_key_sentinel() +constexpr static_set::key_type +static_set::empty_key_sentinel() const noexcept { return empty_key_sentinel_; } template template -auto static_set::ref( +auto static_set::ref( Operators...) const noexcept { static_assert(sizeof...(Operators), "No operators specified"); return ref_type{ cuco::empty_key(empty_key_sentinel_), predicate_, probing_scheme_, storage_.ref()}; } + +template +constexpr auto make_static_set(Args&&... args) +{ + // TODO don't repeat defaults + return static_set{ + detail::find_arg(std::forward(args)...), // required parameter + detail::find_arg>(std::forward(args)...), // required parameter + detail::find_arg::template is_equal_functor_t>( + std::forward(args)..., thrust::equal_to{}), + detail::find_arg(std::forward(args)..., + double_hashing<4, // CG size + cuco::murmurhash3_32, + cuco::murmurhash3_32>{}), + detail::find_arg(std::forward(args)..., + cuco::cuda_allocator{}), + detail::find_arg(std::forward(args)..., + cuco::experimental::aow_storage<1>{}), + detail::find_arg(std::forward(args)..., cuda_stream_ref{})}; +} + } // namespace experimental } // namespace cuco diff --git a/include/cuco/detail/traits.hpp b/include/cuco/detail/traits.hpp new file mode 100644 index 000000000..722b0f9c3 --- /dev/null +++ b/include/cuco/detail/traits.hpp @@ -0,0 +1,118 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include +// #include // FIXME check inheritance +#include + +#include +#include + +namespace cuco::experimental::detail { + +// Trait to check if T is a valid probing scheme type +template +struct is_probing_scheme : std::false_type { +}; + +template +struct is_probing_scheme< + T, + std::enable_if_t< + std::is_base_of_v, T>>> + : std::true_type { +}; + +template +inline constexpr bool is_probing_scheme_v = is_probing_scheme::value; + +template +struct is_storage : std::false_type { +}; + +template +struct is_storage< + T, + std::enable_if_t, T>>> + : std::true_type { +}; + +/* FIXME +std::enable_if_t, + T>>> : std::true_type {}; +*/ + +template +inline constexpr bool is_storage_v = is_storage::value; + +// Trait to check if T is a `cuco::extent` +template +struct is_extent : std::false_type { +}; + +template +struct is_extent< + T, + std::enable_if_t< + std::is_base_of_v, T>>> + : std::true_type { +}; + +template +inline constexpr bool is_extent_v = is_extent::value; + +// Trait to check if T is allocator-like +template > +struct is_allocator : std::false_type { +}; + +template +struct is_allocator().allocate(std::size_t{})), + decltype(std::declval().deallocate( + std::declval(), std::size_t{}))>> + : std::true_type { +}; + +template +inline constexpr bool is_allocator_v = is_allocator::value; + +template +struct key_equal_traits { + template > + struct is_equal_functor : std::false_type { + }; + + template + struct is_equal_functor< + T, + std::void_t>>> + : std::true_type { + }; + + template + using is_equal_functor_t = typename is_equal_functor::type; + + template + static constexpr bool is_equal_functor_v = is_equal_functor::value; +}; + +} // namespace cuco::experimental::detail \ No newline at end of file diff --git a/include/cuco/detail/utils.hpp b/include/cuco/detail/utils.hpp index 513ccd559..50b9034cd 100644 --- a/include/cuco/detail/utils.hpp +++ b/include/cuco/detail/utils.hpp @@ -105,4 +105,44 @@ constexpr ForwardIt lower_bound(ForwardIt first, ForwardIt last, const T& value) } } // namespace detail + +namespace experimental::detail { + +// TODO docs +template