zigzag.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 Zigzag {
    static_assert(std::is_integral<T>::value && std::is_signed<T>::value,
                  "fz::Zigzag<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;
 
    static constexpr int W = sizeof(T) * 8;
 
    FZ_HOST_DEVICE static constexpr UInt encode(SInt x) noexcept {
        // Cast to unsigned before shift to avoid UB on signed overflow,
        // then XOR with the arithmetic right-shift mask.
        return (static_cast<UInt>(x) << 1) ^ static_cast<UInt>(x >> (W - 1));
    }
 
    FZ_HOST_DEVICE static constexpr SInt decode(UInt u) noexcept {
        // (u >> 1) reverses the shift; -(u & 1) reconstructs the sign mask.
        // Cast to signed only at the final XOR to keep arithmetic well-defined.
        return static_cast<SInt>((u >> 1) ^ static_cast<UInt>(-(static_cast<SInt>(u & 1u))));
    }
};
 
// ---------------------------------------------------------------------------
// Convenience aliases for the four standard widths
// ---------------------------------------------------------------------------
using Zigzag8  = Zigzag<int8_t>;
using Zigzag16 = Zigzag<int16_t>;
using Zigzag32 = Zigzag<int32_t>;
using Zigzag64 = Zigzag<int64_t>;
 
}  // namespace fz
 
#undef FZ_HOST_DEVICE