r351747 - [OpenCL] Allow address spaces as method qualifiers.
Anastasia Stulova via cfe-commits
cfe-commits at lists.llvm.org
Mon Jan 21 08:01:39 PST 2019
Author: stulova
Date: Mon Jan 21 08:01:38 2019
New Revision: 351747
URL: http://llvm.org/viewvc/llvm-project?rev=351747&view=rev
Log:
[OpenCL] Allow address spaces as method qualifiers.
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
Added:
cfe/trunk/test/CodeGenOpenCLCXX/method-overload-address-space.cl
cfe/trunk/test/SemaOpenCLCXX/address-space-of-this-class-scope.cl
cfe/trunk/test/SemaOpenCLCXX/method-overload-address-space.cl
Modified:
cfe/trunk/include/clang/AST/Type.h
cfe/trunk/include/clang/Sema/ParsedAttr.h
cfe/trunk/lib/Parse/ParseDecl.cpp
cfe/trunk/lib/Sema/SemaOverload.cpp
cfe/trunk/lib/Sema/SemaType.cpp
cfe/trunk/test/SemaOpenCLCXX/address_space_overloading.cl
Modified: cfe/trunk/include/clang/AST/Type.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/Type.h?rev=351747&r1=351746&r2=351747&view=diff
==============================================================================
--- cfe/trunk/include/clang/AST/Type.h (original)
+++ cfe/trunk/include/clang/AST/Type.h Mon Jan 21 08:01:38 2019
@@ -1982,7 +1982,7 @@ public:
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?
///
@@ -6440,6 +6440,10 @@ inline bool Type::isObjCBuiltinType() co
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); \
Modified: cfe/trunk/include/clang/Sema/ParsedAttr.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/ParsedAttr.h?rev=351747&r1=351746&r2=351747&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/ParsedAttr.h (original)
+++ cfe/trunk/include/clang/Sema/ParsedAttr.h Mon Jan 21 08:01:38 2019
@@ -567,6 +567,25 @@ public:
/// 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;
Modified: cfe/trunk/lib/Parse/ParseDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/ParseDecl.cpp?rev=351747&r1=351746&r2=351747&view=diff
==============================================================================
--- cfe/trunk/lib/Parse/ParseDecl.cpp (original)
+++ cfe/trunk/lib/Parse/ParseDecl.cpp Mon Jan 21 08:01:38 2019
@@ -6177,6 +6177,20 @@ void Parser::ParseFunctionDeclarator(Dec
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,
Modified: cfe/trunk/lib/Sema/SemaOverload.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOverload.cpp?rev=351747&r1=351746&r2=351747&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOverload.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOverload.cpp Mon Jan 21 08:01:38 2019
@@ -1171,16 +1171,14 @@ bool Sema::IsOverload(FunctionDecl *New,
// 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;
}
@@ -5150,6 +5148,16 @@ TryObjectArgumentInitialization(Sema &S,
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);
Modified: cfe/trunk/lib/Sema/SemaType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=351747&r1=351746&r2=351747&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaType.cpp (original)
+++ cfe/trunk/lib/Sema/SemaType.cpp Mon Jan 21 08:01:38 2019
@@ -3915,6 +3915,25 @@ static Attr *createNullabilityAttr(ASTCo
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);
@@ -4822,18 +4841,35 @@ static TypeSourceInfo *GetFullTypeForDec
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);
@@ -5789,19 +5825,9 @@ QualType Sema::BuildAddressSpaceAttr(Qua
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);
}
@@ -5879,34 +5905,14 @@ static void HandleAddressSpaceTypeAttrib
}
} 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);
@@ -7243,9 +7249,12 @@ static void deduceOpenCLImplicitAddrSpac
// 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;
Added: cfe/trunk/test/CodeGenOpenCLCXX/method-overload-address-space.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCLCXX/method-overload-address-space.cl?rev=351747&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCLCXX/method-overload-address-space.cl (added)
+++ cfe/trunk/test/CodeGenOpenCLCXX/method-overload-address-space.cl Mon Jan 21 08:01:38 2019
@@ -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();
+}
Added: cfe/trunk/test/SemaOpenCLCXX/address-space-of-this-class-scope.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCLCXX/address-space-of-this-class-scope.cl?rev=351747&view=auto
==============================================================================
--- cfe/trunk/test/SemaOpenCLCXX/address-space-of-this-class-scope.cl (added)
+++ cfe/trunk/test/SemaOpenCLCXX/address-space-of-this-class-scope.cl Mon Jan 21 08:01:38 2019
@@ -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 at -1{{candidate function not viable: address space mismatch in 1st argument ('decltype(this)' (aka '__global C *')), parameter type must be '__local C *'}}
+// expected-note at -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'}}
+}
Modified: cfe/trunk/test/SemaOpenCLCXX/address_space_overloading.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCLCXX/address_space_overloading.cl?rev=351747&r1=351746&r2=351747&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCLCXX/address_space_overloading.cl (original)
+++ cfe/trunk/test/SemaOpenCLCXX/address_space_overloading.cl Mon Jan 21 08:01:38 2019
@@ -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;
};
@@ -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}}
}
Added: cfe/trunk/test/SemaOpenCLCXX/method-overload-address-space.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCLCXX/method-overload-address-space.cl?rev=351747&view=auto
==============================================================================
--- cfe/trunk/test/SemaOpenCLCXX/method-overload-address-space.cl (added)
+++ cfe/trunk/test/SemaOpenCLCXX/method-overload-address-space.cl Mon Jan 21 08:01:38 2019
@@ -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 at -1{{candidate function}}
+ void m1() __global;
+ //expected-note at -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'}}
+}
More information about the cfe-commits
mailing list