Add hasSideEffects flag to intra-warp shuffle inline assembly

ROCm · Aug 15, 2023 · d0b7793 · d0b7793
1 parent 4215086
commit d0b7793
Showing 1 changed file with 2 additions and 1 deletion.
diff --git a/lib/Conversion/TritonGPUToLLVM/Utility.cpp b/lib/Conversion/TritonGPUToLLVM/Utility.cpp
@@ -139,6 +139,7 @@ Value shflSync(Location loc, ConversionPatternRewriter &rewriter, Value val,
   }
   auto swait = builder.create("s_waitcnt lgkmcnt(0)");
   (*swait)();
+  return builder.launch(rewriter, loc, val.getType(), true);
 #else
   PTXBuilder builder;
   auto &shfl = builder.create("shfl.sync")->o("bfly").o("b32");
@@ -148,8 +149,8 @@ Value shflSync(Location loc, ConversionPatternRewriter &rewriter, Value val,
   auto *cOpr = builder.newConstantOperand("0x1f");
   auto *maskOpr = builder.newConstantOperand("0xffffffff");
   shfl(dOpr, aOpr, bOpr, cOpr, maskOpr);
-#endif
   return builder.launch(rewriter, loc, val.getType(), false);
+#endif
 }
 
 Value addStringToModule(Location loc, ConversionPatternRewriter &rewriter,