[flang][cuda] Use NVVM op for clock64

[clementval]

NVVM op is now available for clock64

[llvmbot]

@llvm/pr-subscribers-flang-fir-hlfir

Author: Valentin Clement (バレンタイン クレメン) (clementval)

Changes

NVVM op is now available for clock64

---

Full diff: https://github.com/llvm/llvm-project/pull/149223.diff

2 Files Affected:

• (modified) flang/lib/Optimizer/Builder/IntrinsicCall.cpp (+1-5)
• (modified) flang/test/Lower/CUDA/cuda-device-proc.cuf (+1-1)

diff --git a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
index 8d0a511744e25..e9a1753a537a1 100644
--- a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
@@ -3563,11 +3563,7 @@ IntrinsicLibrary::genChdir(std::optional<mlir::Type> resultType,
 // CLOCK64
 mlir::Value IntrinsicLibrary::genClock64(mlir::Type resultType,
                                          llvm::ArrayRef<mlir::Value> args) {
-  constexpr llvm::StringLiteral funcName = "llvm.nvvm.read.ptx.sreg.clock64";
-  mlir::MLIRContext *context = builder.getContext();
-  mlir::FunctionType ftype = mlir::FunctionType::get(context, {}, {resultType});
-  auto funcOp = builder.createFunction(loc, funcName, ftype);
-  return builder.create<fir::CallOp>(loc, funcOp, args).getResult(0);
+  return builder.create<mlir::NVVM::Clock64Op>(loc, resultType).getResult();
 }
 
 // CMPLX
diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf
index 42ee7657966e2..0f72a0d89f3e7 100644
--- a/flang/test/Lower/CUDA/cuda-device-proc.cuf
+++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf
@@ -82,7 +82,7 @@ end
 ! CHECK: %{{.*}} = llvm.atomicrmw uinc_wrap  %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32
 ! CHECK: %{{.*}} = llvm.atomicrmw udec_wrap  %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32
 
-! CHECK: fir.call @llvm.nvvm.read.ptx.sreg.clock64()
+! CHECK: %{{.*}} = nvvm.read.ptx.sreg.clock64 : i64
 
 subroutine host1()
   integer, device :: a(32)