bit.hpp

转到此文件的文档。

/*
 * Copyright (c) 2019-2024, NVIDIA CORPORATION.
 *
 * 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
 *
 * https://apache.ac.cn/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.
 */

#pragma once

#include <cudf/types.hpp>

#include <cuda/std/climits>

#include <cassert>

namespace CUDF_EXPORT cudf {
namespace detail {
// @cond
// 解决 NVRTC 中的一个 bug，该 bug 导致 assert() 在 constexpr 函数中编译失败（在 CUDA 11.0 后已修复）
// Work around a bug in NVRTC that fails to compile assert() in constexpr
// functions (fixed after CUDA 11.0)
#if defined __GNUC__
#define LIKELY(EXPR) __builtin_expect(!!(EXPR), 1)
#else
#define LIKELY(EXPR) (!!(EXPR))
#endif

#ifdef NDEBUG
#define constexpr_assert(CHECK) static_cast<void>(0)
#else
#define constexpr_assert(CHECK) (LIKELY(CHECK) ? void(0) : [] { assert(!#CHECK); }())
#endif
// @endcond
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)
{
 constexpr_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>()};
 constexpr_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 // 组结束
} // namespace CUDF_EXPORT cudf