mikael updated this revision to Diff 175646.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D54862/new/
https://reviews.llvm.org/D54862
Files:
lib/AST/DeclCXX.cpp
lib/CodeGen/CGCall.cpp
lib/CodeGen/CGClass.cpp
lib/CodeGen/CGDeclCXX.cpp
lib/CodeGen/CGExprCXX.cpp
lib/Sema/SemaType.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,140 @@
+// 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 c1(c);
+ C c2;
+ c2 = c1;
+}
+
+// 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 invoking 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* %c1 to %class.C* addrspace(4)*
+// CHECK: %1 = bitcast %class.C* addrspace(4)* %0 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)*))
+
+// Test the address space of 'this' when invoking a constructor.
+// CHECK: %2 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)*
+// CHECK: %3 = bitcast %class.C* addrspace(4)* %2 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %3) #4
+
+// Test the address space of 'this' when invoking assignment operator.
+// CHECK: %4 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: %5 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)*
+// CHECK: %6 = bitcast %class.C* addrspace(4)* %5 to %class.C addrspace(4)*
+// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %4)
+
+#define TEST(AS) \
+ __kernel void test##AS() { \
+ AS C c; \
+ int i = c.get(); \
+ C c1(c); \
+ C c2; \
+ c2 = c1; \
+ }
+
+TEST(__local)
+
+// CHECK-LABEL: _Z11test__localv
+// CHECK: @__cxa_guard_acquire
+
+// Test the address space of 'this' when invoking a method.
+// CHECK: call void @_ZN1CC1Ev(%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: %call = call i32 @_ZN1C3getEv(%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 constructor.
+// CHECK: %4 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)*
+// CHECK: %5 = bitcast %class.C* addrspace(4)* %4 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %5)
+
+// Test the address space of 'this' when invoking assignment operator.
+// CHECK: %6 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: %7 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)*
+// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %8, %class.C addrspace(4)* dereferenceable(4) %6)
+
+TEST(__private)
+
+// CHECK-LABEL: @_Z13test__privatev
+
+// Test the address space of 'this' when invoking a method.
+// CHECK: %2 = addrspacecast %class.C* %c to %class.C* addrspace(4)*
+// CHECK: %3 = bitcast %class.C* addrspace(4)* %2 to %class.C addrspace(4)*
+// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %3)
+
+// Test the address space of 'this' when invoking a copy-constructor.
+// CHECK: %5 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %6 = bitcast %class.C* addrspace(4)* %4 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5)
+
+// Test the address space of 'this' when invoking a constructor.
+// CHECK: %7 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)*
+// CHECK: %8 = bitcast %class.C* addrspace(4)* %7 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %8)
+
+// Test the address space of 'this' when invoking a copy-assignment.
+// CHECK: %9 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: %10 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)*
+// CHECK: %11 = bitcast %class.C* addrspace(4)* %10 to %class.C addrspace(4)*
+// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %11, %class.C addrspace(4)* dereferenceable(4) %9)
+
+TEST()
+
+// CHECK-LABEL: @_Z4testv()
+// Test the address space of 'this' when invoking a method.
+// CHECK: %2 = addrspacecast %class.C* %c to %class.C* addrspace(4)*
+// CHECK: %3 = bitcast %class.C* addrspace(4)* %2 to %class.C addrspace(4)*
+// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %3) #4
+
+// Test the address space of 'this' when invoking a copy-constructor.
+// CHECK: %4 = addrspacecast %class.C* %c1 to %class.C* addrspace(4)*
+// CHECK: %5 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %6 = bitcast %class.C* addrspace(4)* %4 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5)
+
+// Test the address space of 'this' when invoking a constructor.
+// CHECK: %7 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)*
+// CHECK: %8 = bitcast %class.C* addrspace(4)* %7 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %8)
+
+// Test the address space of 'this' when invoking a copy-assignment.
+// CHECK: %9 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: %10 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)*
+// CHECK: %11 = bitcast %class.C* addrspace(4)* %10 to %class.C addrspace(4)*
+// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %11, %class.C addrspace(4)* dereferenceable(4) %9)
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -7189,12 +7189,16 @@
bool IsFuncType =
ChunkIndex < D.getNumTypeObjects() &&
D.getTypeObject(ChunkIndex).Kind == DeclaratorChunk::Function;
+ bool IsClassMemberFn = IsFuncType && D.getContext() == DeclaratorContext::MemberContext;
+
if ( // Do not deduce addr space for function return type and function type,
- // otherwise it will fail some sema check.
- IsFuncReturnType || IsFuncType ||
+ // otherwise it will fail some sema check. We want to deduce class member functions.
+ IsFuncReturnType ||
+ (IsFuncType && !IsClassMemberFn) ||
// Do not deduce addr space for member types of struct, except the pointee
- // type of a pointer member type.
- (D.getContext() == DeclaratorContext::MemberContext && !IsPointee) ||
+ // type of a pointer member type. We want to deduce class member functions.
+ (D.getContext() == DeclaratorContext::MemberContext && !IsPointee &&
+ !IsClassMemberFn) ||
// Do not deduce addr space for types used to define a typedef and the
// typedef itself, except the pointee type of a pointer type which is used
// to define the typedef.
@@ -7224,7 +7228,8 @@
// (...)
// Pointers that are declared without pointing to a named address space
// point to the generic address space.
- if (IsPointee) {
+ // Deduce class members functions to be of the generic address space
+ if (IsPointee || IsClassMemberFn) {
ImpAddr = LangAS::opencl_generic;
} else {
if (D.getContext() == DeclaratorContext::FileContext) {
Index: lib/CodeGen/CGExprCXX.cpp
===================================================================
--- lib/CodeGen/CGExprCXX.cpp
+++ lib/CodeGen/CGExprCXX.cpp
@@ -11,12 +11,13 @@
//
//===----------------------------------------------------------------------===//
-#include "CodeGenFunction.h"
#include "CGCUDARuntime.h"
#include "CGCXXABI.h"
#include "CGDebugInfo.h"
#include "CGObjCRuntime.h"
+#include "CodeGenFunction.h"
#include "ConstantEmitter.h"
+#include "TargetInfo.h"
#include "clang/CodeGen/CGFunctionInfo.h"
#include "clang/Frontend/CodeGenOptions.h"
#include "llvm/IR/CallSite.h"
@@ -44,11 +45,26 @@
"Trying to emit a member or operator call expr on a static method!");
ASTContext &C = CGF.getContext();
- // Push the this ptr.
const CXXRecordDecl *RD =
CGF.CGM.getCXXABI().getThisArgumentTypeForMethod(MD);
- Args.add(RValue::get(This),
- RD ? C.getPointerType(C.getTypeDeclType(RD)) : C.VoidPtrTy);
+
+ if (RD) {
+ LangAS AS = MD->getType().getAddressSpace();
+ if (AS != LangAS::Default) {
+ unsigned TargetAS = C.getTargetAddressSpace(AS);
+ llvm::Type *NewType = This->getType()->getPointerTo(TargetAS);
+ This = CGF.getTargetHooks().performAddrSpaceCast(
+ CGF, This,
+ getLangASFromTargetAS(This->getType()->getPointerAddressSpace()), AS,
+ NewType);
+ }
+ // Push the this ptr.
+ Args.add(RValue::get(This), C.getPointerType(C.getAddrSpaceQualType(
+ C.getTypeDeclType(RD), AS)));
+ } else {
+ // Push the this ptr.
+ Args.add(RValue::get(This), C.VoidPtrTy);
+ }
// If there is an implicit parameter (e.g. VTT), emit it.
if (ImplicitParam) {
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/CGClass.cpp
===================================================================
--- lib/CodeGen/CGClass.cpp
+++ lib/CodeGen/CGClass.cpp
@@ -16,6 +16,7 @@
#include "CGDebugInfo.h"
#include "CGRecordLayout.h"
#include "CodeGenFunction.h"
+#include "TargetInfo.h"
#include "clang/AST/CXXInheritance.h"
#include "clang/AST/DeclTemplate.h"
#include "clang/AST/EvaluatedExprVisitor.h"
@@ -2012,8 +2013,18 @@
bool NewPointerIsChecked) {
CallArgList Args;
+ LangAS AS = D->getType().getAddressSpace();
+ llvm::Value *ThisPtr = This.getPointer();
+ if (AS != LangAS::Default) {
+ unsigned TargetAS = getContext().getTargetAddressSpace(AS);
+ llvm::Type *NewType = ThisPtr->getType()->getPointerTo(TargetAS);
+ ThisPtr = getTargetHooks().performAddrSpaceCast(
+ *this, This.getPointer(),
+ getLangASFromTargetAS(This.getAddressSpace()), AS,
+ NewType);
+ }
// Push the this ptr.
- Args.add(RValue::get(This.getPointer()), D->getThisType(getContext()));
+ Args.add(RValue::get(ThisPtr), D->getThisType(getContext()));
// If this is a trivial constructor, emit a memcpy now before we lose
// the alignment information on the argument.
Index: lib/CodeGen/CGCall.cpp
===================================================================
--- lib/CodeGen/CGCall.cpp
+++ lib/CodeGen/CGCall.cpp
@@ -69,9 +69,10 @@
/// Derives the 'this' type for codegen purposes, i.e. ignoring method
/// qualification.
-/// FIXME: address space qualification?
-static CanQualType GetThisType(ASTContext &Context, const CXXRecordDecl *RD) {
+static CanQualType GetThisType(ASTContext &Context, const CXXRecordDecl *RD, const CXXMethodDecl *MD) {
QualType RecTy = Context.getTagDeclType(RD)->getCanonicalTypeInternal();
+ if (MD)
+ RecTy = Context.getAddrSpaceQualType(RecTy, MD->getType().getAddressSpace());
return Context.getPointerType(CanQualType::CreateUnsafe(RecTy));
}
@@ -246,7 +247,7 @@
// Add the 'this' pointer.
if (RD)
- argTypes.push_back(GetThisType(Context, RD));
+ argTypes.push_back(GetThisType(Context, RD, MD));
else
argTypes.push_back(Context.VoidPtrTy);
@@ -302,7 +303,7 @@
SmallVector<CanQualType, 16> argTypes;
SmallVector<FunctionProtoType::ExtParameterInfo, 16> paramInfos;
- argTypes.push_back(GetThisType(Context, MD->getParent()));
+ argTypes.push_back(GetThisType(Context, MD->getParent(), MD));
bool PassParams = true;
@@ -529,7 +530,7 @@
CodeGenTypes::arrangeUnprototypedMustTailThunk(const CXXMethodDecl *MD) {
assert(MD->isVirtual() && "only methods have thunks");
CanQual<FunctionProtoType> FTP = GetFormalType(MD);
- CanQualType ArgTys[] = { GetThisType(Context, MD->getParent()) };
+ CanQualType ArgTys[] = { GetThisType(Context, MD->getParent(), MD) };
return arrangeLLVMFunctionInfo(Context.VoidTy, /*instanceMethod=*/false,
/*chainCall=*/false, ArgTys,
FTP->getExtInfo(), {}, RequiredArgs(1));
@@ -543,7 +544,7 @@
CanQual<FunctionProtoType> FTP = GetFormalType(CD);
SmallVector<CanQualType, 2> ArgTys;
const CXXRecordDecl *RD = CD->getParent();
- ArgTys.push_back(GetThisType(Context, RD));
+ ArgTys.push_back(GetThisType(Context, RD, CD));
if (CT == Ctor_CopyingClosure)
ArgTys.push_back(*FTP->param_type_begin());
if (RD->getNumVBases() > 0)
Index: lib/AST/DeclCXX.cpp
===================================================================
--- lib/AST/DeclCXX.cpp
+++ lib/AST/DeclCXX.cpp
@@ -2185,6 +2185,8 @@
QualType ClassTy = C.getTypeDeclType(getParent());
ClassTy = C.getQualifiedType(ClassTy,
Qualifiers::fromCVRUMask(getTypeQualifiers()));
+
+ ClassTy = C.getAddrSpaceQualType(ClassTy, getType().getAddressSpace());
return C.getPointerType(ClassTy);
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
