FZGPUModules 2.0
GPU-accelerated modular compression pipelines
Loading...
Searching...
No Matches
rle.h
Go to the documentation of this file.
1#pragma once
2
8#include "stage/stage.h"
9#include "fzm_format.h"
10#include "log.h"
11#include "backend/types.h"
12#include <cstdint>
13#include <cstring>
14#include <type_traits>
15
16namespace fz {
17
26template<typename T>
27constexpr size_t rleValuesOffset() {
28 return (alignof(T) > sizeof(uint32_t)) ? alignof(T) : sizeof(uint32_t);
29}
30
37template<typename T>
38constexpr size_t rleChunkedValuesOffset(size_t num_chunks) {
39 const size_t hdr = (num_chunks + 2) * sizeof(uint32_t);
40 return (alignof(T) > sizeof(uint32_t))
41 ? ((hdr + alignof(T) - 1) & ~(alignof(T) - 1))
42 : hdr;
43}
44
81template<typename T = uint16_t>
82class RLEStage : public Stage {
83public:
84 RLEStage() : is_inverse_(false) {}
85 ~RLEStage() override;
86
87 void setInverse(bool inverse) override { is_inverse_ = inverse; }
88 bool isInverse() const override { return is_inverse_; }
89
90 void execute(
91 fz::stream_t stream,
92 MemoryPool* pool,
93 const std::vector<void*>& inputs,
94 const std::vector<void*>& outputs,
95 const std::vector<size_t>& sizes
96 ) override;
97
103 void postStreamSync(fz::stream_t stream) override;
104
105 std::string getName() const override { return "RLE"; }
106 size_t getNumInputs() const override { return 1; }
107 size_t getNumOutputs() const override { return 1; }
108
114 void setChunkSize(size_t bytes) {
115 chunk_size_ = static_cast<uint32_t>(bytes - (bytes % sizeof(T)));
116 }
117 size_t getChunkSize() const { return chunk_size_; }
118 bool isChunked() const { return chunk_size_ >= sizeof(T); }
119
121 size_t getRequiredInputAlignment() const override {
122 return isChunked() ? chunk_size_ : 1;
123 }
124
136 const std::vector<size_t>& input_sizes
137 ) const override {
138 if (is_inverse_ || input_sizes.empty()) return 0;
139 const size_t n = input_sizes[0] / sizeof(T);
140 // is_boundary(1B) + boundary_scan(4B) + boundary_positions(4B)
141 // + values_scratch(sizeof(T)) + lengths_scratch(4B)
142 // Chunked mode reuses the same arrays: boundary_positions holds the
143 // per-chunk run start positions and boundary_scan holds the per-chunk
144 // run counts, so the bound is unchanged (and is_boundary goes unused).
145 return n * (1 + 4 + 4 + sizeof(T) + 4);
146 }
147
148 std::vector<size_t> estimateOutputSizes(
149 const std::vector<size_t>& input_sizes
150 ) const override {
151 if (is_inverse_) {
152 // Use the element count cached from the forward pass (or deserialized
153 // from the file header) for an exact estimate. Falls back to a
154 // conservative 2× bound only when no prior forward pass has run.
155 if (cached_num_elements_ > 0)
156 return {static_cast<size_t>(cached_num_elements_) * sizeof(T)};
157 return {input_sizes[0] * 2};
158 } else {
159 // Compression: worst case is every element is unique.
160 // Wire format: [num_runs:u32][values:T×n, 4B-aligned][lengths:u32×n]
161 // The values section is padded to a 4-byte boundary (matching
162 // rle_pack_kernel), so the estimate must include that padding or
163 // the allocated buffer will be too small and the lengths write OOBs.
164 size_t n = input_sizes[0] / sizeof(T);
165 size_t values_bytes = n * sizeof(T);
166 size_t values_aligned = (values_bytes + 3u) & ~3u;
167 if (isChunked()) {
168 // Same worst case (every element unique) plus the offset table.
169 const size_t nc = numChunks(input_sizes[0]);
170 return {rleChunkedValuesOffset<T>(nc) + values_aligned
171 + n * sizeof(uint32_t)};
172 }
173 return {rleValuesOffset<T>() + values_aligned + n * sizeof(uint32_t)};
174 }
175 }
176
177 std::unordered_map<std::string, size_t> getActualOutputSizesByName() const override {
178 completePendingSync();
179 return {{"output", actual_output_sizes_.empty() ? 0 : actual_output_sizes_[0]}};
180 }
181 size_t getActualOutputSize(int index) const override {
182 completePendingSync();
183 return (index == 0 && !actual_output_sizes_.empty()) ? actual_output_sizes_[0] : 0;
184 }
185
186 uint16_t getStageTypeId() const override {
187 return static_cast<uint16_t>(StageType::RLE);
188 }
189
190 uint8_t getOutputDataType(size_t output_index) const override {
191 (void)output_index;
192 return static_cast<uint8_t>(getDataTypeEnum());
193 }
194
195 uint8_t getInputDataType(size_t /*input_index*/) const override {
196 return static_cast<uint8_t>(getDataTypeEnum());
197 }
198
199 size_t serializeHeader(size_t output_index, uint8_t* header_buffer, size_t max_size) const override {
200 (void)output_index;
201 const size_t needed = sizeof(DataType) + 2 * sizeof(uint32_t);
202 if (max_size < needed) return 0;
203 DataType dt = getDataTypeEnum();
204 std::memcpy(header_buffer, &dt, sizeof(DataType));
205 std::memcpy(header_buffer + sizeof(DataType), &cached_num_elements_, sizeof(uint32_t));
206 std::memcpy(header_buffer + sizeof(DataType) + sizeof(uint32_t),
207 &chunk_size_, sizeof(uint32_t));
208 return needed;
209 }
210
211 void deserializeHeader(const uint8_t* header_buffer, size_t size) override {
212 if (size >= sizeof(DataType) + sizeof(uint32_t))
213 std::memcpy(&cached_num_elements_, header_buffer + sizeof(DataType), sizeof(uint32_t));
214 // chunk_size_ is absent in streams written before chunked mode existed;
215 // leaving it at its default keeps those decoding on the global path.
216 if (size >= sizeof(DataType) + 2 * sizeof(uint32_t))
217 std::memcpy(&chunk_size_, header_buffer + sizeof(DataType) + sizeof(uint32_t),
218 sizeof(uint32_t));
219 }
220
221 size_t getMaxHeaderSize(size_t output_index) const override {
222 (void)output_index;
223 return sizeof(DataType) + 2 * sizeof(uint32_t);
224 }
225
226private:
227 bool is_inverse_;
228
230 uint32_t chunk_size_ = 0;
231
233 uint32_t elemsPerChunk() const {
234 return static_cast<uint32_t>(chunk_size_ / sizeof(T));
235 }
236 size_t numChunks(size_t in_bytes) const {
237 const size_t epc = elemsPerChunk();
238 const size_t n = in_bytes / sizeof(T);
239 return epc ? (n + epc - 1) / epc : 0;
240 }
241
245 uint32_t cached_num_elements_ = 0;
246
247 // ── Persistent forward-path scratch ──────────────────────────────────────
248 // Allocated lazily on the first forward execute(); grown if n increases.
249 uint8_t* d_is_boundary_ = nullptr;
250 uint32_t* d_boundary_scan_ = nullptr;
251 uint32_t* d_boundary_positions_ = nullptr;
252 T* d_values_scratch_ = nullptr;
253 uint32_t* d_lengths_scratch_ = nullptr;
254 size_t fwd_scratch_n_ = 0;
255 MemoryPool* fwd_scratch_pool_ = nullptr;
256 bool fwd_from_pool_ = false;
257
258 // Pinned host buffer for async D2H of num_runs.
259 // mutable so getActualOutputSizesByName() can complete the pending
260 // readback even when called on a const Stage reference.
261 mutable uint32_t* h_num_runs_ = nullptr;
262 mutable bool fwd_sync_pending_ = false;
263 mutable fz::stream_t fwd_last_stream_ = nullptr;
264 mutable std::vector<size_t> actual_output_sizes_;
265
266 // Complete a pending forward-path readback (if any) by syncing the stream
267 // that was used and computing actual_output_sizes_. Safe to call from
268 // const methods; all state it touches is mutable.
269 void completePendingSync() const {
270 if (!fwd_sync_pending_) return;
271 cudaStreamSynchronize(fwd_last_stream_);
272 const uint32_t num_runs = *h_num_runs_;
273 const size_t values_bytes = static_cast<size_t>(num_runs) * sizeof(T);
274 const size_t values_aligned = (values_bytes + 3) & ~3;
275 const size_t values_offset = isChunked()
276 ? rleChunkedValuesOffset<T>(
277 numChunks(static_cast<size_t>(cached_num_elements_) * sizeof(T)))
278 : rleValuesOffset<T>();
279 actual_output_sizes_ = {
280 values_offset + values_aligned + num_runs * sizeof(uint32_t)
281 };
282 fwd_sync_pending_ = false;
283 // Log run count and effective compression ratio.
284 const size_t in_bytes = static_cast<size_t>(cached_num_elements_) * sizeof(T);
285 const size_t out_bytes = actual_output_sizes_[0];
286 const float ratio = in_bytes > 0
287 ? static_cast<float>(in_bytes) / static_cast<float>(out_bytes) : 0.0f;
288 FZ_LOG(DEBUG, "RLE encode: %u runs / %u elems %.1f KB -> %.1f KB ratio %.2fx",
289 num_runs, cached_num_elements_,
290 in_bytes / 1024.0f, out_bytes / 1024.0f, ratio);
291 }
292
293 // Helper to map template type T to DataType enum
294 DataType getDataTypeEnum() const {
295 if (std::is_same<T, uint8_t>::value) return DataType::UINT8;
296 if (std::is_same<T, uint16_t>::value) return DataType::UINT16;
297 if (std::is_same<T, uint32_t>::value) return DataType::UINT32;
298 if (std::is_same<T, uint64_t>::value) return DataType::UINT64;
299 if (std::is_same<T, int8_t>::value) return DataType::INT8;
300 if (std::is_same<T, int16_t>::value) return DataType::INT16;
301 if (std::is_same<T, int32_t>::value) return DataType::INT32;
302 if (std::is_same<T, int64_t>::value) return DataType::INT64;
303 if (std::is_same<T, float>::value) return DataType::FLOAT32;
304 if (std::is_same<T, double>::value) return DataType::FLOAT64;
305 return DataType::UINT8; // Fallback
306 }
307};
308
309extern template class RLEStage<uint8_t>;
310extern template class RLEStage<uint16_t>;
311extern template class RLEStage<uint32_t>;
312extern template class RLEStage<uint64_t>;
313extern template class RLEStage<int8_t>;
314extern template class RLEStage<int16_t>;
315extern template class RLEStage<int32_t>;
316extern template class RLEStage<int64_t>;
317
318} // namespace fz
Definition mempool.h:82
Definition rle.h:82
void setInverse(bool inverse) override
Definition rle.h:87
void setChunkSize(size_t bytes)
Definition rle.h:114
void postStreamSync(fz::stream_t stream) override
void deserializeHeader(const uint8_t *header_buffer, size_t size) override
Definition rle.h:211
std::unordered_map< std::string, size_t > getActualOutputSizesByName() const override
Definition rle.h:177
std::string getName() const override
Definition rle.h:105
size_t serializeHeader(size_t output_index, uint8_t *header_buffer, size_t max_size) const override
Definition rle.h:199
void execute(fz::stream_t stream, MemoryPool *pool, const std::vector< void * > &inputs, const std::vector< void * > &outputs, const std::vector< size_t > &sizes) override
uint16_t getStageTypeId() const override
Definition rle.h:186
size_t getMaxHeaderSize(size_t output_index) const override
Definition rle.h:221
uint8_t getOutputDataType(size_t output_index) const override
Definition rle.h:190
std::vector< size_t > estimateOutputSizes(const std::vector< size_t > &input_sizes) const override
Definition rle.h:148
uint8_t getInputDataType(size_t) const override
Definition rle.h:195
size_t getActualOutputSize(int index) const override
Definition rle.h:181
size_t getRequiredInputAlignment() const override
Chunked mode needs whole chunks; the pipeline zero-pads the input to suit.
Definition rle.h:121
size_t estimateScratchBytes(const std::vector< size_t > &input_sizes) const override
Definition rle.h:135
Definition stage.h:30
FZM binary file format definitions — structs, enums, and helpers.
Logging infrastructure and macros.
#define FZ_LOG(level,...)
Definition log.h:201
Definition algorithms.h:48
constexpr size_t rleChunkedValuesOffset(size_t num_chunks)
Definition rle.h:38
constexpr size_t rleValuesOffset()
Definition rle.h:27
DataType
Element data type identifiers used in buffer and stage descriptors.
Definition fzm_format.h:117
@ DEBUG
Pipeline construction, buffer allocation, data stats.
Base class interface for all compression stages.
Backend-neutral GPU type aliases.