FZGPUModules 1.0
GPU-accelerated modular compression pipeline
Loading...
Searching...
No Matches
zigzag.h
1#pragma once
2
29#include <cstdint>
30#include <type_traits>
31
32#ifdef __CUDACC__
33# define FZ_HOST_DEVICE __host__ __device__
34#else
35# define FZ_HOST_DEVICE
36#endif
37
38namespace fz {
39
40template <typename T>
41struct Zigzag {
42 static_assert(std::is_integral<T>::value && std::is_signed<T>::value,
43 "fz::Zigzag<T>: T must be a signed integer type "
44 "(int8_t, int16_t, int32_t, or int64_t).");
45
46 using SInt = T;
47 using UInt = typename std::make_unsigned<T>::type;
48
49 static constexpr int W = sizeof(T) * 8;
50
57 FZ_HOST_DEVICE static constexpr UInt encode(SInt x) noexcept {
58 // Cast to unsigned before shift to avoid UB on signed overflow,
59 // then XOR with the arithmetic right-shift mask.
60 return (static_cast<UInt>(x) << 1) ^ static_cast<UInt>(x >> (W - 1));
61 }
62
67 FZ_HOST_DEVICE static constexpr SInt decode(UInt u) noexcept {
68 // (u >> 1) reverses the shift; -(u & 1) reconstructs the sign mask.
69 // Cast to signed only at the final XOR to keep arithmetic well-defined.
70 return static_cast<SInt>((u >> 1) ^ static_cast<UInt>(-(static_cast<SInt>(u & 1u))));
71 }
72};
73
74// ---------------------------------------------------------------------------
75// Convenience aliases for the four standard widths
76// ---------------------------------------------------------------------------
77using Zigzag8 = Zigzag<int8_t>;
78using Zigzag16 = Zigzag<int16_t>;
79using Zigzag32 = Zigzag<int32_t>;
80using Zigzag64 = Zigzag<int64_t>;
81
82} // namespace fz
83
84#undef FZ_HOST_DEVICE
Definition fzm_format.h:25