[flang-commits] [flang] WIP: Trying to get a call treated as an intrinsic (PR #120020)

[Renaud Kauffmann via flang-commits]

https://github.com/Renaud-K created https://github.com/llvm/llvm-project/pull/120020

I am trying to get the call to syncthreads1 identified as an intrinsic. I have modelled the changes after ieee_set_rounding_mode.
I am not able to get it to even try to do look-up for a possible intrinsic match. (As the debug messages show)
It works for the ieee call though. I am not sure what I am missing.


```
module mtests
  integer, device :: n(5) = (/ (I, I = 3,7) /)
  interface syncthreads1
    subroutine syncthreads1_0()
    end subroutine
  end interface
  contains
    attributes(global) subroutine testany() bind(c)
      integer i
      i = threadIdx%x
      if (i < 3) then 
        n(i) = 1
        !call syncthreads1()
      endif
    end subroutine
end module mtests
 
program t
use mtests
use ieee_arithmetic
call testany<<<1,5>>> ()
call syncthreads1()
call ieee_set_rounding_mode(ieee_to_zero)
end
```

>From 1b0239b5b7042950cff6afab7ea8766f110e6101 Mon Sep 17 00:00:00 2001
From: Renaud-K <rkauffmann at nvidia.com>
Date: Sun, 15 Dec 2024 13:11:39 -0800
Subject: [PATCH] Trying to get a call treated as an intrinsic

---
 .../flang/Optimizer/Builder/IntrinsicCall.h   |  1 +
 flang/lib/Optimizer/Builder/IntrinsicCall.cpp | 21 ++++++++++++++++++-
 2 files changed, 21 insertions(+), 1 deletion(-)

diff --git a/flang/include/flang/Optimizer/Builder/IntrinsicCall.h b/flang/include/flang/Optimizer/Builder/IntrinsicCall.h
index bc0020e614db24..77683ad4b3c7b1 100644
--- a/flang/include/flang/Optimizer/Builder/IntrinsicCall.h
+++ b/flang/include/flang/Optimizer/Builder/IntrinsicCall.h
@@ -392,6 +392,7 @@ struct IntrinsicLibrary {
   fir::ExtendedValue genSum(mlir::Type, llvm::ArrayRef<fir::ExtendedValue>);
   void genSignalSubroutine(llvm::ArrayRef<fir::ExtendedValue>);
   void genSleep(llvm::ArrayRef<fir::ExtendedValue>);
+  void genSyncThreads(llvm::ArrayRef<fir::ExtendedValue>);
   fir::ExtendedValue genSystem(std::optional<mlir::Type>,
                                mlir::ArrayRef<fir::ExtendedValue> args);
   void genSystemClock(llvm::ArrayRef<fir::ExtendedValue>);
diff --git a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
index 547cebefd2df47..c358c492f66a5d 100644
--- a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
@@ -642,6 +642,7 @@ static constexpr IntrinsicHandler handlers[]{
        {"dim", asValue},
        {"mask", asBox, handleDynamicOptional}}},
      /*isElemental=*/false},
+    {"syncthreads1", &I::genSyncThreads},
     {"system",
      &I::genSystem,
      {{{"command", asBox}, {"exitstat", asBox, handleDynamicOptional}}},
@@ -1639,8 +1640,9 @@ mlir::Value toValue(const fir::ExtendedValue &val, fir::FirOpBuilder &builder,
 //===----------------------------------------------------------------------===//
 
 static bool isIntrinsicModuleProcedure(llvm::StringRef name) {
+  llvm::errs() << "isIntrinsicModuleProcedure: " << name << "\n";
   return name.starts_with("c_") || name.starts_with("compiler_") ||
-         name.starts_with("ieee_") || name.starts_with("__ppc_");
+         name.starts_with("ieee_") || name.starts_with("__ppc_") || name == "syncthreads1";
 }
 
 static bool isCoarrayIntrinsic(llvm::StringRef name) {
@@ -1684,6 +1686,7 @@ lookupIntrinsicHandler(fir::FirOpBuilder &builder,
                        llvm::StringRef intrinsicName,
                        std::optional<mlir::Type> resultType) {
   llvm::StringRef name = genericName(intrinsicName);
+  llvm::errs() << "Looking up " << intrinsicName << " with name " << name << "\n";
   if (const IntrinsicHandler *handler = findIntrinsicHandler(name))
     return std::make_optional<IntrinsicHandlerEntry>(handler);
   bool isPPCTarget = fir::getTargetTriple(builder.getModule()).isPPC();
@@ -7290,6 +7293,22 @@ IntrinsicLibrary::genSum(mlir::Type resultType,
                       resultType, args);
 }
 
+// SYNCTHREADS
+void IntrinsicLibrary::genSyncThreads(llvm::ArrayRef<fir::ExtendedValue> args) {
+  constexpr llvm::StringLiteral funcName = "llvm.nvvm.barrier0";
+  mlir::func::FuncOp funcOp = builder.getNamedFunction(funcName);
+  mlir::MLIRContext *context = builder.getContext();
+  mlir::FunctionType funcType =
+      mlir::FunctionType::get(context, {}, {});
+
+  if (!funcOp) 
+    funcOp = builder.createFunction(loc, funcName, funcType);
+  
+  llvm::SmallVector<mlir::Value> noArgs;
+  builder.create<fir::CallOp>(loc, funcOp, noArgs);
+
+}
+
 // SYSTEM
 fir::ExtendedValue
 IntrinsicLibrary::genSystem(std::optional<mlir::Type> resultType,