summaryrefslogtreecommitdiffstats
path: root/clang
diff options
context:
space:
mode:
Diffstat (limited to 'clang')
-rw-r--r--clang/include/clang/Basic/Specifiers.h4
-rw-r--r--clang/lib/AST/ItaniumMangle.cpp2
-rw-r--r--clang/lib/AST/Type.cpp2
-rw-r--r--clang/lib/AST/TypePrinter.cpp2
-rw-r--r--clang/lib/Basic/Targets.cpp14
-rw-r--r--clang/lib/CodeGen/CGCall.cpp5
-rw-r--r--clang/lib/CodeGen/CGDebugInfo.cpp2
-rw-r--r--clang/lib/CodeGen/CodeGenTypes.h2
-rw-r--r--clang/lib/CodeGen/TargetInfo.cpp13
-rw-r--r--clang/lib/CodeGen/TargetInfo.h3
-rw-r--r--clang/lib/Sema/SemaType.cpp14
-rwxr-xr-xclang/test/CodeGenOpenCL/amdgpu-call-kernel.cl14
-rw-r--r--clang/test/CodeGenOpenCL/amdgpu-calling-conv.cl12
-rw-r--r--clang/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl8
-rw-r--r--clang/tools/libclang/CXType.cpp2
15 files changed, 79 insertions, 20 deletions
diff --git a/clang/include/clang/Basic/Specifiers.h b/clang/include/clang/Basic/Specifiers.h
index b2965782de9..c099d6e6810 100644
--- a/clang/include/clang/Basic/Specifiers.h
+++ b/clang/include/clang/Basic/Specifiers.h
@@ -241,7 +241,7 @@ namespace clang {
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))
@@ -257,7 +257,7 @@ namespace clang {
case CC_X86Pascal:
case CC_X86VectorCall:
case CC_SpirFunction:
- case CC_SpirKernel:
+ case CC_OpenCLKernel:
case CC_Swift:
return false;
default:
diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp
index 8d49c6f586b..902a0ed048b 100644
--- a/clang/lib/AST/ItaniumMangle.cpp
+++ b/clang/lib/AST/ItaniumMangle.cpp
@@ -2161,7 +2161,7 @@ StringRef CXXNameMangler::getCallingConvQualifierName(CallingConv CC) {
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.
diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp
index 6dc585bf296..a0a751e7602 100644
--- a/clang/lib/AST/Type.cpp
+++ b/clang/lib/AST/Type.cpp
@@ -2642,7 +2642,7 @@ StringRef FunctionType::getNameForCallConv(CallingConv CC) {
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";
diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp
index 29a4845d4a9..030afd9cfc5 100644
--- a/clang/lib/AST/TypePrinter.cpp
+++ b/clang/lib/AST/TypePrinter.cpp
@@ -725,7 +725,7 @@ void TypePrinter::printFunctionProtoAfter(const FunctionProtoType *T,
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:
diff --git a/clang/lib/Basic/Targets.cpp b/clang/lib/Basic/Targets.cpp
index 87203f5598c..252f2be7aff 100644
--- a/clang/lib/Basic/Targets.cpp
+++ b/clang/lib/Basic/Targets.cpp
@@ -2137,6 +2137,16 @@ public:
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[] = {
@@ -7927,8 +7937,8 @@ public:
}
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 {
diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index d683b899022..02a85ef0eed 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -30,6 +30,7 @@
#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"
@@ -41,7 +42,7 @@ using namespace CodeGen;
/***/
-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;
@@ -57,7 +58,7 @@ static unsigned ClangCallConvToLLVMCallConv(CallingConv CC) {
// 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;
diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp
index 638b2d46fcf..723f79aee74 100644
--- a/clang/lib/CodeGen/CGDebugInfo.cpp
+++ b/clang/lib/CodeGen/CGDebugInfo.cpp
@@ -848,7 +848,7 @@ static unsigned getDwarfCC(CallingConv CC) {
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:
diff --git a/clang/lib/CodeGen/CodeGenTypes.h b/clang/lib/CodeGen/CodeGenTypes.h
index affa3344103..5796ab8fe5a 100644
--- a/clang/lib/CodeGen/CodeGenTypes.h
+++ b/clang/lib/CodeGen/CodeGenTypes.h
@@ -164,6 +164,8 @@ class CodeGenTypes {
llvm::SmallSet<const Type *, 8> RecordsWithOpaqueMemberPointers;
+ unsigned ClangCallConvToLLVMCallConv(CallingConv CC);
+
public:
CodeGenTypes(CodeGenModule &cgm);
~CodeGenTypes();
diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp
index bbe98df77c9..28651141b20 100644
--- a/clang/lib/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CodeGen/TargetInfo.cpp
@@ -372,6 +372,9 @@ TargetCodeGenInfo::getDependentLibraryOption(llvm::StringRef Lib,
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
@@ -6828,6 +6831,7 @@ public:
: TargetCodeGenInfo(new DefaultABIInfo(CGT)) {}
void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const override;
+ unsigned getOpenCLKernelCallingConv() const override;
};
}
@@ -6856,6 +6860,10 @@ void AMDGPUTargetCodeGenInfo::setTargetAttributes(
}
+unsigned AMDGPUTargetCodeGenInfo::getOpenCLKernelCallingConv() const {
+ return llvm::CallingConv::AMDGPU_KERNEL;
+}
+
//===----------------------------------------------------------------------===//
// SPARC v8 ABI Implementation.
// Based on the SPARC Compliance Definition version 2.4.1.
@@ -7505,6 +7513,7 @@ public:
: TargetCodeGenInfo(new DefaultABIInfo(CGT)) {}
void emitTargetMD(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const override;
+ unsigned getOpenCLKernelCallingConv() const override;
};
} // End anonymous namespace.
@@ -7534,6 +7543,10 @@ void SPIRTargetCodeGenInfo::emitTargetMD(const Decl *D, llvm::GlobalValue *GV,
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);
diff --git a/clang/lib/CodeGen/TargetInfo.h b/clang/lib/CodeGen/TargetInfo.h
index 71f6b0a4c5c..e46382596af 100644
--- a/clang/lib/CodeGen/TargetInfo.h
+++ b/clang/lib/CodeGen/TargetInfo.h
@@ -217,6 +217,9 @@ public:
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
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index d58de8a3822..c0d03a15f31 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -3184,15 +3184,19 @@ getCCForDeclaratorChunk(Sema &S, Declarator &D,
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;
}
}
diff --git a/clang/test/CodeGenOpenCL/amdgpu-call-kernel.cl b/clang/test/CodeGenOpenCL/amdgpu-call-kernel.cl
new file mode 100755
index 00000000000..005793916c6
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/amdgpu-call-kernel.cl
@@ -0,0 +1,14 @@
+// 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);
+}
diff --git a/clang/test/CodeGenOpenCL/amdgpu-calling-conv.cl b/clang/test/CodeGenOpenCL/amdgpu-calling-conv.cl
new file mode 100644
index 00000000000..7da9d7f4d49
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/amdgpu-calling-conv.cl
@@ -0,0 +1,12 @@
+// 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()
+{
+}
diff --git a/clang/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl b/clang/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
index d2ecc7a8c6b..589d00d1eaa 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
@@ -5,23 +5,23 @@
__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
diff --git a/clang/tools/libclang/CXType.cpp b/clang/tools/libclang/CXType.cpp
index fb0e5b5acdf..4fcd8864cd3 100644
--- a/clang/tools/libclang/CXType.cpp
+++ b/clang/tools/libclang/CXType.cpp
@@ -541,7 +541,7 @@ CXCallingConv clang_getFunctionTypeCallingConv(CXType X) {
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
OpenPOWER on IntegriCloud