CC_AAPCS_VFP, // __attribute__((pcs("aapcs-vfp")))
CC_IntelOclBicc, // __attribute__((intel_ocl_bicc))
CC_SpirFunction, // default for OpenCL functions on SPIR target
- CC_SpirKernel, // inferred for OpenCL kernels on SPIR target
+ CC_OpenCLKernel, // inferred for OpenCL kernels
CC_Swift, // __attribute__((swiftcall))
CC_PreserveMost, // __attribute__((preserve_most))
CC_PreserveAll, // __attribute__((preserve_all))
case CC_X86Pascal:
case CC_X86VectorCall:
case CC_SpirFunction:
- case CC_SpirKernel:
+ case CC_OpenCLKernel:
case CC_Swift:
return false;
default:
case CC_AAPCS_VFP:
case CC_IntelOclBicc:
case CC_SpirFunction:
- case CC_SpirKernel:
+ case CC_OpenCLKernel:
case CC_PreserveMost:
case CC_PreserveAll:
// FIXME: we should be mangling all of the above.
case CC_AAPCS_VFP: return "aapcs-vfp";
case CC_IntelOclBicc: return "intel_ocl_bicc";
case CC_SpirFunction: return "spir_function";
- case CC_SpirKernel: return "spir_kernel";
+ case CC_OpenCLKernel: return "opencl_kernel";
case CC_Swift: return "swiftcall";
case CC_PreserveMost: return "preserve_most";
case CC_PreserveAll: return "preserve_all";
OS << " __attribute__((sysv_abi))";
break;
case CC_SpirFunction:
- case CC_SpirKernel:
+ case CC_OpenCLKernel:
// Do nothing. These CCs are not available as attributes.
break;
case CC_Swift:
Opts.cl_khr_3d_image_writes = 1;
}
}
+
+ CallingConvCheckResult checkCallingConvention(CallingConv CC) const override {
+ switch (CC) {
+ default:
+ return CCCR_Warning;
+ case CC_C:
+ case CC_OpenCLKernel:
+ return CCCR_OK;
+ }
+ }
};
const Builtin::Info AMDGPUTargetInfo::BuiltinInfo[] = {
}
CallingConvCheckResult checkCallingConvention(CallingConv CC) const override {
- return (CC == CC_SpirFunction || CC == CC_SpirKernel) ? CCCR_OK
- : CCCR_Warning;
+ return (CC == CC_SpirFunction || CC == CC_OpenCLKernel) ? CCCR_OK
+ : CCCR_Warning;
}
CallingConv getDefaultCallingConv(CallingConvMethodType MT) const override {
#include "clang/Frontend/CodeGenOptions.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/IR/Attributes.h"
+#include "llvm/IR/CallingConv.h"
#include "llvm/IR/CallSite.h"
#include "llvm/IR/DataLayout.h"
#include "llvm/IR/InlineAsm.h"
/***/
-static unsigned ClangCallConvToLLVMCallConv(CallingConv CC) {
+unsigned CodeGenTypes::ClangCallConvToLLVMCallConv(CallingConv CC) {
switch (CC) {
default: return llvm::CallingConv::C;
case CC_X86StdCall: return llvm::CallingConv::X86_StdCall;
// TODO: Add support for __vectorcall to LLVM.
case CC_X86VectorCall: return llvm::CallingConv::X86_VectorCall;
case CC_SpirFunction: return llvm::CallingConv::SPIR_FUNC;
- case CC_SpirKernel: return llvm::CallingConv::SPIR_KERNEL;
+ case CC_OpenCLKernel: return CGM.getTargetCodeGenInfo().getOpenCLKernelCallingConv();
case CC_PreserveMost: return llvm::CallingConv::PreserveMost;
case CC_PreserveAll: return llvm::CallingConv::PreserveAll;
case CC_Swift: return llvm::CallingConv::Swift;
case CC_AAPCS_VFP:
case CC_IntelOclBicc:
case CC_SpirFunction:
- case CC_SpirKernel:
+ case CC_OpenCLKernel:
case CC_Swift:
case CC_PreserveMost:
case CC_PreserveAll:
llvm::SmallSet<const Type *, 8> RecordsWithOpaqueMemberPointers;
+ unsigned ClangCallConvToLLVMCallConv(CallingConv CC);
+
public:
CodeGenTypes(CodeGenModule &cgm);
~CodeGenTypes();
Opt += Lib;
}
+unsigned TargetCodeGenInfo::getOpenCLKernelCallingConv() const {
+ return llvm::CallingConv::C;
+}
static bool isEmptyRecord(ASTContext &Context, QualType T, bool AllowArrays);
/// isEmptyField - Return true iff a the field is "empty", that is it
: TargetCodeGenInfo(new DefaultABIInfo(CGT)) {}
void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const override;
+ unsigned getOpenCLKernelCallingConv() const override;
};
}
}
+unsigned AMDGPUTargetCodeGenInfo::getOpenCLKernelCallingConv() const {
+ return llvm::CallingConv::AMDGPU_KERNEL;
+}
+
//===----------------------------------------------------------------------===//
// SPARC v8 ABI Implementation.
// Based on the SPARC Compliance Definition version 2.4.1.
: TargetCodeGenInfo(new DefaultABIInfo(CGT)) {}
void emitTargetMD(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const override;
+ unsigned getOpenCLKernelCallingConv() const override;
};
} // End anonymous namespace.
OCLVerMD->addOperand(llvm::MDNode::get(Ctx, OCLVerElts));
}
+unsigned SPIRTargetCodeGenInfo::getOpenCLKernelCallingConv() const {
+ return llvm::CallingConv::SPIR_KERNEL;
+}
+
static bool appendType(SmallStringEnc &Enc, QualType QType,
const CodeGen::CodeGenModule &CGM,
TypeStringCache &TSC);
virtual void getDetectMismatchOption(llvm::StringRef Name,
llvm::StringRef Value,
llvm::SmallString<32> &Opt) const {}
+
+ /// Get LLVM calling convention for OpenCL kernel.
+ virtual unsigned getOpenCLKernelCallingConv() const;
};
} // namespace CodeGen
CallingConv CC = S.Context.getDefaultCallingConvention(FTI.isVariadic,
IsCXXInstanceMethod);
- // Attribute AT_OpenCLKernel affects the calling convention only on
- // the SPIR target, hence it cannot be treated as a calling
+ // Attribute AT_OpenCLKernel affects the calling convention for SPIR
+ // and AMDGPU targets, hence it cannot be treated as a calling
// convention attribute. This is the simplest place to infer
- // "spir_kernel" for OpenCL kernels on SPIR.
- if (CC == CC_SpirFunction) {
+ // calling convention for OpenCL kernels.
+ if (S.getLangOpts().OpenCL) {
for (const AttributeList *Attr = D.getDeclSpec().getAttributes().getList();
Attr; Attr = Attr->getNext()) {
if (Attr->getKind() == AttributeList::AT_OpenCLKernel) {
- CC = CC_SpirKernel;
+ llvm::Triple::ArchType arch = S.Context.getTargetInfo().getTriple().getArch();
+ if (arch == llvm::Triple::spir || arch == llvm::Triple::spir64 ||
+ arch == llvm::Triple::amdgcn) {
+ CC = CC_OpenCLKernel;
+ }
break;
}
}
--- /dev/null
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
+// CHECK: define amdgpu_kernel void @test_call_kernel(i32 addrspace(1)* nocapture %out)
+// CHECK: store i32 4, i32 addrspace(1)* %out, align 4
+
+kernel void test_kernel(global int *out)
+{
+ out[0] = 4;
+}
+
+__kernel void test_call_kernel(__global int *out)
+{
+ test_kernel(out);
+}
--- /dev/null
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
+
+// CHECK: define amdgpu_kernel void @calling_conv_amdgpu_kernel()
+kernel void calling_conv_amdgpu_kernel()
+{
+}
+
+// CHECK: define void @calling_conv_none()
+void calling_conv_none()
+{
+}
__attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics
kernel void test_num_vgpr64() {
-// CHECK: define void @test_num_vgpr64() [[ATTR_VGPR64:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_vgpr64() [[ATTR_VGPR64:#[0-9]+]]
}
__attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics
kernel void test_num_sgpr32() {
-// CHECK: define void @test_num_sgpr32() [[ATTR_SGPR32:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_sgpr32() [[ATTR_SGPR32:#[0-9]+]]
}
__attribute__((amdgpu_num_vgpr(64), amdgpu_num_sgpr(32))) // expected-no-diagnostics
kernel void test_num_vgpr64_sgpr32() {
-// CHECK: define void @test_num_vgpr64_sgpr32() [[ATTR_VGPR64_SGPR32:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_vgpr64_sgpr32() [[ATTR_VGPR64_SGPR32:#[0-9]+]]
}
__attribute__((amdgpu_num_sgpr(20), amdgpu_num_vgpr(40))) // expected-no-diagnostics
kernel void test_num_sgpr20_vgpr40() {
-// CHECK: define void @test_num_sgpr20_vgpr40() [[ATTR_SGPR20_VGPR40:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_sgpr20_vgpr40() [[ATTR_SGPR20_VGPR40:#[0-9]+]]
}
__attribute__((amdgpu_num_vgpr(0))) // expected-no-diagnostics
TCALLINGCONV(PreserveMost);
TCALLINGCONV(PreserveAll);
case CC_SpirFunction: return CXCallingConv_Unexposed;
- case CC_SpirKernel: return CXCallingConv_Unexposed;
+ case CC_OpenCLKernel: return CXCallingConv_Unexposed;
break;
}
#undef TCALLINGCONV