Index: cfe/trunk/include/clang/Basic/Specifiers.h =================================================================== --- cfe/trunk/include/clang/Basic/Specifiers.h +++ cfe/trunk/include/clang/Basic/Specifiers.h @@ -241,7 +241,7 @@ 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 @@ case CC_X86Pascal: case CC_X86VectorCall: case CC_SpirFunction: - case CC_SpirKernel: + case CC_OpenCLKernel: case CC_Swift: return false; default: Index: cfe/trunk/lib/AST/ItaniumMangle.cpp =================================================================== --- cfe/trunk/lib/AST/ItaniumMangle.cpp +++ cfe/trunk/lib/AST/ItaniumMangle.cpp @@ -2161,7 +2161,7 @@ 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. Index: cfe/trunk/lib/AST/Type.cpp =================================================================== --- cfe/trunk/lib/AST/Type.cpp +++ cfe/trunk/lib/AST/Type.cpp @@ -2642,7 +2642,7 @@ 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"; Index: cfe/trunk/lib/AST/TypePrinter.cpp =================================================================== --- cfe/trunk/lib/AST/TypePrinter.cpp +++ cfe/trunk/lib/AST/TypePrinter.cpp @@ -725,7 +725,7 @@ 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: Index: cfe/trunk/lib/Basic/Targets.cpp =================================================================== --- cfe/trunk/lib/Basic/Targets.cpp +++ cfe/trunk/lib/Basic/Targets.cpp @@ -2137,6 +2137,16 @@ 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 @@ } 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 { Index: cfe/trunk/lib/CodeGen/CGCall.cpp =================================================================== --- cfe/trunk/lib/CodeGen/CGCall.cpp +++ cfe/trunk/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 @@ /***/ -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 @@ // 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; Index: cfe/trunk/lib/CodeGen/CGDebugInfo.cpp =================================================================== --- cfe/trunk/lib/CodeGen/CGDebugInfo.cpp +++ cfe/trunk/lib/CodeGen/CGDebugInfo.cpp @@ -848,7 +848,7 @@ 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: Index: cfe/trunk/lib/CodeGen/CodeGenTypes.h =================================================================== --- cfe/trunk/lib/CodeGen/CodeGenTypes.h +++ cfe/trunk/lib/CodeGen/CodeGenTypes.h @@ -164,6 +164,8 @@ llvm::SmallSet RecordsWithOpaqueMemberPointers; + unsigned ClangCallConvToLLVMCallConv(CallingConv CC); + public: CodeGenTypes(CodeGenModule &cgm); ~CodeGenTypes(); Index: cfe/trunk/lib/CodeGen/TargetInfo.h =================================================================== --- cfe/trunk/lib/CodeGen/TargetInfo.h +++ cfe/trunk/lib/CodeGen/TargetInfo.h @@ -217,6 +217,9 @@ 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 Index: cfe/trunk/lib/CodeGen/TargetInfo.cpp =================================================================== --- cfe/trunk/lib/CodeGen/TargetInfo.cpp +++ cfe/trunk/lib/CodeGen/TargetInfo.cpp @@ -372,6 +372,9 @@ 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 @@ : 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 @@ } +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 @@ : 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 @@ 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); Index: cfe/trunk/lib/Sema/SemaType.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaType.cpp +++ cfe/trunk/lib/Sema/SemaType.cpp @@ -3184,15 +3184,19 @@ 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; } } Index: cfe/trunk/test/CodeGenOpenCL/amdgpu-call-kernel.cl =================================================================== --- cfe/trunk/test/CodeGenOpenCL/amdgpu-call-kernel.cl +++ cfe/trunk/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); +} Index: cfe/trunk/test/CodeGenOpenCL/amdgpu-calling-conv.cl =================================================================== --- cfe/trunk/test/CodeGenOpenCL/amdgpu-calling-conv.cl +++ cfe/trunk/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() +{ +} Index: cfe/trunk/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl =================================================================== --- cfe/trunk/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl +++ cfe/trunk/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 Index: cfe/trunk/tools/libclang/CXType.cpp =================================================================== --- cfe/trunk/tools/libclang/CXType.cpp +++ cfe/trunk/tools/libclang/CXType.cpp @@ -541,7 +541,7 @@ 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