/* blake3-fanout.cu — BLAKE3 batched-hash CUDA worker, Wave-2 form
 * `cuda-blake3-tree`. Sibling of shake256-fanout.cu — same wire
 * shape, same daemon contract, distinct binary magic (BSB3 / BSR3).
 *
 * One CUDA thread per input. Each thread walks the BLAKE3 chunk-tree
 * serially: per-chunk compression, then parent merges, then a root
 * compression. Captures both single-chunk (input <= 1024 B) and
 * multi-chunk paths.
 *
 * Wire format mirrors shake256-fanout exactly. Binary mode:
 *   request:  "BSB3" | u32 out_bytes | u32 n_items | for each:
 *                                              u32 len | len bytes
 *   response: "BSR3" | u32 n_items   | u32 out_bytes |
 *             for each: out_bytes
 *
 * Daemon protocol on stdin (one line per command):
 *   process <in.portal>     <out.portal>      — S-expr in, S-expr out
 *   process-bin <in.bin>    <out.bin>         — binary in, binary out
 *   quit
 *
 * Stdout response: "done <path>" / "error <reason>" / "bye".
 *
 * License: AGPLv3. BLAKE3 __device__ primitives vendored from
 * Blaze-3/BLAKE3-gpu (MIT, see vendor/blake3-gpu/{LICENSE,NOTICE}).
 * Constants and round function track the BLAKE3 reference spec
 * (https://github.com/BLAKE3-team/BLAKE3 — CC0 / Apache-2.0).
 */

#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <stdint.h>
#include <time.h>

#include "vendor/blake3-gpu/blake3_device.cuh"

using blake3_gpu::u8;
using blake3_gpu::u32;
using blake3_gpu::u64;
using blake3_gpu::IV;
using blake3_gpu::CHUNK_START;
using blake3_gpu::CHUNK_END;
using blake3_gpu::PARENT;
using blake3_gpu::ROOT;
using blake3_gpu::BLOCK_LEN;
using blake3_gpu::CHUNK_LEN;
using blake3_gpu::OUT_LEN;
using blake3_gpu::compress;
using blake3_gpu::words_from_bytes_le;

#define CUDA_CHECK(stmt) do { \
    cudaError_t err = (stmt); \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA error %s at %s:%d: %s\n", \
                #stmt, __FILE__, __LINE__, cudaGetErrorString(err)); \
        exit(1); \
    } \
} while (0)

/* ─── per-thread BLAKE3 over one input ──────────────────────────────── */

/* Produce the 8-word chaining value of a *complete* (1024 B) chunk at
 * `counter`. Used only for chunks 0..N-2 of multi-chunk inputs. */
__device__ static void chunk_cv_complete(const u8 *input, u64 counter,
                                         u32 *out_cv8)
{
    u32 cv[8];
    for (int i = 0; i < 8; i++) cv[i] = IV[i];

    u32 block_words[16];
    u32 state[16];

    /* CHUNK_LEN = 16 * BLOCK_LEN — exactly 16 compressions per full chunk. */
    for (u32 b = 0; b < (CHUNK_LEN / BLOCK_LEN); b++) {
        words_from_bytes_le(input + (size_t)b * BLOCK_LEN, BLOCK_LEN, block_words);
        u32 flags = 0;
        if (b == 0) flags |= CHUNK_START;
        if (b == (CHUNK_LEN / BLOCK_LEN) - 1) flags |= CHUNK_END;
        compress(cv, block_words, counter, BLOCK_LEN, flags, state);
        for (int i = 0; i < 8; i++) cv[i] = state[i];
    }
    for (int i = 0; i < 8; i++) out_cv8[i] = cv[i];
}

/* Build the "pending Output" for the final chunk (or the only chunk). The
 * output is captured as (input_cv, block_words, counter, block_len, flags)
 * — we do not run the final compress yet; finalize-time decides ROOT.
 *
 * `input` / `len` describe the chunk's bytes; `counter` is chunk index;
 * `is_root_chunk` is true ONLY for single-chunk inputs (no parents will
 * be applied later — we still defer the final compress so that the
 * caller adds ROOT during finalize). */
struct Pending {
    u32 input_cv[8];
    u32 block_words[16];
    u64 counter;
    u32 block_len;
    u32 flags;
};

__device__ static void chunk_pending(const u8 *input, u32 len, u64 counter,
                                     Pending &p)
{
    u32 cv[8];
    for (int i = 0; i < 8; i++) cv[i] = IV[i];

    u32 block_words[16];
    u32 state[16];

    u32 offset = 0;
    u32 blocks_compressed = 0;

    /* Compress every full preceding block. Handle empty input (len==0): we
     * enter the post-loop with blocks_compressed==0 and emit one final
     * 0-byte block with CHUNK_START | CHUNK_END. */
    while (offset + BLOCK_LEN < len) {
        words_from_bytes_le(input + offset, BLOCK_LEN, block_words);
        u32 flags = 0;
        if (blocks_compressed == 0) flags |= CHUNK_START;
        compress(cv, block_words, counter, BLOCK_LEN, flags, state);
        for (int i = 0; i < 8; i++) cv[i] = state[i];
        blocks_compressed++;
        offset += BLOCK_LEN;
    }

    /* Final block: bytes [offset, len). Pad inside words_from_bytes_le. */
    u32 final_len = len - offset;
    words_from_bytes_le(input + offset, final_len, p.block_words);
    for (int i = 0; i < 8; i++) p.input_cv[i] = cv[i];
    p.counter   = counter;
    p.block_len = final_len;
    p.flags     = CHUNK_END;
    if (blocks_compressed == 0) p.flags |= CHUNK_START;
}

/* BLAKE3 tree walk for a single input. Stack depth ≤ 54 per spec
 * (2^54 * 1024 B = 2^64). Cap at 56 for slack. Each entry = 8 u32 = 32 B. */
__device__ static void blake3_hash_one(const u8 *input, u32 len, u8 *out32)
{
    constexpr int STACK_MAX = 56;
    u32 stack_cv[STACK_MAX][8];
    int stack_depth = 0;

    /* Walk all chunks except the last. Each complete chunk produces a CV
     * pushed onto the stack with the trailing-zeros merge rule. */
    u64 chunk_idx = 0;
    u32 offset = 0;
    /* `last_chunk_start` = offset of the last (possibly partial) chunk. */
    u32 last_chunk_start;
    if (len == 0) {
        last_chunk_start = 0;
    } else if (len % CHUNK_LEN == 0) {
        last_chunk_start = len - CHUNK_LEN;
    } else {
        last_chunk_start = (len / CHUNK_LEN) * CHUNK_LEN;
    }

    while (offset < last_chunk_start) {
        u32 cv8[8];
        chunk_cv_complete(input + offset, chunk_idx, cv8);

        /* add_chunk_chaining_value: merge while total_chunks (chunk_idx+1)
         * has trailing zero bits. Each merge pops one CV, parents with the
         * incoming cv8, and the result becomes the new cv8. */
        u64 total_chunks = chunk_idx + 1;
        while ((total_chunks & 1ULL) == 0ULL) {
            u32 left[8];
            for (int i = 0; i < 8; i++) left[i] = stack_cv[stack_depth - 1][i];
            stack_depth--;

            u32 block_words[16];
            for (int i = 0; i < 8; i++) block_words[i]     = left[i];
            for (int i = 0; i < 8; i++) block_words[8 + i] = cv8[i];
            u32 parent_cv[8];
            for (int i = 0; i < 8; i++) parent_cv[i] = IV[i];
            u32 state[16];
            compress(parent_cv, block_words, 0ULL, BLOCK_LEN, PARENT, state);
            for (int i = 0; i < 8; i++) cv8[i] = state[i];
            total_chunks >>= 1;
        }
        /* push cv8 */
        for (int i = 0; i < 8; i++) stack_cv[stack_depth][i] = cv8[i];
        stack_depth++;

        chunk_idx++;
        offset += CHUNK_LEN;
    }

    /* Build the pending Output for the final chunk. */
    Pending p;
    chunk_pending(input + last_chunk_start, len - last_chunk_start, chunk_idx, p);

    /* Finalize: fold the stack into Output right-to-left. Each fold computes
     * the chaining value of the current Output (non-root) and uses it as
     * the right child of a parent Output. The final Output (no more stack
     * entries) gets ROOT applied when its compress runs. */
    while (stack_depth > 0) {
        /* compute current Output's chaining value (non-root) */
        u32 child_cv[8];
        {
            u32 state[16];
            compress(p.input_cv, p.block_words, p.counter, p.block_len, p.flags, state);
            for (int i = 0; i < 8; i++) child_cv[i] = state[i];
        }
        /* pop left from stack, build parent Output */
        u32 left[8];
        for (int i = 0; i < 8; i++) left[i] = stack_cv[stack_depth - 1][i];
        stack_depth--;
        for (int i = 0; i < 8; i++) p.block_words[i]     = left[i];
        for (int i = 0; i < 8; i++) p.block_words[8 + i] = child_cv[i];
        for (int i = 0; i < 8; i++) p.input_cv[i]        = IV[i];
        p.counter   = 0ULL;
        p.block_len = BLOCK_LEN;
        p.flags     = PARENT;
    }

    /* Apply ROOT and compress. First 32 B of root output = state[0..8] LE. */
    u32 state[16];
    compress(p.input_cv, p.block_words, 0ULL, p.block_len, p.flags | ROOT, state);
    for (int i = 0; i < 8; i++) {
        u32 w = state[i];
        out32[4 * i + 0] = (u8)(w & 0xFF);
        out32[4 * i + 1] = (u8)((w >> 8) & 0xFF);
        out32[4 * i + 2] = (u8)((w >> 16) & 0xFF);
        out32[4 * i + 3] = (u8)((w >> 24) & 0xFF);
    }
}

__global__ void blake3_fanout_kernel(const u8 *inputs,
                                     const u32 *offsets,
                                     const u32 *lengths,
                                     u8 *outputs,
                                     u32 out_bytes,
                                     u32 n)
{
    u32 t = blockIdx.x * blockDim.x + threadIdx.x;
    if (t >= n) return;
    /* We only realize the first 32 bytes of the root output. */
    u8 digest[32];
    blake3_hash_one(inputs + offsets[t], lengths[t], digest);
    u32 to_copy = out_bytes < 32 ? out_bytes : 32;
    for (u32 i = 0; i < to_copy; i++) {
        outputs[(size_t)t * out_bytes + i] = digest[i];
    }
    for (u32 i = 32; i < out_bytes; i++) {
        /* out_bytes > 32: zero-pad. Future extension can XOF here. */
        outputs[(size_t)t * out_bytes + i] = 0;
    }
}

/* ─── host: portal I/O + driver ─────────────────────────────────────── */

static char *slurp(const char *path, size_t *len_out) {
    FILE *f = fopen(path, "rb");
    if (!f) { perror(path); exit(1); }
    fseek(f, 0, SEEK_END); long n = ftell(f); fseek(f, 0, SEEK_SET);
    char *buf = (char*)malloc(n + 1);
    if (fread(buf, 1, n, f) != (size_t)n) { perror(path); exit(1); }
    buf[n] = 0; fclose(f);
    if (len_out) *len_out = n;
    return buf;
}

static int hex_to_bytes(const char *hex, size_t hex_len, uint8_t *out) {
    if (hex_len % 2 != 0) return 0;
    for (size_t i = 0; i < hex_len; i += 2) {
        unsigned v;
        if (sscanf(hex + i, "%2x", &v) != 1) return 0;
        out[i / 2] = (uint8_t)v;
    }
    return 1;
}

static void bytes_to_hex(const uint8_t *bytes, size_t n, char *out) {
    static const char *H = "0123456789abcdef";
    for (size_t i = 0; i < n; i++) {
        out[2 * i]     = H[bytes[i] >> 4];
        out[2 * i + 1] = H[bytes[i] & 0xF];
    }
    out[2 * n] = 0;
}

static int parse_input_portal(const char *src, uint32_t *out_bytes,
                              char ***hex_inputs, uint32_t *n_inputs)
{
    *out_bytes = 32;
    const char *p = strstr(src, "(output-bytes");
    if (p) sscanf(p, "(output-bytes %u)", out_bytes);
    p = strstr(src, "(inputs");
    if (!p) return 0;
    char **list = NULL;
    uint32_t count = 0, cap = 0;
    p += strlen("(inputs");
    while ((p = strchr(p, '"')) != NULL) {
        const char *end = strchr(p + 1, '"');
        if (!end) break;
        if (count == cap) {
            cap = cap ? cap * 2 : 16;
            list = (char**)realloc(list, cap * sizeof(char*));
        }
        size_t L = end - (p + 1);
        list[count] = (char*)malloc(L + 1);
        memcpy(list[count], p + 1, L); list[count][L] = 0;
        count++;
        p = end + 1;
    }
    *hex_inputs = list;
    *n_inputs = count;
    return count > 0;
}

static int process_one(const char *in_path, const char *out_path);
static int process_one_bin(const char *in_path, const char *out_path);

int main(int argc, char **argv) {
    if (argc >= 2 && !strcmp(argv[1], "--daemon")) {
        int dev_count = 0;
        cudaGetDeviceCount(&dev_count);
        if (dev_count == 0) {
            fprintf(stderr, "no CUDA devices\n"); return 1;
        }
        cudaFree(0);  /* forces context init */
        fprintf(stdout, "ready\n"); fflush(stdout);

        char line[4096];
        while (fgets(line, sizeof(line), stdin)) {
            char *nl = strchr(line, '\n');
            if (nl) *nl = 0;
            if (!strcmp(line, "quit")) {
                fprintf(stdout, "bye\n"); fflush(stdout);
                break;
            }
            int is_bin = 0;
            char *cmd_args = NULL;
            if (!strncmp(line, "process ", 8))            { cmd_args = line + 8;  is_bin = 0; }
            else if (!strncmp(line, "process-bin ", 12))  { cmd_args = line + 12; is_bin = 1; }
            if (cmd_args) {
                char *in_path = cmd_args;
                char *sep = strchr(in_path, ' ');
                if (!sep) {
                    fprintf(stdout, "error bad-args\n"); fflush(stdout);
                    continue;
                }
                *sep = 0;
                char *out_path = sep + 1;
                int r = is_bin ? process_one_bin(in_path, out_path)
                               : process_one(in_path, out_path);
                if (r == 0) fprintf(stdout, "done %s\n", out_path);
                else        fprintf(stdout, "error process-failed\n");
                fflush(stdout);
            } else if (line[0]) {
                fprintf(stdout, "error unknown-command\n"); fflush(stdout);
            }
        }
        return 0;
    }

    if (argc >= 4 && !strcmp(argv[1], "--binary")) {
        return process_one_bin(argv[2], argv[3]);
    }
    if (argc < 3) {
        fprintf(stderr,
            "usage: %s <input.portal> <output.portal>\n"
            "       %s --binary <input.bin> <output.bin>\n"
            "       %s --daemon  (read commands on stdin)\n",
            argv[0], argv[0], argv[0]);
        return 1;
    }
    return process_one(argv[1], argv[2]);
}

/* Binary format (mirrors shake256-fanout layout, distinct magic):
 *   in:  "BSB3" | u32 out_bytes | u32 n | (u32 len | len bytes) × n
 *   out: "BSR3" | u32 n         | u32 out_bytes | (out_bytes bytes) × n
 */
static int process_one_bin(const char *in_path, const char *out_path) {
    struct timespec t0, t1, t2, t3, t4;
    clock_gettime(CLOCK_MONOTONIC, &t0);

    FILE *f = fopen(in_path, "rb");
    if (!f) { perror(in_path); return 1; }
    fseek(f, 0, SEEK_END); long sz = ftell(f); fseek(f, 0, SEEK_SET);
    uint8_t *buf = (uint8_t*)malloc(sz);
    if (fread(buf, 1, sz, f) != (size_t)sz) { perror(in_path); return 1; }
    fclose(f);

    if (sz < 12 || memcmp(buf, "BSB3", 4) != 0) {
        fprintf(stderr, "%s: bad magic (want BSB3)\n", in_path); return 1;
    }
    uint32_t out_bytes, n;
    memcpy(&out_bytes, buf + 4, 4);
    memcpy(&n,         buf + 8, 4);

    uint32_t *offsets = (uint32_t*)malloc(n * sizeof(uint32_t));
    uint32_t *lengths = (uint32_t*)malloc(n * sizeof(uint32_t));
    size_t cur = 12;
    size_t data_off = 0;
    for (uint32_t i = 0; i < n; i++) {
        if (cur + 4 > (size_t)sz) {
            fprintf(stderr, "%s: truncated at input %u\n", in_path, i); return 1;
        }
        uint32_t L; memcpy(&L, buf + cur, 4); cur += 4;
        if (cur + L > (size_t)sz) {
            fprintf(stderr, "%s: bad length at input %u\n", in_path, i); return 1;
        }
        offsets[i] = (uint32_t)data_off;
        lengths[i] = L;
        data_off += L;
        cur += L;
    }
    size_t total = data_off;
    uint8_t *flat = (uint8_t*)malloc(total ? total : 1);
    cur = 12;
    size_t fcur = 0;
    for (uint32_t i = 0; i < n; i++) {
        uint32_t L = lengths[i];
        cur += 4;
        memcpy(flat + fcur, buf + cur, L);
        cur += L; fcur += L;
    }
    free(buf);
    clock_gettime(CLOCK_MONOTONIC, &t1);

    uint8_t *d_in; uint32_t *d_off, *d_len; uint8_t *d_out;
    CUDA_CHECK(cudaMalloc(&d_in,  total ? total : 1));
    CUDA_CHECK(cudaMalloc(&d_off, n * sizeof(uint32_t)));
    CUDA_CHECK(cudaMalloc(&d_len, n * sizeof(uint32_t)));
    CUDA_CHECK(cudaMalloc(&d_out, (size_t)n * out_bytes));
    CUDA_CHECK(cudaMemcpy(d_in,  flat,    total ? total : 1,    cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_off, offsets, n * sizeof(uint32_t), cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_len, lengths, n * sizeof(uint32_t), cudaMemcpyHostToDevice));
    clock_gettime(CLOCK_MONOTONIC, &t2);

    int threads = 64;
    int blocks  = (n + threads - 1) / threads;
    blake3_fanout_kernel<<<blocks, threads>>>(d_in, d_off, d_len, d_out, out_bytes, n);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
    clock_gettime(CLOCK_MONOTONIC, &t3);

    uint8_t *h_out = (uint8_t*)malloc((size_t)n * out_bytes);
    CUDA_CHECK(cudaMemcpy(h_out, d_out, (size_t)n * out_bytes, cudaMemcpyDeviceToHost));

    FILE *fo = fopen(out_path, "wb");
    if (!fo) { perror(out_path); return 1; }
    fwrite("BSR3", 1, 4, fo);
    fwrite(&n, 4, 1, fo);
    fwrite(&out_bytes, 4, 1, fo);
    fwrite(h_out, 1, (size_t)n * out_bytes, fo);
    fclose(fo);
    clock_gettime(CLOCK_MONOTONIC, &t4);

    double ld_ms     = (t1.tv_sec - t0.tv_sec) * 1e3 + (t1.tv_nsec - t0.tv_nsec) / 1e6;
    double launch_ms = (t2.tv_sec - t1.tv_sec) * 1e3 + (t2.tv_nsec - t1.tv_nsec) / 1e6;
    double kern_ms   = (t3.tv_sec - t2.tv_sec) * 1e3 + (t3.tv_nsec - t2.tv_nsec) / 1e6;
    double write_ms  = (t4.tv_sec - t3.tv_sec) * 1e3 + (t4.tv_nsec - t3.tv_nsec) / 1e6;
    fprintf(stderr, "  bin: load=%.1f launch=%.1f kernel=%.2f write=%.1f ms (n=%u, %.0f MB)\n",
            ld_ms, launch_ms, kern_ms, write_ms, n, total / 1e6);

    cudaFree(d_in); cudaFree(d_off); cudaFree(d_len); cudaFree(d_out);
    free(offsets); free(lengths); free(flat); free(h_out);
    return 0;
}

static int process_one(const char *in_path, const char *out_path) {
    struct timespec t0, t1, t2, t3, t4;
    clock_gettime(CLOCK_MONOTONIC, &t0);

    size_t portal_len;
    char *portal = slurp(in_path, &portal_len);
    uint32_t out_bytes = 32;
    char **hex_in = NULL; uint32_t n = 0;
    if (!parse_input_portal(portal, &out_bytes, &hex_in, &n)) {
        fprintf(stderr, "no (inputs ...) found in %s\n", in_path);
        return 1;
    }

    uint32_t *offsets = (uint32_t*)malloc(n * sizeof(uint32_t));
    uint32_t *lengths = (uint32_t*)malloc(n * sizeof(uint32_t));
    size_t total = 0;
    for (uint32_t i = 0; i < n; i++) {
        size_t hex_len = strlen(hex_in[i]);
        if (hex_len % 2 != 0) {
            fprintf(stderr, "input %u has odd hex length\n", i); return 1;
        }
        offsets[i] = total;
        lengths[i] = hex_len / 2;
        total += hex_len / 2;
    }
    uint8_t *flat = (uint8_t*)malloc(total ? total : 1);
    for (uint32_t i = 0; i < n; i++) {
        if (!hex_to_bytes(hex_in[i], strlen(hex_in[i]), flat + offsets[i])) {
            fprintf(stderr, "bad hex at input %u\n", i); return 1;
        }
    }
    clock_gettime(CLOCK_MONOTONIC, &t1);

    uint8_t  *d_in;  uint32_t *d_off, *d_len; uint8_t *d_out;
    CUDA_CHECK(cudaMalloc(&d_in,  total ? total : 1));
    CUDA_CHECK(cudaMalloc(&d_off, n * sizeof(uint32_t)));
    CUDA_CHECK(cudaMalloc(&d_len, n * sizeof(uint32_t)));
    CUDA_CHECK(cudaMalloc(&d_out, (size_t)n * out_bytes));
    CUDA_CHECK(cudaMemcpy(d_in,  flat,    total ? total : 1,    cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_off, offsets, n * sizeof(uint32_t), cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_len, lengths, n * sizeof(uint32_t), cudaMemcpyHostToDevice));
    clock_gettime(CLOCK_MONOTONIC, &t2);

    int threads = 64;
    int blocks  = (n + threads - 1) / threads;
    blake3_fanout_kernel<<<blocks, threads>>>(d_in, d_off, d_len, d_out, out_bytes, n);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
    clock_gettime(CLOCK_MONOTONIC, &t3);

    uint8_t *h_out = (uint8_t*)malloc((size_t)n * out_bytes);
    CUDA_CHECK(cudaMemcpy(h_out, d_out, (size_t)n * out_bytes, cudaMemcpyDeviceToHost));

    FILE *f = fopen(out_path, "w");
    if (!f) { perror(out_path); return 1; }
    fprintf(f, "(cuda-blake3-tree-result\n");
    fprintf(f, "  (n %u)\n", n);
    fprintf(f, "  (output-bytes %u)\n", out_bytes);
    fprintf(f, "  (hashes\n");
    char *hex_buf = (char*)malloc(out_bytes * 2 + 1);
    for (uint32_t i = 0; i < n; i++) {
        bytes_to_hex(h_out + (size_t)i * out_bytes, out_bytes, hex_buf);
        fprintf(f, "    \"%s\"\n", hex_buf);
    }
    fprintf(f, "  )\n");
    clock_gettime(CLOCK_MONOTONIC, &t4);
    double ld_ms     = (t1.tv_sec - t0.tv_sec) * 1e3 + (t1.tv_nsec - t0.tv_nsec) / 1e6;
    double launch_ms = (t2.tv_sec - t1.tv_sec) * 1e3 + (t2.tv_nsec - t1.tv_nsec) / 1e6;
    double kern_ms   = (t3.tv_sec - t2.tv_sec) * 1e3 + (t3.tv_nsec - t2.tv_nsec) / 1e6;
    double write_ms  = (t4.tv_sec - t3.tv_sec) * 1e3 + (t4.tv_nsec - t3.tv_nsec) / 1e6;
    fprintf(f, "  (timing-ms (load %.3f) (launch %.3f) (kernel %.3f) (write %.3f)))\n",
            ld_ms, launch_ms, kern_ms, write_ms);
    fclose(f);

    cudaFree(d_in); cudaFree(d_off); cudaFree(d_len); cudaFree(d_out);
    for (uint32_t i = 0; i < n; i++) free(hex_in[i]);
    free(hex_in); free(offsets); free(lengths); free(flat); free(h_out);
    free(hex_buf); free(portal);
    return 0;
}
