negabinary.h

Go to the documentation of this file.

#pragma once
 
#include <cstdint>
#include <type_traits>
 
// nvcc defines __CUDACC__; hipcc/amdclang++ defines __HIPCC__ instead. Without
// the second arm these helpers stay host-only under HIP and every __global__
// caller fails to resolve them.
#if defined(__CUDACC__) || defined(__HIPCC__)
#  define FZ_HOST_DEVICE __host__ __device__
#else
#  define FZ_HOST_DEVICE
#endif
 
namespace fz {
 
template <typename T>
struct Negabinary {
    static_assert(std::is_integral<T>::value && std::is_signed<T>::value,
                  "fz::Negabinary<T>: T must be a signed integer type "
                  "(int8_t, int16_t, int32_t, or int64_t).");
 
    using SInt = T;
    using UInt = typename std::make_unsigned<T>::type;
 
    // Alternating-bit mask 0xAA…A for the UInt width.
    // Truncating the 64-bit constant to UInt gives:
    //   uint8_t  → 0xAA
    //   uint16_t → 0xAAAA
    //   uint32_t → 0xAAAAAAAA
    //   uint64_t → 0xAAAAAAAAAAAAAAAA
    static constexpr UInt MASK =
        static_cast<UInt>(static_cast<uint64_t>(0xAAAAAAAAAAAAAAAAULL));
 
    FZ_HOST_DEVICE static constexpr UInt encode(SInt n) noexcept {
        return (static_cast<UInt>(n) + MASK) ^ MASK;
    }
 
    FZ_HOST_DEVICE static constexpr SInt decode(UInt u) noexcept {
        return static_cast<SInt>((u ^ MASK) - MASK);
    }
};
 
// ---------------------------------------------------------------------------
// Convenience aliases for the four standard widths
// ---------------------------------------------------------------------------
using Negabinary8  = Negabinary<int8_t>;
using Negabinary16 = Negabinary<int16_t>;
using Negabinary32 = Negabinary<int32_t>;
using Negabinary64 = Negabinary<int64_t>;
 
}  // namespace fz
 
#undef FZ_HOST_DEVICE