Anastasia created this revision.
Anastasia added a reviewer: rjmccall.
Herald added a subscriber: yaxunl.
This patch allows qualifying methods with address spaces and extends some
overloading rules to use those qualifiers.
The main use case is to prevent conversions to generic/default address space
where it's requested by the programmer. More details can be found in:
http://lists.llvm.org/pipermail/cfe-dev/2018-December/060470.html
This patch doesn't enable the feature for C++ yet but it can easily be
generalized if the overall flow is right.
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
@@ -1148,12 +1148,14 @@
if (!getLangOpts().CPlusPlus14 && NewMethod->isConstexpr() &&
!isa<CXXConstructorDecl>(NewMethod))
NewQuals |= Qualifiers::Const;
-
// We do not allow overloading based off of '__restrict'.
OldQuals &= ~Qualifiers::Restrict;
NewQuals &= ~Qualifiers::Restrict;
if (OldQuals != NewQuals)
return true;
+ if (OldMethod->getTypeQualifiers().getAddressSpace() !=
+ NewMethod->getTypeQualifiers().getAddressSpace())
+ return true;
}
// Though pass_object_size is placed on parameters and takes an argument, we
@@ -6654,6 +6656,20 @@
Candidate.Viable = false;
Candidate.FailureKind = ovl_non_default_multiversion_function;
}
+
+ if (Method->getType().getQualifiers().hasAddressSpace() &&
+ !ObjectType.isNull()) {
+ Qualifiers QualsObject =
+ ObjectType->isPointerType() || ObjectType->isReferenceType()
+ ? ObjectType->getPointeeType().getQualifiers()
+ : ObjectType.getQualifiers();
+ Qualifiers QualsMethod = Method->getType().getQualifiers();
+ if (!QualsMethod.isAddressSpaceSupersetOf(QualsObject)) {
+ Candidate.Viable = false;
+ Candidate.FailureKind = ovl_fail_bad_final_conversion;
+ return;
+ }
+ }
}
/// Add a C++ member function template as a candidate to the candidate
@@ -9256,6 +9272,14 @@
if (HasPS1 != HasPS2 && HasPS1)
return true;
+ 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
