|
| 1 | +/** |
| 2 | + * examples/cusz_standalone.cpp |
| 3 | + * |
| 4 | + * LorenzoQuantStage → HuffmanStage driven manually without a Pipeline. |
| 5 | + * |
| 6 | + * Use this pattern when you need precise control over execution order, |
| 7 | + * intermediate buffer placement, or per-stage stream assignment — for |
| 8 | + * instance, to interleave other GPU work between compression stages, |
| 9 | + * or to run stages asynchronously on multiple streams. |
| 10 | + * |
| 11 | + * Demonstrates: |
| 12 | + * - Constructing a caller-managed MemoryPool for stage scratch |
| 13 | + * - HuffmanStage::onFinalize() to pre-allocate persistent scratch buffers |
| 14 | + * - LorenzoQuantStage::execute() — 1 input, 4 outputs |
| 15 | + * (codes, outlier_errors, outlier_indices, outlier_count) |
| 16 | + * - LorenzoQuantStage::postStreamSync() to read the actual outlier count |
| 17 | + * - HuffmanStage::execute() in compress and decompress direction |
| 18 | + * - Stage::setInverse(true) to switch to the decompression pass |
| 19 | + * - LorenzoQuantStage::execute() in decompress mode — 4 inputs, 1 output |
| 20 | + * - Verifying the full compress → decompress round-trip |
| 21 | + * |
| 22 | + * Contrast with the Pipeline DAG path: |
| 23 | + * - Pipeline manages the pool, allocates intermediate buffers, and calls |
| 24 | + * postStreamSync() automatically — you write connect() instead of buffer ptrs. |
| 25 | + * - The standalone path gives you raw pointers but requires manual management |
| 26 | + * of everything the Pipeline normally hides. |
| 27 | + * |
| 28 | + * No external data files required — uses synthetic float data. |
| 29 | + * |
| 30 | + * Usage: |
| 31 | + * ./build/release/bin/examples/cusz_standalone |
| 32 | + * |
| 33 | + * Build: |
| 34 | + * cmake --preset release -DBUILD_EXAMPLES=ON && cmake --build build/release -j$(nproc) |
| 35 | + * Binary: build/release/bin/examples/cusz_standalone |
| 36 | + */ |
| 37 | + |
| 38 | +#include "fzgpumodules.h" |
| 39 | + |
| 40 | +#include <algorithm> |
| 41 | +#include <cmath> |
| 42 | +#include <cstdio> |
| 43 | +#include <vector> |
| 44 | + |
| 45 | +using namespace fz; |
| 46 | + |
| 47 | +int main() { |
| 48 | + // ── 1. Input data ───────────────────────────────────────────────────────── |
| 49 | + static constexpr size_t N = 1 << 18; // 256 K floats = 1 MB |
| 50 | + static constexpr float EB = 1e-4f; |
| 51 | + static constexpr float OUTLIER_CAPACITY = 0.10f; |
| 52 | + static constexpr uint16_t QUANT_RADIUS = 512; |
| 53 | + // Zigzag maps [-511, 511] → [0, 1022]; bklen=1024 covers the full range. |
| 54 | + static constexpr uint16_t BKLEN = 2 * QUANT_RADIUS; |
| 55 | + |
| 56 | + std::vector<float> h_input(N); |
| 57 | + for (size_t i = 0; i < N; ++i) { |
| 58 | + const float t = static_cast<float>(i) / static_cast<float>(N); |
| 59 | + h_input[i] = std::sin(2.0f * 3.14159265f * t) |
| 60 | + + 0.5f * std::cos(6.0f * 3.14159265f * t); |
| 61 | + } |
| 62 | + const size_t input_bytes = N * sizeof(float); |
| 63 | + const size_t codes_bytes = N * sizeof(uint16_t); |
| 64 | + const size_t max_outliers = static_cast<size_t>(N * OUTLIER_CAPACITY) + 1; |
| 65 | + const size_t huf_out_cap = codes_bytes * 2 + 4096; // generous upper bound |
| 66 | + |
| 67 | + float* d_input = nullptr; |
| 68 | + cudaMalloc(&d_input, input_bytes); |
| 69 | + cudaMemcpy(d_input, h_input.data(), input_bytes, cudaMemcpyHostToDevice); |
| 70 | + |
| 71 | + // ── 2. Stage setup ──────────────────────────────────────────────────────── |
| 72 | + LorenzoQuantStage<float, uint16_t> lq; |
| 73 | + lq.setErrorBound(EB); |
| 74 | + lq.setErrorBoundMode(ErrorBoundMode::ABS); |
| 75 | + lq.setQuantRadius(QUANT_RADIUS); |
| 76 | + lq.setOutlierCapacity(OUTLIER_CAPACITY); |
| 77 | + lq.setZigzagCodes(true); |
| 78 | + |
| 79 | + HuffmanStage<uint16_t> huf; |
| 80 | + huf.setBklen(BKLEN); |
| 81 | + |
| 82 | + // ── 3. Pool and persistent scratch ─────────────────────────────────────── |
| 83 | + // |
| 84 | + // The pool serves two roles: |
| 85 | + // a) Huffman persistent scratch: pre-allocated here via onFinalize() |
| 86 | + // so that no allocation happens inside the timed compress() loop. |
| 87 | + // b) LorenzoQuant transient scratch: tiny workspace for the NOA/REL |
| 88 | + // value-range scan (not needed here with ABS mode, but the pool |
| 89 | + // pointer is required by execute() regardless). |
| 90 | + // |
| 91 | + // Size the pool to cover Huffman's full device footprint plus padding. |
| 92 | + MemoryPool pool(MemoryPoolConfig(input_bytes, 4.0f)); |
| 93 | + huf.onFinalize(codes_bytes, &pool); // pre-allocates PHF scratch buffers |
| 94 | + |
| 95 | + // ── 4. Intermediate device buffers ──────────────────────────────────────── |
| 96 | + // |
| 97 | + // Caller owns all of these — cudaFree every one on exit. |
| 98 | + // These are distinct from pool memory (which is internal scratch). |
| 99 | + void* d_codes = nullptr; |
| 100 | + void* d_outlier_errors = nullptr; |
| 101 | + void* d_outlier_indices = nullptr; |
| 102 | + void* d_outlier_count = nullptr; |
| 103 | + void* d_huf_output = nullptr; |
| 104 | + void* d_reconstructed = nullptr; |
| 105 | + |
| 106 | + cudaMalloc(&d_codes, codes_bytes); |
| 107 | + cudaMalloc(&d_outlier_errors, max_outliers * sizeof(float)); |
| 108 | + cudaMalloc(&d_outlier_indices, max_outliers * sizeof(uint32_t)); |
| 109 | + cudaMalloc(&d_outlier_count, sizeof(uint32_t)); |
| 110 | + cudaMalloc(&d_huf_output, huf_out_cap); |
| 111 | + cudaMalloc(&d_reconstructed, input_bytes); |
| 112 | + cudaDeviceSynchronize(); |
| 113 | + |
| 114 | + // ── 5. Compress: LorenzoQuant → Huffman ─────────────────────────────────── |
| 115 | + // |
| 116 | + // LorenzoQuant compress: 1 input → 4 outputs. |
| 117 | + // outputs[3] = outlier_count (device uint32_t) — initialized to 0 by execute(). |
| 118 | + lq.execute(0, &pool, |
| 119 | + {d_input}, |
| 120 | + {d_codes, d_outlier_errors, d_outlier_indices, d_outlier_count}, |
| 121 | + {input_bytes}); |
| 122 | + cudaDeviceSynchronize(); |
| 123 | + |
| 124 | + // postStreamSync() reads the actual outlier count back to host and trims |
| 125 | + // actual_output_sizes_ to the real (not max-capacity) sizes. |
| 126 | + // Must be called after the stream is fully synchronized. |
| 127 | + lq.postStreamSync(0); |
| 128 | + |
| 129 | + const size_t actual_outlier_errors_bytes = lq.getActualOutputSize(1); |
| 130 | + const size_t actual_outlier_indices_bytes = lq.getActualOutputSize(2); |
| 131 | + const size_t n_outliers = actual_outlier_errors_bytes / sizeof(float); |
| 132 | + std::printf("LorenzoQuant compress: %zu elements, %zu outliers (%.2f%%)\n", |
| 133 | + N, n_outliers, 100.0 * n_outliers / N); |
| 134 | + |
| 135 | + // Huffman compress: 1 input (codes) → 1 output (compressed bitstream). |
| 136 | + huf.execute(0, &pool, |
| 137 | + {d_codes}, |
| 138 | + {d_huf_output}, |
| 139 | + {codes_bytes}); |
| 140 | + cudaDeviceSynchronize(); |
| 141 | + |
| 142 | + const size_t compressed_bytes = huf.getActualOutputSize(0); |
| 143 | + std::printf("Huffman compress: %zu bytes → %zu bytes (%.2fx for codes)\n", |
| 144 | + codes_bytes, compressed_bytes, |
| 145 | + static_cast<double>(codes_bytes) / compressed_bytes); |
| 146 | + std::printf("Full ratio (input/compressed_codes): %.2fx\n", |
| 147 | + static_cast<double>(input_bytes) / compressed_bytes); |
| 148 | + |
| 149 | + // ── 6. Decompress: Huffman inverse → LorenzoQuant inverse ───────────────── |
| 150 | + // |
| 151 | + // Both stages reuse the same objects with setInverse(true). |
| 152 | + // The Huffman header embedded in d_huf_output tells the decoder the |
| 153 | + // original length — no external metadata needed for the codes buffer. |
| 154 | + |
| 155 | + // Huffman decompress: compressed bitstream → codes. |
| 156 | + huf.setInverse(true); |
| 157 | + huf.execute(0, &pool, |
| 158 | + {d_huf_output}, |
| 159 | + {d_codes}, // reuse d_codes as output |
| 160 | + {compressed_bytes}); |
| 161 | + cudaDeviceSynchronize(); |
| 162 | + |
| 163 | + // LorenzoQuant decompress: 4 inputs → 1 output (reconstructed floats). |
| 164 | + // sizes[1] = outlier_errors_bytes is used to derive max_outliers inside |
| 165 | + // the kernel, so pass the actual (not max-capacity) sizes. |
| 166 | + lq.setInverse(true); |
| 167 | + lq.execute(0, &pool, |
| 168 | + {d_codes, d_outlier_errors, d_outlier_indices, d_outlier_count}, |
| 169 | + {d_reconstructed}, |
| 170 | + {codes_bytes, |
| 171 | + actual_outlier_errors_bytes, |
| 172 | + actual_outlier_indices_bytes, |
| 173 | + sizeof(uint32_t)}); |
| 174 | + cudaDeviceSynchronize(); |
| 175 | + |
| 176 | + // ── 7. Verify ───────────────────────────────────────────────────────────── |
| 177 | + std::vector<float> h_out(N); |
| 178 | + cudaMemcpy(h_out.data(), d_reconstructed, input_bytes, cudaMemcpyDeviceToHost); |
| 179 | + |
| 180 | + float max_err = 0.0f; |
| 181 | + for (size_t i = 0; i < N; ++i) |
| 182 | + max_err = std::max(max_err, std::abs(h_out[i] - h_input[i])); |
| 183 | + |
| 184 | + std::printf("max_abs_error: %.2e (eb = %.2e)\n", |
| 185 | + max_err, static_cast<double>(EB)); |
| 186 | + |
| 187 | + // ── 8. Cleanup ──────────────────────────────────────────────────────────── |
| 188 | + cudaFree(d_input); |
| 189 | + cudaFree(d_codes); |
| 190 | + cudaFree(d_outlier_errors); |
| 191 | + cudaFree(d_outlier_indices); |
| 192 | + cudaFree(d_outlier_count); |
| 193 | + cudaFree(d_huf_output); |
| 194 | + cudaFree(d_reconstructed); |
| 195 | + // pool is stack-allocated — its destructor frees Huffman's persistent scratch. |
| 196 | + return 0; |
| 197 | +} |
0 commit comments