From d8d49375956fecbedbeec46a2dc2f982368c3478 Mon Sep 17 00:00:00 2001 From: Nikolay Haustov Date: Thu, 30 Jun 2016 09:06:33 +0000 Subject: [PATCH] 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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@274220 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/Basic/Specifiers.h | 4 ++-- lib/AST/ItaniumMangle.cpp | 2 +- lib/AST/Type.cpp | 2 +- lib/AST/TypePrinter.cpp | 2 +- lib/Basic/Targets.cpp | 14 ++++++++++++-- lib/CodeGen/CGCall.cpp | 5 +++-- lib/CodeGen/CGDebugInfo.cpp | 2 +- lib/CodeGen/CodeGenTypes.h | 2 ++ lib/CodeGen/TargetInfo.cpp | 13 +++++++++++++ lib/CodeGen/TargetInfo.h | 3 +++ lib/Sema/SemaType.cpp | 14 +++++++++----- test/CodeGenOpenCL/amdgpu-call-kernel.cl | 14 ++++++++++++++ test/CodeGenOpenCL/amdgpu-calling-conv.cl | 12 ++++++++++++ test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl | 8 ++++---- tools/libclang/CXType.cpp | 2 +- 15 files changed, 79 insertions(+), 20 deletions(-) create mode 100755 test/CodeGenOpenCL/amdgpu-call-kernel.cl create mode 100644 test/CodeGenOpenCL/amdgpu-calling-conv.cl diff --git a/include/clang/Basic/Specifiers.h b/include/clang/Basic/Specifiers.h index b2965782de..c099d6e681 100644 --- a/include/clang/Basic/Specifiers.h +++ b/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/lib/AST/ItaniumMangle.cpp b/lib/AST/ItaniumMangle.cpp index 8d49c6f586..902a0ed048 100644 --- a/lib/AST/ItaniumMangle.cpp +++ b/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/lib/AST/Type.cpp b/lib/AST/Type.cpp index 6dc585bf29..a0a751e760 100644 --- a/lib/AST/Type.cpp +++ b/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/lib/AST/TypePrinter.cpp b/lib/AST/TypePrinter.cpp index 29a4845d4a..030afd9cfc 100644 --- a/lib/AST/TypePrinter.cpp +++ b/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/lib/Basic/Targets.cpp b/lib/Basic/Targets.cpp index 87203f5598..252f2be7af 100644 --- a/lib/Basic/Targets.cpp +++ b/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/lib/CodeGen/CGCall.cpp b/lib/CodeGen/CGCall.cpp index d683b89902..02a85ef0ee 100644 --- a/lib/CodeGen/CGCall.cpp +++ b/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/lib/CodeGen/CGDebugInfo.cpp b/lib/CodeGen/CGDebugInfo.cpp index 638b2d46fc..723f79aee7 100644 --- a/lib/CodeGen/CGDebugInfo.cpp +++ b/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/lib/CodeGen/CodeGenTypes.h b/lib/CodeGen/CodeGenTypes.h index affa334410..5796ab8fe5 100644 --- a/lib/CodeGen/CodeGenTypes.h +++ b/lib/CodeGen/CodeGenTypes.h @@ -164,6 +164,8 @@ class CodeGenTypes { llvm::SmallSet RecordsWithOpaqueMemberPointers; + unsigned ClangCallConvToLLVMCallConv(CallingConv CC); + public: CodeGenTypes(CodeGenModule &cgm); ~CodeGenTypes(); diff --git a/lib/CodeGen/TargetInfo.cpp b/lib/CodeGen/TargetInfo.cpp index bbe98df77c..28651141b2 100644 --- a/lib/CodeGen/TargetInfo.cpp +++ b/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/lib/CodeGen/TargetInfo.h b/lib/CodeGen/TargetInfo.h index 71f6b0a4c5..e46382596a 100644 --- a/lib/CodeGen/TargetInfo.h +++ b/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/lib/Sema/SemaType.cpp b/lib/Sema/SemaType.cpp index d58de8a382..c0d03a15f3 100644 --- a/lib/Sema/SemaType.cpp +++ b/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/test/CodeGenOpenCL/amdgpu-call-kernel.cl b/test/CodeGenOpenCL/amdgpu-call-kernel.cl new file mode 100755 index 0000000000..005793916c --- /dev/null +++ b/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/test/CodeGenOpenCL/amdgpu-calling-conv.cl b/test/CodeGenOpenCL/amdgpu-calling-conv.cl new file mode 100644 index 0000000000..7da9d7f4d4 --- /dev/null +++ b/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/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl b/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl index d2ecc7a8c6..589d00d1ea 100644 --- a/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl +++ b/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/tools/libclang/CXType.cpp b/tools/libclang/CXType.cpp index fb0e5b5acd..4fcd8864cd 100644 --- a/tools/libclang/CXType.cpp +++ b/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 -- 2.40.0