[flang-commits] [flang] 43bfec2 - [flang][cuda] Fix condition in barrier_try_wait lowering (#171916)

[via flang-commits]

Author: Valentin Clement (バレンタイン クレメン)
Date: 2025-12-11T14:48:22-08:00
New Revision: 43bfec29cbecc1ff2e5aa6f8908c4d63e9c896c5

URL: https://github.com/llvm/llvm-project/commit/43bfec29cbecc1ff2e5aa6f8908c4d63e9c896c5
DIFF: https://github.com/llvm/llvm-project/commit/43bfec29cbecc1ff2e5aa6f8908c4d63e9c896c5.diff

LOG: [flang][cuda] Fix condition in barrier_try_wait lowering (#171916)

The condition should have been while 0 then try to wait and not the
opposite.

Added: 
    

Modified: 
    flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
    flang/test/Lower/CUDA/cuda-device-proc.cuf

Removed: 
    


################################################################################
diff  --git a/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp b/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
index 3c86a9d7451f0..fe2db4607f86b 100644
--- a/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
@@ -1010,7 +1010,7 @@ CUDAIntrinsicLibrary::genBarrierTryWait(mlir::Type resultType,
   mlir::Value beforeArg = beforeBlock->addArgument(resultType, loc);
   builder.setInsertionPointToStart(beforeBlock);
   mlir::Value condition = mlir::arith::CmpIOp::create(
-      builder, loc, mlir::arith::CmpIPredicate::ne, beforeArg, zero);
+      builder, loc, mlir::arith::CmpIPredicate::eq, beforeArg, zero);
   mlir::scf::ConditionOp::create(builder, loc, condition, beforeArg);
   mlir::Block *afterBlock = builder.createBlock(&whileOp.getAfter());
   afterBlock->addArgument(resultType, loc);

diff  --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf
index 27ef8e0889627..ca6695c26e5d5 100644
--- a/flang/test/Lower/CUDA/cuda-device-proc.cuf
+++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf
@@ -513,6 +513,8 @@ end subroutine
 
 ! CHECK-LABEL: func.func @_QPtest_barrier_try_wait()
 ! CHECK: scf.while
+! CHECK: %[[COND:.*]] = arith.cmpi eq
+! CHECK: scf.condition(%[[COND]])
 ! CHECK: %{{.*}} = nvvm.inline_ptx "{\0A  .reg .pred p;\0A  mbarrier.try_wait.shared.b64 p, [%{{.*}}], %{{.*}}, %{{.*}};\0A  selp.b32 %{{.*}}, 1, 0, p;\0A}" ro(%{{.*}}, %{{.*}}, %{{.*}} : !llvm.ptr, i64, i32) -> i32
 
 attributes(global) subroutine test_barrier_try_wait_sleep()