-
Notifications
You must be signed in to change notification settings - Fork 0
Expand file tree
/
Copy pathBitSet.h
More file actions
53 lines (44 loc) · 1.59 KB
/
BitSet.h
File metadata and controls
53 lines (44 loc) · 1.59 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
/*
Copyright [2024] [Yao Yao]
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.
*/
//
// Created by yao on 10/21/19.
//
#pragma once
#include <cstddef>
#include <type_traits>
#include <cstdint>
#include <cuda_runtime_api.h>
#include <cuda_utils.h>
template <size_t nbBits>
struct BitSet
{
using StorageElemType = std::conditional_t<nbBits <= 8, uint8_t, std::conditional_t<nbBits <= 16, uint16_t, uint32_t >>;
static constexpr uint32_t bitsPerElem = 8u * sizeof(StorageElemType);
__forceinline__ __device__
void set(uint32_t idx, bool val = true) {
if (val) {
StorageElemType& elem = data[idx / bitsPerElem];
elem |= StorageElemType(1u << idx % bitsPerElem);
} else{
reset(idx);
}
}
__forceinline__ __device__
void reset(uint32_t idx) {
StorageElemType& elem = data[idx / bitsPerElem];
elem &= StorageElemType(~(1u << idx % bitsPerElem));
}
__forceinline__ __device__
bool test(uint32_t idx) const { return (data[idx / bitsPerElem] & (1u << (idx % bitsPerElem))) != 0; }
StorageElemType data[divUp(nbBits, bitsPerElem)];
};