Amir Ghamarian
e21b915381
Fix bs32 tiny tier race condition: promote to small tier
...
The 1-warp 16x16 MFMA tiny tier with kPageBlockSize=32 has a
non-deterministic race condition in block_tile_reduce_sync when
the last KV block has 9+ valid tokens (crossing a half-MFMA
boundary). Promote bs32 decode to the small tier (2-warp, 32x32
MFMA) which uses permlane32_swap and is deterministic.
Made-with: Cursor
2026-03-30 15:10:00 +00:00
..
2026-01-15 07:19:31 -08:00
2026-01-07 16:30:57 +01:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-12-18 07:59:45 +01:00
2025-11-28 13:49:54 -08:00
2025-10-16 03:10:57 -07:00
2026-01-07 16:30:57 +01:00
2026-01-07 16:30:57 +01:00
2026-01-15 16:43:02 +01:00
2026-01-13 07:14:23 +01:00
2026-01-07 16:30:57 +01:00
2025-12-18 07:59:45 +01:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-12-18 07:59:45 +01:00
2025-11-28 13:49:54 -08:00
2026-01-07 16:30:57 +01:00
2026-01-07 16:30:57 +01:00
2026-01-07 16:30:57 +01:00
2026-01-17 08:30:27 +01:00
2025-12-15 07:16:01 -08:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2026-01-17 08:30:27 +01:00
2026-01-07 16:30:57 +01:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2026-01-07 16:30:57 +01:00
2025-11-28 13:49:54 -08:00
2026-01-07 16:30:57 +01:00
2026-01-07 16:30:57 +01:00
2026-01-07 16:30:57 +01:00
2025-12-30 16:25:08 +01:00
2025-11-28 13:49:54 -08:00
2026-01-07 16:30:57 +01:00
2025-11-28 13:49:54 -08:00
2026-01-07 16:30:57 +01:00
2025-11-28 13:49:54 -08:00
2026-01-07 16:30:57 +01:00
2026-01-07 16:30:57 +01:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2025-12-18 13:12:15 -07:00
2025-11-28 13:49:54 -08:00
2025-11-28 13:49:54 -08:00
2026-01-15 16:43:02 +01:00
2025-11-28 13:49:54 -08:00
2026-01-07 16:30:57 +01:00
2026-01-15 16:43:02 +01:00
2026-01-15 16:43:02 +01:00
2026-03-30 15:10:00 +00:00
2026-01-14 07:31:45 -08:00
2024-12-04 00:46:47 +01:00