Skip to content

Commit ab4bc9d

Browse files
tpntpn
committed
Add CUDA threshold-2 bitset helper
1 parent 9e713a2 commit ab4bc9d

6 files changed

Lines changed: 643 additions & 0 deletions

File tree

‎src/PerfectHashCuda/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@ set(target PerfectHashCuda)
66
################################################################################
77
set(Header_Files
88
"Graph.cuh"
9+
"Threshold2Bitset.cuh"
910
)
1011
source_group("Header Files" FILES ${Header_Files})
1112

‎src/PerfectHashCuda/Graph.cuh‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,8 @@ Abstract:
2121
#include <cooperative_groups.h>
2222
#include <cooperative_groups/reduce.h>
2323

24+
#include "Threshold2Bitset.cuh"
25+
2426
namespace cg = cooperative_groups;
2527

2628
#define MAX_NUMBER_OF_SEEDS 8

‎src/PerfectHashCuda/PerfectHashCuda.vcxproj‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -114,6 +114,7 @@
114114
<ClInclude Include="..\..\include\PerfectHash\PerfectHash.h" />
115115
<ClInclude Include="..\..\include\PerfectHash\PerfectHashCuda.h" />
116116
<ClInclude Include="..\..\include\PerfectHash\PerfectHashErrors.h" />
117+
<ClInclude Include="Threshold2Bitset.cuh" />
117118
</ItemGroup>
118119
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
119120
<ImportGroup Label="ExtensionTargets">

‎src/PerfectHashCuda/Threshold2Bitset.cuh‎

Lines changed: 306 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,306 @@
1+
/*++
2+
3+
Copyright (c) 2026 Trent Nelson <trent@trent.me>
4+
5+
Module Name:
6+
7+
Threshold2Bitset.cuh
8+
9+
Abstract:
10+
11+
CUDA helpers for dense threshold-2 histograms represented by two bitsets.
12+
13+
--*/
14+
15+
#pragma once
16+
17+
#include <cuda_runtime.h>
18+
#include <stddef.h>
19+
#include <stdint.h>
20+
#include <type_traits>
21+
22+
enum THRESHOLD2_BITSET_STATE : uint8_t {
23+
Threshold2BitsetZero = 0,
24+
Threshold2BitsetOne = 1,
25+
Threshold2BitsetMany = 2,
26+
};
27+
28+
template <typename WordType>
29+
struct Threshold2BitsetWordTraits
30+
{
31+
static_assert(
32+
std::is_unsigned<WordType>::value,
33+
"WordType must be an unsigned integer type"
34+
);
35+
36+
static_assert(
37+
sizeof(WordType) == sizeof(uint32_t) ||
38+
sizeof(WordType) == sizeof(uint64_t),
39+
"WordType must be a 32-bit or 64-bit unsigned integer type"
40+
);
41+
42+
static constexpr uint32_t BitsPerWord = sizeof(WordType) * 8u;
43+
static constexpr uint32_t WordIndexShift = BitsPerWord == 32u ? 5u : 6u;
44+
static constexpr uint32_t WordBitMask = BitsPerWord - 1u;
45+
};
46+
47+
template <typename OffsetType>
48+
using Threshold2BitsetUnsignedOffsetType =
49+
typename std::make_unsigned<
50+
typename std::remove_cv<OffsetType>::type
51+
>::type;
52+
53+
template <typename OffsetType>
54+
__host__
55+
__device__
56+
constexpr
57+
Threshold2BitsetUnsignedOffsetType<OffsetType>
58+
Threshold2BitsetNormalizeOffset(
59+
OffsetType Offset
60+
)
61+
{
62+
static_assert(
63+
std::is_integral<OffsetType>::value,
64+
"OffsetType must be an integral type"
65+
);
66+
67+
static_assert(
68+
!std::is_same<typename std::remove_cv<OffsetType>::type, bool>::value,
69+
"OffsetType must not be bool"
70+
);
71+
72+
return static_cast<Threshold2BitsetUnsignedOffsetType<OffsetType>>(Offset);
73+
}
74+
75+
template <typename WordType,
76+
typename CountType>
77+
__host__
78+
__device__
79+
constexpr
80+
size_t
81+
Threshold2BitsetWordCount(
82+
CountType NumberOfBins
83+
)
84+
{
85+
using Traits = Threshold2BitsetWordTraits<WordType>;
86+
const auto Count = Threshold2BitsetNormalizeOffset(NumberOfBins);
87+
return static_cast<size_t>(
88+
(Count + Traits::BitsPerWord - 1u) >> Traits::WordIndexShift
89+
);
90+
}
91+
92+
template <typename WordType,
93+
typename OffsetType>
94+
__host__
95+
__device__
96+
constexpr
97+
size_t
98+
Threshold2BitsetWordIndex(
99+
OffsetType Offset
100+
)
101+
{
102+
using Traits = Threshold2BitsetWordTraits<WordType>;
103+
const auto NormalizedOffset = Threshold2BitsetNormalizeOffset(Offset);
104+
return static_cast<size_t>(
105+
NormalizedOffset >> Traits::WordIndexShift
106+
);
107+
}
108+
109+
template <typename WordType,
110+
typename OffsetType>
111+
__host__
112+
__device__
113+
constexpr
114+
WordType
115+
Threshold2BitsetMask(
116+
OffsetType Offset
117+
)
118+
{
119+
using Traits = Threshold2BitsetWordTraits<WordType>;
120+
const auto NormalizedOffset = Threshold2BitsetNormalizeOffset(Offset);
121+
const uint32_t Shift = static_cast<uint32_t>(
122+
NormalizedOffset & Traits::WordBitMask
123+
);
124+
return static_cast<WordType>(static_cast<WordType>(1u) << Shift);
125+
}
126+
127+
template <typename WordType,
128+
typename OffsetType>
129+
__host__
130+
__device__
131+
THRESHOLD2_BITSET_STATE
132+
Threshold2BitsetReadState(
133+
const WordType *Seen,
134+
const WordType *Many,
135+
OffsetType Offset
136+
)
137+
{
138+
const size_t WordIndex = Threshold2BitsetWordIndex<WordType>(Offset);
139+
const WordType Mask = Threshold2BitsetMask<WordType>(Offset);
140+
141+
if ((Many[WordIndex] & Mask) != 0) {
142+
return Threshold2BitsetMany;
143+
} else if ((Seen[WordIndex] & Mask) != 0) {
144+
return Threshold2BitsetOne;
145+
} else {
146+
return Threshold2BitsetZero;
147+
}
148+
}
149+
150+
template <typename WordType>
151+
__device__
152+
__forceinline__
153+
WordType
154+
Threshold2BitsetAtomicOr(
155+
WordType *Address,
156+
WordType Mask
157+
)
158+
{
159+
static_assert(
160+
sizeof(WordType) == sizeof(uint32_t) ||
161+
sizeof(WordType) == sizeof(uint64_t),
162+
"WordType must be a 32-bit or 64-bit unsigned integer type"
163+
);
164+
165+
if constexpr (sizeof(WordType) == sizeof(uint32_t)) {
166+
return static_cast<WordType>(
167+
atomicOr(
168+
reinterpret_cast<unsigned int *>(Address),
169+
static_cast<unsigned int>(Mask)
170+
)
171+
);
172+
} else {
173+
return static_cast<WordType>(
174+
atomicOr(
175+
reinterpret_cast<unsigned long long *>(Address),
176+
static_cast<unsigned long long>(Mask)
177+
)
178+
);
179+
}
180+
}
181+
182+
template <typename WordType,
183+
bool SetSeenWhenMany>
184+
__device__
185+
__forceinline__
186+
THRESHOLD2_BITSET_STATE
187+
Threshold2BitsetMarkMany(
188+
WordType *Seen,
189+
WordType *Many,
190+
WordType Mask
191+
)
192+
{
193+
if constexpr (SetSeenWhenMany) {
194+
Threshold2BitsetAtomicOr(Seen, Mask);
195+
}
196+
197+
Threshold2BitsetAtomicOr(Many, Mask);
198+
return Threshold2BitsetMany;
199+
}
200+
201+
template <typename WordType,
202+
bool SetSeenWhenMany>
203+
__device__
204+
__forceinline__
205+
THRESHOLD2_BITSET_STATE
206+
Threshold2BitsetMarkOne(
207+
WordType *Seen,
208+
WordType *Many,
209+
WordType Mask
210+
)
211+
{
212+
const WordType Previous = Threshold2BitsetAtomicOr(Seen, Mask);
213+
if ((Previous & Mask) != 0) {
214+
return Threshold2BitsetMarkMany<WordType, SetSeenWhenMany>(
215+
Seen,
216+
Many,
217+
Mask
218+
);
219+
}
220+
221+
return Threshold2BitsetOne;
222+
}
223+
224+
__device__
225+
__forceinline__
226+
uint32_t
227+
Threshold2BitsetLaneId()
228+
{
229+
uint32_t LaneId;
230+
asm volatile("mov.u32 %0, %%laneid;" : "=r"(LaneId));
231+
return LaneId;
232+
}
233+
234+
template <typename OffsetType,
235+
typename WordType = uint32_t,
236+
bool SetSeenWhenMany = false,
237+
bool SkipIfAlreadyMany = true>
238+
__device__
239+
__forceinline__
240+
THRESHOLD2_BITSET_STATE
241+
AtomicAggIncThreshold2Bitset(
242+
WordType *Seen,
243+
WordType *Many,
244+
OffsetType Offset
245+
)
246+
{
247+
static_assert(
248+
sizeof(Threshold2BitsetUnsignedOffsetType<OffsetType>) <= sizeof(uint64_t),
249+
"OffsetType must fit in 64 bits"
250+
);
251+
252+
const size_t WordIndex = Threshold2BitsetWordIndex<WordType>(Offset);
253+
const WordType Mask = Threshold2BitsetMask<WordType>(Offset);
254+
WordType *SeenAddress = Seen + WordIndex;
255+
WordType *ManyAddress = Many + WordIndex;
256+
257+
if constexpr (SkipIfAlreadyMany) {
258+
if ((*ManyAddress & Mask) != 0) {
259+
return Threshold2BitsetMany;
260+
}
261+
}
262+
263+
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 700)
264+
return Threshold2BitsetMarkOne<WordType, SetSeenWhenMany>(
265+
SeenAddress,
266+
ManyAddress,
267+
Mask
268+
);
269+
#else
270+
using UnsignedOffsetType = Threshold2BitsetUnsignedOffsetType<OffsetType>;
271+
const UnsignedOffsetType NormalizedOffset =
272+
Threshold2BitsetNormalizeOffset(Offset);
273+
274+
if constexpr (sizeof(UnsignedOffsetType) <= sizeof(uint32_t)) {
275+
const uint32_t ActiveMask = __activemask();
276+
const uint32_t MatchValue = static_cast<uint32_t>(NormalizedOffset);
277+
const uint32_t PeerMask = __match_any_sync(ActiveMask, MatchValue);
278+
const uint32_t LaneId = Threshold2BitsetLaneId();
279+
const uint32_t LeaderLane = static_cast<uint32_t>(__ffs(PeerMask) - 1);
280+
const uint32_t PeerCount = static_cast<uint32_t>(__popc(PeerMask));
281+
282+
if (PeerCount > 1u) {
283+
if (LaneId == LeaderLane) {
284+
Threshold2BitsetMarkMany<WordType, SetSeenWhenMany>(
285+
SeenAddress,
286+
ManyAddress,
287+
Mask
288+
);
289+
}
290+
return Threshold2BitsetMany;
291+
} else {
292+
return Threshold2BitsetMarkOne<WordType, SetSeenWhenMany>(
293+
SeenAddress,
294+
ManyAddress,
295+
Mask
296+
);
297+
}
298+
} else {
299+
return Threshold2BitsetMarkOne<WordType, SetSeenWhenMany>(
300+
SeenAddress,
301+
ManyAddress,
302+
Mask
303+
);
304+
}
305+
#endif
306+
}

‎tests/CMakeLists.txt‎

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -102,6 +102,38 @@ add_test(
102102
)
103103
set_tests_properties(perfecthash.fast.unit PROPERTIES LABELS fast)
104104

105+
if(PERFECTHASH_HAS_CUDA)
106+
add_executable(perfecthash_cuda_threshold2_bitset_tests
107+
PerfectHashCudaThreshold2BitsetTests.cu
108+
)
109+
110+
perfecthash_apply_cuda_settings(perfecthash_cuda_threshold2_bitset_tests)
111+
112+
if(CUDAToolkit_FOUND)
113+
target_link_libraries(
114+
perfecthash_cuda_threshold2_bitset_tests
115+
PRIVATE
116+
CUDA::cudart
117+
)
118+
endif()
119+
120+
set_target_properties(perfecthash_cuda_threshold2_bitset_tests PROPERTIES
121+
CUDA_STANDARD 17
122+
CUDA_STANDARD_REQUIRED YES
123+
)
124+
125+
add_test(
126+
NAME perfecthash.cuda.threshold2_bitset
127+
COMMAND $<TARGET_FILE:perfecthash_cuda_threshold2_bitset_tests>
128+
)
129+
130+
set_tests_properties(
131+
perfecthash.cuda.threshold2_bitset
132+
PROPERTIES
133+
LABELS "cuda"
134+
)
135+
endif()
136+
105137
set(TEST_KEYS_FILE "${PERFECTHASH_ROOT_DIR}/keys/HologramWorld-31016.keys")
106138
set(TEST_KEYS_DIR "${PERFECTHASH_ROOT_DIR}/keys")
107139
set(TEST_OUTPUT_DIR "${CMAKE_CURRENT_BINARY_DIR}/cli-out")

0 commit comments

Comments
 (0)