[llvm-branch-commits] [clang] b33830a - [OpenCL] Restrict addr space conversions in nested pointers
Hans Wennborg via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Wed Feb 12 04:48:44 PST 2020
Author: Anastasia Stulova
Date: 2020-02-12T13:45:17+01:00
New Revision: b33830aea54536f0f03334b2015b03af336ec90c
URL: https://github.com/llvm/llvm-project/commit/b33830aea54536f0f03334b2015b03af336ec90c
DIFF: https://github.com/llvm/llvm-project/commit/b33830aea54536f0f03334b2015b03af336ec90c.diff
LOG: [OpenCL] Restrict addr space conversions in nested pointers
Address space conversion changes pointer representation.
This commit disallows such conversions when they are not
legal i.e. for the nested pointers even with compatible
address spaces. Because the address space conversion in
the nested levels can't be generated to modify the pointers
correctly. The behavior implemented is as follows:
- Any implicit conversions of nested pointers with different
address spaces is rejected.
- Any conversion of address spaces in nested pointers in safe
casts (e.g. const_cast or static_cast) is rejected.
- Conversion in low level C-style or reinterpret_cast is accepted
but with a warning (this aligns with OpenCL C behavior).
Fixes PR39674
Differential Revision: https://reviews.llvm.org/D73360
(cherry picked from commit 6064f426a18304e16b51cc79e74c9c2d55ef5a9c)
Added:
clang/test/SemaOpenCLCXX/address-space-castoperators.cl
Modified:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/Sema/SemaCast.cpp
clang/lib/Sema/SemaOverload.cpp
clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
clang/test/SemaOpenCL/address-spaces.cl
clang/test/SemaOpenCLCXX/address-space-deduction.cl
clang/test/SemaOpenCLCXX/address-space-references.cl
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index ecf959b9077a..2199dfbddc84 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6748,6 +6748,10 @@ def err_bad_cxx_cast_scalar_to_vector_
diff erent_size : Error<
def err_bad_cxx_cast_vector_to_vector_
diff erent_size : Error<
"%select{||reinterpret_cast||C-style cast|}0 from vector %1 "
"to vector %2 of
diff erent size">;
+def warn_bad_cxx_cast_nested_pointer_addr_space : Warning<
+ "%select{reinterpret_cast|C-style cast}0 from %1 to %2 "
+ "changes address space of nested pointers">,
+ InGroup<IncompatiblePointerTypesDiscardsQualifiers>;
def err_bad_lvalue_to_rvalue_cast : Error<
"cannot cast from lvalue of type %1 to rvalue reference type %2; types are "
"not compatible">;
diff --git a/clang/lib/Sema/SemaCast.cpp b/clang/lib/Sema/SemaCast.cpp
index a905ebc67305..7a8cbca1e3f1 100644
--- a/clang/lib/Sema/SemaCast.cpp
+++ b/clang/lib/Sema/SemaCast.cpp
@@ -2311,6 +2311,24 @@ static TryCastResult TryReinterpretCast(Sema &Self, ExprResult &SrcExpr,
return SuccessResult;
}
+ // Diagnose address space conversion in nested pointers.
+ QualType DestPtee = DestType->getPointeeType().isNull()
+ ? DestType->getPointeeType()
+ : DestType->getPointeeType()->getPointeeType();
+ QualType SrcPtee = SrcType->getPointeeType().isNull()
+ ? SrcType->getPointeeType()
+ : SrcType->getPointeeType()->getPointeeType();
+ while (!DestPtee.isNull() && !SrcPtee.isNull()) {
+ if (DestPtee.getAddressSpace() != SrcPtee.getAddressSpace()) {
+ Self.Diag(OpRange.getBegin(),
+ diag::warn_bad_cxx_cast_nested_pointer_addr_space)
+ << CStyle << SrcType << DestType << SrcExpr.get()->getSourceRange();
+ break;
+ }
+ DestPtee = DestPtee->getPointeeType();
+ SrcPtee = SrcPtee->getPointeeType();
+ }
+
// C++ 5.2.10p7: A pointer to an object can be explicitly converted to
// a pointer to an object of
diff erent type.
// Void pointers are not specified, but supported by every compiler out there.
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index 9a9843827b3f..db1884acd349 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -3176,7 +3176,7 @@ static bool isNonTrivialObjCLifetimeConversion(Qualifiers FromQuals,
/// FromType and \p ToType is permissible, given knowledge about whether every
/// outer layer is const-qualified.
static bool isQualificationConversionStep(QualType FromType, QualType ToType,
- bool CStyle,
+ bool CStyle, bool IsTopLevel,
bool &PreviousToQualsIncludeConst,
bool &ObjCLifetimeConversion) {
Qualifiers FromQuals = FromType.getQualifiers();
@@ -3213,11 +3213,15 @@ static bool isQualificationConversionStep(QualType FromType, QualType ToType,
if (!CStyle && !ToQuals.compatiblyIncludes(FromQuals))
return false;
- // For a C-style cast, just require the address spaces to overlap.
- // FIXME: Does "superset" also imply the representation of a pointer is the
- // same? We're assuming that it does here and in compatiblyIncludes.
- if (CStyle && !ToQuals.isAddressSpaceSupersetOf(FromQuals) &&
- !FromQuals.isAddressSpaceSupersetOf(ToQuals))
+ // If address spaces mismatch:
+ // - in top level it is only valid to convert to addr space that is a
+ // superset in all cases apart from C-style casts where we allow
+ // conversions between overlapping address spaces.
+ // - in non-top levels it is not a valid conversion.
+ if (ToQuals.getAddressSpace() != FromQuals.getAddressSpace() &&
+ (!IsTopLevel ||
+ !(ToQuals.isAddressSpaceSupersetOf(FromQuals) ||
+ (CStyle && FromQuals.isAddressSpaceSupersetOf(ToQuals)))))
return false;
// -- if the cv 1,j and cv 2,j are
diff erent, then const is in
@@ -3258,9 +3262,9 @@ Sema::IsQualificationConversion(QualType FromType, QualType ToType,
bool PreviousToQualsIncludeConst = true;
bool UnwrappedAnyPointer = false;
while (Context.UnwrapSimilarTypes(FromType, ToType)) {
- if (!isQualificationConversionStep(FromType, ToType, CStyle,
- PreviousToQualsIncludeConst,
- ObjCLifetimeConversion))
+ if (!isQualificationConversionStep(
+ FromType, ToType, CStyle, !UnwrappedAnyPointer,
+ PreviousToQualsIncludeConst, ObjCLifetimeConversion))
return false;
UnwrappedAnyPointer = true;
}
@@ -4499,7 +4503,7 @@ Sema::CompareReferenceRelationship(SourceLocation Loc,
// If we find a qualifier mismatch, the types are not reference-compatible,
// but are still be reference-related if they're similar.
bool ObjCLifetimeConversion = false;
- if (!isQualificationConversionStep(T2, T1, /*CStyle=*/false,
+ if (!isQualificationConversionStep(T2, T1, /*CStyle=*/false, TopLevel,
PreviousToQualsIncludeConst,
ObjCLifetimeConversion))
return (ConvertedReferent || Context.hasSimilarType(T1, T2))
diff --git a/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl b/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
index 14f7cf3a0e7e..5efea216346a 100644
--- a/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
+++ b/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
@@ -513,17 +513,37 @@ void test_pointer_chains() {
#endif
// Case 3: Corresponded inner pointees has overlapping but not equivalent address spaces.
- // FIXME: Should this really be allowed in C++ mode?
var_as_as_int = var_asc_asc_int;
-#if !__OPENCL_CPP_VERSION__
#ifdef GENERIC
+#if !__OPENCL_CPP_VERSION__
// expected-error at -3 {{assigning '__local int *__local *__private' to '__generic int *__generic *__private' changes address space of nested pointer}}
+#else
+// expected-error at -5 {{assigning to '__generic int *__generic *' from incompatible type '__local int *__local *__private'}}
#endif
#endif
- var_as_as_int = 0 ? var_as_as_int : var_asc_asc_int;
+
+ var_as_as_int = (AS int *AS *)var_asc_asc_int;
+#ifdef GENERIC
#if !__OPENCL_CPP_VERSION__
+// expected-warning at -3 {{casting '__local int *__local *' to type '__generic int *__generic *' discards qualifiers in nested pointer types}}
+#else
+// expected-warning at -5 {{C-style cast from '__local int *__local *' to '__generic int *__generic *' changes address space of nested pointers}}
+#endif
+#endif
+
+ var_as_as_int = (AS int *AS *)var_asc_asn_int;
+#if !__OPENCL_CPP_VERSION__
+// expected-warning-re at -2 {{casting '__{{global|local|constant}} int *__{{local|constant|global}} *' to type '__{{global|constant|generic}} int *__{{global|constant|generic}} *' discards qualifiers in nested pointer types}}
+#else
+// expected-warning-re at -4 {{C-style cast from '__{{global|local|constant}} int *__{{local|constant|global}} *' to '__{{global|constant|generic}} int *__{{global|constant|generic}} *' changes address space of nested pointers}}
+#endif
+
+ var_as_as_int = 0 ? var_as_as_int : var_asc_asc_int;
#ifdef GENERIC
+#if !__OPENCL_CPP_VERSION__
// expected-warning at -3{{pointer type mismatch ('__generic int *__generic *' and '__local int *__local *')}}
+#else
+// expected-error at -5 {{incompatible operand types ('__generic int *__generic *' and '__local int *__local *')}}
#endif
#endif
}
diff --git a/clang/test/SemaOpenCL/address-spaces.cl b/clang/test/SemaOpenCL/address-spaces.cl
index bc9a07d88704..07547ea19680 100644
--- a/clang/test/SemaOpenCL/address-spaces.cl
+++ b/clang/test/SemaOpenCL/address-spaces.cl
@@ -198,9 +198,7 @@ void nested(__global int *g, __global int * __private *gg, __local int *l, __loc
ll = gg; // expected-error {{assigning to '__local int *__private *' from incompatible type '__global int *__private *__private'}}
ll = l; // expected-error {{assigning to '__local int *__private *' from incompatible type '__local int *__private'; take the address with &}}
ll = gg_f; // expected-error {{assigning to '__local int *__private *' from incompatible type '__global float *__private *__private'}}
- // FIXME: The below becomes a reinterpret_cast, and therefore does not emit an error
- // even though the address space mismatches in the nested pointers.
- ll = (__local int * __private *)gg;
+ ll = (__local int *__private *)gg; //expected-warning{{C-style cast from '__global int *__private *' to '__local int *__private *' changes address space of nested pointers}}
gg_f = g; // expected-error {{assigning to '__global float *__private *' from incompatible type '__global int *__private'}}
gg_f = gg; // expected-error {{assigning to '__global float *__private *' from incompatible type '__global int *__private *__private'}}
diff --git a/clang/test/SemaOpenCLCXX/address-space-castoperators.cl b/clang/test/SemaOpenCLCXX/address-space-castoperators.cl
new file mode 100644
index 000000000000..d61a9a72573c
--- /dev/null
+++ b/clang/test/SemaOpenCLCXX/address-space-castoperators.cl
@@ -0,0 +1,12 @@
+//RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
+
+void nester_ptr() {
+ local int * * locgen;
+ constant int * * congen;
+ int * * gengen;
+
+ gengen = const_cast<int**>(locgen); //expected-error{{const_cast from '__local int *__generic *' to '__generic int *__generic *' is not allowed}}
+ gengen = static_cast<int**>(locgen); //expected-error{{static_cast from '__local int *__generic *' to '__generic int *__generic *' is not allowed}}
+// CHECK-NOT: AddressSpaceConversion
+ gengen = reinterpret_cast<int**>(locgen); //expected-warning{{reinterpret_cast from '__local int *__generic *' to '__generic int *__generic *' changes address space of nested pointers}}
+}
diff --git a/clang/test/SemaOpenCLCXX/address-space-deduction.cl b/clang/test/SemaOpenCLCXX/address-space-deduction.cl
index 0275ae57efe4..6a81a8b2d7c7 100644
--- a/clang/test/SemaOpenCLCXX/address-space-deduction.cl
+++ b/clang/test/SemaOpenCLCXX/address-space-deduction.cl
@@ -105,7 +105,7 @@ void t5(float (*(*fYZ))[2]);
__kernel void k() {
__local float x[2];
- __local float(*p)[2];
+ float(*p)[2];
t1(x);
t2(&x);
t3(&x);
diff --git a/clang/test/SemaOpenCLCXX/address-space-references.cl b/clang/test/SemaOpenCLCXX/address-space-references.cl
index a6c7c842da00..068318dfa141 100644
--- a/clang/test/SemaOpenCLCXX/address-space-references.cl
+++ b/clang/test/SemaOpenCLCXX/address-space-references.cl
@@ -13,3 +13,16 @@ int bar(const unsigned int &i);
void foo() {
bar(1) // expected-error{{binding reference of type 'const __global unsigned int' to value of type 'int' changes address space}}
}
+
+// Test addr space conversion with nested pointers
+
+extern void nestptr(int *&); // expected-note {{candidate function not viable: no known conversion from '__global int *__private' to '__generic int *__generic &__private' for 1st argument}}
+extern void nestptr_const(int * const &); // expected-note {{candidate function not viable: cannot pass pointer to address space '__constant' as a pointer to address space '__generic' in 1st argument}}
+int test_nestptr(__global int *glob, __constant int *cons, int* gen) {
+ nestptr(glob); // expected-error{{no matching function for call to 'nestptr'}}
+ // Addr space conversion first occurs on a temporary.
+ nestptr_const(glob);
+ // No legal conversion between disjoint addr spaces.
+ nestptr_const(cons); // expected-error{{no matching function for call to 'nestptr_const'}}
+ return *(*cons ? glob : gen);
+}
More information about the llvm-branch-commits
mailing list