This revision was automatically updated to reflect the committed changes.
Closed by commit rG6064f426a183: [OpenCL] Restrict addr space conversions in
nested pointers (authored by Anastasia).
Herald added a project: clang.
Changed prior to commit:
https://reviews.llvm.org/D73360?vs=242312&id=243132#toc
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D73360/new/
https://reviews.llvm.org/D73360
Files:
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-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/lib/Sema/SemaOverload.cpp
===================================================================
--- clang/lib/Sema/SemaOverload.cpp
+++ clang/lib/Sema/SemaOverload.cpp
@@ -3180,7 +3180,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();
@@ -3217,11 +3217,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
@@ -3262,9 +3266,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;
}
@@ -4508,7 +4512,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
@@ -6760,6 +6760,10 @@
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">,
+ 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">;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
