bit.hpp
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2019-2024, NVIDIA CORPORATION.
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 
17 #pragma once
18 
19 #include <cudf/types.hpp>
20 
21 #include <cuda/std/climits>
22 
23 #include <cassert>
24 
30 namespace CUDF_EXPORT cudf {
31 namespace detail {
32 // @cond
33 // Work around a bug in NVRTC that fails to compile assert() in constexpr
34 // functions (fixed after CUDA 11.0)
35 #if defined __GNUC__
36 #define LIKELY(EXPR) __builtin_expect(!!(EXPR), 1)
37 #else
38 #define LIKELY(EXPR) (!!(EXPR))
39 #endif
40 
41 #ifdef NDEBUG
42 #define constexpr_assert(CHECK) static_cast<void>(0)
43 #else
44 #define constexpr_assert(CHECK) (LIKELY(CHECK) ? void(0) : [] { assert(!#CHECK); }())
45 #endif
46 // @endcond
47 
54 template <typename T>
55 constexpr CUDF_HOST_DEVICE inline std::size_t size_in_bits()
56 {
57  static_assert(CHAR_BIT == 8, "Size of a byte must be 8 bits.");
58  return sizeof(T) * CHAR_BIT;
59 }
60 } // namespace detail
61 
74 constexpr CUDF_HOST_DEVICE inline size_type word_index(size_type bit_index)
75 {
76  return bit_index / detail::size_in_bits<bitmask_type>();
77 }
78 
86 {
87  return bit_index % detail::size_in_bits<bitmask_type>();
88 }
89 
99 CUDF_HOST_DEVICE inline void set_bit_unsafe(bitmask_type* bitmask, size_type bit_index)
100 {
101  assert(nullptr != bitmask);
102  bitmask[word_index(bit_index)] |= (bitmask_type{1} << intra_word_index(bit_index));
103 }
104 
114 CUDF_HOST_DEVICE inline void clear_bit_unsafe(bitmask_type* bitmask, size_type bit_index)
115 {
116  assert(nullptr != bitmask);
117  bitmask[word_index(bit_index)] &= ~(bitmask_type{1} << intra_word_index(bit_index));
118 }
119 
128 CUDF_HOST_DEVICE inline bool bit_is_set(bitmask_type const* bitmask, size_type bit_index)
129 {
130  assert(nullptr != bitmask);
131  return bitmask[word_index(bit_index)] & (bitmask_type{1} << intra_word_index(bit_index));
132 }
133 
144 CUDF_HOST_DEVICE inline bool bit_value_or(bitmask_type const* bitmask,
145  size_type bit_index,
146  bool default_value)
147 {
148  return bitmask != nullptr ? bit_is_set(bitmask, bit_index) : default_value;
149 }
150 
160 {
161  constexpr_assert(0 <= n && n < static_cast<size_type>(detail::size_in_bits<bitmask_type>()));
162  return ((bitmask_type{1} << n) - 1);
163 }
164 
174 {
175  constexpr size_type word_size{detail::size_in_bits<bitmask_type>()};
176  constexpr_assert(0 <= n && n < word_size);
177  return ~((bitmask_type{1} << (word_size - n)) - 1);
178 }
179 
180 #ifdef __CUDACC__
181 
195 __device__ inline void set_bit(bitmask_type* bitmask, size_type bit_index)
196 {
197  assert(nullptr != bitmask);
198  atomicOr(&bitmask[word_index(bit_index)], (bitmask_type{1} << intra_word_index(bit_index)));
199 }
200 
213 __device__ inline void clear_bit(bitmask_type* bitmask, size_type bit_index)
214 {
215  assert(nullptr != bitmask);
216  atomicAnd(&bitmask[word_index(bit_index)], ~(bitmask_type{1} << intra_word_index(bit_index)));
217 }
218 #endif // end of group
220 } // namespace CUDF_EXPORT cudf
constexpr CUDF_HOST_DEVICE std::size_t size_in_bits()
Returns the number of bits the given type can hold.
Definition: bit.hpp:55
CUDF_HOST_DEVICE void set_bit_unsafe(bitmask_type *bitmask, size_type bit_index)
Sets the specified bit to 1
Definition: bit.hpp:99
constexpr CUDF_HOST_DEVICE size_type intra_word_index(size_type bit_index)
Returns the position within a word of the specified bit.
Definition: bit.hpp:85
constexpr CUDF_HOST_DEVICE size_type word_index(size_type bit_index)
Returns the index of the word containing the specified bit.
Definition: bit.hpp:74
constexpr CUDF_HOST_DEVICE bitmask_type set_least_significant_bits(size_type n)
Returns a bitmask word with the n least significant bits set.
Definition: bit.hpp:159
CUDF_HOST_DEVICE bool bit_value_or(bitmask_type const *bitmask, size_type bit_index, bool default_value)
optional-like interface to check if a specified bit of a bitmask is set.
Definition: bit.hpp:144
CUDF_HOST_DEVICE bool bit_is_set(bitmask_type const *bitmask, size_type bit_index)
Indicates whether the specified bit is set to 1
Definition: bit.hpp:128
constexpr CUDF_HOST_DEVICE bitmask_type set_most_significant_bits(size_type n)
Returns a bitmask word with the n most significant bits set.
Definition: bit.hpp:173
CUDF_HOST_DEVICE void clear_bit_unsafe(bitmask_type *bitmask, size_type bit_index)
Sets the specified bit to 0
Definition: bit.hpp:114
int32_t size_type
Row index type for columns and tables.
Definition: types.hpp:95
uint32_t bitmask_type
Bitmask type stored as 32-bit unsigned integer.
Definition: types.hpp:96
cuDF interfaces
Definition: aggregation.hpp:35
Type declarations for libcudf.
#define CUDF_HOST_DEVICE
Indicates that the function or method is usable on host and device.
Definition: types.hpp:32