mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-21 06:40:23 +08:00
Chennes
ed00dfcd7f
* [BugFix] Reject unsupported T.serial access to thread-distributed fragments with clear diagnostics A T.serial loop runs sequentially inside each CUDA thread, but a fragment's elements are distributed across threads, so indexing a fragment by a serial loop variable has no valid thread-ownership mapping. Three fuzzer reports share this root cause: - #2393: serial reduce of a GEMM-output fragment -> unhelpful internal error ("contains inner var j"). - #2395: serial write into a fragment then T.copy out -> nvcc "identifier undefined" (an owner-thread guard built from serial loop vars was emitted outside the loop that binds them). - #2396: serial write of a fragment to global -> silently wrong (threads race-write each cell). AddWrapperForSingleBufStore wraps bare fragment stores in a degenerate T.parallel(1) and only rejected constant non-zero fragment indices, so a dynamic (e.g. serial-loop) index slipped through and lowered to broken code (#2395/#2396). Reject any fragment index that is not the constant all-zero index in this fallback path, with a message pointing to T.Parallel / T.reduce_sum. Rephrase the existing inner-var layout check (#2393) to explain the cause instead of dumping "contains inner var". Also collect every access to a buffer in a statement (not just the last) so the diagnostic is not bypassed when one fragment is accessed at both a constant and a variable index in the same statement; without it such a case falls through to a cryptic downstream StructuralEqual internal error. This is a completeness hardening only, not required by the three issues above. Fixes #2393 Fixes #2395 Fixes #2396 * Generalize fragment diagnostic to non-parallel loops inner_vars_ is populated for every non-parallel For (serial, unroll, vectorized), so the message must not hard-code T.serial. Reword to "non-parallel loop variable" per review feedback. * Narrow fix scope to #2393 only Remove fragment index validation from AddWrapperForSingleBufStore and delete tests for #2395/#2396. These issues cannot be caught at this pass stage because T.grid and T.serial are indistinguishable (both appear as ForKind.SERIAL before LowerTileOp). Keep only the ParallelOpNode diagnostic improvement which successfully catches #2393.