[flang-commits] [flang] [flang][cuda][NFC] Use NVVM VoteBallotOp (PR #134307)

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

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

`llvm.nvvm.vote.ballot.sync` has its own operation so use it in lowering. 

>From 5d7484ec543a2f520cb1d9d6ac338e34d6662891 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 3 Apr 2025 13:46:10 -0700
Subject: [PATCH] [flang][cuda][NFC] Use NVVM VoteBallotOp

---
 flang/lib/Optimizer/Builder/IntrinsicCall.cpp | 7 +++++--
 flang/test/Lower/CUDA/cuda-device-proc.cuf    | 2 +-
 2 files changed, 6 insertions(+), 3 deletions(-)

diff --git a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
index a562d9b7e461c..ffebeb6adc9ca 100644
--- a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
@@ -6542,8 +6542,11 @@ mlir::Value
 IntrinsicLibrary::genVoteBallotSync(mlir::Type resultType,
                                     llvm::ArrayRef<mlir::Value> args) {
   assert(args.size() == 2);
-  return genVoteSync(builder, loc, "llvm.nvvm.vote.ballot.sync",
-                     builder.getI32Type(), args);
+  mlir::Value arg1 =
+      builder.create<fir::ConvertOp>(loc, builder.getI1Type(), args[1]);
+  return builder
+      .create<mlir::NVVM::VoteBallotOp>(loc, resultType, args[0], arg1)
+      .getResult();
 }
 
 // MATCH_ANY_SYNC
diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf
index dbce4a5fa47dd..6a546a5e4b78d 100644
--- a/flang/test/Lower/CUDA/cuda-device-proc.cuf
+++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf
@@ -305,7 +305,7 @@ end subroutine
 ! CHECK-LABEL: func.func @_QPtestvote()
 ! CHECK: fir.call @llvm.nvvm.vote.all.sync
 ! CHECK: fir.call @llvm.nvvm.vote.any.sync
-! CHECK: fir.call @llvm.nvvm.vote.ballot.sync
+! CHECK: %{{.*}} = nvvm.vote.ballot.sync %{{.*}}, %{{.*}} : i32
 
 ! CHECK-DAG: func.func private @__ldca_i4x4_(!fir.ref<!fir.array<4xi32>>, !fir.ref<!fir.array<4xi32>>)
 ! CHECK-DAG: func.func private @__ldcg_i4x4_(!fir.ref<!fir.array<4xi32>>, !fir.ref<!fir.array<4xi32>>)