[flang][cuda] Lower clock() to NNVM op

[clementval]

Also use a same gen function for all NVVM time ops.

[llvmbot]

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

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

Changes

Also use a same gen function for all NVVM time ops.

---

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

4 Files Affected:

• (modified) flang/include/flang/Optimizer/Builder/IntrinsicCall.h (+2-1)
• (modified) flang/lib/Optimizer/Builder/IntrinsicCall.cpp (+10-8)
• (modified) flang/module/cudadevice.f90 (+5)
• (modified) flang/test/Lower/CUDA/cuda-device-proc.cuf (+3)

diff --git a/flang/include/flang/Optimizer/Builder/IntrinsicCall.h b/flang/include/flang/Optimizer/Builder/IntrinsicCall.h
index 363b1d5844d1b..b15dd29d68f65 100644
--- a/flang/include/flang/Optimizer/Builder/IntrinsicCall.h
+++ b/flang/include/flang/Optimizer/Builder/IntrinsicCall.h
@@ -282,7 +282,6 @@ struct IntrinsicLibrary {
                         llvm::ArrayRef<mlir::Value> args);
   mlir::Value genGetUID(mlir::Type resultType,
                         llvm::ArrayRef<mlir::Value> args);
-  mlir::Value genGlobalTimer(mlir::Type, llvm::ArrayRef<mlir::Value>);
   fir::ExtendedValue genHostnm(std::optional<mlir::Type> resultType,
                                llvm::ArrayRef<fir::ExtendedValue> args);
   fir::ExtendedValue genIall(mlir::Type, llvm::ArrayRef<fir::ExtendedValue>);
@@ -377,6 +376,8 @@ struct IntrinsicLibrary {
   fir::ExtendedValue genNorm2(mlir::Type, llvm::ArrayRef<fir::ExtendedValue>);
   mlir::Value genNot(mlir::Type, llvm::ArrayRef<mlir::Value>);
   fir::ExtendedValue genNull(mlir::Type, llvm::ArrayRef<fir::ExtendedValue>);
+  template <typename OpTy>
+  mlir::Value genNVVMTime(mlir::Type, llvm::ArrayRef<mlir::Value>);
   fir::ExtendedValue genPack(mlir::Type, llvm::ArrayRef<fir::ExtendedValue>);
   fir::ExtendedValue genParity(mlir::Type, llvm::ArrayRef<fir::ExtendedValue>);
   void genPerror(llvm::ArrayRef<fir::ExtendedValue>);
diff --git a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
index ddfa27475fa7a..ce767750a56cd 100644
--- a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
@@ -385,6 +385,7 @@ static constexpr IntrinsicHandler handlers[]{
      &I::genChdir,
      {{{"name", asAddr}, {"status", asAddr, handleDynamicOptional}}},
      /*isElemental=*/false},
+    {"clock", &I::genNVVMTime<mlir::NVVM::ClockOp>, {}, /*isElemental=*/false},
     {"clock64", &I::genClock64, {}, /*isElemental=*/false},
     {"cmplx",
      &I::genCmplx,
@@ -503,7 +504,7 @@ static constexpr IntrinsicHandler handlers[]{
     {"getgid", &I::genGetGID},
     {"getpid", &I::genGetPID},
     {"getuid", &I::genGetUID},
-    {"globaltimer", &I::genGlobalTimer, {}, /*isElemental=*/false},
+    {"globaltimer", &I::genNVVMTime<mlir::NVVM::GlobalTimerOp>, {}, /*isElemental=*/false},
     {"hostnm",
      &I::genHostnm,
      {{{"c", asBox}, {"status", asAddr, handleDynamicOptional}}},
@@ -4320,13 +4321,6 @@ mlir::Value IntrinsicLibrary::genGetUID(mlir::Type resultType,
                                fir::runtime::genGetUID(builder, loc));
 }
 
-// GLOBALTIMER
-mlir::Value IntrinsicLibrary::genGlobalTimer(mlir::Type resultType,
-                                             llvm::ArrayRef<mlir::Value> args) {
-  assert(args.size() == 0 && "globalTimer takes no args");
-  return builder.create<mlir::NVVM::GlobalTimerOp>(loc, resultType).getResult();
-}
-
 // GET_COMMAND_ARGUMENT
 void IntrinsicLibrary::genGetCommandArgument(
     llvm::ArrayRef<fir::ExtendedValue> args) {
@@ -7207,6 +7201,14 @@ IntrinsicLibrary::genNull(mlir::Type, llvm::ArrayRef<fir::ExtendedValue> args) {
   return fir::MutableBoxValue(boxStorage, mold->nonDeferredLenParams(), {});
 }
 
+// CLOCK, GLOBALTIMER
+template <typename OpTy>
+mlir::Value IntrinsicLibrary::genNVVMTime(mlir::Type resultType,
+                                          llvm::ArrayRef<mlir::Value> args) {
+  assert(args.size() == 0 && "expect no arguments");
+  return builder.create<OpTy>(loc, resultType).getResult();
+}
+
 // PACK
 fir::ExtendedValue
 IntrinsicLibrary::genPack(mlir::Type resultType,
diff --git a/flang/module/cudadevice.f90 b/flang/module/cudadevice.f90
index 52a619e07165c..05113d86e616f 100644
--- a/flang/module/cudadevice.f90
+++ b/flang/module/cudadevice.f90
@@ -957,6 +957,11 @@ attributes(device) pure integer function atomicxori(address, val)
 
   ! Time function
 
+  interface
+    attributes(device) integer function clock()
+    end function
+  end interface
+
   interface
     attributes(device) integer(8) function clock64()
     end function
diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf
index 888c7961ee2b4..2d6f734670740 100644
--- a/flang/test/Lower/CUDA/cuda-device-proc.cuf
+++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf
@@ -10,6 +10,7 @@ attributes(global) subroutine devsub()
   integer(4) :: ai
   integer(8) :: al
   integer(8) :: time
+  integer :: smalltime
 
   call syncthreads()
   call syncwarp(1)
@@ -45,6 +46,7 @@ attributes(global) subroutine devsub()
   ai = atomicinc(ai, 1_4)
   ai = atomicdec(ai, 1_4)
 
+  smalltime = clock()
   time = clock64()
 
   time = globalTimer()
@@ -84,6 +86,7 @@ end
 ! CHECK: %{{.*}} = llvm.atomicrmw uinc_wrap  %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32
 ! CHECK: %{{.*}} = llvm.atomicrmw udec_wrap  %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32
 
+! CHECK: %{{.*}} = nvvm.read.ptx.sreg.clock : i32
 ! CHECK: fir.call @llvm.nvvm.read.ptx.sreg.clock64()
 ! CHECK: %{{.*}} = nvvm.read.ptx.sreg.globaltimer : i64