FZGPUModules 1.0
GPU-accelerated modular compression pipeline
Loading...
Searching...
No Matches
negabinary.h
1#pragma once
2
37#include <cstdint>
38#include <type_traits>
39
40#ifdef __CUDACC__
41# define FZ_HOST_DEVICE __host__ __device__
42#else
43# define FZ_HOST_DEVICE
44#endif
45
46namespace fz {
47
48template <typename T>
49struct Negabinary {
50 static_assert(std::is_integral<T>::value && std::is_signed<T>::value,
51 "fz::Negabinary<T>: T must be a signed integer type "
52 "(int8_t, int16_t, int32_t, or int64_t).");
53
54 using SInt = T;
55 using UInt = typename std::make_unsigned<T>::type;
56
57 // Alternating-bit mask 0xAA…A for the UInt width.
58 // Truncating the 64-bit constant to UInt gives:
59 // uint8_t → 0xAA
60 // uint16_t → 0xAAAA
61 // uint32_t → 0xAAAAAAAA
62 // uint64_t → 0xAAAAAAAAAAAAAAAA
63 static constexpr UInt MASK =
64 static_cast<UInt>(static_cast<uint64_t>(0xAAAAAAAAAAAAAAAAULL));
65
71 FZ_HOST_DEVICE static constexpr UInt encode(SInt n) noexcept {
72 return (static_cast<UInt>(n) + MASK) ^ MASK;
73 }
74
83 FZ_HOST_DEVICE static constexpr SInt decode(UInt u) noexcept {
84 return static_cast<SInt>((u ^ MASK) - MASK);
85 }
86};
87
88// ---------------------------------------------------------------------------
89// Convenience aliases for the four standard widths
90// ---------------------------------------------------------------------------
91using Negabinary8 = Negabinary<int8_t>;
92using Negabinary16 = Negabinary<int16_t>;
93using Negabinary32 = Negabinary<int32_t>;
94using Negabinary64 = Negabinary<int64_t>;
95
96} // namespace fz
97
98#undef FZ_HOST_DEVICE
Definition fzm_format.h:25