[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 fbf91d3088c7d9aafbe33defe348e13e638ce09e..50271aa8402fe92f8d296ec4e4389364ae56276c 100644 (file)
--- a/include/clang/AST/Type.h
+++ b/include/clang/AST/Type.h
@@ -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() 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); \

diff --git a/include/clang/Sema/ParsedAttr.h b/include/clang/Sema/ParsedAttr.h
index 82488b1700ac67ca3d5f9fc2712764221e913b66..dd598f41ab9137a97bb88af5ecd1ecd879b1c7bf 100644 (file)
--- 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 2ebf757ffa24ea1067e351a0679d4f926e9c6435..bbead42f0315b929c7760f8c816596be1e35973e 100644 (file)
--- 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<CXXRecordDecl>(Actions.CurContext), Q,

diff --git a/lib/Sema/SemaOverload.cpp b/lib/Sema/SemaOverload.cpp
index 7b706cb26251d072f598a2e58e3b73f40b208d8a..17920d00e6f7d3cef9d3101d976f5f1bed35f717 100644 (file)
--- 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<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, 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 538bb6aafd086383816b6334be6c40934eab04c5..93b42829bba95cff78ad01dfb6f71317ad5955ff 100644 (file)
--- 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<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 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 (file)
index 0000000..0864589
--- /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 (file)
index 0000000..ebb7604
--- /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 ccdd5735bb90d3926ebc85a64ef5ceb145675fec..6c1934083a843200b752970424fe9bf83d804993 100644 (file)
--- 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 <<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}}
 }

diff --git a/test/SemaOpenCLCXX/method-overload-address-space.cl b/test/SemaOpenCLCXX/method-overload-address-space.cl
new file mode 100644 (file)
index 0000000..64a2795
--- /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'}}
+}