[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.

Author: clementval

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

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