bit.hpp

Go to the documentation of this file.

/*
 * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: Apache-2.0
 */
 
#pragma once
 
#include <cudf/types.hpp>
 
#include <cuda/std/climits>
 
#include <cassert>
 
namespace CUDF_EXPORT cudf {
namespace detail {
 
template <typename T>
constexpr CUDF_HOST_DEVICE inline std::size_t size_in_bits()
{
  static_assert(CHAR_BIT == 8, "Size of a byte must be 8 bits.");
  return sizeof(T) * CHAR_BIT;
}
}  // namespace detail
 
constexpr CUDF_HOST_DEVICE inline size_type word_index(size_type bit_index)
{
  return bit_index / detail::size_in_bits<bitmask_type>();
}
 
constexpr CUDF_HOST_DEVICE inline size_type intra_word_index(size_type bit_index)
{
  return bit_index % detail::size_in_bits<bitmask_type>();
}
 
CUDF_HOST_DEVICE inline void set_bit_unsafe(bitmask_type* bitmask, size_type bit_index)
{
  assert(nullptr != bitmask);
  bitmask[word_index(bit_index)] |= (bitmask_type{1} << intra_word_index(bit_index));
}
 
CUDF_HOST_DEVICE inline void clear_bit_unsafe(bitmask_type* bitmask, size_type bit_index)
{
  assert(nullptr != bitmask);
  bitmask[word_index(bit_index)] &= ~(bitmask_type{1} << intra_word_index(bit_index));
}
 
CUDF_HOST_DEVICE inline bool bit_is_set(bitmask_type const* bitmask, size_type bit_index)
{
  assert(nullptr != bitmask);
  return bitmask[word_index(bit_index)] & (bitmask_type{1} << intra_word_index(bit_index));
}
 
CUDF_HOST_DEVICE inline bool bit_value_or(bitmask_type const* bitmask,
                                          size_type bit_index,
                                          bool default_value)
{
  return bitmask != nullptr ? bit_is_set(bitmask, bit_index) : default_value;
}
 
constexpr CUDF_HOST_DEVICE inline bitmask_type set_least_significant_bits(size_type n)
{
  assert(0 <= n && n < static_cast<size_type>(detail::size_in_bits<bitmask_type>()));
  return ((bitmask_type{1} << n) - 1);
}
 
constexpr CUDF_HOST_DEVICE inline bitmask_type set_most_significant_bits(size_type n)
{
  constexpr size_type word_size{detail::size_in_bits<bitmask_type>()};
  assert(0 <= n && n < word_size);
  return ~((bitmask_type{1} << (word_size - n)) - 1);
}
 
#ifdef __CUDACC__
 
__device__ inline void set_bit(bitmask_type* bitmask, size_type bit_index)
{
  assert(nullptr != bitmask);
  atomicOr(&bitmask[word_index(bit_index)], (bitmask_type{1} << intra_word_index(bit_index)));
}
 
__device__ inline void clear_bit(bitmask_type* bitmask, size_type bit_index)
{
  assert(nullptr != bitmask);
  atomicAnd(&bitmask[word_index(bit_index)], ~(bitmask_type{1} << intra_word_index(bit_index)));
}
#endif  // end of group
}  // namespace CUDF_EXPORT cudf