r360152 - [OpenCL] Prevent mangling kernel functions.
Anastasia Stulova via cfe-commits
cfe-commits at lists.llvm.org
Tue May 7 07:22:34 PDT 2019
Author: stulova
Date: Tue May 7 07:22:34 2019
New Revision: 360152
URL: http://llvm.org/viewvc/llvm-project?rev=360152&view=rev
Log:
[OpenCL] Prevent mangling kernel functions.
Kernel function names have to be preserved as in the original
source to be able to access them from the host API side.
This commit also adds restriction to kernels that prevents them
from being used in overloading, templates, etc.
Differential Revision: https://reviews.llvm.org/D60454
Added:
cfe/trunk/test/SemaOpenCLCXX/kernel_invalid.cl
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/AST/Decl.cpp
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl
cfe/trunk/test/CodeGenOpenCLCXX/local_addrspace_init.cl
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=360152&r1=360151&r2=360152&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Tue May 7 07:22:34 2019
@@ -8582,6 +8582,10 @@ def err_invalid_astype_of_different_size
"invalid reinterpretation: sizes of %0 and %1 must match">;
def err_static_kernel : Error<
"kernel functions cannot be declared static">;
+def err_method_kernel : Error<
+ "kernel functions cannot be class members">;
+def err_template_kernel : Error<
+ "kernel functions cannot be used in a template declaration, instantiation or specialization">;
def err_opencl_ptrptr_kernel_param : Error<
"kernel parameter cannot be declared as a pointer to a pointer">;
def err_kernel_arg_address_space : Error<
Modified: cfe/trunk/lib/AST/Decl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/Decl.cpp?rev=360152&r1=360151&r2=360152&view=diff
==============================================================================
--- cfe/trunk/lib/AST/Decl.cpp (original)
+++ cfe/trunk/lib/AST/Decl.cpp Tue May 7 07:22:34 2019
@@ -2961,6 +2961,8 @@ bool FunctionDecl::isExternC() const {
}
bool FunctionDecl::isInExternCContext() const {
+ if (hasAttr<OpenCLKernelAttr>())
+ return true;
return getLexicalDeclContext()->isExternCContext();
}
Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=360152&r1=360151&r2=360152&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue May 7 07:22:34 2019
@@ -9214,18 +9214,9 @@ Sema::ActOnFunctionDeclarator(Scope *S,
MarkUnusedFileScopedDecl(NewFD);
- if (getLangOpts().CPlusPlus) {
- if (FunctionTemplate) {
- if (NewFD->isInvalidDecl())
- FunctionTemplate->setInvalidDecl();
- return FunctionTemplate;
- }
- if (isMemberSpecialization && !NewFD->isInvalidDecl())
- CompleteMemberSpecialization(NewFD, Previous);
- }
- if (NewFD->hasAttr<OpenCLKernelAttr>()) {
+ if (getLangOpts().OpenCL && NewFD->hasAttr<OpenCLKernelAttr>()) {
// OpenCL v1.2 s6.8 static is invalid for kernel functions.
if ((getLangOpts().OpenCLVersion >= 120)
&& (SC == SC_Static)) {
@@ -9245,7 +9236,30 @@ Sema::ActOnFunctionDeclarator(Scope *S,
llvm::SmallPtrSet<const Type *, 16> ValidTypes;
for (auto Param : NewFD->parameters())
checkIsValidOpenCLKernelParameter(*this, D, Param, ValidTypes);
+
+ if (getLangOpts().OpenCLCPlusPlus) {
+ if (DC->isRecord()) {
+ Diag(D.getIdentifierLoc(), diag::err_method_kernel);
+ D.setInvalidType();
+ }
+ if (FunctionTemplate) {
+ Diag(D.getIdentifierLoc(), diag::err_template_kernel);
+ D.setInvalidType();
+ }
+ }
}
+
+ if (getLangOpts().CPlusPlus) {
+ if (FunctionTemplate) {
+ if (NewFD->isInvalidDecl())
+ FunctionTemplate->setInvalidDecl();
+ return FunctionTemplate;
+ }
+
+ if (isMemberSpecialization && !NewFD->isInvalidDecl())
+ CompleteMemberSpecialization(NewFD, Previous);
+ }
+
for (const ParmVarDecl *Param : NewFD->parameters()) {
QualType PT = Param->getType();
Modified: cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl?rev=360152&r1=360151&r2=360152&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl (original)
+++ cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl Tue May 7 07:22:34 2019
@@ -83,7 +83,7 @@ __kernel void test__global() {
// EXPL-LABEL: @__cxx_global_var_init()
// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
-// COMMON-LABEL: @_Z12test__globalv()
+// COMMON-LABEL: @test__global()
// Test the address space of 'this' when invoking a method.
// COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
@@ -152,19 +152,19 @@ __kernel void test__global() {
TEST(__local)
-// COMMON-LABEL: _Z11test__localv
+// COMMON-LABEL: @test__local
// Test that we don't initialize an object in local address space.
-// EXPL-NOT: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+// EXPL-NOT: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
// Test the address space of 'this' when invoking a method.
-// COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+// COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
// Test the address space of 'this' when invoking copy-constructor.
// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// EXPL: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+// EXPL: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
// IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8*
-// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}addrspacecast (i8 addrspace(3)* bitcast (%class.C addrspace(3)* @_ZZ11test__localvE1c to i8 addrspace(3)*) to i8 addrspace(4)*), i32 4, i1 false)
+// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}addrspacecast (i8 addrspace(3)* bitcast (%class.C addrspace(3)* @_ZZ11test__localE1c to i8 addrspace(3)*) to i8 addrspace(4)*), i32 4, i1 false)
// Test the address space of 'this' when invoking a constructor.
// EXPL: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
@@ -180,7 +180,7 @@ TEST(__local)
TEST(__private)
-// CHECK-LABEL: @_Z13test__privatev
+// CHECK-LABEL: @test__private
// Test the address space of 'this' when invoking a constructor for an object in non-default address space
// EXPL: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
Modified: cfe/trunk/test/CodeGenOpenCLCXX/local_addrspace_init.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCLCXX/local_addrspace_init.cl?rev=360152&r1=360151&r2=360152&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCLCXX/local_addrspace_init.cl (original)
+++ cfe/trunk/test/CodeGenOpenCLCXX/local_addrspace_init.cl Tue May 7 07:22:34 2019
@@ -1,8 +1,8 @@
// RUN: %clang_cc1 %s -triple spir -cl-std=c++ -emit-llvm -O0 -o - | FileCheck %s
// Test that we don't initialize local address space objects.
-//CHECK: @_ZZ4testvE1i = internal addrspace(3) global i32 undef
-//CHECK: @_ZZ4testvE2ii = internal addrspace(3) global %class.C undef
+//CHECK: @_ZZ4testE1i = internal addrspace(3) global i32 undef
+//CHECK: @_ZZ4testE2ii = internal addrspace(3) global %class.C undef
class C {
int i;
};
Added: cfe/trunk/test/SemaOpenCLCXX/kernel_invalid.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCLCXX/kernel_invalid.cl?rev=360152&view=auto
==============================================================================
--- cfe/trunk/test/SemaOpenCLCXX/kernel_invalid.cl (added)
+++ cfe/trunk/test/SemaOpenCLCXX/kernel_invalid.cl Tue May 7 07:22:34 2019
@@ -0,0 +1,17 @@
+// RUN: %clang_cc1 %s -cl-std=c++ -pedantic -verify -fsyntax-only
+
+struct C {
+ kernel void m(); //expected-error{{kernel functions cannot be class members}}
+};
+
+template <typename T>
+kernel void templ(T par) { //expected-error{{kernel functions cannot be used in a template declaration, instantiation or specialization}}
+}
+
+template <int>
+kernel void bar(int par) { //expected-error{{kernel functions cannot be used in a template declaration, instantiation or specialization}}
+}
+
+kernel void foo(int); //expected-note{{previous declaration is here}}
+
+kernel void foo(float); //expected-error{{conflicting types for 'foo'}}
More information about the cfe-commits
mailing list