[flang-commits] [flang] 89d9a83 - [flang][cuda] Use NVVM op for barrier0 intrinsic (#140947)

[via flang-commits]

Author: Valentin Clement (バレンタイン クレメン)
Date: 2025-05-21T13:05:14-07:00
New Revision: 89d9a83b704a8f6b5bd64dac93095a9228c601d5

URL: https://github.com/llvm/llvm-project/commit/89d9a83b704a8f6b5bd64dac93095a9228c601d5
DIFF: https://github.com/llvm/llvm-project/commit/89d9a83b704a8f6b5bd64dac93095a9228c601d5.diff

LOG: [flang][cuda] Use NVVM op for barrier0 intrinsic (#140947)

The simple form of `Barrier0Op` is available in the NVVM dialect. It is
needed to use it instead of the string version since
https://github.com/llvm/llvm-project/pull/140615

Added: 
    

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

Removed: 
    


################################################################################
diff  --git a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
index 1ac0627da9524..178b6770d6b53 100644
--- a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
@@ -8332,12 +8332,7 @@ IntrinsicLibrary::genSum(mlir::Type resultType,
 
 // SYNCTHREADS
 void IntrinsicLibrary::genSyncThreads(llvm::ArrayRef<fir::ExtendedValue> args) {
-  constexpr llvm::StringLiteral funcName = "llvm.nvvm.barrier0";
-  mlir::FunctionType funcType =
-      mlir::FunctionType::get(builder.getContext(), {}, {});
-  auto funcOp = builder.createFunction(loc, funcName, funcType);
-  llvm::SmallVector<mlir::Value> noArgs;
-  builder.create<fir::CallOp>(loc, funcOp, noArgs);
+  builder.create<mlir::NVVM::Barrier0Op>(loc);
 }
 
 // SYNCTHREADS_AND

diff  --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf
index 8f5e6dd36da4e..42ee7657966e2 100644
--- a/flang/test/Lower/CUDA/cuda-device-proc.cuf
+++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf
@@ -49,7 +49,7 @@ attributes(global) subroutine devsub()
 end
 
 ! CHECK-LABEL: func.func @_QPdevsub() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
-! CHECK: fir.call @llvm.nvvm.barrier0() fastmath<contract> : () -> ()
+! CHECK: nvvm.barrier0
 ! CHECK: fir.call @llvm.nvvm.bar.warp.sync(%c1{{.*}}) fastmath<contract> : (i32) -> ()
 ! CHECK: fir.call @llvm.nvvm.membar.gl() fastmath<contract> : () -> ()
 ! CHECK: fir.call @llvm.nvvm.membar.cta() fastmath<contract> : () -> ()
@@ -106,7 +106,7 @@ end
 
 ! CHECK-LABEL: func.func @_QPhost1()
 ! CHECK: cuf.kernel
-! CHECK: fir.call @llvm.nvvm.barrier0() fastmath<contract> : () -> ()
+! CHECK: nvvm.barrier0
 ! CHECK: fir.call @llvm.nvvm.bar.warp.sync(%c1{{.*}}) fastmath<contract> : (i32) -> ()
 ! CHECK: fir.call @llvm.nvvm.barrier0.and(%c1{{.*}}) fastmath<contract> : (i32) -> i32
 ! CHECK: fir.call @llvm.nvvm.barrier0.popc(%c1{{.*}}) fastmath<contract> : (i32) -> i32