Anastasia updated this revision to Diff 178904.
Anastasia marked an inline comment as done.
Anastasia added a comment.

Address review comments.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D55850/new/

https://reviews.llvm.org/D55850

Files:
  lib/Parse/ParseDecl.cpp
  lib/Sema/SemaOverload.cpp
  lib/Sema/SemaType.cpp
  test/CodeGenOpenCLCXX/method-overload-address-space.cl
  test/SemaOpenCLCXX/method-overload-address-space.cl

Index: test/SemaOpenCLCXX/method-overload-address-space.cl
===================================================================
--- /dev/null
+++ 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'}}
+}
Index: test/CodeGenOpenCLCXX/method-overload-address-space.cl
===================================================================
--- /dev/null
+++ 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();
+}
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -3921,6 +3921,14 @@
   llvm_unreachable("unknown NullabilityKind");
 }
 
+/// IsClassMember - Determines whether a state belongs to a class member.
+static bool IsClassMember(TypeProcessingState &State) {
+  return (!State.getDeclarator().getCXXScopeSpec().isEmpty() &&
+          State.getDeclarator().getCXXScopeSpec().getScopeRep()->getKind() ==
+              NestedNameSpecifier::TypeSpec) ||
+         State.getDeclarator().getContext() == DeclaratorContext::MemberContext;
+}
+
 static TypeSourceInfo *
 GetTypeSourceInfoForDeclarator(TypeProcessingState &State,
                                QualType T, TypeSourceInfo *ReturnTypeInfo);
@@ -4825,18 +4833,45 @@
                                       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();
+            IsClassMember(state)) {
+          LangAS ASIdx = LangAS::Default;
+          // Take address space attr if any and mark as invalid to avoid adding
+          // them later while creating QualType.
+          for (ParsedAttr &attr : DeclType.getAttrs()) {
+            switch (attr.getKind()) {
+            case ParsedAttr::AT_OpenCLConstantAddressSpace:
+              ASIdx = LangAS::opencl_constant;
+              attr.setInvalid();
+              break;
+            case ParsedAttr::AT_OpenCLLocalAddressSpace:
+              ASIdx = LangAS::opencl_local;
+              attr.setInvalid();
+              break;
+            case ParsedAttr::AT_OpenCLGlobalAddressSpace:
+              ASIdx = LangAS::opencl_global;
+              attr.setInvalid();
+              break;
+            case ParsedAttr::AT_OpenCLPrivateAddressSpace:
+              ASIdx = LangAS::opencl_private;
+              attr.setInvalid();
+              break;
+            case ParsedAttr::AT_OpenCLGenericAddressSpace:
+              ASIdx = LangAS::opencl_generic;
+              attr.setInvalid();
+              break;
+            default:
+              break;
+            }
+            if (ASIdx != LangAS::Default)
+              break;
+          }
           // 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);
           T = state.getSema().Context.getAddrSpaceQualType(T, AS);
@@ -5834,7 +5869,10 @@
 
   // ISO/IEC TR 18037 S5.3 (amending C99 6.7.3): "A function type shall not be
   // qualified by an address-space qualifier."
-  if (Type->isFunctionType()) {
+  // Allow qualifying methods as an extension.
+  // FIXME: For now only enabled for OpenCL.
+  if (Type->isFunctionType() &&
+      !(S.getLangOpts().OpenCLCPlusPlus && IsClassMember(State))) {
     S.Diag(Attr.getLoc(), diag::err_attribute_address_function_type);
     Attr.setInvalid();
     return;
Index: lib/Sema/SemaOverload.cpp
===================================================================
--- lib/Sema/SemaOverload.cpp
+++ lib/Sema/SemaOverload.cpp
@@ -1142,16 +1142,14 @@
     // 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;
   }
@@ -5120,6 +5118,16 @@
     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);
@@ -9256,6 +9264,15 @@
   if (HasPS1 != HasPS2 && HasPS1)
     return true;
 
+  // Always prefer a specific address space to a generic or default one.
+  if (S.getLangOpts().OpenCLCPlusPlus && Cand1.Function && Cand2.Function) {
+    LangAS CandAS1 = Cand1.Function->getType().getAddressSpace();
+    LangAS CandAS2 = Cand2.Function->getType().getAddressSpace();
+    if ((CandAS2 == LangAS::opencl_generic || CandAS2 == LangAS::Default) &&
+        (CandAS1 != LangAS::opencl_generic && CandAS1 != LangAS::Default))
+      return true;
+  }
+
   return isBetterMultiversionCandidate(Cand1, Cand2);
 }
 
Index: lib/Parse/ParseDecl.cpp
===================================================================
--- lib/Parse/ParseDecl.cpp
+++ lib/Parse/ParseDecl.cpp
@@ -6143,6 +6143,24 @@
         RestrictQualifierLoc = DS.getRestrictSpecLoc();
       }
 
+      // Propagate address space attributes with a language semantic.
+      if (getLangOpts().OpenCLCPlusPlus) {
+        for (ParsedAttr &attr : DS.getAttributes()) {
+          switch (attr.getKind()) {
+          case ParsedAttr::AT_OpenCLConstantAddressSpace:
+          case ParsedAttr::AT_OpenCLLocalAddressSpace:
+          case ParsedAttr::AT_OpenCLGlobalAddressSpace:
+          case ParsedAttr::AT_OpenCLPrivateAddressSpace:
+          case ParsedAttr::AT_OpenCLGenericAddressSpace:
+            FnAttrs.addNew(attr.getName(), attr.getLoc(), nullptr,
+                           attr.getLoc(), nullptr, 0, ParsedAttr::AS_Keyword);
+            break;
+          default:
+            break;
+          }
+        }
+      }
+
       // Parse ref-qualifier[opt].
       if (ParseRefQualifier(RefQualifierIsLValueRef, RefQualifierLoc))
         EndLoc = RefQualifierLoc;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to