[PR40778][PR41157][OpenCL] Prevent implicit initialization of local address space objects

[Anastasia]

Summary

In OpenCL C variables in local address space can't be initialized (s6.5.2). However in C++ default initialization can be performed even then there is no initializer explicitly provided. This won't work for local address space objects because they are not the same as regular global or local variables.

Current solution is to disable implicit default initialization. We still need some work to be done on user defined initialization and also construction of objects in local address space.

[bader]

lib/Sema/SemaDecl.cpp

11648–11652

Shouldn't we invalidate Var declaration?

test/CodeGenOpenCLCXX/addrspace-of-this.cl

154–155

onbject -> object

test/CodeGenOpenCLCXX/local_addrspace_init.cl

13

I guess this declaration should be disallowed for non-POD types. Can we add a check for that to some test/Sema* test?

[Anastasia]

Fixed typo in the comment.

lib/Sema/SemaDecl.cpp

11648–11652

This early exit is to prevent adding default initializer implicitly to the declarations that are valid.

test/CodeGenOpenCLCXX/local_addrspace_init.cl

13

What should be disallowed specifically? Declaring non-POD types in the local address space?

[bader]

test/CodeGenOpenCLCXX/local_addrspace_init.cl

13

Something like

class C {
  int i = 0;
};

kernel void test() {
  __local C c; // error
}

[Anastasia]

test/CodeGenOpenCLCXX/local_addrspace_init.cl

13

Yes, potentially we should disallow those or specify that the initialization is to be ignored. I plan to look into all of those. But for this change I am just fixing the issue reported in the bug and adding a FIXME for the future work.

Thanks!

[rjmccall]

lib/Sema/SemaDecl.cpp

11648–11652

"In OpenCL, we can't initialize objects in the __local address space, even implicitly, so don't synthesize an implicit initializer."

I think it's important to add the underscores on __local to make it obvious that we're talking about OpenCL's "local" address space, not the address space that local variables appear in.

[Anastasia]

Improved comment about initializers in __local addr space.

[rjmccall]

LGTM.

lib/Sema/SemaDecl.cpp

if (getLangOpts().CPlusPlus && Var->hasLocalStorage()) {
         CXXRecordDecl *CXXRecord = cast<CXXRecordDecl>(Record->getDecl());
         // Mark the function (if we're in one) for further checking even if the
         // looser rules of C++11 do not require such checks, so that we can
         // diagnose incompatibilities with C++98.
         if (!CXXRecord->isPOD())
           setFunctionHasBranchProtectedScope();
       }
     }
-
+    // In OpenCL we don't allow to initialize objects in local address space.
+    if (getLangOpts().OpenCL &&
+        Var->getType().getAddressSpace() == LangAS::opencl_local)
+      return;
     // C++03 [dcl.init]p9:
     //   If no initializer is specified for an object, and the
     //   object is of (possibly cv-qualified) non-POD class type (or
     //   array thereof), the object shall be default-initialized; if
     //   the object is of const-qualified type, the underlying class
     //   type shall have a user-declared default
     //   constructor. Otherwise, if no initializer is specified for
     //   a non- static object, the object and its subobjects, if
     //   any, have an indeterminate initial value); if the object

test/CodeGenOpenCLCXX/addrspace-of-this.cl

__kernel void test##AS() { \
     C c1(c);                 \
     C c2;                    \
     c2 = c1;                 \
   }
 
 TEST(__local)
 
 // COMMON-LABEL: _Z11test__localv
-// EXPL: @__cxa_guard_acquire
 
-// Test the address space of 'this' when invoking a constructor for an object in non-default address space
-// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+// 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)*))
 
 // Test the address space of 'this' when invoking a method.
 // COMMON: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c 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)*))
 // 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)
 
 // Test the address space of 'this' when invoking a constructor.
 // EXPL: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*

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
+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();
+}