Anastasia created this revision.
Anastasia added reviewers: rjmccall, jeroen.dobbelaere, bevinh.
Herald added subscribers: ebevhan, yaxunl.
This patch is fixing the issue reported in:
http://lists.llvm.org/pipermail/cfe-dev/2020-January/064273.html
and in the bug:
https://bugs.llvm.org/show_bug.cgi?id=39674
Since address space conversion changes pointer representation it is not legal
to do it for the nested pointers even with compatible address spaces. Because
the address space conversion in the nested levels can't be generated. The
behavior implemented by this patch is as follows:
- Reject any implicit conversions of nested pointers with address spaces.
- Reject address space of nested pointers conversions in safe casts e.g.
`const_cast` or `static_cast.`
- Allow conversion in low level C-style or `reinterpret_cast` but with a
warning (this aligns with OpenCL C behavior).
https://reviews.llvm.org/D73360
Files:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/Sema/SemaCast.cpp
clang/lib/Sema/SemaOverload.cpp
clang/test/Misc/warning-flags.c
clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
clang/test/SemaOpenCL/address-spaces.cl
clang/test/SemaOpenCLCXX/address-space-castoperators.cl
clang/test/SemaOpenCLCXX/address-space-deduction.cl
clang/test/SemaOpenCLCXX/address-space-references.cl
Index: clang/test/SemaOpenCLCXX/address-space-references.cl
===================================================================
--- clang/test/SemaOpenCLCXX/address-space-references.cl
+++ clang/test/SemaOpenCLCXX/address-space-references.cl
@@ -13,3 +13,16 @@
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);
+}
Index: clang/test/SemaOpenCLCXX/address-space-deduction.cl
===================================================================
--- clang/test/SemaOpenCLCXX/address-space-deduction.cl
+++ clang/test/SemaOpenCLCXX/address-space-deduction.cl
@@ -105,7 +105,7 @@
__kernel void k() {
__local float x[2];
- __local float(*p)[2];
+ float(*p)[2];
t1(x);
t2(&x);
t3(&x);
Index: clang/test/SemaOpenCLCXX/address-space-castoperators.cl
===================================================================
--- /dev/null
+++ 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}}
+}
Index: clang/test/SemaOpenCL/address-spaces.cl
===================================================================
--- clang/test/SemaOpenCL/address-spaces.cl
+++ clang/test/SemaOpenCL/address-spaces.cl
@@ -198,9 +198,7 @@
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'}}
Index: clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
===================================================================
--- clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
+++ clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
@@ -513,17 +513,37 @@
#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@-3 {{assigning '__local int *__local *__private' to '__generic int *__generic *__private' changes address space of nested pointer}}
+#else
+// expected-error@-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@-3 {{casting '__local int *__local *' to type '__generic int *__generic *' discards qualifiers in nested pointer types}}
+#else
+// expected-warning@-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@-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@-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@-3{{pointer type mismatch ('__generic int *__generic *' and '__local int *__local *')}}
+#else
+// expected-error@-5 {{incompatible operand types ('__generic int *__generic *' and '__local int *__local *')}}
#endif
#endif
}
Index: clang/test/Misc/warning-flags.c
===================================================================
--- clang/test/Misc/warning-flags.c
+++ clang/test/Misc/warning-flags.c
@@ -18,7 +18,7 @@
The list of warnings below should NEVER grow. It should gradually shrink to 0.
-CHECK: Warnings without flags (74):
+CHECK: Warnings without flags (75):
CHECK-NEXT: ext_excess_initializers
CHECK-NEXT: ext_excess_initializers_in_char_array_initializer
CHECK-NEXT: ext_expected_semi_decl_list
@@ -38,6 +38,7 @@
CHECK-NEXT: warn_accessor_property_type_mismatch
CHECK-NEXT: warn_arcmt_nsalloc_realloc
CHECK-NEXT: warn_asm_label_on_auto_decl
+CHECK-NEXT: warn_bad_cxx_cast_nested_pointer_addr_space
CHECK-NEXT: warn_c_kext
CHECK-NEXT: warn_call_wrong_number_of_arguments
CHECK-NEXT: warn_case_empty_range
Index: clang/lib/Sema/SemaOverload.cpp
===================================================================
--- clang/lib/Sema/SemaOverload.cpp
+++ clang/lib/Sema/SemaOverload.cpp
@@ -3176,7 +3176,7 @@
/// 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 @@
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 different, then const is in
@@ -3258,9 +3262,9 @@
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;
}
@@ -4504,7 +4508,7 @@
// 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))
Index: clang/lib/Sema/SemaCast.cpp
===================================================================
--- clang/lib/Sema/SemaCast.cpp
+++ clang/lib/Sema/SemaCast.cpp
@@ -2311,6 +2311,24 @@
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 different type.
// Void pointers are not specified, but supported by every compiler out there.
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6679,6 +6679,9 @@
def err_bad_cxx_cast_vector_to_vector_different_size : Error<
"%select{||reinterpret_cast||C-style cast|}0 from vector %1 "
"to vector %2 of different 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">;
def err_bad_lvalue_to_rvalue_cast : Error<
"cannot cast from lvalue of type %1 to rvalue reference type %2; types are "
"not compatible">;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
