Skip to content

Commit

Permalink
[OpenCL] Allow address spaces as method qualifiers.
Browse files Browse the repository at this point in the history
Methods can now be qualified with address spaces to prevent
undesirable conversions to generic or to provide custom 
implementation to be used if the object is located in certain
memory segments.

This commit extends parsing and standard C++ overloading to
work for an address space of a method (i.e. implicit 'this'
parameter).

Differential Revision: https://reviews.llvm.org/D55850

llvm-svn: 351747
  • Loading branch information
Anastasia Stulova committed Jan 21, 2019
1 parent cfa434a commit 5cffa45
Show file tree
Hide file tree
Showing 9 changed files with 187 additions and 60 deletions.
6 changes: 5 additions & 1 deletion clang/include/clang/AST/Type.h
Expand Up @@ -1982,7 +1982,7 @@ class Type : public ExtQualsTypeCommonBase {
bool isObjCQualifiedClassType() const; // Class<foo>
bool isObjCObjectOrInterfaceType() const;
bool isObjCIdType() const; // id

bool isDecltypeType() const;
/// Was this type written with the special inert-in-ARC __unsafe_unretained
/// qualifier?
///
Expand Down Expand Up @@ -6440,6 +6440,10 @@ inline bool Type::isObjCBuiltinType() const {
return isObjCIdType() || isObjCClassType() || isObjCSelType();
}

inline bool Type::isDecltypeType() const {
return isa<DecltypeType>(this);
}

#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
inline bool Type::is##Id##Type() const { \
return isSpecificBuiltinType(BuiltinType::Id); \
Expand Down
19 changes: 19 additions & 0 deletions clang/include/clang/Sema/ParsedAttr.h
Expand Up @@ -567,6 +567,25 @@ class ParsedAttr final
/// parsed attribute does not have a semantic equivalent, or would not have
/// a Spelling enumeration, the value UINT_MAX is returned.
unsigned getSemanticSpelling() const;

/// If this is an OpenCL addr space attribute returns its representation
/// in LangAS, otherwise returns default addr space.
LangAS asOpenCLLangAS() const {
switch (getKind()) {
case ParsedAttr::AT_OpenCLConstantAddressSpace:
return LangAS::opencl_constant;
case ParsedAttr::AT_OpenCLGlobalAddressSpace:
return LangAS::opencl_global;
case ParsedAttr::AT_OpenCLLocalAddressSpace:
return LangAS::opencl_local;
case ParsedAttr::AT_OpenCLPrivateAddressSpace:
return LangAS::opencl_private;
case ParsedAttr::AT_OpenCLGenericAddressSpace:
return LangAS::opencl_generic;
default:
return LangAS::Default;
}
}
};

class AttributePool;
Expand Down
14 changes: 14 additions & 0 deletions clang/lib/Parse/ParseDecl.cpp
Expand Up @@ -6177,6 +6177,20 @@ void Parser::ParseFunctionDeclarator(Declarator &D,
Qualifiers Q = Qualifiers::fromCVRUMask(DS.getTypeQualifiers());
if (D.getDeclSpec().isConstexprSpecified() && !getLangOpts().CPlusPlus14)
Q.addConst();
// FIXME: Collect C++ address spaces.
// If there are multiple different address spaces, the source is invalid.
// Carry on using the first addr space for the qualifiers of 'this'.
// The diagnostic will be given later while creating the function
// prototype for the method.
if (getLangOpts().OpenCLCPlusPlus) {
for (ParsedAttr &attr : DS.getAttributes()) {
LangAS ASIdx = attr.asOpenCLLangAS();
if (ASIdx != LangAS::Default) {
Q.addAddressSpace(ASIdx);
break;
}
}
}

Sema::CXXThisScopeRAII ThisScope(
Actions, dyn_cast<CXXRecordDecl>(Actions.CurContext), Q,
Expand Down
22 changes: 15 additions & 7 deletions clang/lib/Sema/SemaOverload.cpp
Expand Up @@ -1171,16 +1171,14 @@ bool Sema::IsOverload(FunctionDecl *New, FunctionDecl *Old,
// function yet (because we haven't yet resolved whether this is a static
// or non-static member function). Add it now, on the assumption that this
// is a redeclaration of OldMethod.
// FIXME: OpenCL: Need to consider address spaces
unsigned OldQuals = OldMethod->getTypeQualifiers().getCVRUQualifiers();
unsigned NewQuals = NewMethod->getTypeQualifiers().getCVRUQualifiers();
auto OldQuals = OldMethod->getTypeQualifiers();
auto NewQuals = NewMethod->getTypeQualifiers();
if (!getLangOpts().CPlusPlus14 && NewMethod->isConstexpr() &&
!isa<CXXConstructorDecl>(NewMethod))
NewQuals |= Qualifiers::Const;

NewQuals.addConst();
// We do not allow overloading based off of '__restrict'.
OldQuals &= ~Qualifiers::Restrict;
NewQuals &= ~Qualifiers::Restrict;
OldQuals.removeRestrict();
NewQuals.removeRestrict();
if (OldQuals != NewQuals)
return true;
}
Expand Down Expand Up @@ -5150,6 +5148,16 @@ TryObjectArgumentInitialization(Sema &S, SourceLocation Loc, QualType FromType,
return ICS;
}

if (FromTypeCanon.getQualifiers().hasAddressSpace()) {
Qualifiers QualsImplicitParamType = ImplicitParamType.getQualifiers();
Qualifiers QualsFromType = FromTypeCanon.getQualifiers();
if (!QualsImplicitParamType.isAddressSpaceSupersetOf(QualsFromType)) {
ICS.setBad(BadConversionSequence::bad_qualifiers,
FromType, ImplicitParamType);
return ICS;
}
}

// Check that we have either the same type or a derived type. It
// affects the conversion rank.
QualType ClassTypeCanon = S.Context.getCanonicalType(ClassType);
Expand Down
107 changes: 58 additions & 49 deletions clang/lib/Sema/SemaType.cpp
Expand Up @@ -3915,6 +3915,25 @@ static Attr *createNullabilityAttr(ASTContext &Ctx, ParsedAttr &Attr,
llvm_unreachable("unknown NullabilityKind");
}

// Diagnose whether this is a case with the multiple addr spaces.
// Returns true if this is an invalid case.
// ISO/IEC TR 18037 S5.3 (amending C99 6.7.3): "No type shall be qualified
// by qualifiers for two or more different address spaces."
static bool DiagnoseMultipleAddrSpaceAttributes(Sema &S, LangAS ASOld,
LangAS ASNew,
SourceLocation AttrLoc) {
if (ASOld != LangAS::Default) {
if (ASOld != ASNew) {
S.Diag(AttrLoc, diag::err_attribute_address_multiple_qualifiers);
return true;
}
// Emit a warning if they are identical; it's likely unintended.
S.Diag(AttrLoc,
diag::warn_attribute_address_multiple_identical_qualifiers);
}
return false;
}

static TypeSourceInfo *
GetTypeSourceInfoForDeclarator(TypeProcessingState &State,
QualType T, TypeSourceInfo *ReturnTypeInfo);
Expand Down Expand Up @@ -4822,18 +4841,35 @@ static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state,
Exceptions,
EPI.ExceptionSpec);

const auto &Spec = D.getCXXScopeSpec();
// FIXME: Set address space from attrs for C++ mode here.
// OpenCLCPlusPlus: A class member function has an address space.
if (state.getSema().getLangOpts().OpenCLCPlusPlus &&
((!Spec.isEmpty() &&
Spec.getScopeRep()->getKind() == NestedNameSpecifier::TypeSpec) ||
state.getDeclarator().getContext() ==
DeclaratorContext::MemberContext)) {
LangAS CurAS = EPI.TypeQuals.getAddressSpace();
auto IsClassMember = [&]() {
return (!state.getDeclarator().getCXXScopeSpec().isEmpty() &&
state.getDeclarator()
.getCXXScopeSpec()
.getScopeRep()
->getKind() == NestedNameSpecifier::TypeSpec) ||
state.getDeclarator().getContext() ==
DeclaratorContext::MemberContext;
};

if (state.getSema().getLangOpts().OpenCLCPlusPlus && IsClassMember()) {
LangAS ASIdx = LangAS::Default;
// Take address space attr if any and mark as invalid to avoid adding
// them later while creating QualType.
if (FTI.MethodQualifiers)
for (ParsedAttr &attr : FTI.MethodQualifiers->getAttributes()) {
LangAS ASIdxNew = attr.asOpenCLLangAS();
if (DiagnoseMultipleAddrSpaceAttributes(S, ASIdx, ASIdxNew,
attr.getLoc()))
D.setInvalidType(true);
else
ASIdx = ASIdxNew;
}
// If a class member function's address space is not set, set it to
// __generic.
LangAS AS =
(CurAS == LangAS::Default ? LangAS::opencl_generic : CurAS);
(ASIdx == LangAS::Default ? LangAS::opencl_generic : ASIdx);
EPI.TypeQuals.addAddressSpace(AS);
}
T = Context.getFunctionType(T, ParamTys, EPI);
Expand Down Expand Up @@ -5789,19 +5825,9 @@ QualType Sema::BuildAddressSpaceAttr(QualType &T, Expr *AddrSpace,
LangAS ASIdx =
getLangASFromTargetAS(static_cast<unsigned>(addrSpace.getZExtValue()));

// If this type is already address space qualified with a different
// address space, reject it.
// ISO/IEC TR 18037 S5.3 (amending C99 6.7.3): "No type shall be qualified
// by qualifiers for two or more different address spaces."
if (T.getAddressSpace() != LangAS::Default) {
if (T.getAddressSpace() != ASIdx) {
Diag(AttrLoc, diag::err_attribute_address_multiple_qualifiers);
return QualType();
} else
// Emit a warning if they are identical; it's likely unintended.
Diag(AttrLoc,
diag::warn_attribute_address_multiple_identical_qualifiers);
}
if (DiagnoseMultipleAddrSpaceAttributes(*this, T.getAddressSpace(), ASIdx,
AttrLoc))
return QualType();

return Context.getAddrSpaceQualType(T, ASIdx);
}
Expand Down Expand Up @@ -5879,34 +5905,14 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
}
} else {
// The keyword-based type attributes imply which address space to use.
switch (Attr.getKind()) {
case ParsedAttr::AT_OpenCLGlobalAddressSpace:
ASIdx = LangAS::opencl_global; break;
case ParsedAttr::AT_OpenCLLocalAddressSpace:
ASIdx = LangAS::opencl_local; break;
case ParsedAttr::AT_OpenCLConstantAddressSpace:
ASIdx = LangAS::opencl_constant; break;
case ParsedAttr::AT_OpenCLGenericAddressSpace:
ASIdx = LangAS::opencl_generic; break;
case ParsedAttr::AT_OpenCLPrivateAddressSpace:
ASIdx = LangAS::opencl_private; break;
default:
ASIdx = Attr.asOpenCLLangAS();
if (ASIdx == LangAS::Default)
llvm_unreachable("Invalid address space");
}

// If this type is already address space qualified with a different
// address space, reject it.
// ISO/IEC TR 18037 S5.3 (amending C99 6.7.3): "No type shall be qualified by
// qualifiers for two or more different address spaces."
if (Type.getAddressSpace() != LangAS::Default) {
if (Type.getAddressSpace() != ASIdx) {
S.Diag(Attr.getLoc(), diag::err_attribute_address_multiple_qualifiers);
Attr.setInvalid();
return;
} else
// Emit a warning if they are identical; it's likely unintended.
S.Diag(Attr.getLoc(),
diag::warn_attribute_address_multiple_identical_qualifiers);
if (DiagnoseMultipleAddrSpaceAttributes(S, Type.getAddressSpace(), ASIdx,
Attr.getLoc())) {
Attr.setInvalid();
return;
}

Type = S.Context.getAddrSpaceQualType(Type, ASIdx);
Expand Down Expand Up @@ -7243,9 +7249,12 @@ static void deduceOpenCLImplicitAddrSpace(TypeProcessingState &State,
// Do not deduce addr space of the void type, e.g. in f(void), otherwise
// it will fail some sema check.
(T->isVoidType() && !IsPointee) ||
// Do not deduce address spaces for dependent types because they might end
// Do not deduce addr spaces for dependent types because they might end
// up instantiating to a type with an explicit address space qualifier.
T->isDependentType())
T->isDependentType() ||
// Do not deduce addr space of decltype because it will be taken from
// its argument.
T->isDecltypeType())
return;

LangAS ImpAddr = LangAS::Default;
Expand Down
35 changes: 35 additions & 0 deletions clang/test/CodeGenOpenCLCXX/method-overload-address-space.cl
@@ -0,0 +1,35 @@
//RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -O0 -o - | FileCheck %s

struct C {
void foo() __local;
void foo() __global;
void foo();
void bar();
};

__global C c1;

__kernel void k() {
__local C c2;
C c3;
__global C &c_ref = c1;
__global C *c_ptr;

// CHECK: call void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
c1.foo();
// CHECK: call void @_ZNU3AS31C3fooEv(%struct.C addrspace(3)*
c2.foo();
// CHECK: call void @_ZNU3AS41C3fooEv(%struct.C addrspace(4)*
c3.foo();
// CHECK: call void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
c_ptr->foo();
// CHECK: void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
c_ref.foo();

// CHECK: call void @_ZNU3AS41C3barEv(%struct.C addrspace(4)* addrspacecast (%struct.C addrspace(1)* @c1 to %struct.C addrspace(4)*))
c1.bar();
//FIXME: Doesn't compile yet
//c_ptr->bar();
// CHECK: call void @_ZNU3AS41C3barEv(%struct.C addrspace(4)* addrspacecast (%struct.C addrspace(1)* @c1 to %struct.C addrspace(4)*))
c_ref.bar();
}
18 changes: 18 additions & 0 deletions clang/test/SemaOpenCLCXX/address-space-of-this-class-scope.cl
@@ -0,0 +1,18 @@
//RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -pedantic -verify

struct C {
auto fGlob() __global -> decltype(this);
auto fGen() -> decltype(this);
auto fErr() __global __local -> decltype(this); //expected-error{{multiple address spaces specified for type}}
};

void bar(__local C*);
// expected-note@-1{{candidate function not viable: address space mismatch in 1st argument ('decltype(this)' (aka '__global C *')), parameter type must be '__local C *'}}
// expected-note@-2{{candidate function not viable: address space mismatch in 1st argument ('decltype(this)' (aka 'C *')), parameter type must be '__local C *'}}

__global C Glob;
void foo(){
bar(Glob.fGlob()); // expected-error{{no matching function for call to 'bar'}}
// FIXME: AS of 'this' below should be correctly deduced to generic
bar(Glob.fGen()); // expected-error{{no matching function for call to 'bar'}}
}
6 changes: 3 additions & 3 deletions clang/test/SemaOpenCLCXX/address_space_overloading.cl
@@ -1,12 +1,12 @@
// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=c++

// expected-no-diagnostics
// FIXME: This test shouldn't trigger any errors.

struct RetGlob {
int dummy;
};

struct RetGen {
struct RetGen { //expected-error{{binding value of type '__generic RetGen' to reference to type 'RetGen' drops <<ERROR>> qualifiers}}
char dummy;
};

Expand All @@ -19,5 +19,5 @@ void kernel k() {
__local int *ArgLoc;
RetGlob TestGlob = foo(ArgGlob);
RetGen TestGen = foo(ArgGen);
TestGen = foo(ArgLoc);
TestGen = foo(ArgLoc); //expected-note{{in implicit copy assignment operator for 'RetGen' first required here}}
}
20 changes: 20 additions & 0 deletions clang/test/SemaOpenCLCXX/method-overload-address-space.cl
@@ -0,0 +1,20 @@
//RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -pedantic -verify

struct C {
void m1() __local __local; //expected-warning{{multiple identical address spaces specified for type}}
//expected-note@-1{{candidate function}}
void m1() __global;
//expected-note@-1{{candidate function}}
void m2() __global __local; //expected-error{{multiple address spaces specified for type}}
};

__global C c_glob;

__kernel void bar() {
__local C c_loc;
C c_priv;

c_glob.m1();
c_loc.m1();
c_priv.m1(); //expected-error{{no matching member function for call to 'm1'}}
}

0 comments on commit 5cffa45

Please sign in to comment.