- Generated by
1.9.8
|
FZGPUModules 2.0
GPU-accelerated modular compression pipelines
|
Header: modules/coders/bitpack/bitpack_stage.h
Class: fz::BitpackStage<T>
Category: Coder (lossless)
Packs each element using only its low nbits bits into a dense byte stream.
T[] → uint8_t[] — ceil(n × nbits / 8) bytes.uint8_t[] → T[] — unpacks elements, zero-extending to full width.Most useful after a Lorenzo predictor stage where small delta values only need a few bits of representation.
| Parameter | Constraint |
|---|---|
T | Unsigned integer type (see available instantiations below) |
Only these types are compiled and linked:
BitpackStage<uint8_t>BitpackStage<uint16_t>BitpackStage<uint32_t>Using any other type will result in a linker error. Most common: BitpackStage<uint32_t> (to match typical quantizer code width).
| Setting | Purpose | Notes |
|---|---|---|
setNBits(nbits) | Bits per element | Power of two, 1..8 * sizeof(T); ignored when auto-detect is on |
setAutoDetect(bool) | GPU scan to pick nbits automatically | Disables CUDA Graph compatibility while active |
setBase(base) | Frame-of-reference offset subtracted before packing | Removes dead high bits; always lossless |
setShift(shift) | Right shift applied after the base subtraction | Removes dead low bits; lossy unless those bits are zero |
setAutoBase(bool) | GPU min-reduce to pick base automatically | Disables CUDA Graph compatibility while active |
setAutoShift(bool) | GPU OR-reduce to pick the largest lossless shift | Disables CUDA Graph compatibility while active |
setAdaptive(bool) | Convenience: auto base + auto shift + auto nbits | Fully adaptive lossless mode |
nbits must be a power of two in [1, 8 × sizeof(T)]. Default is 8 * sizeof(T) (identity, no compression).
T | Allowed nbits |
|---|---|
uint8_t | 1, 2, 4, 8 |
uint16_t | 1, 2, 4, 8, 16 |
uint32_t | 1, 2, 4, 8, 16, 32 |
Violations throw std::invalid_argument at setNBits() time.
When enabled, forward execute scans the input for its maximum value using cub::DeviceReduce::Max and selects the smallest valid power-of-two nbits that covers it. The chosen nbits is written into the compressed header so the inverse pass unpacks correctly without any out-of-band configuration.
After compress(), getNBits() reflects the detected value.
Scratch buffers for the scan are allocated through the pipeline's memory pool (with a transparent cudaMalloc fallback in vGPU / pool-fallback mode), so all device memory remains tracked by the pipeline.
CUDA Graph incompatibility: auto-detect requires a device-to-host transfer and stream synchronization to read the max value, making it incompatible with CUDA Graph capture. isGraphCompatible() returns false while auto-detect is on. If you know the bit-width ahead of time, use setNBits() instead to keep graph capture available.
Output size estimate: estimateOutputSizes() returns the worst-case (full input size) when auto-detect is enabled, so PREALLOCATE mode reserves sufficient space regardless of the detected nbits.
nbits alone only helps when the useful information already sits in the low bits of each word. The shift transform moves it there first:
The two knobs attack opposite ends of the word and compose:
| Knob | Removes | Lossless? |
|---|---|---|
base | dead high bits — values clustered far from zero | always |
shift | dead low bits — values that are all multiples of 1 << shift | only if those bits are zero |
Both default to 0, which is the identity: behaviour is unchanged unless you opt in.
**
setShift()is lossy in general.** The inverse restores(packed << shift) + base, so any low bits you drop are gone — values come back rounded down to a multiple of1 << shift. UsesetAutoShift(true)if you want the largest shift that is provably lossless for your data.
Forward execute then runs three GPU scans, in this order, and writes all three results into the compressed header:
base ← cub::DeviceReduce::Min over the input.shift ← trailing-zero count of an OR-reduce of every (v - base). This is the largest shift that drops no information, so adaptive mode is always lossless.nbits ← smallest power of two covering (max - base) >> shift.After compress(), getBase(), getShift(), and getNBits() reflect the detected values.
Worked example — values of the form 1000 + k*16 for k in [0, 15]:
| value | bits needed | |
|---|---|---|
raw uint16_t | 1000 … 1240 | 16 |
after base = 1000 | 0 … 240 | 8 |
after shift = 4 | 0 … 15 | 4 |
A 4× gain that neither knob delivers on its own. The individual setters are still available if you want only one half (e.g. setAutoBase(true) with a hand-set nbits, keeping the shift at 0).
CUDA Graph incompatibility: each auto mode needs a device-to-host readback, so isGraphCompatible() returns false while any of setAutoDetect, setAutoBase, or setAutoShift is on. Hand-set base/shift/nbits stay graph-capturable.
File format: the serialized header is 15 bytes — the previous 10 plus shift (1 byte) and base (4 bytes, zero-extended). Archives written before this option decode with shift = base = 0.
1.9.8