What happens
Jump threading already knows convergent operations are special: it refuses to
duplicate any block that contains one. But there's a second way the pass can
hurt them that has no check.
When the pass threads an edge, it makes some paths skip a redundant branch and
jump directly to their known destination. The destination block itself is not
copied or changed but it gains a new incoming edge from a different part
of the function. Nothing checks what's inside the destination block, so this
is currently allowed even when it contains convergent operations.
That's a problem on GPUs. Instructions like NVPTX's bar.sync and ldmatrix
need all threads of a warp to arrive together, and (without convergence
tokens) the places where threads re-join are worked out from the shape of the
CFG, specifically from the merge points that all paths funnel through. The
new edge lets some paths enter the block from elsewhere, bypassing that merge
point. After that, there is no valid spot left to re-join the warp before the
convergent instruction, so it executes with only part of the warp present.
How to see it
Attached unopt.ll is real unoptimized IR from a Triton kernel. Just run:
opt -passes='default<O3>' unopt.ll -S -o out.ll
In the output, the block with the bar.sync + ldmatrix has become
.thread312 : it now has an extra predecessor coming from a completely
different region, added by jump threading (confirmed by bisecting with
-opt-bisect-limit: the CFG changes exactly at the jump-threading pass).
The pattern in the kernel is simple: the same opaque (runtime) condition is
tested twice, so on one path the second test is redundant, and the pass
threads that path straight into the block containing the barrier.
On real hardware this miscompiles. The ldmatrix runs with a partially
converged warp.
Downstream reports: triton-lang/triton#11148, triton-lang/triton#11153;
Triton is currently working around it by disabling the pass for functions
with convergent ops (triton-lang/triton#11156) which isn't ideal.
unopt.ll.txt
What happens
Jump threading already knows convergent operations are special: it refuses to
duplicate any block that contains one. But there's a second way the pass can
hurt them that has no check.
When the pass threads an edge, it makes some paths skip a redundant branch and
jump directly to their known destination. The destination block itself is not
copied or changed but it gains a new incoming edge from a different part
of the function. Nothing checks what's inside the destination block, so this
is currently allowed even when it contains convergent operations.
That's a problem on GPUs. Instructions like NVPTX's
bar.syncandldmatrixneed all threads of a warp to arrive together, and (without convergence
tokens) the places where threads re-join are worked out from the shape of the
CFG, specifically from the merge points that all paths funnel through. The
new edge lets some paths enter the block from elsewhere, bypassing that merge
point. After that, there is no valid spot left to re-join the warp before the
convergent instruction, so it executes with only part of the warp present.
How to see it
Attached
unopt.llis real unoptimized IR from a Triton kernel. Just run:In the output, the block with the
bar.sync+ldmatrixhas become.thread312: it now has an extra predecessor coming from a completelydifferent region, added by jump threading (confirmed by bisecting with
-opt-bisect-limit: the CFG changes exactly at thejump-threadingpass).The pattern in the kernel is simple: the same opaque (runtime) condition is
tested twice, so on one path the second test is redundant, and the pass
threads that path straight into the block containing the barrier.
On real hardware this miscompiles. The
ldmatrixruns with a partiallyconverged warp.
Downstream reports: triton-lang/triton#11148, triton-lang/triton#11153;
Triton is currently working around it by disabling the pass for functions
with convergent ops (triton-lang/triton#11156) which isn't ideal.
unopt.ll.txt