This revision was automatically updated to reflect the committed changes.
Closed by commit rL274220: AMDGPU: Set amdgpu_kernel calling convention for
OpenCL kernels. (authored by nhaustov).
Changed prior to commit:
http://reviews.llvm.org/D21367?vs=62198&id=62343#toc
Repository:
rL LLVM
http://reviews.llvm.org/D21367
Files:
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-call-kernel.cl
cfe/trunk/test/CodeGenOpenCL/amdgpu-calling-conv.cl
cfe/trunk/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
cfe/trunk/tools/libclang/CXType.cpp
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/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/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/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/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/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/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/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/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/CodeGen/CodeGenTypes.h
===================================================================
--- cfe/trunk/lib/CodeGen/CodeGenTypes.h
+++ cfe/trunk/lib/CodeGen/CodeGenTypes.h
@@ -164,6 +164,8 @@
llvm::SmallSet<const Type *, 8> RecordsWithOpaqueMemberPointers;
+ unsigned ClangCallConvToLLVMCallConv(CallingConv CC);
+
public:
CodeGenTypes(CodeGenModule &cgm);
~CodeGenTypes();
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/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/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/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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
