From: Anastasia Stulova Date: Mon, 21 Jan 2019 16:01:38 +0000 (+0000) Subject: [OpenCL] Allow address spaces as method qualifiers. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=710b2573329fdf3f6180e159812b4f4fc368fc79;p=clang [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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@351747 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/include/clang/AST/Type.h b/include/clang/AST/Type.h index fbf91d3088..50271aa840 100644 --- a/include/clang/AST/Type.h +++ b/include/clang/AST/Type.h @@ -1982,7 +1982,7 @@ public: bool isObjCQualifiedClassType() const; // Class 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() const { return isObjCIdType() || isObjCClassType() || isObjCSelType(); } +inline bool Type::isDecltypeType() const { + return isa(this); +} + #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ inline bool Type::is##Id##Type() const { \ return isSpecificBuiltinType(BuiltinType::Id); \ diff --git a/include/clang/Sema/ParsedAttr.h b/include/clang/Sema/ParsedAttr.h index 82488b1700..dd598f41ab 100644 --- a/include/clang/Sema/ParsedAttr.h +++ b/include/clang/Sema/ParsedAttr.h @@ -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; diff --git a/lib/Parse/ParseDecl.cpp b/lib/Parse/ParseDecl.cpp index 2ebf757ffa..bbead42f03 100644 --- a/lib/Parse/ParseDecl.cpp +++ b/lib/Parse/ParseDecl.cpp @@ -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(Actions.CurContext), Q, diff --git a/lib/Sema/SemaOverload.cpp b/lib/Sema/SemaOverload.cpp index 7b706cb262..17920d00e6 100644 --- a/lib/Sema/SemaOverload.cpp +++ b/lib/Sema/SemaOverload.cpp @@ -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(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, 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); diff --git a/lib/Sema/SemaType.cpp b/lib/Sema/SemaType.cpp index 538bb6aafd..93b42829bb 100644 --- a/lib/Sema/SemaType.cpp +++ b/lib/Sema/SemaType.cpp @@ -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); @@ -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); @@ -5789,19 +5825,9 @@ QualType Sema::BuildAddressSpaceAttr(QualType &T, Expr *AddrSpace, LangAS ASIdx = getLangASFromTargetAS(static_cast(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 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); @@ -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; diff --git a/test/CodeGenOpenCLCXX/method-overload-address-space.cl b/test/CodeGenOpenCLCXX/method-overload-address-space.cl new file mode 100644 index 0000000000..0864589b05 --- /dev/null +++ b/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(); +} diff --git a/test/SemaOpenCLCXX/address-space-of-this-class-scope.cl b/test/SemaOpenCLCXX/address-space-of-this-class-scope.cl new file mode 100644 index 0000000000..ebb76043bb --- /dev/null +++ b/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'}} +} diff --git a/test/SemaOpenCLCXX/address_space_overloading.cl b/test/SemaOpenCLCXX/address_space_overloading.cl index ccdd5735bb..6c1934083a 100644 --- a/test/SemaOpenCLCXX/address_space_overloading.cl +++ b/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 <> 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}} } diff --git a/test/SemaOpenCLCXX/method-overload-address-space.cl b/test/SemaOpenCLCXX/method-overload-address-space.cl new file mode 100644 index 0000000000..64a279549c --- /dev/null +++ b/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'}} +}