CUDA kernel contract: the Qwen 3.5 decode attention cases launch tf_q35_attn_partial and tf_q35_attn_partial8 at 192 threads, past their __launch_bounds__(128), so every gate fails with CUDA_ERROR_INVALID_VALUE #1267

Activity

1. Bizuayeu commented on Oct 25, 2026

   

   Bizuayeu

   on Oct 25, 2026

   ContributorAuthor

   More actions

   I read the case against the kernels and the engine on 3d219d21. Three things the case launches have moved since a56d0ce2:

   1. The partials' launch. attn_one takes its row, chunk and key head from blockIdx.x, .y and .z and is written for 128 threads (qwen35_attn_tc.cuh:152-153, :172), and the bf16 partial keeps its tiles in attn_tiles, 65,536 bytes of dynamic shared memory that the engine opts into (cuda_engine.zig:225). The engine launches both partials at grid (rows, chunks, key heads) with 128 threads (cuda_forward.zig:81, :84); the case launches them at (rows, key heads, chunks) with 192 threads and no dynamic shared memory.
   2. A long stream's rows. A partial returns for a stream of GROUP_LEN (8,192) keys or more (qwen35_attn_tc.cuh:156), and tf_q35_attn_rows / tf_q35_attn_rows8 take those rows, at grid (row groups, chunks, key heads) with 256 threads (cuda_forward.zig:82, :85). The case's fourth served shape, 16 rows at 9,000 keys, is such a stream, so with the partials alone its rows would have nothing to merge.
   3. A prompt chunk. One stream's prompt chunk goes through cuda_prompt.zig's tiles instead of the decode partials when the query heads per key head are even (cuda_forward.zig:73), as the case's 24 over 4 are; its third shape, a 128-row chain, is that call in the engine.

   So the fix is in the case: launch what the engine launches, the partial at (rows, chunks, key heads) with 128 threads (and attn_tiles with the opt-in for bf16), the rows kernel at (row groups, chunks, key heads) with 256 threads for the long stream, then the merge as now, and keep the prompt-chunk shape on the partials or move it to the prompt path as the gate's coverage asks. The control that drops each row's last chunk goes through the same launches.

   Until the case follows the kernels: the refused launches do not lose the context, so the rest of the contract runs either way, and a run that needs a clean SUMMARY can leave these three registrations out with one line in contract_tests.zig (on GB10 together with the W4A8 line, #1264):

       // for (&q35s) |*x| try kernels.append(a, x.kernel());

   I have not run either on GB10 yet. I can run the kernel contract with the change and post the result here.
2. Bizuayeu commented on Oct 25, 2026

   

   Bizuayeu

   on Oct 25, 2026

   ContributorAuthor

   More actions

   I ran the kernel contract on GB10 at 464824fd4 (main with 50fa3c7a0, "cuda contract: match Qwen TC attention launches"), with the Qwen 3.5 registrations in: zig build kernel-contract -Dnvcc=/usr/local/cuda/bin/nvcc -Dsm=121, nvcc 13.3.73, driver 580.178.04, Zig 0.17.0, sm_121. The launches are no longer refused, and the run reaches its SUMMARY:

   SUMMARY 336 kernels: 334 pass gates 1-4, 2 fail, 3 known findings; 89 controls: 89 red as planted, 0 missed; 291 limit problems; gates 1873218 ms, limits 1020243 ms
   

   Both decode attention registrations pass repeat, shapes and sums, their nonfinite gate gives the known finding, and the control that drops each row's last chunk, which the refused launches had left missed, goes red at shapes:

   FAIL nonfinite cuda qwen35 decode attention, bf16 cache (tf_q35_attn_partial, tf_q35_attn_merge): head 18 element 0: finite where float64 gives nan
   KNOWN nonfinite cuda qwen35 decode attention, bf16 cache (tf_q35_attn_partial, tf_q35_attn_merge): a NaN key in one of a row's several chunks leaves the row finite: tf_q35_attn_merge skips that chunk
   PASS repeat cuda qwen35 decode attention, bf16 cache (tf_q35_attn_partial, tf_q35_attn_merge): five runs byte-identical (196608 bytes), inputs and guards untouched, no overlap
   PASS shapes cuda qwen35 decode attention, bf16 cache (tf_q35_attn_partial, tf_q35_attn_merge): 13 shapes bit-equal to the kernel it replaces (worst 0.000 of the bound), 0 edges refused by the kernel
   PASS sums cuda qwen35 decode attention, bf16 cache (tf_q35_attn_partial, tf_q35_attn_merge): exact sums at k 9000, where fp16 sums would miss by infx and bf16 by 2107142x
   FAIL nonfinite cuda qwen35 decode attention, fp8 cache (tf_q35_attn_partial8, tf_q35_attn_merge): head 18 element 0: finite where float64 gives nan
   KNOWN nonfinite cuda qwen35 decode attention, fp8 cache (tf_q35_attn_partial8, tf_q35_attn_merge): a NaN key in one of a row's several chunks leaves the row finite: tf_q35_attn_merge skips that chunk
   PASS repeat cuda qwen35 decode attention, fp8 cache (tf_q35_attn_partial8, tf_q35_attn_merge): five runs byte-identical (196608 bytes), inputs and guards untouched, no overlap
   PASS shapes cuda qwen35 decode attention, fp8 cache (tf_q35_attn_partial8, tf_q35_attn_merge): 13 shapes bit-equal to the kernel it replaces (worst 0.000 of the bound), 0 edges refused by the kernel
   PASS sums cuda qwen35 decode attention, fp8 cache (tf_q35_attn_partial8, tf_q35_attn_merge): exact sums at k 9000, where fp16 sums would miss by infx and bf16 by 1512858x
   FAIL shapes control qwen35 attention dropping each row's last chunk: at 16x24x2100 its bytes differ from the chunk-major grid (each row alone)'s from byte 0
   CONTROL control qwen35 attention dropping each row's last chunk: red at shapes, as planted
   

   The limits check finds the five kernels' launch resources as recorded:

   PASS limits live qwen35 tf_q35_attn_partial: 128 threads launched, the driver allows 128; 168 registers, 11648 + 65536 shared bytes (recorded 168, 11648)
   PASS limits live qwen35 tf_q35_attn_rows: 256 threads launched, the driver allows 256; 246 registers, 6912 + 65536 shared bytes (recorded 246, 6912)
   PASS limits live qwen35 tf_q35_attn_merge: 32 threads launched, the driver allows 32; 42 registers, 0 + 0 shared bytes (recorded 42, 0)
   PASS limits live qwen35_kv8 tf_q35_attn_partial8: 128 threads launched, the driver allows 128; 197 registers, 44672 + 0 shared bytes (recorded 197, 44672)
   PASS limits live qwen35_kv8 tf_q35_attn_rows8: 256 threads launched, the driver allows 256; 245 registers, 39936 + 0 shared bytes (recorded 245, 39936)
   

   Two of the three known findings are these; the third is GLM's tf_glm_scores. The two fails are W4A8's nonfinite gate (#1264).
3. ashhart commented on Oct 25, 2026

   

   ashhart

   on Oct 25, 2026

   Owner

   More actions

   Fixed on main in 50fa3c7 (81b476d..2ff0191). The Qwen 3.5 decode attention contract adapter now launches exactly what the engine launches for the tensor-core kernels: (rows, chunks, key head) grid order, 128 threads on the short-row entry and 256 on the long-row entry, and the 64 KiB dynamic shared opt-in for bf16. A recording-driver host test pins the symbols, grids, blocks, shared bytes and arguments. On a GB10 the full contract now passes all 362 kernels, and both decode attention cases leave only their documented known finding (a NaN key in one chunk leaves the row finite, because the merge skips that chunk). Thank you for the precise report.

   - Synthema, Ash's AI Assistant.

   👍React with 👍1Reacted by ビーズちゃん
4. ashhart

   closed this as completedon Oct 25, 2026