[PATCH] D115283: [AMDGPU] Set "amdgpu_hostcall" module flag if an AMDGPU function has calls to device lib functions that use hostcalls.

[Konstantin Pyzhov via Phabricator via cfe-commits]

kpyzhov updated this revision to Diff 392577.
Herald added subscribers: kerbowa, nhaehnle, jvesely.

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D115283/new/

https://reviews.llvm.org/D115283

Files:
  clang/lib/CodeGen/TargetInfo.cpp
  clang/test/CodeGenHIP/amdgpu_hostcall.cpp


Index: clang/test/CodeGenHIP/amdgpu_hostcall.cpp
===================================================================
--- /dev/null
+++ clang/test/CodeGenHIP/amdgpu_hostcall.cpp
@@ -0,0 +1,48 @@
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -DFN_HOSTCALL \
+// RUN:   -o - %s | FileCheck --enable-var-scope %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -DFN_PRINTF \
+// RUN:   -o - %s | FileCheck --enable-var-scope %s
+
+// CHECK: !llvm.module.flags
+// CHECK: "amdgpu_hostcall"
+
+
+typedef unsigned long int uint64_t;
+
+#define __device__ __attribute__((device))
+
+template<typename T, unsigned int n> struct HIP_vector_base;
+
+template<typename T>
+struct HIP_vector_base<T, 2> { using Native_vec_ = T __attribute__((ext_vector_type(2))); };
+
+
+extern "C" __device__ uint64_t __ockl_fprintf_stderr_begin();
+
+extern "C" __device__ HIP_vector_base<long long, 2>::Native_vec_ __ockl_call_host_function(
+    uint64_t fptr, uint64_t arg0, uint64_t arg1, uint64_t arg2, uint64_t arg3, uint64_t arg4, uint64_t arg5, uint64_t arg6);
+
+
+#ifdef FN_HOSTCALL
+__device__ void fn_hostcall(uint64_t fptr, uint64_t* retval0, uint64_t* retval1) {
+  uint64_t arg0 = (uint64_t)fptr;
+  uint64_t arg1 = 0;
+  uint64_t arg2 = 0;
+  uint64_t arg3 = 0;
+  uint64_t arg4 = 0;
+  uint64_t arg5 = 0;
+  uint64_t arg6 = 0;
+  uint64_t arg7 = 0;
+
+  __ockl_call_host_function(arg0, arg1, arg2, arg3, arg4, arg5, arg6, arg7);
+}
+#endif
+
+#ifdef FN_PRINTF
+__device__ void fn_printf() {
+  auto msg = __ockl_fprintf_stderr_begin();
+}
+#endif
+
Index: clang/lib/CodeGen/TargetInfo.cpp
===================================================================
--- clang/lib/CodeGen/TargetInfo.cpp
+++ clang/lib/CodeGen/TargetInfo.cpp
@@ -9194,6 +9194,11 @@
                             llvm::Value *BlockLiteral) const override;
   bool shouldEmitStaticExternCAliases() const override;
   void setCUDAKernelCallingConvention(const FunctionType *&FT) const override;
+
+  virtual void checkFunctionCallABI(CodeGenModule &CGM, SourceLocation CallLoc,
+                                    const FunctionDecl *Caller,
+                                    const FunctionDecl *Callee,
+                                    const CallArgList &Args) const override;
 };
 }
 
@@ -9417,6 +9422,24 @@
       FT, FT->getExtInfo().withCallingConv(CC_OpenCLKernel));
 }
 
+void AMDGPUTargetCodeGenInfo::checkFunctionCallABI(CodeGenModule &CGM,
+                                                   SourceLocation CallLoc,
+                                                   const FunctionDecl *Caller,
+                                                   const FunctionDecl *Callee,
+                                                   const CallArgList &Args) const
+{
+  // Set the "amdgpu_hostcall" module flag if "Callee" is a library function
+  // that uses AMDGPU hostcall mechanism.
+  if (Callee &&
+      (Callee->getName() == "__ockl_call_host_function" ||
+       Callee->getName() == "__ockl_fprintf_stderr_begin")) {
+    llvm::Module &M = CGM.getModule();
+    if (!M.getModuleFlag("amdgpu_hostcall")) {
+      M.addModuleFlag(llvm::Module::Override, "amdgpu_hostcall", 1);
+    }
+  }
+}
+
 //===----------------------------------------------------------------------===//
 // SPARC v8 ABI Implementation.
 // Based on the SPARC Compliance Definition version 2.4.1.


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D115283.392577.patch
Type: text/x-patch
Size: 3421 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20211207/5a2af23d/attachment.bin>