nhaustov updated this revision to Diff 61247.
nhaustov added a comment.
Add test for calling OpenCL kernel from kernel.
http://reviews.llvm.org/D21367
Files:
include/clang/AST/Type.h
include/clang/Basic/Specifiers.h
lib/AST/ItaniumMangle.cpp
lib/AST/Type.cpp
lib/AST/TypePrinter.cpp
lib/Basic/Targets.cpp
lib/CodeGen/CGCall.cpp
lib/CodeGen/CGDebugInfo.cpp
lib/Sema/SemaType.cpp
test/CodeGenOpenCL/amdgpu-call-kernel.cl
test/CodeGenOpenCL/amdgpu-calling-conv.cl
test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
tools/libclang/CXType.cpp
Index: tools/libclang/CXType.cpp
===================================================================
--- tools/libclang/CXType.cpp
+++ tools/libclang/CXType.cpp
@@ -542,6 +542,7 @@
TCALLINGCONV(PreserveAll);
case CC_SpirFunction: return CXCallingConv_Unexposed;
case CC_SpirKernel: return CXCallingConv_Unexposed;
+ case CC_AMDGPUKernel: return CXCallingConv_Unexposed;
break;
}
#undef TCALLINGCONV
Index: test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
===================================================================
--- test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
+++ 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: test/CodeGenOpenCL/amdgpu-calling-conv.cl
===================================================================
--- /dev/null
+++ 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: test/CodeGenOpenCL/amdgpu-call-kernel.cl
===================================================================
--- /dev/null
+++ 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: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -3182,15 +3182,20 @@
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) {
+ CC = CC_SpirKernel;
+ } else if (arch == llvm::Triple::amdgcn) {
+ CC = CC_AMDGPUKernel;
+ }
break;
}
}
Index: lib/CodeGen/CGDebugInfo.cpp
===================================================================
--- lib/CodeGen/CGDebugInfo.cpp
+++ lib/CodeGen/CGDebugInfo.cpp
@@ -859,6 +859,7 @@
case CC_Swift:
case CC_PreserveMost:
case CC_PreserveAll:
+ case CC_AMDGPUKernel:
return 0;
}
return 0;
Index: lib/CodeGen/CGCall.cpp
===================================================================
--- lib/CodeGen/CGCall.cpp
+++ lib/CodeGen/CGCall.cpp
@@ -61,6 +61,7 @@
case CC_PreserveMost: return llvm::CallingConv::PreserveMost;
case CC_PreserveAll: return llvm::CallingConv::PreserveAll;
case CC_Swift: return llvm::CallingConv::Swift;
+ case CC_AMDGPUKernel: return llvm::CallingConv::AMDGPU_KERNEL;
}
}
Index: lib/Basic/Targets.cpp
===================================================================
--- lib/Basic/Targets.cpp
+++ lib/Basic/Targets.cpp
@@ -2118,6 +2118,17 @@
Opts.cl_khr_3d_image_writes = 1;
Opts.cl_khr_gl_msaa_sharing = 1;
}
+
+
+ CallingConvCheckResult checkCallingConvention(CallingConv CC) const override {
+ switch (CC) {
+ default:
+ return CCCR_Warning;
+ case CC_C:
+ case CC_AMDGPUKernel:
+ return CCCR_OK;
+ }
+ }
};
const Builtin::Info AMDGPUTargetInfo::BuiltinInfo[] = {
Index: lib/AST/TypePrinter.cpp
===================================================================
--- lib/AST/TypePrinter.cpp
+++ lib/AST/TypePrinter.cpp
@@ -726,6 +726,7 @@
break;
case CC_SpirFunction:
case CC_SpirKernel:
+ case CC_AMDGPUKernel:
// Do nothing. These CCs are not available as attributes.
break;
case CC_Swift:
Index: lib/AST/Type.cpp
===================================================================
--- lib/AST/Type.cpp
+++ lib/AST/Type.cpp
@@ -2646,6 +2646,7 @@
case CC_Swift: return "swiftcall";
case CC_PreserveMost: return "preserve_most";
case CC_PreserveAll: return "preserve_all";
+ case CC_AMDGPUKernel: return "amdgpu_kernel";
}
llvm_unreachable("Invalid calling convention.");
Index: lib/AST/ItaniumMangle.cpp
===================================================================
--- lib/AST/ItaniumMangle.cpp
+++ lib/AST/ItaniumMangle.cpp
@@ -2144,6 +2144,7 @@
case CC_SpirKernel:
case CC_PreserveMost:
case CC_PreserveAll:
+ case CC_AMDGPUKernel:
// FIXME: we should be mangling all of the above.
return "";
Index: include/clang/Basic/Specifiers.h
===================================================================
--- include/clang/Basic/Specifiers.h
+++ include/clang/Basic/Specifiers.h
@@ -245,6 +245,7 @@
CC_Swift, // __attribute__((swiftcall))
CC_PreserveMost, // __attribute__((preserve_most))
CC_PreserveAll, // __attribute__((preserve_all))
+ CC_AMDGPUKernel, // OpenCL kernel on AMDGPU target
};
/// \brief Checks whether the given calling convention supports variadic
@@ -259,6 +260,7 @@
case CC_SpirFunction:
case CC_SpirKernel:
case CC_Swift:
+ case CC_AMDGPUKernel:
return false;
default:
return true;
Index: include/clang/AST/Type.h
===================================================================
--- include/clang/AST/Type.h
+++ include/clang/AST/Type.h
@@ -1377,7 +1377,7 @@
/// Extra information which affects how the function is called, like
/// regparm and the calling convention.
- unsigned ExtInfo : 9;
+ unsigned ExtInfo : 10;
/// Used only by FunctionProtoType, put here to pack with the
/// other bitfields.
@@ -2902,14 +2902,14 @@
// Type::FunctionTypeBitfields.
// | CC |noreturn|produces|regparm|
- // |0 .. 3| 4 | 5 | 6 .. 8|
+ // |0 .. 4| 5 | 6 | 7 .. 9|
//
// regparm is either 0 (no regparm attribute) or the regparm value+1.
- enum { CallConvMask = 0xF };
- enum { NoReturnMask = 0x10 };
- enum { ProducesResultMask = 0x20 };
+ enum { CallConvMask = 0x1F };
+ enum { NoReturnMask = 0x20 };
+ enum { ProducesResultMask = 0x40 };
enum { RegParmMask = ~(CallConvMask | NoReturnMask | ProducesResultMask),
- RegParmOffset = 6 }; // Assumed to be the last field
+ RegParmOffset = 7 }; // Assumed to be the last field
uint16_t Bits;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
