r274220 - AMDGPU: Set amdgpu_kernel calling convention for OpenCL kernels.
Nikolay Haustov via cfe-commits
cfe-commits at lists.llvm.org
Thu Jun 30 02:06:34 PDT 2016
Author: nhaustov
Date: Thu Jun 30 04:06:33 2016
New Revision: 274220
URL: http://llvm.org/viewvc/llvm-project?rev=274220&view=rev
Log:
AMDGPU: Set amdgpu_kernel calling convention for OpenCL kernels.
Summary:
Summary:
Change Clang calling convention SpirKernel to OpenCLKernel.
Set calling convention OpenCLKernel for amdgcn as well.
Add virtual method .getOpenCLKernelCallingConv() to TargetCodeGenInfo
and use it to set target calling convention for AMDGPU and SPIR.
Update tests.
Reviewers: rsmith, tstellarAMD, Anastasia, yaxunl
Subscribers: kzhuravl, cfe-commits
Differential Revision: http://reviews.llvm.org/D21367
Added:
cfe/trunk/test/CodeGenOpenCL/amdgpu-call-kernel.cl (with props)
cfe/trunk/test/CodeGenOpenCL/amdgpu-calling-conv.cl
Modified:
cfe/trunk/include/clang/Basic/Specifiers.h
cfe/trunk/lib/AST/ItaniumMangle.cpp
cfe/trunk/lib/AST/Type.cpp
cfe/trunk/lib/AST/TypePrinter.cpp
cfe/trunk/lib/Basic/Targets.cpp
cfe/trunk/lib/CodeGen/CGCall.cpp
cfe/trunk/lib/CodeGen/CGDebugInfo.cpp
cfe/trunk/lib/CodeGen/CodeGenTypes.h
cfe/trunk/lib/CodeGen/TargetInfo.cpp
cfe/trunk/lib/CodeGen/TargetInfo.h
cfe/trunk/lib/Sema/SemaType.cpp
cfe/trunk/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
cfe/trunk/tools/libclang/CXType.cpp
Modified: cfe/trunk/include/clang/Basic/Specifiers.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Specifiers.h?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Specifiers.h (original)
+++ cfe/trunk/include/clang/Basic/Specifiers.h Thu Jun 30 04:06:33 2016
@@ -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:
Modified: cfe/trunk/lib/AST/ItaniumMangle.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ItaniumMangle.cpp?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ItaniumMangle.cpp (original)
+++ cfe/trunk/lib/AST/ItaniumMangle.cpp Thu Jun 30 04:06:33 2016
@@ -2161,7 +2161,7 @@ StringRef CXXNameMangler::getCallingConv
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.
Modified: cfe/trunk/lib/AST/Type.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/Type.cpp?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/lib/AST/Type.cpp (original)
+++ cfe/trunk/lib/AST/Type.cpp Thu Jun 30 04:06:33 2016
@@ -2642,7 +2642,7 @@ StringRef FunctionType::getNameForCallCo
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";
Modified: cfe/trunk/lib/AST/TypePrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/TypePrinter.cpp?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/lib/AST/TypePrinter.cpp (original)
+++ cfe/trunk/lib/AST/TypePrinter.cpp Thu Jun 30 04:06:33 2016
@@ -725,7 +725,7 @@ void TypePrinter::printFunctionProtoAfte
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:
Modified: cfe/trunk/lib/Basic/Targets.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets.cpp?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets.cpp (original)
+++ cfe/trunk/lib/Basic/Targets.cpp Thu Jun 30 04:06:33 2016
@@ -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 {
Modified: cfe/trunk/lib/CodeGen/CGCall.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGCall.cpp?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGCall.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGCall.cpp Thu Jun 30 04:06:33 2016
@@ -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 ClangCallConvToLLVMCallC
// 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;
Modified: cfe/trunk/lib/CodeGen/CGDebugInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDebugInfo.cpp?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGDebugInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGDebugInfo.cpp Thu Jun 30 04:06:33 2016
@@ -848,7 +848,7 @@ static unsigned getDwarfCC(CallingConv C
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:
Modified: cfe/trunk/lib/CodeGen/CodeGenTypes.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenTypes.h?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenTypes.h (original)
+++ cfe/trunk/lib/CodeGen/CodeGenTypes.h Thu Jun 30 04:06:33 2016
@@ -164,6 +164,8 @@ class CodeGenTypes {
llvm::SmallSet<const Type *, 8> RecordsWithOpaqueMemberPointers;
+ unsigned ClangCallConvToLLVMCallConv(CallingConv CC);
+
public:
CodeGenTypes(CodeGenModule &cgm);
~CodeGenTypes();
Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.cpp?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.cpp Thu Jun 30 04:06:33 2016
@@ -372,6 +372,9 @@ TargetCodeGenInfo::getDependentLibraryOp
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::setTargetA
}
+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
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);
Modified: cfe/trunk/lib/CodeGen/TargetInfo.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.h?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.h (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.h Thu Jun 30 04:06:33 2016
@@ -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
Modified: cfe/trunk/lib/Sema/SemaType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaType.cpp (original)
+++ cfe/trunk/lib/Sema/SemaType.cpp Thu Jun 30 04:06:33 2016
@@ -3184,15 +3184,19 @@ getCCForDeclaratorChunk(Sema &S, Declara
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;
}
}
Added: cfe/trunk/test/CodeGenOpenCL/amdgpu-call-kernel.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/amdgpu-call-kernel.cl?rev=274220&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/amdgpu-call-kernel.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/amdgpu-call-kernel.cl Thu Jun 30 04:06:33 2016
@@ -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);
+}
Propchange: cfe/trunk/test/CodeGenOpenCL/amdgpu-call-kernel.cl
------------------------------------------------------------------------------
svn:executable = *
Added: cfe/trunk/test/CodeGenOpenCL/amdgpu-calling-conv.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/amdgpu-calling-conv.cl?rev=274220&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/amdgpu-calling-conv.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/amdgpu-calling-conv.cl Thu Jun 30 04:06:33 2016
@@ -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()
+{
+}
Modified: cfe/trunk/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl Thu Jun 30 04:06:33 2016
@@ -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
Modified: cfe/trunk/tools/libclang/CXType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/tools/libclang/CXType.cpp?rev=274220&r1=274219&r2=274220&view=diff
==============================================================================
--- cfe/trunk/tools/libclang/CXType.cpp (original)
+++ cfe/trunk/tools/libclang/CXType.cpp Thu Jun 30 04:06:33 2016
@@ -541,7 +541,7 @@ CXCallingConv clang_getFunctionTypeCalli
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
More information about the cfe-commits
mailing list