Add support for Qwen3.8 27B model CPU + GPU - #150
Open
orionpapadakis wants to merge 110 commits into
Open
Conversation
Qwen3.8-27B declares general.architecture=qwen35: a hybrid stack where only every fourth layer is attention and the remaining 48 of 64 are Gated Delta Net linear attention, plus one MTP/NextN block. Records the inventory (metadata, tensor table, both layer geometries, the quantization mix), walks the extension points, and justifies the four OperationKind values the mixer needs -- L2_NORM, CAUSAL_CONV_1D, DELTA_RULE_UPDATE, GATED_NORM -- against the alternatives, including the family-shaped composite the porting skill warns about. Also records why the GPU path is blocked on memory rather than kernels: materializing this file's Q4_0 weights as Q8_0 needs ~28GB.
Qwen3.8-27B-Q4_0.gguf is Q4_0 throughout except the first eight layers' ffn_down, which the quantizer emitted as Q4_1 -- so the file could not be loaded at all without it. Q4_1 is Q4_0's affine sibling: 32 values to a block, unsigned nibbles, and a per-block minimum alongside the scale, reconstructing d*q + m. The vectorized dot product distributes over that form, carrying the running products and the running activations and scaling each by its own block parameter rather than reconstructing weights. Format-decoded like Q4_0: no target materializes it, and the GPU path maps it to Q8_0 at load, which ForwardPlanFactory now states by name. Also clarifies in the port proposal that the architecture is qwen35 while the model is Qwen3.8 -- the GGUF's declared architecture is what names the family, and llama.cpp's LLM_ARCH_QWEN35 covers 3.5/3.6/3.8.
Adds the qwen35 architecture: 64 trunk layers of which only every fourth attends, the other 48 mixing with a Gated Delta Net recurrence, plus the MTP block the file carries but the trunk does not execute. Four operations enter the shared vocabulary, because the arithmetic is not this family's -- every recurrent architecture needs it: L2_NORM unit-length scaling, which RMS_NORM is not CAUSAL_CONV_1D depthwise causal convolution over a retained window DELTA_RULE_UPDATE decay, correct, accumulate, read back GATED_NORM rms_norm(x, w) * silu(gate) RoPE gains a partial form: this model states rope.dimension_count = 64 against a 256-wide head, so three quarters of every head is unrotated. Partial rotation is a parameter of RoPE, not a second scheme. Recognition needed no case: the file declares general.architecture = qwen35 and nothing else claims that name, so a provider and a service line were the whole of it. The recurrent state is session state, not KV storage -- fixed size, unpageable, and with no position mask, which is why State gains resetSequenceState(): a delta-net matrix that survives a reset silently conditions the new sequence on the old one. Key/value caches are allocated only for the layers that attend, sparing 48 unused ones. No accelerator claims the architecture. Both the loader and the model refuse a device by name rather than falling back to the host path, which would report GPU throughput for CPU work. Verified on Qwen3.8-27B-Q4_0.gguf: 'What is the capital of France? Answer in one word.' answers Paris.
The delta rule has 48 value heads against 16 key heads. The reference
repeats the key heads with ggml_repeat_4d, which *tiles* -- cycling
0,1,...,15,0,1,... -- and the fused kernel states the same mapping
directly as iv1 % nek1. Dividing instead blocks them 0,0,0,1,1,1,...,
which pairs every value head with the wrong key.
The symptom is the one this defect class always has: short answers stay
correct ('capital of France' still answered Paris) while long output
decays into a repetition loop, because the error compounds through the
recurrence rather than failing outright.
The equivalence test passed through all of this, because its reference
carried the same misreading -- exactly the trap the porting skill names:
an invalid fixture fails both paths identically and looks like
agreement. Both sides are corrected, and the reference now says in a
comment which ordering is right and why, so the next reader cannot
re-derive the wrong one from the divisibility alone.
l2Norm also now floors the divisor at eps rather than adding eps under
the root, which is what ggml_l2_norm does.
Adds the checks that were missing:
- Q4_1FloatTensorTest, against hand-encoded blocks, pinning the
unsigned nibble and the affine reconstruction
- Qwen35ConfigurationTest, over Qwen3.8-27B's real metadata, pinning
the derived delta-net widths and that the head width is stated
rather than derived
- the fixture's SHA-256 in GoldenFixture
Verified: 'Write three sentences about the Roman Empire' now produces a
coherent reasoning block and three correct sentences, and stops.
Qwen3.8-27B carries a NextN block past its 64 trunk layers: a complete attention decoder block whose input is not the previous layer's output but the pair (the trunk's hidden state at position p, the token chosen for p+1), each through its own norm, concatenated and projected back to dim. It predicts the token at p+2. Neither nextn.embed_tokens nor nextn.shared_head_head is present in this file, so both fall back to the trunk's, as llama.cpp does when they are absent. Generation through the head is behind -Dllama.qwen35.speculative, off by default, and the reason is stated rather than hedged: an accepted draft only saves work on a backend that verifies several positions in one forward pass, and this host path verifies them one at a time. What the default-off path does deliver is a measurement -- -Dllama.qwen35.speculative.stats reports how often the head agreed with the trunk, which is the only way to tell a correctly fed draft head from a badly fed one: the trunk's token is what gets emitted either way, so no amount of fluent output would show the difference. The head gets its own logits buffer and its own residual stream. Sharing the trunk's would let a draft overwrite the prediction the loop is about to commit, and would move a stream that a question about the future must leave where it is. Qwen35MtpTest asserts both, plus the block itself against an independently written reference -- the concatenation order and which hidden state is consumed being the two things that change nothing about any shape when they are wrong.
Adding four OperationKind values and a DataType left five gates
failing, each of them correctly:
- Operation is sealed and every kind must have exactly one type, so
L2Norm, CausalConv1d, DeltaRuleUpdate and GatedNorm now exist as
descriptions, not only as host implementations. DeltaRuleUpdate
takes keyHeads rather than a ratio on purpose: a ratio invites the
division that produces the wrong head pairing.
- CausalConv1d and DeltaRuleUpdate declare their state as an output as
well as an input. A description that did not would let a backend
advance a window twice, or not at all.
- OperationSupport lists Q4_1 among the representations the host reads
weights in, and lists the four recurrent kinds as unsupported on the
GPU by name rather than leaving them absent.
- The dependency allowlists gain Qwen35 (TornadoVMMasterPlan in the
generateTokensGPU signature, like every other model), Q4_1FloatTensor
(LlamaApp and GGMLType, like Q4_0FloatTensor) and ModelType$12.
549 tests pass; the shaded jar still carries all seven service files,
with the two new providers in them.
verification.md records what was verified and, at more length, the two
defects that produced fluent output -- the tiled key-head mapping whose
test reference shared its misreading, and why the MTP head's acceptance
rate is measured rather than assumed.
These two files were untracked in the working tree before this branch and were swept in by a git add -A. They are not part of this port; this restores them to untracked so they stay the user's own working files.
standalone-inference.yml asserts a resolved backend and a real execution_path on every row; qwen35 has neither, and the smallest release of it is far larger than any fixture that matrix carries. Saying so is the point -- an omission with no stated cause reads as an oversight, and the porting checklist asks for the rows.
…fits Qwen 3.5 abandoned the JSON body Qwen 3 put inside <tool_call> for nested pseudo-XML, one element per argument: <tool_call><function=get_weather><parameter=location>Boston</parameter> </function></tool_call> Reusing Qwen3ChatFormat meant prompting the model for a format it was not trained on and parsing it for one it does not emit -- wrong in both directions, and invisible, because a model that emits calls we cannot read looks exactly like a model that chose not to call anything. Qwen35ToolCalls does the translation at the boundary; nothing above it sees anything but a name and a JSON object. The format erases types -- the template writes strings verbatim and everything else through tojson, so 'Boston' and '42' occupy the same position -- so they are recovered by shape on the way back. That is a heuristic and the class says so: a tool whose string argument is literally '42' round-trips as a number. Quoting everything instead would break every numeric and boolean parameter, which is far commoner. Separately: a caller who asks for no context length now gets 8192 rather than the 262144 this family declares. Sixteen of its layers attend with a 1024-wide key and value, so the declared maximum is 34 GB of host arrays allocated eagerly at session construction -- the public API's ModelOptions.defaults() died on it with an OutOfMemoryError before generating a token, where the CLI survived only because it passes its own default. Scoped to this loader, and the comment says why: the facade's 'the model's own maximum' default is optimistic for every long-context family, but this is the first model where it cannot run at all, and changing that rule for everyone is not this port's call.
The chat template puts consecutive tool results in one user turn; ConversationEncoder calls encodeToolResultTurn once per result, so they become several. No difference for a single result. Merging them means a batched entry point on the shared encoder, which changes every family -- so it is recorded rather than worked around here.
examples.ToolCalling through the public API: the model emitted this
family's pseudo-XML, the call parsed as get_weather({"city":"Athens"}),
the assistant turn replayed, and the final answer used the tool's data.
The one check a unit test cannot stand in for. A model prompted for the
wrong tool-call format answers in prose, which is indistinguishable from
one that decided a tool was unnecessary -- so the parser being correct
says nothing about whether the model is being asked correctly.
Q4_0 was materialized as Q8_0 at load, taking 4.5 bits per weight to
8.5 and roughly doubling what a model occupies on the device. That is
the same problem retaining Q4_K solved for Devstral, and this is the
same solution for the representation that is far commoner.
Measured on one file, switching only -Dllama.q4_0.retain, so the
comparison isolates the representation rather than the quantization:
retained as Q4_0 1570 MiB 172.4 tok/s
materialized Q8_0 2060 MiB 136.0 tok/s
Faster as well as smaller: single-token decode is bandwidth-bound, so
fewer bytes per weight is fewer bytes read per token.
- Q4_0TornadoTensor wraps the file's bytes shallowly, as its Q4_K and
Q8_0 siblings do; nothing is converted at load.
- TransformerComputeKernelsQ4_0 decodes inside the dot product, shaped
like the Q4_K kernels so a layer differs only in which method
reference it names.
- LlamaQ4_0FFNLayers extends the Q8_0 layer and replaces four tasks.
Its feed-forward block applies the RMS norm as its own task rather
than folding it into the gate/up projection -- one more task per
layer, one fewer kernel to keep correct, which is the trade the Q4_K
path already makes.
- The loader retains Q4_0 only when *every* per-layer weight is Q4_0.
All or nothing, because these layers have no per-tensor dispatch: a
retained tensor in a plan whose kernels read Q8_0 blocks would read
18-byte blocks as 34-byte ones, and produce fluent, wrong text.
- The decision follows the per-layer weights, not the output
projection. A Q4_0 file leaves token_embd as Q6_K, so the output
weight alone would say Q8_0 and select the wrong plan.
Q4_0 has single-token kernels only. TornadoPlanRegistry now refuses the
prefill and batch modes for it by name instead of failing on a cast into
an interface the components do not implement -- a general fix, since any
provider may support a mode for one representation and not another.
Q4_0DecodeTest holds the device decode against the host tensor on random
bytes and against the specification on hand-built ones. Agreement alone
would not be enough: the two could agree and both be a different format.
Verified on CUDA with execution_combination llama/Q4_0/STANDARD, and the
Q8_0 and F16 paths are unchanged.
The earlier note said a qwen35 GPU path was blocked by needing ~28GB against 24GB of VRAM. With Q4_0 retained that becomes roughly 17GB -- this file is mostly Q4_0 -- so memory is no longer the objection it was. The delta-net kernels are, and the retention would still have to be wired into that family's loader. Also records why the Q4_0 measurement is an A/B on one file rather than a comparison between a Q4_0 and a Q8_0 model: the easier measurement would have measured the quantization as well as the residency.
The four kernels the qwen35 mixer needs -- depthwise causal convolution with its rolling window, per-head L2 norm, the gated delta rule, and the gated norm -- plus the small one that turns the raw alpha/beta projections into a decay and a write strength. Two decisions shape the file. Every kernel body is a static method taking an explicit lane index, and the kernel is a two-line wrapper passing context.globalIdx. A body written directly against KernelContext cannot be called on the host, so it can only be exercised by running a model on a device, where an indexing mistake surfaces as slightly wrong text rather than as a failure. Lifting the arithmetic out is what lets the parity test run every lane on the host against CpuOperations. TornadoVM inlines it, so it costs nothing at run time. The delta rule needs no barrier and no cross-lane reduction, which is not obvious for what looks like a matrix-vector product per head. Give a lane one value column of a head's state and every quantity it needs is its own: the decayed column, the prediction for that column, the correction, the rank-one update, the readout. The existing state layout then makes it coalesced as well as correct -- per lane the access is strided, but across the lanes of a head it is contiguous, which is what a GPU is paid for. It is the same layout the host uses, deliberately: two layouts would mean the parity test compares a transpose against a transpose and proves nothing about either. Four of six comparisons are bit-exact, including the delta rule's state and the convolution's window. The other two are equal to float rounding because the host evaluates sqrt, log and the logistic in double; that is stated in the test rather than smoothed over with a blanket tolerance. Establishes the arithmetic and the addressing only. That these compile and run on a device, and that a layer graph binds them correctly, are separate gates and are not met.
The three things that separate a qwen35 attention layer from Qwen3's,
none of them a different algorithm:
- the query projection is twice as wide as the query, carrying an
interleaved output gate, split here into contiguous halves so that
the per-head norm, the rotation and attention can all address
head * headDim as they already do;
- the rotary width is 64 of a 256-wide head, so three quarters of
every head passes through unrotated -- a parameter of the rotation,
not a second scheme;
- the attention result is scaled by the logistic of its gate. A
logistic, not a SiLU: reusing the SwiGLU kernel would multiply in an
extra factor of the gate and stay plausible.
The key/value append is its own kernel rather than fused into the
rotation, which is what Llama and Qwen3 do. They can fuse because they
rotate the whole head, so the lanes that rotate are the lanes that must
be written; here the rotation covers a quarter of each head and the
append covers all of it.
The rotation reads the host's precomputed frequency tables rather than
recomputing them with pow and cos, which is what makes it bit-identical
to the host rather than merely close.
Two of the six checks assert a property directly instead of against the
host, because both guard a wrong reading that is also well-formed: the
whole-buffer query/gate split, and a paged append with the wrong stride
landing inside another layer's slice. A comparison in which both sides
made the same assumption would pass.
Also corrects a claim in the class comment. fusedQKRmsNorm is expected
to serve a 256-wide head unchanged, but it reduces through local memory
and cannot run on the host, so nothing has exercised it at that width.
Saying it is reused would have been stating an untested assumption as a
result.
…recurrent Three things about this family's device memory are unlike any other's, and each of them is a size question rather than a correctness one -- which is why they are asserted rather than left to read correctly. Key/value storage covers a quarter of the blocks. Only the attending layers write to it, so they address it by a dense index and the store is sized by them; sizing it by the block count would cost nearly four times as much, which is gigabytes at any useful context. State gains a fillKvFields overload taking that count -- the kernels need no change, since a layer is already just a stride multiplier to them. The recurrent layers hold state that is neither cache nor scratch. The convolution windows and delta-net matrices persist across tokens and are updated in place, so they live in one array per kind addressed by a per-layer offset; 48 device buffers per kind would be 48 transfers to arrange and keep resident. The delta-net and convolution kernels take that offset now, and the parity tests cover a non-zero one -- a lane that dropped it, or applied it to one of its two passes and not the other, would read one layer's state while writing another's and look fine for a while. That state must start at zero on whichever path is running, and a reset must clear both representations. A key/value cache needs neither: attention reads no further than the current position. A recurrence has no such mask. Whether the device arrays are allocated is decided by the use.tornadovm property, which is the facade's default and not the same question as which backend a session resolved. That is adequate only because no plan provider claims this architecture, so nothing can disagree with it; the comment says what has to replace it and when. Zeroing moved behind TornadoWorkspaces rather than calling init on the arrays from the state: naming a TornadoVM type outside the backend is Rule 1, whose allowlist is empty and stays empty. Also records that the memory preflight is materialization-only. It refuses Qwen3.8-27B at 27760 MiB against a 14 GB budget, which is right today and becomes wrong as soon as qwen35 retains Q4_0 -- the retained figure is roughly 17 GB.
…ined Qwen35TornadoWeights and the loader that fills it. Like the host weights, it implements Weights directly rather than extending TornadoWeights, whose fields assume every layer has a query, key and value projection -- here three in four have none. Retention matters more for this model than for any other. Qwen3.8-27B is 16GB and almost entirely Q4_0; materialized as Q8_0 it costs about 28GB, which does not fit on a 24GB device, where retained it is about 17GB, which does. The rest of the file -- ssm_out in Q5_K, eight ffn_down in Q4_1, output in Q6_K -- has no kernel and is still materialized, so the model is mixed and each weight is read by the kernel matching its type. Retention is all-or-nothing over the projections, not per tensor: the graphs dispatch per tensor, but one layer's q/k/v are read by a single fused kernel and must agree. Mixing block layouts inside one kernel reads 18-byte blocks as 34-byte ones and produces fluent, wrong text. The memory preflight predicted every quantized weight at its Q8_0 size, which was already wrong for Llama's Q4_0 files and would have refused qwen35 at 28GB once it retains. It now asks the plan provider which representations that family reads as they lie -- not a new declaration, just supportedDataTypes, which already answers exactly this. Measured on Llama-3.2-1B-Instruct-Q4_0: per-layer weights 547MB retained against 1034MB materialized, the 34/18 block ratio. The preflight decides per tensor where a loader may decide per model, so a file mixing Q4_0 with another quantization in its layers would be predicted smaller than it loads. No quantizer produces one. The direction is the tolerable one: this prediction refuses loads and a refusal cannot be overruled, so an over-estimate blocks a configuration that would have run where an under-estimate reaches the backend's own error. None of this makes qwen35 runnable on a device. Nothing consumes these weights yet, and the preflight's retention does not apply to a family with no provider -- so the model is still predicted at 27760 MiB, which is the right answer while nothing can build it a plan.
Completes the set of quantized representations the engine can hold on a device in the file's own layout. Q4_0, Q4_K, Q6_K and Q8_0 already had wrappers and kernels; Q4_1 and Q5_K did not, and Qwen3.8-27B needs both -- eight Q4_1 ffn_down tensors and forty-eight Q5_K ssm_out. Q4_1 is 32 weights in 20 bytes, d * q + m with an unsigned nibble. Q5_K is Q4_K with a fifth bit held in a separate 32-byte plane, indexed by position within the pair's half and by which nibble the element came from -- not by the element's index in the super-block. Both kernel files are shaped like their siblings, so a layer differs only in which method reference it names, and each decode lives in a package-private helper the parity test can call on the host. QuantizedDeviceDecodeParityTest covers all six against their CPU tensors on adversarial blocks: each block's scale fields walk the half-precision corners -- both zeros, both ends of the subnormal range, unity of both signs, both extremes of the normal range -- and payloads walk all-zero, all-ones and both alternations, with every fourth block random. Bit equality, not a tolerance. It also asserts its own sensitivity. Four deliberate faults -- the wrong nibble half, Q4_0's recentring applied to Q4_1, Q5_K's fifth bit from the wrong bit, K-quant scales read without the straddled high bits -- must each be seen to disagree with the host. A parity test that could not detect them would prove less than it appears to. Writing it found a fixture bug worth recording rather than a kernel one: Q6_K's scale is at offset 208, after ql, qh and sixteen signed sub-block scales, so injecting halves at offset 0 corrupted ql and left the scale random. Random scales reach Inf and NaN, which TransformerComputeKernelsQ6_K states it does not handle because a quantized block scale is neither -- a constraint of well-formed GGUF, not an oversight, and the alternative is two branches in the innermost loop of every Q6_K weight. The offsets are now passed per format and the divergence is recorded in the test instead of asserted away. Also corrects the port proposal's architecture model: quantized storage is retained on every backend, decoding during compute is an implementation property rather than a CPU-only one, materialization to Q8_0 is not the normal accelerator fallback, support is declared per operation and dtype, fusion stays backend-owned, mixed quantization between tensors is legal, and every operand fused into one kernel must have a layout combination that kernel explicitly supports.
DataType conflated three questions, and the conflation is what made the engine believe quantized weights were CPU-only: 1. must arithmetic decode a block to read a value 2. can a backend store the representation 3. does a given operation have a kernel for it Only the first is a property of a representation, and it is true on every backend -- a device kernel decodes inside its dot product exactly as a host one does. The second belongs to the backend's storage vocabulary; the third to OperationSupport, per operation and per target. 'The GPU cannot do Q5_K' was never a fact about Q5_K. So isFormatDecoded is gone rather than renamed, and materializedFallback becomes narrowedFallback with one case left: BF16 to F16, which is a real loss of mantissa bits for want of BF16 arithmetic rather than a capability gap dressed up as a representation. It no longer answers Q8_0 for anything, which is what used to double a 4-bit model on a device. The Q8_0 promotion still exists for families on the older loading path, but it now lives in DataTypeMapping as legacyDevicePromotion, named and documented as a property of that path rather than of the type. Families with native kernels use ModelLoader.loadTornadoTensorNative, which keeps every representation as the file gave it and refuses one the device cannot store instead of promoting it. OperationSupport now declares what is actually true: matrix-vector and vocabulary projection read all six quantizations on the GPU; matrix-matrix has tensor-core kernels for F16 and Q8_0 only; the device embedding gather covers F16 and Q8_0, other representations being gathered on the host. FusedOperandSupport is the safety net the mixed model needs. Mixed quantization between tensors is legal, but a fused kernel that decodes its operands with one block layout must reject a mixture -- giving it a Q5_K operand where it expects Q4_0 reads 176-byte super-blocks as ten Q4_0 blocks and produces fluent, wrong text. Refused at plan construction with the operands and their representations named. Five tests asserted the retired premise and are replaced by two that state the new one. One further test was removed by an over-eager edit and is restored, reworded.
A model is not one representation. Qwen3.8-27B holds Q4_0 projections and token embeddings, eight Q4_1 ffn_down, forty-eight Q5_K ssm_out, a Q6_K vocabulary projection, a Q8_0 MTP projection and F32 norms, SSM parameters and convolution kernels. A prediction derived from one model-wide dtype is wrong for every tensor that is not that dtype, and here the error runs to twelve gigabytes. DeviceRetention is asked per tensor, by name as well as representation. The name matters because support can differ by role: a family may have a matrix-vector kernel for a representation and no vocabulary-projection kernel for it, and only the name distinguishes them. Most policies will not need it; it is there so the ones that do are expressible rather than approximated. Measured on the fixture: retained per-layer 13.306 GiB global 1.637 GiB total 14.944 GiB converted per-layer 24.590 GiB global 2.516 GiB total 27.107 GiB The retained figure is the file's own weight bytes, arrived at independently from the GGUF tensor inventory rather than recorded from a run. Qwen35WeightFootprintTest validates each representation separately: retaining exactly one dtype at a time isolates its contribution, and the saving must be a whole number of that format's blocks times the difference from Q8_0's block size. A predictor with one format's block size wrong would pass a whole-model check and fail this. Q8_0 and F32 are asserted unchanged either way.
QuantizedDeviceDecodeParityTest settles the arithmetic on the host. It
cannot settle whether a kernel compiles, or whether its row addressing
survives a launch, and both have bitten this backend before.
NativeQuantizedKernelAccelTest compiles and executes Q4_0, Q4_1, Q4_K,
Q5_K and Q6_K matrix-vector kernels on CUDA and compares each against
its CPU tensor over the same bytes, on a rectangular shape (512x37) that
would expose a row/column transposition.
Four passed immediately. Q5_K did not, and finding out why took seven
formulations:
element-strided loop, variable shift deopt scaffolding the CUDA
backend cannot declare --
identifier 'context' is
undefined, identifier 'slots'
is undefined
scale/min extracted to a helper same
mask + conditional move assertion inside
CUDALIRGenerator.emitIntegerTestMove
mask + if branch deopt scaffolding
index-derived branch decomposition deopt scaffolding
shuffle reduction, no local memory deopt scaffolding -- so 'slots'
was the deopt frame, never the
local array
sub-block-wise nested loop deopt scaffolding
The decode compiled and ran correctly outside a loop throughout, which
is what narrowed it: the problem was ByteArray.getHalfFloat inlined into
a loop. TransformerComputeKernelsQ6_K already documents that call as one
TornadoVM's sketcher chokes on, and assembles its half from two byte
loads instead. Q5_K now does the same and compiles.
Two of those attempts are kept because they are better code regardless:
the scale/min helper, and walking whole 32-element sub-blocks so a
sub-block's scale, minimum, nibble plane and fifth-bit position are
computed once rather than thirty-two times.
decode() uses the same byte-assembled half as the kernel. It had been
left on getHalfFloat, which would have meant the parity test covering
arithmetic the device never runs.
The host parity tests run every lane on the CPU. That settles the arithmetic and says nothing about whether TornadoVM can compile the kernel -- a distinction that is not academic, since Q5_K's matvec passed its host parity test and then failed to compile in seven formulations. These kernels use nested loops, a retained state array written in place, and TornadoMath transcendentals, none of which the host tests exercise as device code. All five groups compile and run on CUDA and match the host operations: causal convolution, and the window it advances the delta rule, and the 48 x 128 x 128 state it updates in place gated norm and per-head L2 norm decay and write strength, through exp/log/logistic attention: query/gate split, partial RoPE, output gate Dimensions are the 27B's own where it matters -- a 128-wide delta-net head, 48 value heads against 16 key heads, a 256-wide attention head with a 64-wide rotary.
…sor dtype Two shared paths still assumed a model-wide representation, which is wrong for every mixed file and blocks a Qwen3.5 device plan. The embedding lookup staged and converted by the model-wide quantization string. Qwen3.5 reports Q8_0 model-wide while holding its token embeddings as retained Q4_0, so 18-byte blocks would have been staged and read as 34-byte ones -- a plausible activation and wrong output. Activation now dispatches on the embedding tensor's own dataType, which is what TornadoForwardPass already did for staging; this is the other half. convertQ4_0toFP32 is the kernel it needed, assembling its half from two byte loads rather than getHalfFloat, which TornadoVM's sketcher rejects inside a kernel. The vocabulary projection was hard-wired to the Q8_0 kernel. It now selects by what the output projection actually holds -- Qwen3.5's is Q6_K where its layers are Q4_0 -- across all six quantizations, and refuses an unsupported one by name rather than converting it, which would double what it occupies and hide a missing kernel behind a memory cost. Qwen35TornadoWeights now extends TornadoWeights, reversing an earlier choice. Implementing Weights directly kept the base class's per-layer arrays honest, but Activation, AbstractLogitsTaskGraph and TornadoForwardPass are all written against TornadoWeights, and staying outside it meant changing each of them for one family. Fitting the shape costs a documented convention -- per-layer arrays indexed by absolute block, null where the block is of the other kind, exactly as the host weights already do -- and the class says so. 587 tests pass; Llama Q4_0/Q8_0 and Qwen3-0.6B unchanged on CUDA.
…ssion supportedDataTypes answers whether a provider can build a plan for a model whose weights report one representation. A mixed model reports one and holds several, so the memory preflight cannot read that set as what each tensor occupies on the device without under- or over-predicting every tensor whose representation is not the model's. nativeTensorTypes is that second declaration, defaulting to the first so a family whose model is uniform states it once. The preflight reads it. The qwen35 loader now retains every tensor in the file's own representation rather than materializing all but Q4_0 as Q8_0, and reports the projections' shared representation as the model-wide one, failing by name when they disagree instead of picking a representative tensor.
A delta-net mixer splits its convolved projection into q | k | v of unequal widths and scales the queries before the recurrence. splitQKV assumes a key and a value of equal width, which is true of attention and false here: on the 27B the widths are 2048 | 2048 | 6144, and an equal-halves split takes the value slice from inside the keys. Both new kernels are stated per operation rather than per family, and both carry a lane helper so their addressing is checked on the host as well as on the device. scaleInPlace moves out of Gemma4Kernels for the same reason: the arithmetic belongs to the operation, and a second copy named after another family is how one kernel becomes two. Adds device coverage for the F32 matrix-vector kernel at an ssm_alpha shape and for the fused Q4_0 gate/up SwiGLU, neither of which had been launched.
A graph per trunk layer: 20 tasks for a recurrent block, 17 for an attention block, each following the host branch operation for operation. Both kinds run the same dense SwiGLU feed-forward and the same two normalizations, and only the mixer differs. Every task that reads a weight is bound to a kernel chosen from that tensor's own representation at construction, so the compiler sees fixed block addressing rather than a dtype switch inside a K-loop. The one task reading two weights states that they must agree and refuses a mixture by name. A block whose weights do not match the kind the metadata declares fails naming the layer, the role and what is missing, rather than dereferencing null inside graph construction. Key/value storage is addressed by a dense index because the store is sized by the layers that attend; the convolution window and delta-net state are per-layer slices of one array each, uploaded once and updated in place. The topology test builds the smallest model with both layer kinds and pins the graph count, the task counts, the mixer selection, the absence of the draft head, and that changing one tensor's representation changes exactly one task's dispatch.
Adds the plan components, the provider and its service entry, and points the model's GPU generation at the shared single-token loop. The provider declares STANDARD only, and Q4_0 admission — the representation this family's trunk projections share — while declaring separately the eight representations its tasks decode per tensor, which is what the memory preflight needs to predict 14.944 GiB rather than 27. A quantized row must be a whole number of blocks, and the dispatch now says so by name: every block-decoding kernel addresses a row by its block offset, so a partial row would silently read the next row's blocks. The synthetic parity test runs a whole forward pass on the device against the same one on the host, over the same weight bytes, on a model with both layer kinds and the same mixture of representations as the real file. Two positions, because the first says nothing about whether the recurrence carried.
A session's device arrays were gated on use.tornadovm, which the CLI sets and nothing else does. A caller that loaded device weights without it got a state whose device arrays were all null, and the failure arrived from inside TornadoVM as "null object passed into streamIn()" in the activation graph rather than anywhere that named the cause. The model decides now, from whether its weights are device weights, handed to the state for one construction. Attention uses the single-workgroup online-softmax kernel rather than the split-KV one: that kernel fixes its query staging at 128 floats per head and this family's head is 256 wide, which reads past the array and faults with an illegal address that surfaces as an unrelated allocation failure. The load message reports the representation the weights were actually built in, rather than predicting one from the output tensor's type before loading. Adds the real-fixture parity gate. Qwen3.8-27B against the CPU reference on CUDA, teacher-forced over 63 rows of 248320 logits: no elementwise violations against the tightest bounds in the suite, worst ratio 0.074, argmax agreement 63/63, cosine 1.000000.
…mily holds The memory model sized the key/value cache by the layer count, which is right for a stack that attends throughout and four times too large for one that attends in a layer of four. It also had no term for state that is neither cache nor scratch: a recurrent layer keeps its history in a fixed-size convolution window and delta-net matrix, updated in place and independent of the context length — 149.6 MiB on Qwen3.8-27B, previously counted nowhere. Both are asked of the configuration, which is where the answer differs by family, and both default to what an attention-only stack would say. Measured on the real fixture by bisecting the device budget: 15455 MiB succeeds and 15450 MiB fails, against a prediction of 15453.6 MiB. Adds two gates: the retained representation of every role at runtime, and the public route on the device through generate, reset, close and use-after-close, asserting the resolved execution path so a silent host fallback cannot pass.
The measured baseline on the real fixture — 15.0 tok/s against 0.77 on the CPU, where the time goes per kernel, and the memory prediction against the measured minimum budget — plus the two shape facts the port established: a 256-wide head does not fit the split-KV kernel's fixed local arrays, and a quantized row must be a whole number of blocks. Also records an optimization that was tried and rejected. Walking 32-element groups in the Q6_K projection hoists an fp16 scale out of the inner loop and is slower, because the element-strided loop it replaced had consecutive lanes reading consecutive bytes. Coalescing beat the arithmetic saving, and reverting reproduced the original timing to within 0.05%.
… and the batched envelope is settled
Two things that belong together: the Q6_K projection becomes automatic, and the
four batched parity tests that had been failing since ffn_down was packed get
their own envelope rather than a widened shared one.
The projection is chosen from three facts and no preference: the output tensor
is Q6_K, the device lowers dp4a, and this session's state carries the
quantization scratch. Any other representation, device or family keeps the
kernel it had. It is one call site, quantizing the final normalized activation
in the same graph that reads it and nowhere else, which is why it needs no
provenance flag -- there is no second consumer to confuse it with.
tg128 -d 0 24.18, 24.14 against 21.55, 21.47 (+12.4%)
tg128 -d 381 17.21 against 15.85, 15.77 (+8.8%)
pp381 80.24 against 80.38 (prefill skips logits)
Correctness, reported apart from the drift against the floating-point kernel:
- unpacking, one-hot so each output is one decoded weight held against
Q6_KFloatTensor: 256 rows covering every group, both nibbles, every qh shift,
every scale index and both signs of byte scale, raw values walking 0..63;
- accumulation, dense: 96 rows of 1024 inputs over four super-blocks, weights
alternating about the midpoint so terms cancel rather than average, scales of
both signs, a zero activation block and magnitudes three orders apart. Worst
0.000541 against a largest output of 3649.6, which is 1.5e-07 relative.
Q8_0_PACKED_DECODE takes ceilingPerRms 0.70, relL2 0.12, minCosine 0.9955, and
nothing else changes. Engineering regression limits for an accepted arithmetic,
not calibrated quality thresholds: they say this path computes what it computed
yesterday, not that what it computes is good. A future failure there is to be
investigated, not widened away. Q8_0_PACKED_ACTIVATION is left exactly as it
was and now has no users.
Six synthetic parity cases failed first, and they were right to. They pin
themselves to the floating-point path because their subject is addressing and
they compare exactly; the new projection ignored that escape hatch. It honours
it now.
61 tests, 0 failures, 1 skipped -- the skip is Qwen35CrossWidthCaptureAccelTest,
which is a capture tool and does nothing unless driven twice with different
widths. Driven that way on this build, widths 32 and 64 are byte-identical.
The previous commit message said widths 32 and 64 were "byte-identical" on the strength of a cmp over the whole capture file. That was the wrong check: the file begins with a header carrying the width itself, so cmp reports a difference at byte 4 whatever the logits say, and it did. Compared properly, skipping the header: the token sequences match and all 15,644,160 logits are bit-identical between the two widths, with a partial final chunk at both. The claim stands; the way it was checked did not.
…ad of a barrier tree
Only the reduction tail changes. The one-workgroup-per-row mapping, the block
walk, the nibble packing, the dp4a pairing, the -8 * sum correction, the
activation quantization and the dispatch are all untouched, and no new
representation, option or capability is introduced.
Each 32-lane warp reduces its own partial with five simdShuffleDown steps, one
lane per warp writes to shared memory, and a single barrier separates that from
the combine. At the 128-lane width this repository dispatches that is one
barrier where the tree emitted seven. All three packed forms take it: the plain
projection, the residual one, and the fused gate/up, whose two accumulators are
reduced independently and whose SiLU still sees the gate summed across every
warp -- applying it per warp would be a different function.
The reduction rides on PACKED_INTEGER_DOT, which is granted on CUDA alone.
Shuffles miscompile on OpenCL, and no device without that capability reaches a
packed kernel at all, so the floating-point kernels and every other backend are
unaffected. WARP_SHUFFLE stays ungranted and is unrelated to this.
Interleaved whole-model A/B, FP16 KV and CUDA graphs, RTX 5090 Laptop, order
alternated across six runs:
tg128 b32 24.69, 24.70, 24.63 against 24.31, 24.14, 24.02 (+2.1%)
tg128@d381 b32 17.39, 17.35 against 17.24, 17.13 (+1.2%)
Screening first, on three projection shapes with both plans built, warmed and
resident and the rounds alternating -- the flaw this method replaces timed each
plan once in sequence and measured the ordering. Best-of-seven rounds of 200
iterations, twice: attn_q 1.04-1.07x, ssm_qkv 1.16-1.22x, gate/up 1.02-1.08x.
Screening evidence; the whole-model number is the one that decided it.
Q4_0Dp4aReductionAccelTest is about the tail, because that is what changed and
because a reduction that keeps one warp and drops three still returns plausible
finite numbers. Exact integer cases pin it: unit block scales and integer
activations at 128, 2048, 4096 and 5120 inputs, so the answer is a specific
integer that a missing warp changes; large alternating terms that cancel, so a
lost warp shows against a total far smaller than any partial; the residual form
accumulated onto a destination of a million, where a per-warp accumulate cannot
hide; and the fused form against silu(full gate) * up. Then zero activation
blocks and zero block scales, scales spanning four orders of magnitude with
mixed signs, finiteness, and an inspection of the generated code asserting five
shuffles and exactly one barrier.
Reduction order changes, so the floating-point total may round differently; the
integer dot products are exact either way. No numerical bound moved. The parity
tests report what they reported before -- worst relL2 0.0429 against a bound of
0.12, minCosine 0.99939 against 0.9955, argmax 0/63 -- and the NLL screen,
which is a reused development screen and not held-out validation, reads pooled
1.436257 against 1.437495 at HEAD over the same five passages, identical
sha256s.
68 tests, 0 failures, 1 skipped. The skip is Qwen35CrossWidthCaptureAccelTest,
which is a capture tool and asserts nothing unless driven twice.
… instruction The packed kernels reduce with simdShuffleDown, and nothing said so where the grant is written. The capability reads as though it were about dp4a alone, which would make granting it on OpenCL -- where dp4a lowers and the shuffle miscompiles -- look merely unmeasured rather than wrong. No new capability and no dispatch change: this states the dependency the code already has, in the two places a reader decides whether to extend the grant.
ssm_out is the largest single kernel in decode at 22.2%, and the only one of the
large ones still decoding every weight to a float. This gives it the treatment
the Q4_0 and Q6_K projections already have: the weights stay Q5_K in the file's
own representation, the activation is quantized to eight-bit blocks of 32, and
the dot product runs on dp4a.
Its algebra is not Q4_0's and nothing of Q6_K's carries over. A Q5_K weight is
d*sc(sub)*q - dmin*m(sub) with q five bits and never recentred, so a sub-block's
contribution splits into d*sc times the integer dot product and dmin*m times the
sum of the activation's quants. That second term is what this representation
adds: it reads the same scratch array the Q4_0 path multiplies by the constant
8, and multiplies it instead by the sub-block's own six-bit minimum. Getting
that wrong produces plausible finite numbers. llama.cpp computes the sum on the
fly with a dp4a against 0x01010101; ours is precomputed by
quantizeActivationQ8Blocks, so the inner loop is eight instructions per
sub-block rather than sixteen. Because the quants are 0..31 and unsigned, every
packed byte is a valid signed byte and nothing here approaches the
unsigned-recentring defect in docs/architecture/tornadovm-issues.
The activation is the delta-net readout, which is none of the three already
quantized in a layer, so it gets its own task after ssm_gated_norm and its own
provenance flag -- and the scratch it borrows is checked for capacity and
alignment rather than assumed: 6144 against arrays sized for 17408 here, a fact
about this configuration and not about the code.
**Off by default**, behind -Dllama.qwen35.q5kPackedSsmOut. That property is a
temporary A/B control, not an option: it either becomes unconditional under the
existing packed-path gate or it and the kernel go.
Correctness, none of it depending on an envelope:
- one-hot decoding of all 256 positions of eight adversarial super-blocks
against Q5_KFloatTensor -- both nibble halves, all eight sub-blocks including
the straddling branch of the six-bit scale packing, and every bit position of
the fifth-bit plane;
- dense at ssm_out's own 6144 against a same-quantization reference computed in
doubles from an independent reading of the bytes, with mixed signs, zero d and
zero dmin blocks, zero activation blocks and alternating-sign cancellation:
worst 9.27e-06 against a largest output of 122.0, which is 7.6e-08 relative;
- a zero activation contributing exactly nothing, minimum term included;
- the residual added once per row;
- the generated CUDA carrying dp4a and five shuffles.
Measured with the temporary flag, interleaved, FP16 KV and CUDA graphs, RTX 5090
Laptop:
tg128 b32 28.26, 28.26, 28.14 against 24.94, 24.76, 24.73 (+13.8%)
tg128@d381 b32 19.36, 18.39 against 17.46, 16.45
pp381 b32 75.80, 75.77 against 75.92, 75.78 (prefill is the
batched path, untouched)
Screened first at the loaded model's own 6144 -> 5120 with both plans resident,
the rounds alternating and the candidate paying for its quantization: medians
2.54x and 2.48x over two runs. Medians rather than a best-of-N, which reports
the luckiest round.
No bound moved. With the candidate on, parity reads relL2 0.0365 against 0.0750,
minCosine 0.99933 against 0.99750, argmax 0/63, elementwise 46.23% against a 50%
budget. The NLL screen moves by about a thousandth of a nat in both directions
depending on the passage -- pooled -0.000678 on STANDARD and -0.001098 batched,
worst single passage +0.21% perplexity -- which is movement on a reused
development screen, not evidence about quality.
80 tests, 0 failures, 1 skipped at the default. With the flag on,
Qwen35GraphTopologyAccelTest fails on the task count it pins, 24 against 23:
that is the added ssm_out_quantize, and it is how the execution path is
asserted rather than assumed.
…a mismatch A task-count mismatch is not an assertion. With the candidate enabled the topology test failed on 24 against 23, which said the plan had changed but nothing about whether it had changed correctly, and which would have gone on failing for any wrong extra task. Now both modes are asserted, neither weakened. The count expects the readout's quantization exactly when the packed Q5_K path is enabled; a new case asserts that the task is present in every recurrent layer and in no attention layer, and absent everywhere when it is disabled. Placement is asserted through consequences rather than through a task order the grid scheduler does not preserve. The readout's quantization overwrites the same three scratch arrays the branch's projections read, and emitting it clears the provenance flag those projections consult, so a quantization emitted before ssm_qkv_proj and ssm_gate_proj shows up as those two losing the packed path, and one emitted after ssm_out_proj shows up as ssm_out_proj not taking it. Requiring all three at once pins the task to the window between them. That required fixing what the inventory records: matVec computed the Q5_K condition inline for the kernel choice while the Dispatch it recorded still said false, so the inventory described a plan that had not been built. It is one boolean now, used for both. Qwen35GraphTopologyAccelTest: 10 tests, 0 failures, in both modes.
The temporary switch is gone. Eligibility is what it was for every other packed
projection: the device grants PACKED_INTEGER_DOT, the tensor is Q5_K, and the
buffer holds the activation this layer quantized for it -- plus the scratch
length check the readout needs, since its width is not one the shared arrays are
sized for.
One decision drives both the dispatch and the inventory: ssmActivationQuantized
is set only where the quantization task was actually emitted, which is only
under the packed gate, only in a recurrent layer, and only when the scratch is
long enough and 32-aligned. matVec reads that one flag to choose the kernel and
to record what it chose, so the inventory cannot describe a plan that was not
built. The floating-point kernel stays for every device and configuration where
any of that does not hold, and -Dllama.qwen35.packedIntegerDot=false still puts
the whole packed path back on floats for the exact-comparison tests.
Measured before this commit, with the switch, and unchanged by removing it --
the selected computation is the same:
tg128 b32 28.26, 28.26, 28.14 against 24.94, 24.76, 24.73 (+13.8%)
tg128@d381 b32 19.36, 18.39 against 17.46, 16.45
pp381 b32 75.80, 75.77 against 75.92, 75.78
No bound moved. Against their own unchanged envelopes: Q8_0_FULLY_PACKED reads
relL2 0.0365 against 0.0750 and minCosine 0.99933 against 0.99750;
Q8_0_PACKED_DECODE reads relL2 0.0386 against 0.1200 and minCosine 0.99926
against 0.99550; argmax is 0/63 in both. The elementwise violation fraction is
46.23% against a 50% budget -- passing, and closer to its limit than any other
metric here.
The topology tests now assert the default rather than a mismatch, in both
directions: the readout's quantization is expected in every recurrent layer and
in no attention layer, and under the packedIntegerDot escape hatch the same case
asserts the floating-point plan -- nothing packed, no quantization tasks -- so
the fallback is asserted rather than skipped.
The attribution in this file predates the MMA projections and is superseded: ssm_out is no longer an untiled per-row matvec at 25.6%, and nothing decodes weights in a float kernel any more. The reasoning is kept; the numbers are replaced. Two captures, kept apart because they are not the same run: the TornadoVM profiler with graphs off for per-task attribution, and Nsight Systems with graphs on and node tracing for the timeline. Neither is the throughput number, and the three figures -- 71.33, 80.44 and 75.8-76.2 t/s -- differ in more than one thing each, so none of them prices the instrumentation. Attribution counts the last 24 chunk executions only: the two measured passes, 762 tokens, 1,536 layer graph executions, with the warm-up pass and the compilation chunk excluded. ffn_gate_up is 46.11% of it. Activation preparation is 0.11%, per-step transfers are the batch-info holders, and cuGraphLaunch is 0.4% of kernel time. The partial chunk costs a full chunk: mmaEligible tests the configured batch size rather than the live chunk, so the 29-row chunk launches the 32-row grid and the padding rows are computed, stored and never read. 402.9 and 403.3 ms against 385-406 ms for the full ones. What the numbers say: cost tracks panels decoded per block iteration, not activation bytes. Single-panel iterations cost 1.38-1.63 ns whatever the shape; the fused two-panel iteration costs 4.10 ns, and ffn_gate_up and ffn_down run the identical 696,320 block iterations per call while differing 2.65x. Why the fused form costs 2.6x rather than 2x is a hypothesis -- shared memory and accumulator fragments -- and is labelled as one. Nsight Compute still cannot arbitrate it here (ERR_NVGPUCTRPERM).
The fused two-panel MMA task is replaced by two projectionMMAQ4_0 calls at the
same M, N and K, writing the same two buffers, followed by the same SwiGLU task.
The kernel is the one every other batched projection already uses. No new
kernel, no option, no precision or layout change, and the scalar fallback for an
ineligible shape is untouched; both new tasks are registered as tensor-core
tasks so they receive the MMA grid rather than the matrix-vector one.
Equivalence first, before any timing: at the real 32x17408x5120, on a full chunk
and on a 29-row chunk with zeroed padding rows, gate and up came back
bit-identical to the fused kernel -- 557,056 of 557,056 values each, all finite.
The whole-model cross-width capture then returned the same logits as before the
change: 15,644,160 values and 63 token ids identical at widths 32 and 64, SHA-256
e17b0f731220525c, which is the same digest the fused default produced.
Interleaved, warmed, graphs on, tensor cores on, FP16 KV, two repetitions each:
pp381 b32 85.65, 85.56, 92.00 against 76.21, 75.92, 75.89 (+12.8%
on medians; +12.3% on the
conservative pairing)
tg128 b32 28.41, 28.35 against 28.42, 28.33 (unchanged)
Per layer-chunk, from the TornadoVM profiler on the measured passes only: 985.1
+ 1047.6 = 2032.7 us against the fused kernel's 2855.0 us, a 28.8% reduction at
the same shape, and the window's kernel total falls from 9511.5 ms to 8184.4 ms.
What this establishes is that the two-panel form cost more than two single-panel
forms of the same shape on this device and geometry, and that staging the
activation tile twice is cheaper here than whatever the fused body paid for. It
does not establish why: registers, shared memory, occupancy, scheduling and
intra-kernel synchronization were not measured, and Nsight Compute remains
unavailable here (ERR_NVGPUCTRPERM).
No bound moved. The four batched parity tests, the sequential and STANDARD
parity, the batched FP16-KV lifecycle and the topology assertions pass with the
same numbers as before -- as they must, the logits being identical: 24 tests, 0
failures.
The fused kernel had no other production caller and is removed. Its host-parity
case is covered by the single-panel one; the unsigned-nibble guard it hosted now
runs against two projectionMMAQ4_0 tasks and asserts the same values.
…projection Same capture as the A/B, attributed over the same window: the last 24 chunk executions, two measured passes, 762 tokens, compilation and warm-up excluded. Gate and up are 18.49% and 19.66%, 38.15% together, against 46.11% as one fused task; projectionMMAQ4_0 is now 77.86% of prefill kernel time across six task names. Each projection's dtype, M/N/K, workgroup size and grid are recorded with it, since one warp per m16n8k16 tile is what makes those grids what they are. Inspection of that kernel's generated CUDA finds one concrete candidate: the Q4_0 block scale is decoded by hand, ten branches per call, once per lane per block, where the DP4A kernel in the same dump reads the same field through getHalfFloat and gets a single __half2float. Both fragments are quoted. Two further observations are recorded and deliberately not proposed: four lanes duplicate each column's scale decode, which would be a restructure rather than a small change, and the loop-invariant nibble and tile decisions are re-tested per element in the emitted C, which ptxas may hoist. Neither instruction counts nor this inspection identify a hardware bottleneck, and Nsight Compute is still unavailable here.
…ssor
One line, one call site. projectionMMAQ4_0 read its block scale with
halfFromBytes, which rebuilds the half by hand and lowers to a ten-branch
software expansion; it now reads w.getHalfFloat(base).getFloat32(), which lowers
to a single __half2float. Geometry, staging order, synchronization, the A-tile
path, the Q4_1 and Q5_K MMA kernels and halfFromBytes itself are unchanged.
Same bytes, same interpretation: getHalfFloat takes a byte index, requires
two-byte alignment and reads a native-endian short; Q4_0's block stride is 18
bytes with the scale at offset 0, so every scale is aligned, and halfFromBytes
assembles the same little-endian pair.
Checked before it was timed:
- every finite half encoding, read from a Q4_0 block-scale position, agrees bit
for bit with Float.float16ToFloat -- 63,488 encodings, 31,744 negative, 2,046
subnormal, both zeros. The case is in HalfFloatConversionAccelTest;
- the whole projection output at 32x17408x5120, captured from both builds and
compared byte for byte: identical on a full chunk and on a 29-row chunk with
zeroed padding rows, 1,114,112 values, all finite;
- the emitted CUDA carries one __half2float and none of the software expansion.
Interleaved, warmed, graphs on, tensor cores on, FP16 KV, two repetitions each:
pp381 b32 88.46, 88.39, 88.27 against 85.54, 85.31, 85.40
medians 88.39 against 85.40, +3.50%, ranges disjoint
The speedup is not inferred from the instruction count: the generated code says
the work existed, the A/B says what removing it was worth, and neither says what
the hardware was doing. No bound moved -- 28 focused MMA, topology, parity and
lifecycle tests pass with the parity numbers unchanged, which is what
bit-identical outputs require.
The emitted CUDA compiled the way TornadoVM compiles it -- sm_120 cubin, no other options -- and disassembled, so the record separates what the source repeats, what the generated C repeats, and what the hardware is actually issued. 61 registers, no spills of either kind, 1536 bytes of shared memory, one barrier. No occupancy claim follows: residency limits were not measured and no counters are available. The per-element comparisons in the generated C do not survive: 2 BRA and 3 ISETP in the whole kernel. What survives is predication, and in one place a duplicated store -- 16 STS.U16 to write 8 B values and 12 STS to write the A tile's 8 ints, because ptxas predicates both destinations rather than selecting one address between two distinct shared allocations. Loads are individual: eight LD.E.U8 at consecutive offsets for the eight packed weight bytes, sixteen LD.E.U16 for the A tile where each adjacent half pair could be one 32-bit load, one LD.E.U16 for the scale. No wide loads anywhere. The scale is loaded and converted once per lane, so the four lanes that share a column do it four times on the same bytes. Also settles, retrospectively, that the hand-written half decode 5e68ffc removed was not being optimized away: 9 BRA, 9 ISETP, 5 FSEL and 472 instructions before against 2, 3, 0 and 368 after. That checks the premise; the +3.50% still comes from the whole-model A/B, not from instruction counts.
This file claimed merging the two B tiles had no small implementation because the swizzled store and load take an array and not an array plus offset. That was wrong. The installed SDK carries mmaLoadA/mmaLoadB/mmaLoadBSwizzled overloads with a byte offset and a matching mmaStoreBSwizzled, verified with javap against tornado-api-6.0.1-jdk21-dev.jar, and tornado-examples uses exactly that shape. This repository already declares B_SUBTILE_BYTES = 256 and uses it nowhere. Where the offset is applied was verified in the emitters rather than assumed: both the store and the swizzled load compute the address, apply the XOR permutation, and only then add the byte offset. The offset lands after the swizzle on both sides, so one allocation holds the two panels with the layout they have today. Wider reads of the weights are a different answer: ByteArray offers one-byte and two-byte accessors and nothing else, and the one wider path, cp.async, needs a four-byte-aligned source, which Q4_0's 18-byte stride gives for only half the blocks -- the quant bytes sit at 18B + 2. What is missing is named: a four-byte read at two-byte alignment. The same cp.async does suit the activation tile, whose addresses are four-byte aligned by construction. Also corrected in the SASS section: it is an NVCC reconstruction rather than the runtime NVRTC cubin, and the predicated store pairs are static instructions where each thread executes one -- the tile is written once per value, so the cost is issue slots and not store traffic.
The two B panels become one HalfFloat[2 * PANEL * BK] -- 512 bytes -- with the
first at byte offset 0 and the second at B_SUBTILE_BYTES, staged with
mmaStoreBSwizzled and read with the offset-aware mmaLoadBSwizzled. The per-lane
choice of panel is now a byte offset rather than a choice between two arrays.
A tiles, geometry, grid, reduction order, activation staging, barriers, weight
layout and precision are unchanged, and no other kernel is touched. This uses
API the SDK already ships; B_SUBTILE_BYTES was declared here and unused.
Both sides apply the swizzle to the in-panel address and add the offset
afterwards, so each panel keeps the layout it had as its own array. A panel's
swizzled address stays inside its own 256 bytes -- the XOR permutes bit 4 from
bit 7, within the window -- so the panels cannot overlap, and both offsets are
multiples of 16 for ldmatrix.
Verified before timing: the whole projection output at 32x17408x5120, captured
from both builds into NaN-poisoned buffers and compared byte for byte, is
identical on a full chunk and on a 29-row padded chunk -- 1,114,112 values, all
finite. The cross-width capture returns the same logits as before, SHA-256
e17b0f731220525c at widths 32 and 64.
The generated CUDA now has one shared half array and a single staging store with
the offset added after the swizzle, no if/else. In the NVCC reconstruction --
same source and ptxas, not the runtime NVRTC image -- predicated STS falls from
24 to 8 and STS.U16 from 16 to 8, while registers rise 61 to 64, instructions 368
to 392 and IADD 59 to 78, with shared memory unchanged at 1536 bytes and no
spills either way. The kernel got larger, not smaller.
pp381 b32 92.91, 93.19, 93.07 against 88.72, 88.93, 88.66
medians 93.07 against 88.72, +4.90%, ranges disjoint
A static instruction count that rose while measured time fell is why this
repository does not price changes from listings.
No bound moved. 26 focused MMA, topology, batched-parity, sequential and
STANDARD parity and lifecycle tests pass with the parity numbers unchanged, and
the cross-width comparison was driven explicitly rather than left to the suite's
capture-only skip.
The same representation change as the B merge, on the other operand. The two A
panels become one int[2 * BM * BK / 2] -- 256 ints, 1024 bytes -- the first at
byte offset 0 and the second at A_SUBTILE_BYTES. Staging writes
aTile[half * (BM * BK / 2) + j] where it used to choose between two arrays, and
the reads become mmaLoadA(aTile, BK, 0) and mmaLoadA(aTile, BK, A_SUBTILE_BYTES).
Merged B, geometry, grid, activation reads and packing, arithmetic, barriers,
staging order, weight representation and bounds are unchanged; no other kernel is
touched.
Each load sees the layout it saw before. The A load is Variant.X4 -- trans=false,
swizzle=false -- so there is no permutation to preserve, and its per-lane address
(__row << 5) + __col has __row in [0,15] and __col in {0,16}, reaching at most
byte 496 and fitting inside 512. The offset is added after that address is
formed: the emitted code shows __bo += 0 for the first load and __bo += 512 for
the second, so aTile[128 + j] is exactly what aTileHi[j] was and the panels
cannot overlap.
Verified before timing: the projection output at 32x17408x5120 is bit-identical
to the accepted default on a full chunk and on a 29-row padded chunk, from
NaN-poisoned buffers, 1,114,112 values, all finite; the cross-width capture
returns the same logits at widths 32 and 64, SHA-256 e17b0f731220525c.
The emission now has one shared int array and unpredicated stores. In the NVCC
reconstruction -- same source and ptxas, not the runtime NVRTC image --
predicated STS falls from 8 to 0, total STS from 20 to 16, instructions from 392
to 344 and IADD from 78 to 67, with registers at 64, shared memory at 1536 bytes
and no spills either way. That describes the listing and prices nothing.
pp381 b32 117.63, 117.46, 117.16 against 92.80, 92.69, 92.55
medians 117.46 against 92.69, +26.7%, ranges disjoint
No bound moved. 26 focused MMA, topology, batched-parity, sequential and STANDARD
parity and lifecycle tests pass with the parity numbers unchanged, and the
cross-width comparison was driven explicitly.
Same window as every capture here: the last 24 chunk executions, two measured passes, 762 tokens, compilation chunk and warm-up excluded. Kernel time in that window falls from 8,184.4 ms to 6,047.9 ms. Gate and up are 17.05% and 17.32%, 34.37% combined; projectionMMAQ4_0 across its eight task names is 66.01%. Q4_0 and Q4_1 ffn_down are listed separately, at 14.53% and 3.71%, since only the first was affected by the recent work. The per-call comparison carries its own controls: Q5_K ssm_out, Q4_1 ffn_down, attention, attn_output and the delta-rule scan all use kernels nothing touched and move by at most 1.6%, while the six projectionMMAQ4_0 tasks fall 29-38%. Task times only -- no occupancy, bandwidth or hardware cause is claimed, and no counters exist on this host.
…ojections projectionMMAQ5_K and projectionMMAQ4_1 carried the structure projectionMMAQ4_0 had before 3d35094 and ad44666: four shared allocations and a branch per element choosing between them on both operands. Each now uses one allocation per operand with the second panel at a byte offset. Decoding, scale and minimum arithmetic, geometry, activation reads, staging order, barriers and numerical bounds are unchanged; Q4_0 is untouched, and the two kernels remain separable hunks. Bounds were checked per kernel rather than inherited. Both use the same BM, BK and PANEL and the same store arguments, so a B panel's in-panel address is at most 254 before the swizzle -- which permutes within the same 256 bytes -- and an A panel's per-lane address reaches at most 496 of its 512. Emission for both shows one shared int array and one shared half array, loads at __bo += 0, += 256 and += 512, and unpredicated stores. Verified before timing, per kernel, against its accepted implementation at its production shape: Q5_K ssm_out 32x5120x6144 and Q4_1 ffn_down 32x5120x17408, full chunk and 29-row padded chunk, NaN-poisoned destinations -- bit-identical, 327,680 values each, all finite. Isolated screens, warmed, resident weights, medians of nine rounds, two alternating rounds per build: Q5_K ssm_out 0.4587, 0.4510 ms against 0.5086, 0.5199 ms Q4_1 ffn_down 0.8666, 0.8700 ms against 1.3646, 1.2967 ms Whole model, pp381 b32, interleaved, four runs each, all preserved: accepted 120.08, 121.03, 117.74, 117.98 median 119.03, mean 119.21 both merges 123.58, 122.73, 124.34, 120.68 median 123.16, mean 122.83 That is an observed improvement of about 3% -- +3.47% on medians, +3.04% on means -- with a few percent of run-to-run variation in both builds and overlapping ranges. It is not a proven minimum gain, and the isolated screens are the cleaner evidence that each kernel improved. No bound moved. The focused gate passes: 27 tests over both MMA dtypes, topology, batched parity at widths 2, 7, 32 and 64, sequential and STANDARD parity, and both lifecycle tests, with the parity numbers unchanged; the cross-width comparison was driven explicitly and returns identical logits at both widths.
The manual sequence -- two 2-byte global loads, a shift-and-or, a shared store,
per slot -- becomes one asyncCopyToLocal per slot, lowering to
cp.async.ca.shared.global [dst],[src],4 with a commit and a wait before the
barrier that publishes the tiles. The merged A/B allocations, geometry, weight
decoding, arithmetic, MMA order and barriers are unchanged, B staging is
untouched, and Q5_K and Q4_1 keep their manual staging.
The sm_80 floor is read from TornadoVM's compiler, not inferred from a
successful sm_120 run: CUDATensorCoreSupportPhase enforces compute capability
8.0 and gates the cp.async nodes in the same set as the MMA nodes, which this
kernel already contains. Any device that can run this kernel has already cleared
that floor, so no device selection changes and no new capability, gate or option
is introduced; ineligible devices and shapes keep the scalar tiled kernels.
Alignment holds against the kernel's actual reduction dimension, not a variable
name: at every call site the first argument of mmaEligible -- the one checked
% 32 == 0 -- is the value passed as k. So 32 divides k, k is even, and
base = (blockRow + row) * k + kBase + half * BK + kk is even for every copy,
making the source byte address header + 2 * base four-byte aligned.
Each lane issues eight copies to its own destinations: i = lane + slot * 32 over
lane < 32 and slot < 8 takes every value in [0,256) exactly once, and the
destination half * 128 + j equals i, so the 256-int tile is covered completely
with no overlap. Every lane commits and waits before the publishing barrier, and
the trailing barrier still separates a round's reads from the next round's
writes.
Bit-identical to the accepted kernel at 32x17408x5120 on a full chunk and a
29-row padded chunk, from NaN-poisoned destinations, 1,114,112 values, all
finite.
pp381 b32 139.94, 139.69, 139.51 against 121.21, 120.80, 120.57
medians 139.69 against 120.80, +15.6%, ranges disjoint,
within-build spread 0.53% and 0.31%
What produced that is not established. The change replaces a load-pack-store
sequence with one copy instruction per slot; whether any transfer overlapped the
B staging that follows was not measured, and there is no cross-iteration
pipelining here.
No bound moved. The focused gate passes: 27 tests over the MMA host references,
dispatch and topology, batched parity at widths 2, 7, 32 and 64, synthetic
batched parity, sequential and STANDARD parity and both lifecycle tests, with the
parity numbers unchanged; the cross-width payloads are identical at both widths.
…and Q4_1 The A-staging replacement that works in projectionMMAQ4_0 was applied unchanged to projectionMMAQ5_K and projectionMMAQ4_1 and reverted. Every prerequisite holds for both -- the reduction dimension is the argument mmaEligible checks, so 32 divides k and every source address is four-byte aligned; the staging loop is the same one, so each lane copies to its own destinations and the 256-int tile is covered exactly once; commit and wait precede the publishing barrier, and the trailing barrier still separates a round's reads from the next round's writes; and the sm_80 floor is the gate their MMA nodes already pass. NVRTC rejects the result: identifier "half" is undefined, on the shared tile declaration. CUDACompilationResultBuilder prepends cuda_fp16.h only when the emitted source contains __half, half2 or 2half, and with cp.async staging these two kernels emit none of those spellings: no __half_as_ushort from the A loads, a bare ((half *) ...) B store, and byte-arithmetic scale decoding. Q4_0 survives the same change only because 5e68ffc made its scale read emit __half2float. Both keep their manual A staging. The alternatives are a TornadoVM change to that scan or an fp16 spelling added to the kernel to satisfy a text match; neither was pursued, and no timing was taken. Also corrects the Q4_0 comment to say that each lane issues its own eight copies, with the covering argument, rather than "the same eight copies".
…a reproducer A kernel that stores into a half-precision shared tile fails to compile unless some other construct happens to spell fp16 in the generated text. Both offending lines are emitted by the backend itself -- the tile by allocateHalfFloatLocalArray, the store by mmaStoreBSwizzled -- while the include is decided by scanning the source for __half, half2 or 2half, which does not cover the bare half the backend writes. MissingFp16IncludeRepro.java is standalone: one kernel, four arrays, one m16n8k16 step, no model and nothing from this engine. It has two modes that differ by a single getHalfFloat read; the first fails with identifier "half" is undefined, the second compiles and prints a result. A compile-time constant value hides the defect, since new HalfFloat(0.5f) emits __float2half, so the reproducer stores a value computed at run time. The record carries the environment, the exact build and run commands, expected against actual output, the generated excerpt, the source location of the condition and a suspected fix that has not been applied. No TornadoVM change and no issue filed. It also states that the engine-side move of the Q5_K and Q4_1 scale reads to getHalfFloat is an avoidance rather than a fix: the defect is unchanged and any kernel that stores into a half tile without another fp16 spelling still fails.
…finding 5 The record named the version string only, and 6.0.1-jdk21-dev is a development build: it does not identify the source the quoted include condition and its line numbers came from. The build in use is ae7152e20797b13902590183e3f07e06fc76843b on develop, git describe build-lock-v6.0.0-29-gae7152e20.
projectionMMAQ5_K and projectionMMAQ4_1 assembled their fp16 block fields from two byte loads through halfFromBytes, which lowers to a ten-branch software expansion. They now read them with ByteArray.getHalfFloat, which lowers to a single __half2float -- the same change projectionMMAQ4_0 took in 5e68ffc. Manual A staging, scale and minimum algebra, geometry, barriers, MMA order and numerical bounds are unchanged. The complete field inventory, checked per field rather than inherited from the Q4_0 case: Q4_1 scale offset 0 of a 20-byte block -> 20k even Q4_1 minimum offset 2 of a 20-byte block -> 20k+2 even Q5_K d offset 0 of a 176-byte block -> 176k even Q5_K dmin offset 2 of a 176-byte block -> 176k+2 even Q5_K six-bit sub-block scales at K_SCALES_OFFSET are not halves and stay byte reads getHalfFloat requires a two-byte-aligned byte index, which every one of those addresses satisfies. It reads a native-endian short where halfFromBytes composed the pair little-endian; the two agree on this little-endian CUDA target, which is where the equivalence was verified, and no claim is made for any other. HalfFloatConversionAccelTest gains a case that reads every finite half encoding at each of those layouts -- strides 18, 20 and 176, offsets 0 and 2 -- and holds the value against Float.float16ToFloat: 63,488 encodings per layout, bit for bit, including 31,744 negatives, 2,046 subnormals and both signed zeros. Infinities and NaN are excluded, as elsewhere in this repository, because a quantized block field is neither. Outputs are bit-identical to the accepted kernels at the production shapes -- Q5_K ssm_out 32x5120x6144 and Q4_1 ffn_down 32x5120x17408 -- on a full chunk and a 29-row padded chunk, from NaN-poisoned destinations, 327,680 values each, all finite. Isolated kernel screens, medians of nine rounds, warmed, resident weights, two alternating rounds per build: Q5_K ssm_out 0.4224, 0.4142 ms against 0.4331, 0.4330 ms Q4_1 ffn_down 0.7538, 0.7550 ms against 0.8315, 0.8373 ms This change also happens to restore the fp16 spelling TornadoVM's include scan looks for, which is what unblocks the async staging that follows. That is an avoidance of the defect recorded in docs/architecture/tornadovm-issues, not a fix for it.
The manual A-tile sequence in projectionMMAQ5_K and projectionMMAQ4_1 -- two 2-byte global loads, a shift-and-or, a shared store, per slot -- becomes one asyncCopyToLocal per slot, with a commit and a wait before the barrier that publishes the tiles, exactly as projectionMMAQ4_0 does since 8b04c83. Merged tiles, geometry, weight decoding, scale and minimum algebra, B staging, MMA order and numerical bounds are unchanged, and there is no double buffering. Verified per kernel rather than by analogy. Alignment: the dispatch guard checks the value passed as k -- mmaEligible(valueDim, dim) for ssm_out and mmaEligible(hiddenDim, dim) for the early ffn_down -- so 32 divides k, k is even, and base = (blockRow + row) * k + kBase + half * BK + kk is even for every copy, making the source byte address header + 2 * base four-byte aligned. Coverage: i = lane + slot * 32 over lane < 32 and slot < 8 takes every value in [0,256) exactly once and the destination half * 128 + j equals i, so each lane copies to its own slots and the tile is written completely, without overlap. Ordering: every lane commits and waits before the publishing barrier, and the trailing barrier still separates a round's MMA reads from the next round's copies. Architecture: CUDATensorCoreSupportPhase gates the cp.async nodes in the same set as the MMA nodes at compute capability 8.0, which these kernels already require. This compiles only because the preceding commit restored an fp16 spelling in both kernels; without it NVRTC rejects them with identifier "half" is undefined. That defect is recorded with a standalone reproducer in docs/architecture/tornadovm-issues and is avoided here, not fixed. Bit-identical to the accepted kernels at the production shapes -- Q5_K ssm_out 32x5120x6144 and Q4_1 ffn_down 32x5120x17408 -- on a full chunk and a 29-row padded chunk, from NaN-poisoned destinations, 327,680 values each, all finite. The emission carries eight cp.async copies, one commit and one wait per round. Isolated screens, medians of nine rounds, two alternating rounds per build: Q5_K ssm_out 0.3670, 0.3653 ms against 0.4331, 0.4330 ms Q4_1 ffn_down 0.6315, 0.6325 ms against 0.8315, 0.8373 ms Whole model, pp381 b32, interleaved, identical settings: accepted 147.15, 146.51, 146.36 median 146.51 both stages 149.93, 150.03, 149.60 median 149.93 +2.33% observed within that session, ranges disjoint, and not a guaranteed minimum; the absolute numbers are well above the previous session's on the same code path, so only the interleaved comparison carries. No bound moved. The focused gate passes: 29 tests over the MMA host references, the half-conversion layouts, dispatch and topology, batched parity at widths 2, 7, 32 and 64, synthetic batched parity, sequential and STANDARD parity and the FP16-KV lifecycle, with the parity numbers unchanged; the cross-width payloads are identical at both widths.
The stage-two paragraph was written before its own commit existed and carried a placeholder hash.
Same window as every capture here -- the last 24 chunk executions, two measured passes, 762 tokens, 1,536 layer graph executions, compilation chunk and warm-up excluded. Window kernel time falls 6,047.9 to 4,979.2 ms; task-graph 5,855.2 ms, copy-in 122.3 ms. Gate and up are 16.67% and 16.53%, 33.21% combined. By kernel: projectionMMAQ4_0 63.55% over eight task names, projectionMMAQ5_K 7.91%, projectionMMAQ4_1 2.41%. Non-projection work is 26.14%: delta-rule scan 8.89%, attention 4.83%, attn_output_proj 4.12% on the scalar tiled kernel, the two RMS reductions 3.67%, and activation preparation 0.67%. The per-call comparison against the ad44666 capture carries its own controls: the delta-rule scan, attention, attn_output_proj and batchedRmsReduce use kernels nothing touched and move by at most 2%, while the Q4_0 projections fall 19-22%, Q5_K ssm_out 28% and Q4_1 ffn_down 47%. Task times only; no hardware cause is claimed and no counters exist on this host.
…e old gates proved Qwen35BatchPrefillLayers reads llama.qwen35.tensorCores into a static final at class initialization. No parity, lifecycle or cross-width class set it, so every one of those runs built the scalar batched plan -- matrixVectorTiledBatchQ4_0 and its siblings, not projectionMMAQ4_0, projectionMMAQ5_K or projectionMMAQ4_1. Harmless while the MMA changes were bit-identical; not coverage of them. The claim that those gates validated the tensor-core path is withdrawn in both records. The direct kernel comparisons and the performance runs are unaffected: they used the MMA kernels themselves, or --tensor-cores. New coverage, each class in its own JVM under reuseForks=false, with the property set in a static initializer before anything touches the layer class: - Qwen35MmaBatchedPrefillParityAccelTest, width 32, which against the 63-token fixture is also the partial-final-chunk case; - Qwen35MmaBatchedPrefillWidth64ParityAccelTest, a width wider than the prompt; - Qwen35MmaBatchedFp16KvLifecycleAccelTest, the FP16 key/value lifecycle and its session reset over the MMA plan; - Qwen35MmaCrossWidthCaptureAccelTest, the capture tool built against the MMA kernels, still requiring two runs to assert anything; - theBatchedPlanSelectsTensorCoresPerWidth in the topology test, which builds the plan at widths 32, 64 and an ineligible 8, asserts the MMA task names and their one-warp-per-tile grids at the first two and the scalar tasks at the third. A property being set is not execution evidence, so the parity classes assert Qwen35BatchPrefillLayers.tensorCoresSelected() -- a reporting accessor added for this, consulted by nothing in production -- which returns the value of the static the plan in that process was built with. The scalar classes are untouched and still cover the ineligible widths. 16 tests, 0 failures, 1 skipped on the accepted build; the skip is the cross-width capture, which asserts nothing when driven once.
The same passages and helper as Qwen35NllScreenAccelTest -- the unscored prefix ingested through the batched path, the identical teacher-forced continuation positions scored -- with llama.qwen35.tensorCores set in a static initializer so the prefix goes through the MMA plan rather than the scalar one. Drive it with -Dllama.nllScreen.batch=32. It is the same reused development screen as its parent: five passages from this repository's own files, one register, one domain, already used for several accept/reject decisions. Not independent quality validation.
mmaEligible excluded attn_output_proj because it folds its residual, which
left it as the one batched projection still on
matrixVectorTiledBatchWithResidualQ4_0. It now takes the same three-task
shape ffn_down and ssm_out take -- convertToFP16, projectionMMAQ4_0,
residualAdd -- within the existing gate: no change to the --tensor-cores
default, no new switch, no bounds moved.
Buffers are reused from this graph's task order, not from capacity alone.
wrapHbFP16BatchMMA stages the converted activation and is rewritten by
ffn_down_fp16 later in the same layer, with nothing reading it in between;
wrapFFNDownBatch holds exactly this projection's output width and is read
only by the residual pass that follows each write; wrapXBatch, the residual
destination, is touched by the add alone.
FP16 multiplicands are not bit-identical to the scalar kernel, so this moves
the arithmetic. Against the unchanged bounds, at widths 32 and 64: relL2
0.0400938 (bound .12), minCosine 0.99920286 (bound .99550), maxAbs 0.835752
(ceiling 2.29034), violations 45.30% (budget 50%), argmax 0/63. Interleaved
pp381 b32 with tensor cores, FP16 KV and CUDA graphs: 143.57/143.33/142.92
to 147.03/146.77/146.48 tok/s, median +2.40%.
Coverage that shows the plan is the MMA one. Setting llama.qwen35.tensorCores
in a static initializer selects the path; reading the layer class's folded
value back says only that it was asked for, so tensorCoresSelected() is gone
and each MMA class now reads the grid scheduler of the plan it built --
PlanDispatchEvidence, the seam ProgramIdentity.gridEntries and
ResolvedPolicyReachesThePlanAccelTest already use -- asserting per attention
layer the conversion and residual tasks only the MMA branch adds and the
projection's grid of one warp per BM x BN output tile. That is the configured
dispatch and geometry of the plan that ran, from its own process; it is not a
trace of individual task executions. The parity, lifecycle, cross-width and
NLL harnesses each check their own plan, the last three through a
verifyDispatch hook, and the lifecycle one through a package-private
DelegatingSession.planIfBuilt() because the facade hides the plan. An
unreachable plan shape records no scheduler rather than failing the generic
capture; a missing one fails the MMA assertions.
15 tests, 0 failures, 0 skipped on Qwen35Mma{BatchedPrefill,Width64,Fp16Kv}
and the topology test; cross-width driven at 32 and 64 compares 15,644,160
logits and 63 token ids identical, digest ccd4c4fc.
…ivity Two records against the accepted build, neither changing any code. The double-buffering experiment is closed as rejected. It staged only the A activation tile through two buffers -- prologue, cp.async.wait_group 1 in the steady state so the next block's copy stays in flight, wait_group 0 in the drain -- for 1024 more bytes of shared memory per workgroup, with geometry, weight decode, arithmetic and accumulation order unchanged. Interleaved pp381 b32 with CUDA graphs off, which is what llama.cudaGraphs defaults to and what that invocation ran: 143.97 and 143.74 t/s baseline against 84.77 candidate. The cause was not measured and the run was stopped rather than diagnosed. Outputs were bit-identical to the accepted kernel on the shapes and inputs tested, including single-block and partial-chunk cases; that is agreement where it was checked, not proof for every shape and not the exclusion of a compiler defect. The batch-width screen is configuration sensitivity, not a speedup: at pp1024, widths 32, 64 and 128 give 136.06, 134.39 and 140.12 t/s over two interleaved rounds, all asserting qwen35/Q4_0/BATCH_PREFILL_DECODE and all dispatching the same MMA plan -- attn_output_fp16 in 16 attention layers, ffn_down_fp16 in 64 -- with no allocation failure at a 22 GB budget. The ordering is non-monotonic and was not explained. The default width stays where it is.
…needs The row-pair candidate gave one warp two stacked m16n8k16 tiles against a single staged B panel, so the weight panel both tiles read was decoded once rather than once per tile: BN, the one-warp workgroup, the merged tiles, the native scale reads and the cp.async staging all unchanged, A grown from two panels to four, a second accumulator, two stores, and eligibility at batchSize % 32 == 0 with 16 and 48 keeping the accepted kernel and its grid. It was bit-identical to the accepted kernel wherever it was compared -- poisoned outputs, row-distinct activations, k of 32, 64, 96, 512 and 6144, m=64, and both production shapes -- and it lost 19.1%: 145.71 and 145.97 baseline against 117.91 candidate at pp381 b32 with tensor cores, FP16 KV and CUDA graphs, run interleaved with a discarded warm-up. What that establishes is narrow, and the record says so. The combined implementation lost; it isolates neither decode volume nor occupancy nor latency hiding, no counter was read, and the generated code was never dumped. None of the three rejected experiments -- wide tile, A double buffering, row pairing -- has an established cause. Their jars and logs are preserved outside the tree; future candidates get their source and emission captured before the revert, not after. Finding 3 is also corrected. The integer MMQ capability is not "scale a fragment": each s32 block partial has to be converted to floating point, multiplied by its own weight and activation scales and added into a persistent fp32 sum, with the integer accumulator reset or consumed every block because it cannot carry a differently-scaled partial forward. An in-place fragment scale would not by itself supply that, and which primitive is the right upstream request is not settled here.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.