[ifcnv] Don't duplicate blocks that contain convergent instructions.

[jlebar]

Summary

Also add a comment briefly explaining what ifcnv is.

It's unsafe to duplicate blocks that contain convergent instructions during ifcnv. See the patch for details.

[jlebar]

Friendly ping. Based on the discussion in D17430, this does seem confusing enough to be worth a comment.

[rnk]

I'm not sure I understand the comment. This transform is basically turning ifs into predicated instructions, right? And we're basically saying that a predicated convergent operation (threadsync or what have you) is the same as having the operation in a conditional basic block?

[jlebar]

In D17518#387824, @rnk wrote:

I'm not sure I understand the comment. This transform is basically turning
ifs into predicated instructions, right? And we're basically saying that a
predicated convergent operation (threadsync or what have you) is the same as
having the operation in a conditional basic block?

Not quite. The simple case of converting

if (pred) __syncthreads();
==>
(predicated on pred) __syncthreads();

is trivially safe, like I think you're saying. But ifconv can also perform
transformations which duplicate the newly-predicated instruction. The question
is whether or not it's safe to duplicate a convergent instruction, in the
process of if-conversion.

A simple example is

   BB0
  /   \
 BB1  BB2
 |\_ _/
 | | |
 | TBB --> exit
 |
FBB

TBB forms ifconv's "simple" shape with its predecessors. Can we move TBB's
contents into BB1 and BB2 (as predicated instructions) if TBB contains a
convergent op?

I'd argued in D17430 that this was safe, but now I think it is not. If we
think concretely in terms of CUDA, a necessary condition for correctness is
that

• if all threads in the warp executed TBB "together" (i.e., convergently) before the transformation,
• then all threads in the warp must execute *the same copy* of TBB's instructions after the transformation.

But clearly that's not necessarily true in this example.

In terms of the LLVM spec, we say that you're not allowed to "add" a
control-flow dependency to a convergent op. In this case, if BB1 switches on
"cond", then the original CFG dependency expression for TBB is

(BB1 && cond) || BB2

But after predication, we have two copies of the convergent operation:

BB1 && cond
BB2

I think my mistake earlier was concluding from this that we "removed" CFG
dependencies, rather than adding them. But let me write this in an equivalent
way:

BB1 && cond && !BB2
BB2 && !BB1

Now it's suddenly clear that we added a CFG dependency to each of these new
instructions.

Put another way, there's a set of control flows which result in us reaching an
instruction. If I add a CFG dependency, I make that set *smaller*. That's
what we're not allowed to do with convergent operations, and that's what we're
doing here.

@resistor, I'm going to revert this back to what we originally had in D17430.
Reid, thank you for being confused by this. :)

[jlebar]

Go back to disallowing this conversion; see comment in code.

[hfinkel]

lib/CodeGen/IfConversion.cpp

691

instr -> instruction

693

I understand (I think) what you're saying here, but everything from here down seems confusing. I think this is because 'cond' is not in the diagram. I think you can either:

1. Put 'cond' in the diagram, and then say that you've gone from having no control dependencies on the convergent instruction to have a control dependency on 'cond' or '!cond' depending on into which block the instruction is duplicated.
2. If I'm reading this correctly and BB1 and BB2 below indicate the control dependencies of their respective blocks, then you can do the same as above, but say that you've gone from having no relevant control dependencies to have a dependency 'BB1' or 'BB2'.

Feel free to tell me I'm missing something ;)

[jlebar]

Update per Hal's review.

[jlebar]

Thank you for the review, Hal.

This is really pushing my ASCII art skills. lmk if this makes any more sense.

[hfinkel]

In D17518#401665, @jlebar wrote:

Thank you for the review, Hal.

This is really pushing my ASCII art skills. lmk if this makes any more sense.

I like this much better, thanks! LGTM.

[jlebar]

Excellent. Thank you for the review!

lib/CodeGen/IfConversion.cpp

void IfConverter::ScanInstructions(BBInfo &BBI) {
   BBI.NonPredSize = 0;
   BBI.ExtraCost = 0;
   BBI.ExtraCost2 = 0;
   BBI.ClobbersPred = false;
   for (auto &MI : *BBI.BB) {
     if (MI.isDebugValue())
       continue;
 
-    if (MI.isNotDuplicable())
+    // It's unsafe to duplicate convergent instructions in this context, so set
+    // BBI.CannotBeCopied to true if MI is convergent.  Consider the following
+    // CFG, which is subject to our "simple" transformation.
+    //
+    //   BB1  BB2
+    //   |\_ _/
+    //   | | |
+    //   | TBB --> exit
+    //   |
+    //   FBB
+    //
+    // Suppose we want to move TBB's contents up into BB1 and BB2, and suppose
+    // TBB contains a convergent instruction.  This is safe iff doing so does
+    // not add a control-flow dependency to the convergent instruction -- i.e.,
+    // it's safe if the set of control flows that leads us to the convergent
+    // instr does not get smaller after the transformation.
+    //
+    // Assuming BB1 switches on "cond", originally we got to TBB if
+    //
+    //   (BB1 && cond) || BB2.
+    //
+    // After the transformation, there are two copies of TBB's instructions:
+    //
+    //   BB1 && cond, and
+    //   BB2.
+    //
+    // There are clearly fewer ways to satisfy the condition "BB1 && cond" than
+    // "(BB1 && cond) || BB2".  Since we've shrunk the set of control flows
+    // which lead to our convergent instruction, the transformation is unsafe.
+    if (MI.isNotDuplicable() || MI.isConvergent())
       BBI.CannotBeCopied = true;
 
     bool isPredicated = TII->isPredicated(MI);
     bool isCondBr = BBI.IsBrAnalyzable && MI.isConditionalBranch();
 
     // A conditional branch is not predicable, but it may be eliminated.
     if (isCondBr)
       continue;