Anastasia updated this revision to Diff 181825.
Anastasia added a comment.
Rebased on top of recent related changes and addressed remaining review
comments.
This now depends on: https://reviews.llvm.org/D56735
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D55850/new/
https://reviews.llvm.org/D55850
Files:
include/clang/Sema/ParsedAttr.h
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
@@ -3916,6 +3916,25 @@
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;
+ } else
+ // 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);
@@ -4823,18 +4842,35 @@
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);
@@ -5790,19 +5826,9 @@
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);
}
@@ -5880,34 +5906,14 @@
}
} 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);
Index: lib/Sema/SemaOverload.cpp
===================================================================
--- lib/Sema/SemaOverload.cpp
+++ lib/Sema/SemaOverload.cpp
@@ -1172,16 +1172,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;
}
@@ -5148,6 +5146,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);
Index: lib/Parse/ParseDecl.cpp
===================================================================
--- lib/Parse/ParseDecl.cpp
+++ lib/Parse/ParseDecl.cpp
@@ -6159,6 +6159,16 @@
Qualifiers Q = Qualifiers::fromCVRUMask(DS.getTypeQualifiers());
if (D.getDeclSpec().isConstexprSpecified() && !getLangOpts().CPlusPlus14)
Q.addConst();
+ // FIXME: Collect C++ address spaces.
+ 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,
Index: include/clang/Sema/ParsedAttr.h
===================================================================
--- include/clang/Sema/ParsedAttr.h
+++ include/clang/Sema/ParsedAttr.h
@@ -568,6 +568,25 @@
/// 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 attribute return its representation in LangAS,
+ /// otherwise return 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;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
