[LoopUnswitch] Deny loop unswitch if its body contains volatile inline asm

[kovdan01]

Summary

Volatile inline asm might contain convergent insructions.
Unswitching such loops breaks code correctness for targets like NVPTX.

[fhahn]

If inline assembly calls can be convergent in general, should CallBase::isConvergent handle this? otherwise we might run into similar mis-compiles in other transforms.

[kovdan01]

In D117303#3243350, @fhahn wrote:

If inline assembly calls can be convergent in general, should CallBase::isConvergent handle this? otherwise we might run into similar mis-compiles in other transforms.

I suppose that CallBase::isConvergent can’t handle this. Without the change, test with volatile inline asm fails (the unswitching is performed) – it indicates that marking inline asm as volatile does not set convergent attribute on corresponding CallBase.

[nikic]

The frontend is responsible for adding the convergent attribute to inline asm if necessary, see https://github.com/llvm/llvm-project/blob/e58e401b798845ebaf131528d286b39dadfda914/clang/lib/CodeGen/CGStmt.cpp#L2245-L2250.

[kovdan01]

In D117303#3243390, @nikic wrote:

The frontend is responsible for adding the convergent attribute to inline asm if necessary, see https://github.com/llvm/llvm-project/blob/e58e401b798845ebaf131528d286b39dadfda914/clang/lib/CodeGen/CGStmt.cpp#L2245-L2250.

Thank you for your review! Didn't know that. Appropriate tests already exist (for example, clang/test/CodeGenCUDA/convergent.cu), so close this revision.

llvm/lib/Transforms/Scalar/LoopUnswitch.cpp

 #include "llvm/IR/Attributes.h"
 #include "llvm/IR/BasicBlock.h"
 #include "llvm/IR/Constant.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/IR/DerivedTypes.h"
 #include "llvm/IR/Dominators.h"
 #include "llvm/IR/Function.h"
 #include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/InlineAsm.h"
 #include "llvm/IR/InstrTypes.h"
 #include "llvm/IR/Instruction.h"
 #include "llvm/IR/Instructions.h"
 #include "llvm/IR/IntrinsicInst.h"
 #include "llvm/IR/Intrinsics.h"
 #include "llvm/IR/Module.h"
 #include "llvm/IR/Type.h"
 #include "llvm/IR/User.h"
bool LoopUnswitch::processCurrentLoop() {
 
   for (const auto BB : CurrentLoop->blocks()) {
     for (auto &I : *BB) {
       auto *CB = dyn_cast<CallBase>(&I);
       if (!CB)
         continue;
       if (CB->isConvergent())
         return Changed;
+      // Volatile inline asm may contain convergent instructions
+      // e.g., PTX shfl.sync, so rejecting them for correctness
+      if (auto *Asm = dyn_cast<InlineAsm>(CB->getCalledOperand()))
+        if (Asm->hasSideEffects())
+          return Changed;
       if (auto *II = dyn_cast<InvokeInst>(&I))
         if (!II->getUnwindDest()->canSplitPredecessors())
           return Changed;
       if (auto *II = dyn_cast<IntrinsicInst>(&I))
         if (II->getIntrinsicID() == Intrinsic::experimental_guard)
           Guards.push_back(II);
     }
   }

llvm/test/Transforms/LoopUnswitch/convergent-hoist-modified.ll

 ; RUN: opt -loop-unswitch -enable-new-pm=0 %s -S | FileCheck %s
 
 ; When hoisting simple values out from a loop, and not being able to unswitch
 ; the loop due to the convergent call, the pass would return an incorrect
 ; Modified status. This was caught by the pass return status check that is
 ; hidden under EXPENSIVE_CHECKS.
 
-; CHECK-LABEL: entry:
-; CHECK-NEXT: %0 = call i32 @llvm.objectsize.i32.p0i8(i8* bitcast (%struct.anon* @b to i8*), i1 false, i1 false, i1 false)
-; CHECK-NEXT: %1 = icmp uge i32 %0, 1
-; CHECK-NEXT: br label %for.cond
-
 %struct.anon = type { i16 }
 
 @b = global %struct.anon zeroinitializer, align 1
 
+; Function Attrs: nounwind readnone speculatable willreturn
+declare i32 @llvm.objectsize.i32.p0i8(i8*, i1 immarg, i1 immarg, i1 immarg) #1
+
+declare void @conv() convergent
+
 ; Function Attrs: nounwind
 define i16 @c() #0 {
+; CHECK-LABEL: entry:
+; CHECK-NEXT: %0 = call i32 @llvm.objectsize.i32.p0i8(i8* bitcast (%struct.anon* @b to i8*), i1 false, i1 false, i1 false)
+; CHECK-NEXT: %1 = icmp uge i32 %0, 1
+; CHECK-NEXT: br label %for.cond
 entry:
   br label %for.cond
 
 for.cond:                                         ; preds = %cont, %entry
   br label %for.inc
 
 for.inc:                                          ; preds = %for.cond
   %0 = call i32 @llvm.objectsize.i32.p0i8(i8* bitcast (%struct.anon* @b to i8*), i1 false, i1 false, i1 false)
   %1 = icmp uge i32 %0, 1
   br i1 %1, label %cont, label %cont
 
 cont:                                             ; preds = %for.inc
   call void @conv() convergent
   %2 = load i16, i16* getelementptr inbounds (%struct.anon, %struct.anon* @b, i32 0, i32 0), align 1
   br label %for.cond
 }
 
-; Function Attrs: nounwind readnone speculatable willreturn
-declare i32 @llvm.objectsize.i32.p0i8(i8*, i1 immarg, i1 immarg, i1 immarg) #1
+; Function Attrs: nounwind
+define i16 @d() #0 {
+; CHECK-LABEL: entry:
+; CHECK-NEXT: %0 = call i32 @llvm.objectsize.i32.p0i8(i8* bitcast (%struct.anon* @b to i8*), i1 false, i1 false, i1 false)
+; CHECK-NEXT: %1 = icmp uge i32 %0, 1
+; CHECK-NEXT: br label %for.cond
+entry:
+  br label %for.cond
 
-declare void @conv() convergent
+for.cond:                                         ; preds = %cont, %entry
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.cond
+  %0 = call i32 @llvm.objectsize.i32.p0i8(i8* bitcast (%struct.anon* @b to i8*), i1 false, i1 false, i1 false)
+  %1 = icmp uge i32 %0, 1
+  br i1 %1, label %cont, label %cont
+
+cont:                                             ; preds = %for.inc
+  call void asm sideeffect "; some convergent instructions here", ""()
+  %2 = load i16, i16* getelementptr inbounds (%struct.anon, %struct.anon* @b, i32 0, i32 0), align 1
+  br label %for.cond
+}
 
 attributes #0 = { nounwind }
 attributes #1 = { nounwind readnone speculatable willreturn }