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

[Valentin Clement バレンタイン クレメン via flang-commits]

https://github.com/clementval created https://github.com/llvm/llvm-project/pull/171916

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



>From c38d758ed0ddc80c53536f4ff1f2e5c49f607ae6 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 11 Dec 2025 13:52:16 -0800
Subject: [PATCH] [flang][cuda] Fix condition in barrier_try_wait lowering

---
 flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp | 2 +-
 flang/test/Lower/CUDA/cuda-device-proc.cuf        | 2 ++
 2 files changed, 3 insertions(+), 1 deletion(-)

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()