On Thu, Jun 30, 2016 at 09:06:34AM -0000, Nikolay Haustov via cfe-commits wrote:
> 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) {
Was this change left over from a previous version of the patch? This
patch doesn't seem to require that this be a member function.
-Tom
> 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
>
>
> _______________________________________________
> cfe-commits mailing list
> cfe-commits@lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
