[clang] [AMDGPU][clang] provide device implementation for __builtin_logb and … (PR #129347)

[via cfe-commits]

https://github.com/choikwa created https://github.com/llvm/llvm-project/pull/129347

…__builtin_scalbn

Clang generates library calls for __builtin_* functions which can be a problem for GPUs that cannot handle them. This patch generates a call to device implementation for __builtin_logb and an ldexp intrinsic for __builtin_scalbn.

>From 86deca6179c6dce4490e859ffd536f553d8d6c0d Mon Sep 17 00:00:00 2001
From: Kevin Choi <kevin.choi at amd.com>
Date: Fri, 28 Feb 2025 16:52:03 -0600
Subject: [PATCH] [AMDGPU][clang] provide device implementation for
 __builtin_logb and __builtin_scalbn

Clang generates library calls for __builtin_* functions which can be a problem for GPUs that cannot handle them.
This patch generates a call to device implementation for __builtin_logb and an ldexp intrinsic for __builtin_scalbn.
---
 clang/lib/CodeGen/CGBuiltin.cpp   | 49 ++++++++++++++++++++++++++++++-
 clang/lib/CodeGen/CodeGenModule.h |  5 ++++
 2 files changed, 53 insertions(+), 1 deletion(-)

diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 03b8d16b76e0d..6a0497df7acfb 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -259,6 +259,26 @@ Value *readX18AsPtr(CodeGenFunction &CGF) {
   return CGF.Builder.CreateIntToPtr(X18, CGF.Int8PtrTy);
 }
 
+llvm::Constant *CodeGenModule::getDeviceLibFunction(const FunctionDecl *FD,
+                                                    unsigned BuiltinID) {
+  GlobalDecl D(FD);
+  llvm::SmallString<64> Name;
+  if (getTarget().getTriple().isAMDGCN()) {
+    switch (BuiltinID) {
+    default: return nullptr;
+    case Builtin::BIlogb:
+    case Builtin::BI__builtin_logb:
+      Name = "__ocml_logb_f64";
+    }
+  }
+  if (Name.empty())
+    return nullptr;
+
+  llvm::FunctionType *Ty =
+    cast<llvm::FunctionType>(getTypes().ConvertType(FD->getType()));
+  return GetOrCreateLLVMFunction(Name, Ty, D, /*ForVTable*/false);
+}
+
 /// getBuiltinLibFunction - Given a builtin id for a function like
 /// "__builtin_fabsf", return a Function* for "fabsf".
 llvm::Constant *CodeGenModule::getBuiltinLibFunction(const FunctionDecl *FD,
@@ -6579,10 +6599,32 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
   }
   }
 
+  // Some targets like GPUs do not support library call and must provide
+  // device overload implementation.
+  if (getTarget().getTriple().isAMDGCN())
+    // Emit library call to device-lib implementation
+    if (auto *DevLibFunc = CGM.getDeviceLibFunction(FD, BuiltinID))
+      return emitLibraryCall(*this, FD, E, DevLibFunc);
+
+  // These will be emitted as Intrinsic later.
+  auto NeedsDeviceOverloadToIntrin = [&](unsigned BuiltinID) {
+    if (getTarget().getTriple().isAMDGCN()) {
+      switch (BuiltinID) {
+      default:
+        return false;
+      case Builtin::BIscalbn:
+      case Builtin::BI__builtin_scalbn:
+        return true;
+      }
+    }
+    return false;
+  };
+
   // If this is an alias for a lib function (e.g. __builtin_sin), emit
   // the call using the normal call path, but using the unmangled
   // version of the function name.
-  if (getContext().BuiltinInfo.isLibFunction(BuiltinID))
+  if (!NeedsDeviceOverloadToIntrin(BuiltinID) &&
+      getContext().BuiltinInfo.isLibFunction(BuiltinID))
     return emitLibraryCall(*this, FD, E,
                            CGM.getBuiltinLibFunction(FD, BuiltinID));
 
@@ -20804,6 +20846,11 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_s_prefetch_data:
     return emitBuiltinWithOneOverloadedType<2>(
         *this, E, Intrinsic::amdgcn_s_prefetch_data);
+  case Builtin::BIscalbn:
+  case Builtin::BI__builtin_scalbn:
+    return emitBinaryExpMaybeConstrainedFPBuiltin(
+        *this, E, Intrinsic::ldexp,
+        Intrinsic::experimental_constrained_ldexp);
   default:
     return nullptr;
   }
diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 4a269f622ece4..890dc8556cc1c 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -1231,6 +1231,11 @@ class CodeGenModule : public CodeGenTypeCache {
       llvm::FunctionType *FnType = nullptr, bool DontDefer = false,
       ForDefinition_t IsForDefinition = NotForDefinition);
 
+  /// Given a builtin id for a function, return a Function* for device
+  /// overload implementation.
+  llvm::Constant *getDeviceLibFunction(const FunctionDecl *FD,
+                                       unsigned BuiltinID);
+
   /// Given a builtin id for a function like "__builtin_fabsf", return a
   /// Function* for "fabsf".
   llvm::Constant *getBuiltinLibFunction(const FunctionDecl *FD,