FZGPUModules 2.0
GPU-accelerated modular compression pipelines
Loading...
Searching...
No Matches
zigzag.h
Go to the documentation of this file.
1#pragma once
2
30#include <cstdint>
31#include <type_traits>
32
33#ifdef __CUDACC__
34# define FZ_HOST_DEVICE __host__ __device__
35#else
36# define FZ_HOST_DEVICE
37#endif
38
39namespace fz {
40
53template <typename T>
54struct Zigzag {
55 static_assert(std::is_integral<T>::value && std::is_signed<T>::value,
56 "fz::Zigzag<T>: T must be a signed integer type "
57 "(int8_t, int16_t, int32_t, or int64_t).");
58
59 using SInt = T;
60 using UInt = typename std::make_unsigned<T>::type;
61
62 static constexpr int W = sizeof(T) * 8;
63
70 FZ_HOST_DEVICE static constexpr UInt encode(SInt x) noexcept {
71 // Cast to unsigned before shift to avoid UB on signed overflow,
72 // then XOR with the arithmetic right-shift mask.
73 return (static_cast<UInt>(x) << 1) ^ static_cast<UInt>(x >> (W - 1));
74 }
75
80 FZ_HOST_DEVICE static constexpr SInt decode(UInt u) noexcept {
81 // (u >> 1) reverses the shift; -(u & 1) reconstructs the sign mask.
82 // Cast to signed only at the final XOR to keep arithmetic well-defined.
83 return static_cast<SInt>((u >> 1) ^ static_cast<UInt>(-(static_cast<SInt>(u & 1u))));
84 }
85};
86
87// ---------------------------------------------------------------------------
88// Convenience aliases for the four standard widths
89// ---------------------------------------------------------------------------
90using Zigzag8 = Zigzag<int8_t>;
91using Zigzag16 = Zigzag<int16_t>;
92using Zigzag32 = Zigzag<int32_t>;
93using Zigzag64 = Zigzag<int64_t>;
94
95} // namespace fz
96
97#undef FZ_HOST_DEVICE
Definition zigzag.h:54
static FZ_HOST_DEVICE constexpr SInt decode(UInt u) noexcept
Definition zigzag.h:80
static FZ_HOST_DEVICE constexpr UInt encode(SInt x) noexcept
Definition zigzag.h:70