Skip to content

Commit 1cbf8bb

Browse files
sleeepyjacksleeepyjack
committed
Address review comments
1 parent fa595c1 commit 1cbf8bb

7 files changed

Lines changed: 43 additions & 28 deletions

File tree

‎benchmarks/benchmark_defaults.hpp‎

Lines changed: 19 additions & 6 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,24 @@
2526

2627
namespace cuco::benchmark::defaults {
2728

28-
using KEY_TYPE_RANGE = nvbench::type_list<nvbench::int32_t, nvbench::int64_t, __int128_t>;
29-
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>>; //,
29+
using KEY_TYPE_RANGE = nvbench::type_list<nvbench::int32_t,
30+
nvbench::int64_t
31+
#if defined(CUCO_HAS_128BIT_ATOMICS)
32+
,
33+
__int128_t
34+
#endif
35+
>;
36+
using VALUE_TYPE_RANGE = nvbench::type_list<nvbench::int32_t,
37+
nvbench::int64_t
38+
#if defined(CUCO_HAS_128BIT_ATOMICS)
39+
,
40+
__int128_t
41+
#endif
42+
>;
43+
using HASH_RANGE = nvbench::type_list<cuco::identity_hash<char>,
44+
cuco::xxhash_32<char>,
45+
cuco::xxhash_64<char>,
46+
cuco::murmurhash3_32<char>>; //,
3447
// cuco::murmurhash3_x86_128<char>,
3548
// cuco::murmurhash3_x64_128<char>>; // TODO handle tuple-like hash value
3649

‎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/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: 11 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,16 +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
18591859
}
1860+
// Unreachable: the class-level max_key_size static_assert rejects 16-byte
1861+
// keys in key-only containers when CUCO_HAS_128BIT_ATOMICS is not defined.
1862+
__builtin_unreachable();
18601863
}
18611864

18621865
/**
@@ -1889,12 +1892,12 @@ class open_addressing_ref_impl {
18891892
return packed_cas(address, expected, desired);
18901893
}
18911894
#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.");
1895+
else if constexpr (has_payload) {
18961896
return cas_dependent_write(address, expected, desired);
18971897
}
1898+
// Unreachable: the class-level max_key_size static_assert rejects 16-byte
1899+
// keys in key-only containers when CUCO_HAS_128BIT_ATOMICS is not defined.
1900+
__builtin_unreachable();
18981901
}
18991902

19001903
/**

‎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

0 commit comments

Comments
 (0)