mikael created this revision.
mikael added a reviewer: Anastasia.
Herald added subscribers: cfe-commits, yaxunl.
Address spaces are cast into generic before invoking the constructor.
Repository:
rC Clang
https://reviews.llvm.org/D54862
Files:
lib/AST/DeclCXX.cpp
lib/CodeGen/CGCall.cpp
lib/CodeGen/CGDeclCXX.cpp
test/CodeGenOpenCLCXX/addrspace-of-this.cl
Index: test/CodeGenOpenCLCXX/addrspace-of-this.cl
===================================================================
--- /dev/null
+++ test/CodeGenOpenCLCXX/addrspace-of-this.cl
@@ -0,0 +1,102 @@
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s
+// expected-no-diagnostics
+
+// Test that the 'this' pointer is in the __generic address space.
+
+// FIXME: Add support for __constant address space.
+
+class C {
+public:
+ int v;
+ C() { v = 2; }
+ C(const C &c) { v = c.v; }
+ C &operator=(const C &c) {
+ v = c.v;
+ return *this;
+ }
+ int get() { return v; }
+};
+
+__global C c;
+
+__kernel void test__global() {
+ int i = c.get();
+ C c2 = c;
+ C c3(c);
+}
+
+// CHECK-LABEL: @__cxx_global_var_init()
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) #4
+
+// Test that the address space is __generic for the constructor
+// CHECK-LABEL: @_ZN1CC1Ev(%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 @_ZN1CC2Ev(%class.C addrspace(4)* %this1) #4
+// CHECK: ret void
+
+// CHECK-LABEL: @_Z12test__globalv()
+
+// Test the address space of 'this' when accessing a method.
+// CHECK: %call = call i32 @_ZN1C3getEv(%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: %0 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %0, %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 assignment operator.
+// CHECK: %1 = addrspacecast %class.C* %c3 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %1, %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+
+The following checkers are testing similar things like the __global test but for other address spaces.
+
+#define TEST(AS) \
+ __kernel void test##AS() { \
+ AS C c; \
+ int i = c.get(); \
+ C c2 = c; \
+ C c3(c); \
+ }
+
+TEST(__local)
+
+// CHECK-LABEL: _Z11test__localv
+// CHECK: @__cxa_guard_acquire
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+// CHECK: %2 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %2, %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+// CHECK: %3 = addrspacecast %class.C* %c3 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %3, %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+
+TEST(__private)
+
+// CHECK-LABEL: @_Z13test__privatev
+// CHECK-NOT: @__cxa_guard_acquire
+// CHECK: %0 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %0)
+// CHECK: %1 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %1)
+// CHECK: %2 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %3, %class.C addrspace(4)* dereferenceable(4) %2)
+// CHECK: %4 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %5 = addrspacecast %class.C* %c3 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %5, %class.C addrspace(4)* dereferenceable(4) %4)
+
+TEST()
+
+// CHECK-LABEL: @_Z4testv()
+// CHECK: %0 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %0)
+// CHECK: %1 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %1)
+// CHECK: %2 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %3, %class.C addrspace(4)* dereferenceable(4) %2)
+// CHECK: %4 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %5 = addrspacecast %class.C* %c3 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %5, %class.C addrspace(4)* dereferenceable(4) %4)
+
Index: lib/CodeGen/CGDeclCXX.cpp
===================================================================
--- lib/CodeGen/CGDeclCXX.cpp
+++ lib/CodeGen/CGDeclCXX.cpp
@@ -26,7 +26,10 @@
static void EmitDeclInit(CodeGenFunction &CGF, const VarDecl &D,
ConstantAddress DeclPtr) {
- assert(D.hasGlobalStorage() && "VarDecl must have global storage!");
+ assert(
+ (D.hasGlobalStorage() ||
+ (D.hasLocalStorage() && CGF.getContext().getLangOpts().OpenCLCPlusPlus)) &&
+ "VarDecl must have global or local (in the case of OpenCL) storage!");
assert(!D.getType()->isReferenceType() &&
"Should not call EmitDeclInit on a reference!");
Index: lib/CodeGen/CGCall.cpp
===================================================================
--- lib/CodeGen/CGCall.cpp
+++ lib/CodeGen/CGCall.cpp
@@ -72,6 +72,13 @@
/// FIXME: address space qualification?
static CanQualType GetThisType(ASTContext &Context, const CXXRecordDecl *RD) {
QualType RecTy = Context.getTagDeclType(RD)->getCanonicalTypeInternal();
+
+ if (Context.getLangOpts().OpenCLCPlusPlus)
+ // For OpenCL add generic address space for 'this' pointer. This allows to
+ // instantiate classes in various address spaces (except for constant) to be
+ // used with the same version of generated operators.
+ RecTy = Context.getAddrSpaceQualType(RecTy, LangAS::opencl_generic);
+
return Context.getPointerType(CanQualType::CreateUnsafe(RecTy));
}
@@ -4019,11 +4026,13 @@
V->getType()->isIntegerTy())
V = Builder.CreateZExt(V, ArgInfo.getCoerceToType());
- // If the argument doesn't match, perform a bitcast to coerce it. This
- // can happen due to trivial type mismatches.
+ // If the argument doesn't match, perform a bitcast or addrspacecast to
+ // coerce it. This can happen due to trivial type mismatches.
if (FirstIRArg < IRFuncTy->getNumParams() &&
- V->getType() != IRFuncTy->getParamType(FirstIRArg))
- V = Builder.CreateBitCast(V, IRFuncTy->getParamType(FirstIRArg));
+ V->getType() != IRFuncTy->getParamType(FirstIRArg)) {
+ auto DestTy = IRFuncTy->getParamType(FirstIRArg);
+ V = Builder.CreatePointerBitCastOrAddrSpaceCast(V, DestTy);
+ }
IRCallArgs[FirstIRArg] = V;
break;
Index: lib/AST/DeclCXX.cpp
===================================================================
--- lib/AST/DeclCXX.cpp
+++ lib/AST/DeclCXX.cpp
@@ -2185,6 +2185,13 @@
QualType ClassTy = C.getTypeDeclType(getParent());
ClassTy = C.getQualifiedType(ClassTy,
Qualifiers::fromCVRUMask(getTypeQualifiers()));
+
+ // For OpenCL add generic address space for 'this' pointer. This allows to
+ // instantiate classes in various address spaces (except for constant) to be
+ // used with the same version of generated operators.
+ if (C.getLangOpts().OpenCLCPlusPlus)
+ ClassTy = C.getAddrSpaceQualType(ClassTy, LangAS::opencl_generic);
+
return C.getPointerType(ClassTy);
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
