87 void setInverse(
bool inverse)
override { is_inverse_ = inverse; }
88 bool isInverse()
const override {
return is_inverse_; }
93 const std::vector<void*>& inputs,
94 const std::vector<void*>& outputs,
95 const std::vector<size_t>& sizes
105 std::string
getName()
const override {
return "RLE"; }
106 size_t getNumInputs()
const override {
return 1; }
107 size_t getNumOutputs()
const override {
return 1; }
115 chunk_size_ =
static_cast<uint32_t
>(bytes - (bytes %
sizeof(T)));
117 size_t getChunkSize()
const {
return chunk_size_; }
118 bool isChunked()
const {
return chunk_size_ >=
sizeof(T); }
122 return isChunked() ? chunk_size_ : 1;
136 const std::vector<size_t>& input_sizes
138 if (is_inverse_ || input_sizes.empty())
return 0;
139 const size_t n = input_sizes[0] /
sizeof(T);
145 return n * (1 + 4 + 4 +
sizeof(T) + 4);
149 const std::vector<size_t>& input_sizes
155 if (cached_num_elements_ > 0)
156 return {
static_cast<size_t>(cached_num_elements_) *
sizeof(T)};
157 return {input_sizes[0] * 2};
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;
169 const size_t nc = numChunks(input_sizes[0]);
170 return {rleChunkedValuesOffset<T>(nc) + values_aligned
171 + n *
sizeof(uint32_t)};
173 return {rleValuesOffset<T>() + values_aligned + n *
sizeof(uint32_t)};
178 completePendingSync();
179 return {{
"output", actual_output_sizes_.empty() ? 0 : actual_output_sizes_[0]}};
182 completePendingSync();
183 return (index == 0 && !actual_output_sizes_.empty()) ? actual_output_sizes_[0] : 0;
187 return static_cast<uint16_t
>(StageType::RLE);
192 return static_cast<uint8_t
>(getDataTypeEnum());
196 return static_cast<uint8_t
>(getDataTypeEnum());
199 size_t serializeHeader(
size_t output_index, uint8_t* header_buffer,
size_t max_size)
const override {
201 const size_t needed =
sizeof(
DataType) + 2 *
sizeof(uint32_t);
202 if (max_size < needed)
return 0;
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));
212 if (size >=
sizeof(
DataType) +
sizeof(uint32_t))
213 std::memcpy(&cached_num_elements_, header_buffer +
sizeof(
DataType),
sizeof(uint32_t));
216 if (size >=
sizeof(
DataType) + 2 *
sizeof(uint32_t))
217 std::memcpy(&chunk_size_, header_buffer +
sizeof(
DataType) +
sizeof(uint32_t),
223 return sizeof(
DataType) + 2 *
sizeof(uint32_t);
230 uint32_t chunk_size_ = 0;
233 uint32_t elemsPerChunk()
const {
234 return static_cast<uint32_t
>(chunk_size_ /
sizeof(T));
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;
245 uint32_t cached_num_elements_ = 0;
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;
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_;
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)))
279 actual_output_sizes_ = {
280 values_offset + values_aligned + num_runs *
sizeof(uint32_t)
282 fwd_sync_pending_ =
false;
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);
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;
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