The GitHub Actions job "Lint" on tvm.git/main has failed. Run started by GitHub user tqchen (triggered by tqchen).
Head commit for run: ece32bfaee86dc21920ee0e68e0e4a3681afbc49 / Hongyi Jin <[email protected]> [TIRx][CUDA] Preserve local mbarrier predicate and count (#20250) ## Motivation A KDA kernel-evolution experiment introduced a count-1 local handoff barrier. The producer intentionally selected one lane: ```python pred = K.cuda.elect_sync() bar.arrive(0, pred=pred) ``` However, the local branch of `MBarrier.arrive` discarded `pred` and emitted an unconditional `mbarrier.arrive.shared.b64`. Synccheck and NumSim consequently observed 32 arrivals from the producer warp where the barrier expected 1. The kernel had to work around the lowering bug by placing an unconditional `arrive()` inside an explicit elected-lane guard. The same local branch also discarded an explicit `count`, even though both arguments are part of the public `MBarrier.arrive` contract and are already preserved by the remote branches. ## Summary - preserve explicit predicate and arrival count for local `MBarrier.arrive` calls - keep local arrivals on the direct shared-memory path while leaving remote address mapping unchanged - cover both the TIR intrinsic operands and generated predicated PTX ## Testing - `python -m pytest -q tests/python/tirx/codegen/test_codegen_blackwell.py` - 14 passed - `pre-commit run --files python/tvm/backend/cuda/lang/pipeline.py tests/python/tirx/codegen/test_codegen_blackwell.py` - passed Report URL: https://github.com/apache/tvm/actions/runs/33504194151 With regards, GitHub Actions via GitBox --------------------------------------------------------------------- To unsubscribe, e-mail: [email protected] For additional commands, e-mail: [email protected]
