Skip to content

Commit 6a2a1d2

Browse files
sleeepyjacksleeepyjack
committed
Address review comments
1 parent fa595c1 commit 6a2a1d2

13 files changed

Lines changed: 49 additions & 41 deletions

File tree

‎benchmarks/benchmark_defaults.hpp‎

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@
1616

1717
#pragma once
1818

19+
#include <cuco/detail/__config>
1920
#include <cuco/hash_functions.cuh>
2021

2122
#include <nvbench/nvbench.cuh>
@@ -25,12 +26,17 @@
2526

2627
namespace cuco::benchmark::defaults {
2728

29+
#if defined(CUCO_HAS_128BIT_ATOMICS)
2830
using KEY_TYPE_RANGE = nvbench::type_list<nvbench::int32_t, nvbench::int64_t, __int128_t>;
2931
using VALUE_TYPE_RANGE = nvbench::type_list<nvbench::int32_t, nvbench::int64_t, __int128_t>;
30-
using HASH_RANGE = nvbench::type_list<cuco::identity_hash<char>,
31-
cuco::xxhash_32<char>,
32-
cuco::xxhash_64<char>,
33-
cuco::murmurhash3_32<char>>; //,
32+
#else
33+
using KEY_TYPE_RANGE = nvbench::type_list<nvbench::int32_t, nvbench::int64_t>;
34+
using VALUE_TYPE_RANGE = nvbench::type_list<nvbench::int32_t, nvbench::int64_t>;
35+
#endif
36+
using HASH_RANGE = nvbench::type_list<cuco::identity_hash<char>,
37+
cuco::xxhash_32<char>,
38+
cuco::xxhash_64<char>,
39+
cuco::murmurhash3_32<char>>; //,
3440
// cuco::murmurhash3_x86_128<char>,
3541
// cuco::murmurhash3_x64_128<char>>; // TODO handle tuple-like hash value
3642

‎benchmarks/benchmark_utils.hpp‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@
1616

1717
#pragma once
1818

19+
#include <cuco/detail/__config>
1920
#include <cuco/detail/error.hpp>
2021
#include <cuco/utility/key_generator.cuh>
2122

@@ -93,4 +94,6 @@ NVBENCH_DECLARE_TYPE_STRINGS(cuco::utility::distribution::gaussian,
9394
"GAUSSIAN",
9495
"distribution::gaussian");
9596

97+
#if defined(CUCO_HAS_128BIT_ATOMICS)
9698
NVBENCH_DECLARE_TYPE_STRINGS(__int128_t, "I128", "__int128_t");
99+
#endif

‎include/cuco/detail/__config‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -73,9 +73,12 @@ inline constexpr std::size_t max_key_size =
7373
#endif
7474

7575
/// Maximum supported payload/mapped type size (in bytes) for open-addressing containers.
76+
/// Tied to `max_key_size`: a slot stores at most a key plus an equally-sized payload.
7677
inline constexpr std::size_t max_payload_size = max_key_size;
7778

7879
/// Maximum supported slot size (in bytes) for open-addressing containers.
80+
/// Tied to `max_key_size`: a slot stores at most a key plus an equally-sized payload
81+
/// (i.e., `sizeof(pair<Key, Value>) <= 2 * max_key_size`).
7982
inline constexpr std::size_t max_slot_size = 2 * max_key_size;
8083

8184
/// Checks if the given size is a valid mapped_type size for packed CAS operations.

‎include/cuco/detail/bitwise_compare.cuh‎

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@
1616

1717
#pragma once
1818

19+
#include <cuco/detail/__config>
1920
#include <cuco/utility/traits.hpp>
2021

2122
#include <cuda/functional>
@@ -63,15 +64,16 @@ struct bitwise_compare_impl<8> {
6364
}
6465
};
6566

67+
#ifdef CUCO_HAS_INT128
6668
template <>
6769
struct bitwise_compare_impl<16> {
6870
__host__ __device__ inline static bool compare(char const* lhs, char const* rhs)
6971
{
70-
return *reinterpret_cast<uint64_t const*>(lhs) == *reinterpret_cast<uint64_t const*>(rhs) and
71-
*reinterpret_cast<uint64_t const*>(lhs + 8) ==
72-
*reinterpret_cast<uint64_t const*>(rhs + 8);
72+
return *reinterpret_cast<unsigned __int128 const*>(lhs) ==
73+
*reinterpret_cast<unsigned __int128 const*>(rhs);
7374
}
7475
};
76+
#endif
7577

7678
/**
7779
* @brief Gives value to use as alignment for a type that is at least the

‎include/cuco/detail/open_addressing/open_addressing_impl.cuh‎

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -69,12 +69,6 @@ template <class Key,
6969
class Allocator,
7070
class Storage>
7171
class open_addressing_impl {
72-
static_assert(sizeof(Key) <= cuco::detail::max_key_size,
73-
"Key size exceeds the maximum supported size (8 bytes, or 16 with sm_90+).");
74-
75-
static_assert(sizeof(Value) <= cuco::detail::max_slot_size,
76-
"Slot size exceeds the maximum supported size (16 bytes, or 32 with sm_90+).");
77-
7872
static_assert(
7973
cuco::is_bitwise_comparable_v<Key>,
8074
"Key type must have unique object representations or have been explicitly declared as safe for "

‎include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh‎

Lines changed: 13 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -92,6 +92,9 @@ class open_addressing_ref_impl {
9292
static_assert(sizeof(Key) <= cuco::detail::max_key_size,
9393
"Key size exceeds the maximum supported size (8 bytes, or 16 with sm_90+).");
9494

95+
static_assert(sizeof(typename StorageRef::value_type) <= cuco::detail::max_slot_size,
96+
"Slot size exceeds the maximum supported size (16 bytes, or 32 with sm_90+).");
97+
9598
static_assert(
9699
cuco::is_bitwise_comparable_v<Key>,
97100
"Key type must have unique object representations or have been explicitly declared as safe for "
@@ -1847,15 +1850,16 @@ class open_addressing_ref_impl {
18471850
return packed_cas(address, expected, desired);
18481851
}
18491852
#endif
1850-
else {
1851-
static_assert(
1852-
has_payload,
1853-
"16-byte key types in key-only containers require sm_90+ for 128-bit atomic CAS.");
1853+
else if constexpr (has_payload) {
18541854
#if (__CUDA_ARCH__ < 700)
18551855
return cas_dependent_write(address, expected, desired);
18561856
#else
18571857
return back_to_back_cas(address, expected, desired);
18581858
#endif
1859+
} else {
1860+
static_assert(cuco::dependent_false<Value>,
1861+
"No valid atomic CAS path: 16-byte key in a key-only container must be "
1862+
"packable (have unique object representations) and target sm_90+.");
18591863
}
18601864
}
18611865

@@ -1889,11 +1893,12 @@ class open_addressing_ref_impl {
18891893
return packed_cas(address, expected, desired);
18901894
}
18911895
#endif
1892-
else {
1893-
static_assert(
1894-
has_payload,
1895-
"16-byte key types in key-only containers require sm_90+ for 128-bit atomic CAS.");
1896+
else if constexpr (has_payload) {
18961897
return cas_dependent_write(address, expected, desired);
1898+
} else {
1899+
static_assert(cuco::dependent_false<Value>,
1900+
"No valid atomic CAS path: 16-byte key in a key-only container must be "
1901+
"packable (have unique object representations) and target sm_90+.");
18971902
}
18981903
}
18991904

‎include/cuco/detail/pair/pair.inl‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -23,14 +23,14 @@ namespace cuco {
2323

2424
template <typename First, typename Second>
2525
__host__ __device__ constexpr pair<First, Second>::pair(First const& f, Second const& s)
26-
: first(f), second(s)
26+
: first{f}, second{s}
2727
{
2828
}
2929

3030
template <typename First, typename Second>
3131
template <typename F, typename S>
3232
__host__ __device__ constexpr pair<First, Second>::pair(pair<F, S> const& p)
33-
: first(p.first), second(p.second)
33+
: first{p.first}, second{p.second}
3434
{
3535
}
3636

‎include/cuco/detail/static_map/static_map.inl‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -55,7 +55,7 @@ constexpr static_map<Key, T, Extent, Scope, KeyEqual, ProbingScheme, Allocator,
5555
probing_scheme,
5656
alloc,
5757
stream)},
58-
empty_value_sentinel_(empty_value_sentinel)
58+
empty_value_sentinel_{empty_value_sentinel}
5959
{
6060
}
6161

@@ -85,7 +85,7 @@ constexpr static_map<Key, T, Extent, Scope, KeyEqual, ProbingScheme, Allocator,
8585
probing_scheme,
8686
alloc,
8787
stream)},
88-
empty_value_sentinel_(empty_value_sentinel)
88+
empty_value_sentinel_{empty_value_sentinel}
8989
{
9090
}
9191

@@ -115,7 +115,7 @@ constexpr static_map<Key, T, Extent, Scope, KeyEqual, ProbingScheme, Allocator,
115115
probing_scheme,
116116
alloc,
117117
stream)},
118-
empty_value_sentinel_(empty_value_sentinel)
118+
empty_value_sentinel_{empty_value_sentinel}
119119
{
120120
}
121121

‎tests/static_map/insert_or_apply_test.cu‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -164,10 +164,9 @@ TEMPLATE_TEST_CASE_SIG(
164164
(int64_t, int64_t, cuco::test::probe_sequence::linear_probing, 2)
165165
#if defined(CUCO_HAS_128BIT_ATOMICS)
166166
,
167-
(__int128_t, __int128_t, cuco::test::probe_sequence::double_hashing, 1),
168167
(__int128_t, __int128_t, cuco::test::probe_sequence::double_hashing, 2),
169-
(__int128_t, __int128_t, cuco::test::probe_sequence::linear_probing, 1),
170-
(__int128_t, __int128_t, cuco::test::probe_sequence::linear_probing, 2)
168+
(__int128_t, int64_t, cuco::test::probe_sequence::double_hashing, 1),
169+
(int32_t, __int128_t, cuco::test::probe_sequence::linear_probing, 2)
171170
#endif
172171
)
173172
{

‎tests/static_map/insert_or_assign_test.cu‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -96,10 +96,9 @@ TEMPLATE_TEST_CASE_SIG(
9696
(int64_t, int64_t, cuco::test::probe_sequence::linear_probing, 2)
9797
#if defined(CUCO_HAS_128BIT_ATOMICS)
9898
,
99-
(__int128_t, __int128_t, cuco::test::probe_sequence::double_hashing, 1),
10099
(__int128_t, __int128_t, cuco::test::probe_sequence::double_hashing, 2),
101-
(__int128_t, __int128_t, cuco::test::probe_sequence::linear_probing, 1),
102-
(__int128_t, __int128_t, cuco::test::probe_sequence::linear_probing, 2)
100+
(__int128_t, int64_t, cuco::test::probe_sequence::double_hashing, 1),
101+
(int32_t, __int128_t, cuco::test::probe_sequence::linear_probing, 2)
103102
#endif
104103
)
105104
{

0 commit comments

Comments
 (0)