[OpenCL] Prevent mangling kernel functions

[Anastasia]

Summary

Kernel function names have to be preserved as in the original source to be able to access them from the host API side.

This patch adds restriction to kernels that prevents them to be used in overloading, templates, etc. This is implemented by enclosing the kernels implicitly into C linkage clause. Therefore all C linkage restrictions will apply to the kernels.

[rjmccall]

I think it would be more reasonable to just change getDeclLanguageLinkage to check for a kernel function.

[Anastasia]

Extended FunctionDecl exten C logic to take kernel function into account.

[rjmccall]

I think it would be more reasonable to just change getDeclLanguageLinkage to check for a kernel function.

lib/AST/Decl.cpp

2961

Both of these changes should be unnecessary because they ultimately defer to isInExternCContext.

I assume that OpenCL bans making a class member function a kernel?

[Anastasia]

Removed redundant C linkage checks!

[Anastasia]

In D60454#1488348, @rjmccall wrote:

I think it would be more reasonable to just change getDeclLanguageLinkage to check for a kernel function.

I tried to change getDeclLanguageLinkage only but Sema calls isInExternCContext separately while giving diagnostics. getDeclLanguageLinkage also calls it through isFirstInExternCContext. So it seems we just need to check in isInExternCContext and that is a common part for all various program paths. Do you agree?

lib/AST/Decl.cpp

2961

I assume that OpenCL bans making a class member function a kernel?

Yep, that's right! I am adding a diagnostic in this patch.

[rjmccall]

LGTM.

include/clang/Basic/DiagnosticSemaKinds.td

 def err_asm_naked_parm_ref : Error<
   "parameter references not allowed in naked functions">;
 
 // OpenCL warnings and errors.
 def err_invalid_astype_of_different_size : Error<
   "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<
   "pointer arguments to kernel functions must reside in '__global', "
   "'__constant' or '__local' address space">;
 def err_opencl_ext_vector_component_invalid_length : Error<
   "vector component access has invalid length %0.  Supported: 1,2,3,4,8,16.">;
 def err_opencl_function_variable : Error<

lib/AST/Decl.cpp

 }
 
 LanguageLinkage FunctionDecl::getLanguageLinkage() const {
   return getDeclLanguageLinkage(*this);
 }
 
 bool FunctionDecl::isExternC() const {
   return isDeclExternC(*this);
 }
 
 bool FunctionDecl::isInExternCContext() const {
+  if (hasAttr<OpenCLKernelAttr>())
+    return true;
   return getLexicalDeclContext()->isExternCContext();
 }
 
 bool FunctionDecl::isInExternCXXContext() const {
   return getLexicalDeclContext()->isExternCXXContext();
 }
 
 bool FunctionDecl::isGlobal() const {

lib/Sema/SemaDecl.cpp

if (!getLangOpts().CUDAAllowVariadicFunctions && NewFD->isVariadic() &&
         !(II && II->isStr("printf") && NewFD->isExternC() &&
           !D.isFunctionDefinition())) {
       Diag(NewFD->getLocation(), diag::err_variadic_device_fn);
     }
   }
 
   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)) {
       Diag(D.getIdentifierLoc(), diag::err_static_kernel);
       D.setInvalidType();
     }
 
     // OpenCL v1.2, s6.9 -- Kernels can only have return type void.
     if (!NewFD->getReturnType()->isVoidType()) {
       SourceRange RTRange = NewFD->getReturnTypeSourceRange();
       Diag(D.getIdentifierLoc(), diag::err_expected_kernel_void_return_type)
           << (RTRange.isValid() ? FixItHint::CreateReplacement(RTRange, "void")
                                 : FixItHint());
       D.setInvalidType();
     }
 
     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();
 
     // OpenCL 2.0 pipe restrictions forbids pipe packet types to be non-value
     // types.
     if (getLangOpts().OpenCLVersion >= 200) {
       if(const PipeType *PipeTy = PT->getAs<PipeType>()) {
         QualType ElemTy = PipeTy->getElementType();

test/CodeGenOpenCLCXX/addrspace-of-this.cl

 // EXPL: @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %this
 // EXPL: @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %this
 // EXPL: @_ZNU3AS4R1CaSEOU3AS4S_(%class.C addrspace(4)* %this
 // COMMON: @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* %this)
 
 // 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)*))
 // Test the address space of 'this' when invoking a method using a pointer to the object.
 // COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
 
 // Test the address space of 'this' when invoking a method that is declared in the file contex.
 // COMMON: call i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
__kernel void test##AS() { \
     int i = c.get();         \
     C c1(c);                 \
     C c2;                    \
     c2 = c1;                 \
   }
 
 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)*
 // EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
 
 // Test the address space of 'this' when invoking assignment operator.
 // COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
 // COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // EXPL: call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
 // IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)*
 // IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)*
 // IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]]
 
 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)*
 // EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]])
 
 // Test the address space of 'this' when invoking a method.
 // COMMON: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
 // COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]])

test/CodeGenOpenCLCXX/local_addrspace_init.cl

 // 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;
 };
 
 kernel void test() {
   __local int i;
   __local C ii;
   // FIXME: In OpenCL C we don't accept initializers for local
   // address space variables. User defined initialization could
   // make sense, but would it mean that all work items need to
   // execute it? Potentially disallowing any initialization would
   // make things easier and assingments can be used to set specific
   // values. This rules should make it consistent with OpenCL C.
   //__local C c();
 }

test/SemaOpenCLCXX/kernel_invalid.cl

+// 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'}}