mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-21 06:40:23 +08:00
431c85f100
* fix: T.Parallel + call_extern + address_of correctness in warp-specialized pipelines Two bugs caused incorrect results when T.Parallel loops use T.call_extern with T.address_of(shared_buffer) inside pipelined, warp-specialized kernels: Bug 1 (lower_tile_op.cc): has_non_local_store did not recognize address_of(shared_buffer) inside call_extern as a non-local access, so parallel_loop was set to false and thread partitioning was skipped entirely — every thread ran the full serial loop. Bug 3 (producer_consumer_ws.cc): wait_insert_pos was only computed for kTmaProducer statements. SIMT/cp.async producers (e.g. T.Parallel global→shared copies) were skipped, so the consumer read from SIMT-produced shared buffers before the barrier wait — a data race. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * lin * fix ws async producer wait placement --------- Co-authored-by: Rachmanino <wutong.1109@bytedance.com> Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>