From b198751157296c7f49c52d7af5b3dc8b8bf9f078 Mon Sep 17 00:00:00 2001 From: Anastasia Stulova Date: Wed, 30 Jan 2019 11:18:08 +0000 Subject: [PATCH] [OpenCL] Add generic addr space to the return of implicit assignment. When creating the prototype of implicit assignment operators the returned reference to the class should be qualified with the same addr space as 'this' (i.e. __generic in OpenCL). Differential Revision: https://reviews.llvm.org/D57101 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@352617 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Sema/SemaDeclCXX.cpp | 7 +- test/CodeGenOpenCLCXX/addrspace-of-this.cl | 216 ++++++++++-------- .../address_space_overloading.cl | 7 +- 3 files changed, 126 insertions(+), 104 deletions(-) diff --git a/lib/Sema/SemaDeclCXX.cpp b/lib/Sema/SemaDeclCXX.cpp index 031e0620b3..fd6eaa8cb6 100644 --- a/lib/Sema/SemaDeclCXX.cpp +++ b/lib/Sema/SemaDeclCXX.cpp @@ -11813,14 +11813,13 @@ CXXMethodDecl *Sema::DeclareImplicitCopyAssignment(CXXRecordDecl *ClassDecl) { return nullptr; QualType ArgType = Context.getTypeDeclType(ClassDecl); + if (Context.getLangOpts().OpenCLCPlusPlus) + ArgType = Context.getAddrSpaceQualType(ArgType, LangAS::opencl_generic); QualType RetType = Context.getLValueReferenceType(ArgType); bool Const = ClassDecl->implicitCopyAssignmentHasConstParam(); if (Const) ArgType = ArgType.withConst(); - if (Context.getLangOpts().OpenCLCPlusPlus) - ArgType = Context.getAddrSpaceQualType(ArgType, LangAS::opencl_generic); - ArgType = Context.getLValueReferenceType(ArgType); bool Constexpr = defaultedSpecialMemberIsConstexpr(*this, ClassDecl, @@ -12138,6 +12137,8 @@ CXXMethodDecl *Sema::DeclareImplicitMoveAssignment(CXXRecordDecl *ClassDecl) { // constructor rules. QualType ArgType = Context.getTypeDeclType(ClassDecl); + if (Context.getLangOpts().OpenCLCPlusPlus) + ArgType = Context.getAddrSpaceQualType(ArgType, LangAS::opencl_generic); QualType RetType = Context.getLValueReferenceType(ArgType); ArgType = Context.getRValueReferenceType(ArgType); diff --git a/test/CodeGenOpenCLCXX/addrspace-of-this.cl b/test/CodeGenOpenCLCXX/addrspace-of-this.cl index 83af3fbcb3..0e76523f38 100644 --- a/test/CodeGenOpenCLCXX/addrspace-of-this.cl +++ b/test/CodeGenOpenCLCXX/addrspace-of-this.cl @@ -1,25 +1,26 @@ -// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL" +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - -DDECL -DUSE_DEFLT | FileCheck %s --check-prefixes="COMMON,IMPL" +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s --check-prefixes="COMMON,IMPL" // expected-no-diagnostics // Test that the 'this' pointer is in the __generic address space. -// FIXME: Add support for __constant address space. +#ifdef USE_DEFLT +#define DEFAULT =default +#else +#define DEFAULT +#endif class C { public: int v; - C() { v = 2; } - C(C &&c) { v = c.v; } - C(const C &c) { v = c.v; } - C &operator=(const C &c) { - v = c.v; - return *this; - } - C &operator=(C &&c) & { - v = c.v; - return *this; - } - +#ifdef DECL + C() DEFAULT; + C(C &&c) DEFAULT; + C(const C &c) DEFAULT; + C &operator=(const C &c) DEFAULT; + C &operator=(C &&c) & DEFAULT; +#endif C operator+(const C& c) { v += c.v; return *this; @@ -30,6 +31,24 @@ public: int outside(); }; +#if defined(DECL) && !defined(USE_DEFLT) +C::C() { v = 2; }; + +C::C(C &&c) { v = c.v; } + +C::C(const C &c) { v = c.v; } + +C &C::operator=(const C &c) { + v = c.v; + return *this; +} + +C &C::operator=(C &&c) & { + v = c.v; + return *this; +} +#endif + int C::outside() { return v; } @@ -49,58 +68,76 @@ __kernel void test__global() { C c5 = foo(); } -// CHECK-LABEL: @__cxx_global_var_init() -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) +// Test that the address space is __generic for all members +// EXPL: @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this) +// EXPL: @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %this) +// EXPL: @_ZNU3AS41CC2EOU3AS4S_(%class.C addrspace(4)* %this +// EXPL: @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* %this +// EXPL: @_ZNU3AS41CC2ERU3AS4KS_(%class.C addrspace(4)* %this +// 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) -// Test that the address space is __generic for the constructor -// CHECK-LABEL: @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %this) -// CHECK: entry: -// CHECK: %this.addr = alloca %class.C addrspace(4)*, align 4 -// CHECK: store %class.C addrspace(4)* %this, %class.C addrspace(4)** %this.addr, align 4 -// CHECK: %this1 = load %class.C addrspace(4)*, %class.C addrspace(4)** %this.addr, align 4 -// CHECK: call void @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this1) -// CHECK: ret void +// 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)*)) -// CHECK-LABEL: @_Z12test__globalv() +// COMMON-LABEL: @_Z12test__globalv() // Test the address space of 'this' when invoking a method. -// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) +// COMMON: %call = 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. -// CHECK: %call1 = call i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) +// COMMON: %call1 = 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. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: 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)*)) +// 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. -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) +// 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. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: %call2 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) +// 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 the address space of 'this' when invoking the operator+ -// CHECK: [[C3GEN:%[0-9]+]] = addrspacecast %class.C* %c3 to %class.C addrspace(4)* -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret %ref.tmp, %class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[C2GEN]]) -// CHECK: [[REFGEN:%[0-9]+]] = addrspacecast %class.C* %ref.tmp to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C3GEN]], %class.C addrspace(4)* dereferenceable(4) [[REFGEN]]) +// COMMON: [[C3GEN:%[0-9]+]] = addrspacecast %class.C* %c3 to %class.C addrspace(4)* +// 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)* +// COMMON: call void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret %ref.tmp, %class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[C2GEN]]) +// COMMON: [[REFGEN:%[0-9]+]] = addrspacecast %class.C* %ref.tmp to %class.C addrspace(4)* +// EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C3GEN]], %class.C addrspace(4)* dereferenceable(4) [[REFGEN]]) +// IMPL: [[C3VOID:%[0-9]+]] = bitcast %class.C* %c3 to i8* +// IMPL: [[REFGENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[REFGEN]] to i8 addrspace(4)* +// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C3VOID]], i8 addrspace(4)* {{.*}}[[REFGENVOID]] // Test the address space of 'this' when invoking the move constructor -// CHECK: [[C4GEN:%[0-9]+]] = addrspacecast %class.C* %c4 to %class.C addrspace(4)* -// CHECK: %call3 = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov() -// CHECK: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C4GEN]], %class.C addrspace(4)* dereferenceable(4) %call3) +// COMMON: [[C4GEN:%[0-9]+]] = addrspacecast %class.C* %c4 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)* [[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 -// CHECK: [[C5GEN:%[0-9]+]] = addrspacecast %class.C* %c5 to %class.C addrspace(4)* -// CHECK: %call4 = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov() #5 -// CHECK: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C5GEN:%[0-9]+]], %class.C addrspace(4)* dereferenceable(4) %call4) - - +// 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) +// 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; \ @@ -112,77 +149,62 @@ __kernel void test__global() { TEST(__local) -// CHECK-LABEL: _Z11test__localv -// CHECK: @__cxa_guard_acquire +// 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 -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) +// EXPL: 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. -// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) +// 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. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: 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)*)) +// 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: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) +// 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. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) +// 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 // Test the address space of 'this' when invoking a constructor for an object in non-default address space -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]]) - -// Test the address space of 'this' when invoking a method. -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]]) - -// Test the address space of 'this' when invoking a copy-constructor. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]]) - -// Test the address space of 'this' when invoking a constructor. -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) - -// Test the address space of 'this' when invoking a copy-assignment. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) - -TEST() - -// CHECK-LABEL: @_Z4testv() - -// Test the address space of 'this' when invoking a constructor for an object in non-default address space -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]]) +// 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. -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]]) +// COMMON: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* +// COMMON: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]]) // Test the address space of 'this' when invoking a copy-constructor. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]]) +// 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. -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) +// 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. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) +// 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]]) +// 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]] diff --git a/test/SemaOpenCLCXX/address_space_overloading.cl b/test/SemaOpenCLCXX/address_space_overloading.cl index 6c1934083a..6458ade562 100644 --- a/test/SemaOpenCLCXX/address_space_overloading.cl +++ b/test/SemaOpenCLCXX/address_space_overloading.cl @@ -1,12 +1,11 @@ // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=c++ - -// FIXME: This test shouldn't trigger any errors. +// expected-no-diagnostics struct RetGlob { int dummy; }; -struct RetGen { //expected-error{{binding value of type '__generic RetGen' to reference to type 'RetGen' drops <> qualifiers}} +struct RetGen { char dummy; }; @@ -19,5 +18,5 @@ void kernel k() { __local int *ArgLoc; RetGlob TestGlob = foo(ArgGlob); RetGen TestGen = foo(ArgGen); - TestGen = foo(ArgLoc); //expected-note{{in implicit copy assignment operator for 'RetGen' first required here}} + TestGen = foo(ArgLoc); } -- 2.40.0