[PR41674] [OpenCL] Fix initialisation of this via pointer

[kpet]

Summary

When the expression used to initialise this has a pointer type,
check the address space of the pointee type instead of the pointer
type to decide whether an address space cast is required.
It is the pointee type that carries the address space qualifier.

[Anastasia]

I would suggest an IR test, especially since what you want to achieve is addrspacecast in IR output. We use AST dump only if there is no other good way to test something. In fact we already have a test with similar logic that tests various OpenCL address spaces: test/CodeGenOpenCLCXX/addrspace-of-this.cl. I think it should be reasonable enough to extend it.

[kpet]

Ok, I'll replace the test with an IR-level test.

[kpet]

Replaced the AST dump test with an IR test.

[Anastasia]

test/CodeGenOpenCLCXX/addrspace-of-this.cl

94

Ok I think we should regex name of call registers or may be even prevent checking it if it's not required. The same applies to other places btw.

136

Shouldn't %call5 just be [[CALL]]? If not I see little value in checking the name...

[kpet]

Most of these are pre-existing issues with the tests but I agree they're worth fixing. I'll update the tests.

[kpet]

Clean up the tests.

[Anastasia]

LGTM! Thanks!

lib/Sema/SemaOverload.cpp

if (ICS.Standard.Second == ICK_Derived_To_Base) {
     ExprResult FromRes =
       PerformObjectMemberConversion(From, Qualifier, FoundDecl, Method);
     if (FromRes.isInvalid())
       return ExprError();
     From = FromRes.get();
   }
 
   if (!Context.hasSameType(From->getType(), DestType)) {
-    if (From->getType().getAddressSpace() != DestType.getAddressSpace())
-      From = ImpCastExprToType(From, DestType, CK_AddressSpaceConversion,
-                             From->getValueKind()).get();
+    CastKind CK;
+    if (FromRecordType.getAddressSpace() != DestType.getAddressSpace())
+      CK = CK_AddressSpaceConversion;
     else
-      From = ImpCastExprToType(From, DestType, CK_NoOp,
-                             From->getValueKind()).get();
+      CK = CK_NoOp;
+    From = ImpCastExprToType(From, DestType, CK, From->getValueKind()).get();
   }
   return From;
 }
 
 /// TryContextuallyConvertToBool - Attempt to contextually convert the
 /// expression From to bool (C++0x [conv]p3).
 static ImplicitConversionSequence
 TryContextuallyConvertToBool(Sema &S, Expr *From) {

test/CodeGenOpenCLCXX/addrspace-of-this.cl

 }
 
 extern C&& foo();
 
 __global C c;
 
 __kernel void test__global() {
   int i = c.get();
-  int i2 = c.outside();
+  int i2 = (&c)->get();
+  int i3 = c.outside();
   C c1(c);
   C c2;
   c2 = c1;
   C c3 = c1 + c2;
   C c4(foo());
   C c5 = foo();
 }
 

 // 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()
 
 // Test the address space of 'this' when invoking a method.
-// COMMON: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+// 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: %call1 = call i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+// COMMON: call i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c 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)*
 // 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(1)* bitcast (%class.C addrspace(1)* @c to i8 addrspace(1)*) to i8 addrspace(4)*)
 // EXPL: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
 
 // Test the address space of 'this' when invoking a constructor.

 // EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C4GEN]], %class.C addrspace(4)* dereferenceable(4) [[CALL]])
 // IMPL: [[C4VOID:%[0-9]+]] = bitcast %class.C* %c4 to i8*
 // IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)*
 // IMPL:  call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C4VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]]
 
 // Test the address space of 'this' when invoking the move assignment
 // COMMON: [[C5GEN:%[0-9]+]] = addrspacecast %class.C* %c5 to %class.C addrspace(4)*
 // COMMON: [[CALL:%call[0-9]+]] = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov()
-// EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C5GEN:%[0-9]+]], %class.C addrspace(4)* dereferenceable(4) %call4)
+// EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C5GEN:%[0-9]+]], %class.C addrspace(4)* dereferenceable(4) [[CALL]])
 // IMPL: [[C5VOID:%[0-9]+]] = bitcast %class.C* %c5 to i8*
 // IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)*
 // IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C5VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]]
 
 // Tests address space of inline members
 //COMMON: @_ZNU3AS41C3getEv(%class.C addrspace(4)* %this)
 //COMMON: @_ZNU3AS41CplERU3AS4KS_(%class.C* noalias sret %agg.result, %class.C addrspace(4)* %this
 #define TEST(AS)             \
   __kernel void test##AS() { \
     AS C c;                  \
     int i = c.get();         \
     C c1(c);                 \
     C c2;                    \
     c2 = c1;                 \
   }
 
 TEST(__local)
 
 // COMMON-LABEL: _Z11test__localv
 
 // 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)*))
+// COMMON: 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.

 // CHECK-LABEL: @_Z13test__privatev
 
 // 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 = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]])
+// COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]])
 
 // Test the address space of 'this' when invoking a copy-constructor.
 // COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
 // COMMON: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
 // EXPL: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]])
 // IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8*
 // IMPL: [[CGENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CGEN]] to i8 addrspace(4)*
 // IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}[[CGENVOID]]
 
 // 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 a copy-assignment.
 // 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: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
+// 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]]