tlopex opened a new pull request, #19741: URL: https://github.com/apache/tvm/pull/19741
This fixes a correctness issue in predicated `cp.async` generated from `if_then_else(predicate, global_load, 0)`. Previously, when the predicate was false, the generated predicated `cp.async` skipped the copy but did not zero-fill the shared-memory destination, losing the original `if_then_else(..., 0)` semantics. This also fixes the async wait count for interleaved async producers. The steady-state logical wait count was `5`, but the prologue can coalesce the A/B copies into fewer physical CUDA commit groups. At the first body iteration, `wait_group_5` may therefore not wait for the oldest prologue group, causing the consumer to read incomplete shared memory. The wait count is adjusted to account for that prologue/body boundary case. The test no longer relies on the full generated CUDA source snapshot or the dead CSE variable shape. Instead, it checks the relevant codegen invariants and runs the generated CUDA kernel to validate the actual output. -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected] --------------------------------------------------------------------- To unsubscribe, e-mail: [email protected] For additional commands, e-mail: [email protected]
