FP16 KV cache with packed half2 split-KV attention

[orionpapadakis]

Description

Adds an opt-in half-precision (FP16) KV cache for the NVIDIA/CUDA decode path,
halving KV-cache bandwidth while keeping accumulation in FP32. On top of it, a
packed-half2 split-KV attention path ("deep-half2") stages the query as a
__half2 local-memory tile and consumes each K/V pair with a single __hfma2,
converting to FP32 once per row (llama.cpp fattn-vec style) instead of per pair.

Everything is flag-gated and additive: the FP32 KV cache remains the default and
the fallback for non-NVIDIA backends and unsupported models, so existing behaviour
is unchanged unless a flag is set.

Scope

• Models: Llama and Qwen3.
• Path: single-token decode (and the decode/batch-decode task graphs).
• Backend: NVIDIA/CUDA only — the packed kernels rely on half2 codegen in
  the CUDA backend; other backends transparently keep the FP32 cache.
• FP32 KV cache is still allocated and used as the fallback; Q8_0 is unaffected.

Flags

Flag	Default	Effect
-Dllama.kvcache.fp16=true	off	Allocate and use the FP16 KV cache on the CUDA decode path.
-Dllama.kvcache.fp16.scalar=true	off	Read the FP16 cache with scalar half loads instead of packed half2 (isolates the packed-load gain).
-Dllama.attention.deepHalf2=true	off	Keep the K·Q score accumulation packed (__hfma2); requires the FP16 cache.
-Dllama.attention.splitKv.count=N	8	Split-KV (flash-decoding) partition count.
-Dllama.bench.ignoreEos=true	off	Benchmarking aid: keep decoding past the stop token for fixed-length runs.

Dependency

The deep-half2 path uses the packed half2 local-memory API
(allocateHalf2LocalArray, Half2.fma/lowFloat/highFloat) added in the
TornadoVM feat/packed-half2 branch, so pom.xml targets
tornadovm.base.version = 5.1.1-jdk21-dev.

This PR does not build against a released TornadoVM until that branch is merged
and published. Please treat as dependent on the corresponding TornadoVM PR.

Testing

Built with JDK 21 against a local TornadoVM 5.1.1-jdk21-dev (CUDA backend) and
run on an RTX 5090 (Linux).

Baseline = the implementation before FP16 KV / half2. Enable the path with:

# baseline (FP32 KV)
llama-tornado --model <llama-or-qwen3-model> --prompt "..." --max-tokens 2048

# FP16 KV + packed split-KV attention
llama-tornado --model <llama-or-qwen3-model> --prompt "..." --max-tokens 2048 \
    -Dllama.kvcache.fp16=true -Dllama.attention.deepHalf2=true

Output remains coherent with the flags enabled; FP32 remains the default when they
are not set.

Throughput

Standard single-token decode path (no prefill/decode), 2048 tokens per run at equal
length (-Dllama.bench.ignoreEos=true), fixed seed and prompt, RTX 5090 Laptop.
Baseline = same binary with the flags unset (FP32 KV cache). Decode throughput (tok/s):

Model	baseline	-Dllama.kvcache.fp16	+ -Dllama.attention.deepHalf2
Llama-3.2-1B	47.4	72.8 (+53.6%)	71.8 (+51.5%)
Llama-3.2-3B	13.9	19.3 (+39.1%)	19.0 (+37.0%)
Qwen3-0.6B	62.1	61.2 (−1.4%)	62.7 (+0.9%)
Qwen3-1.7B	51.1	52.8 (+3.2%)	54.4 (+6.4%)
Qwen3-4B	32.2	34.0 (+5.4%)	34.5 (+6.9%)

• Llama sees a large gain (+39–54%): KV-cache bandwidth is a real decode bottleneck
  there, so halving it helps directly.
• Qwen3 gains are modest (~0–7%, growing with model size): its decode is heavier per
  token (QK-norm, always-on split-KV attention), so the KV-cache byte share is smaller;
  0.6B is essentially launch-overhead-bound.
• The packed deep-half2 accumulation tracks plain FP16-KV within run-to-run noise at
  this depth — slightly below on Llama, slightly above on Qwen3 — so it is kept as an
  opt-in refinement rather than folded into the default FP16-KV path.

The batched-prefill/decode path (--with-prefill-decode --batch-prefill-size 32,
decode-phase tok/s) shows the same behaviour:

Model	baseline	-Dllama.kvcache.fp16	+ -Dllama.attention.deepHalf2
Llama-3.2-1B	45.2	68.1 (+50.6%)	69.2 (+52.9%)
Llama-3.2-3B	13.2	18.3 (+39.0%)	18.5 (+40.2%)
Qwen3-0.6B	56.5	59.9 (+6.1%)	60.5 (+7.2%)
Qwen3-1.7B	50.9	52.6 (+3.3%)	51.4 (+1.0%)
Qwen3-4B	31.3	32.5 (+3.6%)	32.7 (+4.4%)

(Qwen3-4B batched needs a larger device-memory budget, e.g. --gpu-memory 22GB, than
the 14 GB default; the standard-path run fits in the default.)