Upload the frame path's metadata in one transfer

[iderex]

What & why

Closes #134.

DecodeAndAssemble allocated four device metadata arrays and filled them with
four separate blocking cudaMemcpy calls. They now live in one allocation as
four equal-stride sections and go up in one transfer, which is the layout
src/stream.cpp already uses (kMetaStride / kMetaBytes at lines 44-48, the
single copy at 295-298). The frame path predates that work and never picked it
up.

No behavioural change. The same pointers, sizes and capacities reach
cudec_lz4_decompress_batch in the same order; only the number of submissions
moves, from four to one.

The means is C++ inside the existing host translation unit, because the thing
being changed is the transfer shape of that function and the layout it copies
is three files away in the same language.

Type of change

• Performance
• Decode kernel / device code
• Format support (LZ4 / Snappy / GDeflate / Zstd)
• API surface
• Bug fix
• Refactor / code quality
• Build / CI / supply chain
• Docs

Engineering checklist

• Fail-closed preserved. The n != 0 guard and the n * block_max
  overflow check are untouched. Packing multiplies the metadata request by
  four, so that product gets the same treatment before the driver sees it.
• Oracle coverage: frame_twin decodes against liblz4's frame API and is
  green, unchanged.
• Determinism preserved. determinism_gpu green.
• Every CUDA call goes through FRAME_CUDA, as before. No exceptions cross
  the ABI.
• Adversarial review: read against the four lenses by one reader. See the
  disclosure at the end.
• Review record below.

The new guard n > SIZE_MAX / (4 * sizeof(void*)) is stated as unproven rather
than as proven. No input the frame parser admits can reach it: n counts
compressed blocks in a frame already bounded by the caller's byte count, so a
value above SIZE_MAX / 32 cannot be constructed and no test makes it bite.
It sits beside the identical unreachable guard on n * block_max that was
already there, for the same arithmetic reason, and both are here so the
multiplication is checked rather than assumed.

static_assert(sizeof(size_t) == sizeof(void*)) carries the alignment
constraint the issue asked for: every section gets a void*-sized stride,
including the two holding size_t, so a host where the two differ is a
compile error rather than a silent misalignment of the two size sections.
It is unproven for the same class of reason as the guard above and the reason
is worth being exact about, because it is not the reason it looks like.
Rewriting the stride expression does not trip it: the assertion compares two
sizeof results and says nothing about how the stride is spelled. What trips
it is a host where the two types differ in width, and no compiler in this
gate is one. What it is worth is stated as what it is, a refusal that fires on
a port rather than a guard this suite can redden.

Review record

One reader, four lenses, over the whole diff. CONFIRMED 0 / PLAUSIBLE 2.

• Correctness. The four sections are written into the staging buffer in the
  order the four casts read them back out, checked element by element against
  the pre-image. Section offsets are multiples of 8 and the allocation base is
  a cudaMalloc result, so every section is naturally aligned. No finding.
• Robustness. PLAUSIBLE, disposition DECLINE with the reason above: a guard
  that cannot be shown to bite. Declined rather than removed because the
  alternative is an unchecked multiplication.
• Performance. PLAUSIBLE: the host now does four memcpy calls into a staging
  vector that it did not do before, and that vector is a fresh allocation per
  decode. Disposition DECLINE, measured: the four host copies move at most
  1248 bytes each at the largest block count in the sweep, against the 71 to 86
  microseconds of submission cost removed. The numbers are below.
• Integration. The ABI is untouched, the kernel arguments are identical, and
  dd_res stays its own allocation because it is a device-to-host destination
  and not metadata. No finding.

GPU sanitizer gate

• Not applicable: this change touches no device code. No .cu file, no
  kernel and no launch configuration is modified.

Separately, and so the empty block is not read as an answered one: #258 records
that no route to a device the sanitizer can attach to is available on this
machine. Under WSL2 the device is in WDDM mode, the debugger interface is
absent, and all four tools report the same two initialization errors on a
program with no fault in it.

Performance checklist

• Measured, not reasoned.
• No regression against docs/BENCHMARKS.md. That document records the M1
  block path and the M2 batch path; it carries no frame-path baseline, so
  there is no recorded number for this change to move.

Both sides were built and run in the pinned container against the local RTX
3080, one after the other:

docker run --rm --gpus all -v "<tree>:/w" -v "<corpora>:/corpora:ro" \
  -w /w nvidia/cuda:12.6.2-devel-ubuntu24.04 sh -c \
  "apt-get update -q >/dev/null && apt-get install -yq cmake >/dev/null 2>&1 \
   && cmake -B build-cuda -DCUDEC_ENABLE_CUDA=ON && cmake --build build-cuda -j 12 \
   && ./build-cuda/bench/bench_lz4 --frame --runs 30 --warmup 3 /corpora/silesia/dickens"

RTX 3080 (sm_86), driver 12.6, runtime 12.6, host AMD Ryzen 9 5950X, liblz4
1.10.0, 3 warmup plus 30 measured runs, wall clock around the whole
synchronous cudec_lz4f_decompress call.

p50 wall per run, before at 3fb2c55 and after at this branch's head:

corpus	block max	blocks	before	after	delta
dickens	64 KB	156	21.364 ms	21.734 ms	+1.7%
dickens	256 KB	39	31.331 ms	32.460 ms	+3.6%
dickens	1 MB	10	82.912 ms	82.787 ms	-0.2%
xml	64 KB	82	12.725 ms	12.056 ms	-5.3%
xml	256 KB	21	21.600 ms	20.905 ms	-3.2%
xml	1 MB	6	55.154 ms	53.606 ms	-2.8%

Three rungs faster, three slower, and the signs do not follow the block count.
That is the honest reading: this sweep cannot resolve the change. Its own
within-run spread says so, p50 to p99 running from 2.0% to 14.8% across the
twelve report blocks, on one binary and one set of bytes, which is wider than
every delta in the table.

So the mechanism was measured where it is visible, on the same device, with a
program that does the two transfer shapes and nothing else. Source and
invocation, so the numbers can be rebuilt rather than trusted:

/* four separate H2D copies of n*8 bytes each, against one copy of 4*n*8,
   at the block counts the sweep decodes; the two shapes swap order every
   iteration so a warm-up artifact cannot read as the win; 2000 iterations,
   median reported. */
for (size_t n : {6, 10, 21, 39, 82, 156}) {
    const size_t stride = n * sizeof(void*);
    std::vector<unsigned char> host(4 * stride, 0x5a);
    void* dev; cudaMalloc(&dev, 4 * stride);
    /* timed A: four cudaMemcpy(d + s*stride, host + s*stride, stride, H2D) */
    /* timed B: one  cudaMemcpy(d, host, 4*stride, H2D)                     */
}

nvcc -O2 -arch=sm_86 -o submit submit.cu && ./submit

blocks	four copies	one copy	removed
6	113.86 us	27.57 us	86.29 us
10	95.10 us	23.81 us	71.29 us
21	110.88 us	25.54 us	85.34 us
39	106.54 us	23.38 us	83.16 us
82	104.19 us	23.64 us	80.55 us
156	103.84 us	23.60 us	80.24 us

Three submissions, 71 to 86 microseconds, near enough flat in the block count
because it is per-call cost and not per-byte cost. Against the walls in the
first table that is 0.10% at the slowest rung and 0.71% at the fastest, which
is why the sweep swallows it.

One correction to the expectation written on the issue, which said the win is
largest as a fraction on few-block frames. On this harness it is the other way
round. Fewer blocks means less parallelism and a longer wall, so the 1 MB rungs
are the slowest and the fixed saving is the smallest fraction of them. The
fraction is largest on the many-block 64 KB rungs.

Quality checklist

• Minimal: one function, four allocations to one, four transfers to one.
• Self-documenting: the comment says why the stride is written in units of
  void* and what the assertion is holding, not what the memcpy does.
• Conformance tests pass. The change establishes no new structural property
  to lock.

Verification

Full gate in the pinned container against the local RTX 3080, at this
branch's head:

cmake -B build-cuda -DCUDEC_ENABLE_CUDA=ON && cmake --build build-cuda -j 12
ctest --test-dir build-cuda --output-on-failure

100% tests passed, 0 tests failed out of 36
Label Time Summary:
gpu    =   5.20 sec*proc (8 tests)
Total Test time (real) =   9.17 sec

• Builds clean, -Wall -Wextra -Werror and the strict device flags.
• ctest green, 36 of 36, including the eight GPU-labelled entries
  (frame_twin, stream_twin, determinism_gpu, termination_gpu,
  bench_frame_selfcheck among them).
• Prettier clean. No .md, .yml or .yaml file is touched.
• Docs synced. No behaviour, API or recorded number moves.

The same gate was run at 3fb2c55 for the before column, also 36 of 36.

Notes

This change had no second reader. The review record above is one reader over
four lenses, and the pasted gate output and the two measurements are the
evidence in place of one.