diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/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. diff --git a/clang/test/CodeGenHIP/amdgpu_hostcall.cpp b/clang/test/CodeGenHIP/amdgpu_hostcall.cpp new file mode 100644 --- /dev/null +++ b/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 struct HIP_vector_base; + +template +struct HIP_vector_base { using Native_vec_ = T __attribute__((ext_vector_type(2))); }; + + +extern "C" __device__ uint64_t __ockl_fprintf_stderr_begin(); + +extern "C" __device__ HIP_vector_base::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 +