r355609 - [Sema] Change addr space diagnostics in casts to follow C++ style.
Anastasia Stulova via cfe-commits
cfe-commits at lists.llvm.org
Thu Mar 7 09:06:30 PST 2019
Author: stulova
Date: Thu Mar 7 09:06:30 2019
New Revision: 355609
URL: http://llvm.org/viewvc/llvm-project?rev=355609&view=rev
Log:
[Sema] Change addr space diagnostics in casts to follow C++ style.
This change adds a new diagnostic for mismatching address spaces
to be used for C++ casts (only enabled in C style cast for now,
the rest will follow!).
The change extends C-style cast rules to account for address spaces.
It also adds a separate function for address space cast checking that
can be used to map from a separate address space cast operator
addrspace_cast (to be added as a follow up patch).
Note, that after this change clang will no longer allows arbitrary
address space conversions in reinterpret_casts because they can lead
to accidental errors. The implicit safe conversions would still be
allowed.
Differential Revision: https://reviews.llvm.org/D58346
Added:
cfe/trunk/test/CodeGenOpenCLCXX/address-space-castoperators.cpp
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/Sema/SemaCast.cpp
cfe/trunk/test/SemaCXX/address-space-conversion.cpp
cfe/trunk/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
cfe/trunk/test/SemaOpenCL/address-spaces.cl
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=355609&r1=355608&r2=355609&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Thu Mar 7 09:06:30 2019
@@ -6271,6 +6271,10 @@ def err_bad_cxx_cast_bitfield : Error<
def err_bad_cxx_cast_qualifiers_away : Error<
"%select{const_cast|static_cast|reinterpret_cast|dynamic_cast|C-style cast|"
"functional-style cast}0 from %1 to %2 casts away qualifiers">;
+def err_bad_cxx_cast_addr_space_mismatch : Error<
+ "%select{const_cast|static_cast|reinterpret_cast|dynamic_cast|C-style cast|"
+ "functional-style cast}0 from %1 to %2 converts between mismatching address"
+ " spaces">;
def ext_bad_cxx_cast_qualifiers_away_incoherent : ExtWarn<
"ISO C++ does not allow "
"%select{const_cast|static_cast|reinterpret_cast|dynamic_cast|C-style cast|"
Modified: cfe/trunk/lib/Sema/SemaCast.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCast.cpp?rev=355609&r1=355608&r2=355609&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCast.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCast.cpp Thu Mar 7 09:06:30 2019
@@ -2212,7 +2212,15 @@ static TryCastResult TryReinterpretCast(
/*CheckObjCLifetime=*/CStyle))
SuccessResult = getCastAwayConstnessCastKind(CACK, msg);
- if (IsLValueCast) {
+ if (IsAddressSpaceConversion(SrcType, DestType)) {
+ Kind = CK_AddressSpaceConversion;
+ assert(SrcType->isPointerType() && DestType->isPointerType());
+ if (!CStyle &&
+ !DestType->getPointeeType().getQualifiers().isAddressSpaceSupersetOf(
+ SrcType->getPointeeType().getQualifiers())) {
+ SuccessResult = TC_Failed;
+ }
+ } else if (IsLValueCast) {
Kind = CK_LValueBitCast;
} else if (DestType->isObjCObjectPointerType()) {
Kind = Self.PrepareCastToObjCObjectPointer(SrcExpr);
@@ -2222,8 +2230,6 @@ static TryCastResult TryReinterpretCast(
} else {
Kind = CK_BitCast;
}
- } else if (IsAddressSpaceConversion(SrcType, DestType)) {
- Kind = CK_AddressSpaceConversion;
} else {
Kind = CK_BitCast;
}
@@ -2278,6 +2284,41 @@ static TryCastResult TryReinterpretCast(
return SuccessResult;
}
+static TryCastResult TryAddressSpaceCast(Sema &Self, ExprResult &SrcExpr,
+ QualType DestType, bool CStyle,
+ unsigned &msg) {
+ if (!Self.getLangOpts().OpenCL)
+ // FIXME: As compiler doesn't have any information about overlapping addr
+ // spaces at the moment we have to be permissive here.
+ return TC_NotApplicable;
+ // Even though the logic below is general enough and can be applied to
+ // non-OpenCL mode too, we fast-path above because no other languages
+ // define overlapping address spaces currently.
+ auto SrcType = SrcExpr.get()->getType();
+ auto SrcPtrType = SrcType->getAs<PointerType>();
+ if (!SrcPtrType)
+ return TC_NotApplicable;
+ auto DestPtrType = DestType->getAs<PointerType>();
+ if (!DestPtrType)
+ return TC_NotApplicable;
+ auto SrcPointeeType = SrcPtrType->getPointeeType();
+ auto DestPointeeType = DestPtrType->getPointeeType();
+ if (SrcPointeeType.getAddressSpace() == DestPointeeType.getAddressSpace())
+ return TC_NotApplicable;
+ if (!DestPtrType->isAddressSpaceOverlapping(*SrcPtrType)) {
+ msg = diag::err_bad_cxx_cast_addr_space_mismatch;
+ return TC_Failed;
+ }
+ auto SrcPointeeTypeWithoutAS =
+ Self.Context.removeAddrSpaceQualType(SrcPointeeType.getCanonicalType());
+ auto DestPointeeTypeWithoutAS =
+ Self.Context.removeAddrSpaceQualType(DestPointeeType.getCanonicalType());
+ return Self.Context.hasSameType(SrcPointeeTypeWithoutAS,
+ DestPointeeTypeWithoutAS)
+ ? TC_Success
+ : TC_NotApplicable;
+}
+
void CastOperation::checkAddressSpaceCast(QualType SrcType, QualType DestType) {
// In OpenCL only conversions between pointers to objects in overlapping
// addr spaces are allowed. v2.0 s6.5.5 - Generic addr space overlaps
@@ -2372,30 +2413,35 @@ void CastOperation::CheckCXXCStyleCast(b
// listed above, the interpretation that appears first in the list is used,
// even if a cast resulting from that interpretation is ill-formed.
// In plain language, this means trying a const_cast ...
+ // Note that for address space we check compatibility after const_cast.
unsigned msg = diag::err_bad_cxx_cast_generic;
TryCastResult tcr = TryConstCast(Self, SrcExpr, DestType,
- /*CStyle*/true, msg);
+ /*CStyle*/ true, msg);
if (SrcExpr.isInvalid())
return;
if (isValidCast(tcr))
Kind = CK_NoOp;
- Sema::CheckedConversionKind CCK
- = FunctionalStyle? Sema::CCK_FunctionalCast
- : Sema::CCK_CStyleCast;
+ Sema::CheckedConversionKind CCK =
+ FunctionalStyle ? Sema::CCK_FunctionalCast : Sema::CCK_CStyleCast;
if (tcr == TC_NotApplicable) {
- // ... or if that is not possible, a static_cast, ignoring const, ...
- tcr = TryStaticCast(Self, SrcExpr, DestType, CCK, OpRange,
- msg, Kind, BasePath, ListInitialization);
+ tcr = TryAddressSpaceCast(Self, SrcExpr, DestType, /*CStyle*/ true, msg);
if (SrcExpr.isInvalid())
return;
-
if (tcr == TC_NotApplicable) {
- // ... and finally a reinterpret_cast, ignoring const.
- tcr = TryReinterpretCast(Self, SrcExpr, DestType, /*CStyle*/true,
- OpRange, msg, Kind);
+ // ... or if that is not possible, a static_cast, ignoring const, ...
+ tcr = TryStaticCast(Self, SrcExpr, DestType, CCK, OpRange, msg, Kind,
+ BasePath, ListInitialization);
if (SrcExpr.isInvalid())
return;
+
+ if (tcr == TC_NotApplicable) {
+ // ... and finally a reinterpret_cast, ignoring const.
+ tcr = TryReinterpretCast(Self, SrcExpr, DestType, /*CStyle*/ true,
+ OpRange, msg, Kind);
+ if (SrcExpr.isInvalid())
+ return;
+ }
}
}
@@ -2426,8 +2472,6 @@ void CastOperation::CheckCXXCStyleCast(b
}
}
- checkAddressSpaceCast(SrcExpr.get()->getType(), DestType);
-
if (isValidCast(tcr)) {
if (Kind == CK_BitCast)
checkCastAlign();
Added: cfe/trunk/test/CodeGenOpenCLCXX/address-space-castoperators.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCLCXX/address-space-castoperators.cpp?rev=355609&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCLCXX/address-space-castoperators.cpp (added)
+++ cfe/trunk/test/CodeGenOpenCLCXX/address-space-castoperators.cpp Thu Mar 7 09:06:30 2019
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -O0 -o - | FileCheck %s
+
+void test_reinterpret_cast(){
+__private float x;
+__private float& y = x;
+// We don't need bitcast to cast pointer type and
+// address space at the same time.
+//CHECK: addrspacecast float* %x to i32 addrspace(4)*
+//CHECK: [[REG:%[0-9]+]] = load float*, float** %y
+//CHECK: addrspacecast float* [[REG]] to i32 addrspace(4)*
+//CHECK-NOT: bitcast
+__generic int& rc1 = reinterpret_cast<__generic int&>(x);
+__generic int& rc2 = reinterpret_cast<__generic int&>(y);
+}
Modified: cfe/trunk/test/SemaCXX/address-space-conversion.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCXX/address-space-conversion.cpp?rev=355609&r1=355608&r2=355609&view=diff
==============================================================================
--- cfe/trunk/test/SemaCXX/address-space-conversion.cpp (original)
+++ cfe/trunk/test/SemaCXX/address-space-conversion.cpp Thu Mar 7 09:06:30 2019
@@ -131,24 +131,24 @@ void test_dynamic_cast(A_ptr ap, A_ptr_1
void test_reinterpret_cast(void_ptr vp, void_ptr_1 vp1, void_ptr_2 vp2,
A_ptr ap, A_ptr_1 ap1, A_ptr_2 ap2,
B_ptr bp, B_ptr_1 bp1, B_ptr_2 bp2,
- const void __attribute__((address_space(1))) *cvp1) {
- // reinterpret_cast can be used to cast to a different address space.
- (void)reinterpret_cast<A_ptr>(ap1);
- (void)reinterpret_cast<A_ptr>(ap2);
+ const void __attribute__((address_space(1))) * cvp1) {
+ // reinterpret_cast can't be used to cast to a different address space unless they are matching (i.e. overlapping).
+ (void)reinterpret_cast<A_ptr>(ap1); // expected-error{{reinterpret_cast from 'A_ptr_1' (aka '__attribute__((address_space(1))) A *') to 'A_ptr' (aka 'A *') is not allowed}}
+ (void)reinterpret_cast<A_ptr>(ap2); // expected-error{{reinterpret_cast from 'A_ptr_2' (aka '__attribute__((address_space(2))) A *') to 'A_ptr' (aka 'A *') is not allowed}}
(void)reinterpret_cast<A_ptr>(bp);
- (void)reinterpret_cast<A_ptr>(bp1);
- (void)reinterpret_cast<A_ptr>(bp2);
+ (void)reinterpret_cast<A_ptr>(bp1); // expected-error{{reinterpret_cast from 'B_ptr_1' (aka '__attribute__((address_space(1))) B *') to 'A_ptr' (aka 'A *') is not allowed}}
+ (void)reinterpret_cast<A_ptr>(bp2); // expected-error{{reinterpret_cast from 'B_ptr_2' (aka '__attribute__((address_space(2))) B *') to 'A_ptr' (aka 'A *') is not allowed}}
(void)reinterpret_cast<A_ptr>(vp);
- (void)reinterpret_cast<A_ptr>(vp1);
- (void)reinterpret_cast<A_ptr>(vp2);
- (void)reinterpret_cast<A_ptr_1>(ap);
- (void)reinterpret_cast<A_ptr_1>(ap2);
- (void)reinterpret_cast<A_ptr_1>(bp);
+ (void)reinterpret_cast<A_ptr>(vp1); // expected-error{{reinterpret_cast from 'void_ptr_1' (aka '__attribute__((address_space(1))) void *') to 'A_ptr' (aka 'A *') is not allowed}}
+ (void)reinterpret_cast<A_ptr>(vp2); // expected-error{{reinterpret_cast from 'void_ptr_2' (aka '__attribute__((address_space(2))) void *') to 'A_ptr' (aka 'A *') is not allowed}}
+ (void)reinterpret_cast<A_ptr_1>(ap); // expected-error{{reinterpret_cast from 'A_ptr' (aka 'A *') to 'A_ptr_1' (aka '__attribute__((address_space(1))) A *') is not allowed}}
+ (void)reinterpret_cast<A_ptr_1>(ap2); // expected-error{{reinterpret_cast from 'A_ptr_2' (aka '__attribute__((address_space(2))) A *') to 'A_ptr_1' (aka '__attribute__((address_space(1))) A *') is not allowed}}
+ (void)reinterpret_cast<A_ptr_1>(bp); // expected-error{{reinterpret_cast from 'B_ptr' (aka 'B *') to 'A_ptr_1' (aka '__attribute__((address_space(1))) A *') is not allowed}}
(void)reinterpret_cast<A_ptr_1>(bp1);
- (void)reinterpret_cast<A_ptr_1>(bp2);
- (void)reinterpret_cast<A_ptr_1>(vp);
+ (void)reinterpret_cast<A_ptr_1>(bp2); // expected-error{{reinterpret_cast from 'B_ptr_2' (aka '__attribute__((address_space(2))) B *') to 'A_ptr_1' (aka '__attribute__((address_space(1))) A *') is not allowed}}
+ (void)reinterpret_cast<A_ptr_1>(vp); // expected-error{{reinterpret_cast from 'void_ptr' (aka 'void *') to 'A_ptr_1' (aka '__attribute__((address_space(1))) A *') is not allowed}}
(void)reinterpret_cast<A_ptr_1>(vp1);
- (void)reinterpret_cast<A_ptr_1>(vp2);
+ (void)reinterpret_cast<A_ptr_1>(vp2); // expected-error{{reinterpret_cast from 'void_ptr_2' (aka '__attribute__((address_space(2))) void *') to 'A_ptr_1' (aka '__attribute__((address_space(1))) A *') is not allowed}}
// ... but don't try to cast away constness!
(void)reinterpret_cast<A_ptr_2>(cvp1); // expected-error{{casts away qualifiers}}
Modified: cfe/trunk/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl?rev=355609&r1=355608&r2=355609&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl (original)
+++ cfe/trunk/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl Thu Mar 7 09:06:30 2019
@@ -129,27 +129,47 @@ void test_conversion(__global int *arg_g
AS int *var_cast1 = (AS int *)arg_glob;
#ifdef CONSTANT
-// expected-error at -2{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -3{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error at -5{{C-style cast from '__global int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
#endif
AS int *var_cast2 = (AS int *)arg_loc;
#ifndef GENERIC
-// expected-error-re at -2{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re at -3{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re at -5{{C-style cast from '__local int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+#endif
#endif
AS int *var_cast3 = (AS int *)arg_const;
#ifndef CONSTANT
-// expected-error-re at -2{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re at -3{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
+#else
+// expected-error-re at -5{{C-style cast from '__constant int *' to '__{{global|generic}} int *' converts between mismatching address spaces}}
+#endif
#endif
AS int *var_cast4 = (AS int *)arg_priv;
#ifndef GENERIC
-// expected-error-re at -2{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re at -3{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re at -5{{C-style cast from 'int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+#endif
#endif
AS int *var_cast5 = (AS int *)arg_gen;
#ifdef CONSTANT
-// expected-error at -2{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -3{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error at -5{{C-style cast from '__generic int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
#endif
AS int *var_impl;
@@ -200,27 +220,47 @@ void test_conversion(__global int *arg_g
var_cast1 = (AS int *)arg_glob;
#ifdef CONSTANT
-// expected-error at -2{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -3{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error at -5{{C-style cast from '__global int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
#endif
var_cast2 = (AS int *)arg_loc;
#ifndef GENERIC
-// expected-error-re at -2{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re at -3{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re at -5{{C-style cast from '__local int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+#endif
#endif
var_cast3 = (AS int *)arg_const;
#ifndef CONSTANT
-// expected-error-re at -2{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re at -3{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
+#else
+// expected-error-re at -5{{C-style cast from '__constant int *' to '__{{global|generic}} int *' converts between mismatching address spaces}}
+#endif
#endif
var_cast4 = (AS int *)arg_priv;
#ifndef GENERIC
-// expected-error-re at -2{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re at -3{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re at -5{{C-style cast from 'int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+#endif
#endif
var_cast5 = (AS int *)arg_gen;
#ifdef CONSTANT
-// expected-error at -2{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -3{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error at -5{{C-style cast from '__generic int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
#endif
AS int *var_cmp;
Modified: cfe/trunk/test/SemaOpenCL/address-spaces.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/address-spaces.cl?rev=355609&r1=355608&r2=355609&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/address-spaces.cl (original)
+++ cfe/trunk/test/SemaOpenCL/address-spaces.cl Thu Mar 7 09:06:30 2019
@@ -26,24 +26,96 @@ __kernel void foo(__global int *gip) {
}
void explicit_cast(__global int *g, __local int *l, __constant int *c, __private int *p, const __constant int *cc) {
- g = (__global int *)l; // expected-error {{casting '__local int *' to type '__global int *' changes address space of pointer}}
- g = (__global int *)c; // expected-error {{casting '__constant int *' to type '__global int *' changes address space of pointer}}
- g = (__global int *)cc; // expected-error {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}}
- g = (__global int *)p; // expected-error {{casting 'int *' to type '__global int *' changes address space of pointer}}
-
- l = (__local int *)g; // expected-error {{casting '__global int *' to type '__local int *' changes address space of pointer}}
- l = (__local int *)c; // expected-error {{casting '__constant int *' to type '__local int *' changes address space of pointer}}
- l = (__local int *)cc; // expected-error {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}}
- l = (__local int *)p; // expected-error {{casting 'int *' to type '__local int *' changes address space of pointer}}
-
- c = (__constant int *)g; // expected-error {{casting '__global int *' to type '__constant int *' changes address space of pointer}}
- c = (__constant int *)l; // expected-error {{casting '__local int *' to type '__constant int *' changes address space of pointer}}
- c = (__constant int *)p; // expected-error {{casting 'int *' to type '__constant int *' changes address space of pointer}}
-
- p = (__private int *)g; // expected-error {{casting '__global int *' to type 'int *' changes address space of pointer}}
- p = (__private int *)l; // expected-error {{casting '__local int *' to type 'int *' changes address space of pointer}}
- p = (__private int *)c; // expected-error {{casting '__constant int *' to type 'int *' changes address space of pointer}}
- p = (__private int *)cc; // expected-error {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
+ g = (__global int *)l;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting '__local int *' to type '__global int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from '__local int *' to '__global int *' converts between mismatching address spaces}}
+#endif
+ g = (__global int *)c;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting '__constant int *' to type '__global int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from '__constant int *' to '__global int *' converts between mismatching address spaces}}
+#endif
+ g = (__global int *)cc;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from 'const __constant int *' to '__global int *' converts between mismatching address spaces}}
+#endif
+ g = (__global int *)p;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting 'int *' to type '__global int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from 'int *' to '__global int *' converts between mismatching address spaces}}
+#endif
+ l = (__local int *)g;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting '__global int *' to type '__local int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from '__global int *' to '__local int *' converts between mismatching address spaces}}
+#endif
+ l = (__local int *)c;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting '__constant int *' to type '__local int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from '__constant int *' to '__local int *' converts between mismatching address spaces}}
+#endif
+ l = (__local int *)cc;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from 'const __constant int *' to '__local int *' converts between mismatching address spaces}}
+#endif
+ l = (__local int *)p;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting 'int *' to type '__local int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from 'int *' to '__local int *' converts between mismatching address spaces}}
+#endif
+ c = (__constant int *)g;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting '__global int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from '__global int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
+ c = (__constant int *)l;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting '__local int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from '__local int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
+ c = (__constant int *)p;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting 'int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from 'int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
+ p = (__private int *)g;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting '__global int *' to type 'int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from '__global int *' to 'int *' converts between mismatching address spaces}}
+#endif
+ p = (__private int *)l;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting '__local int *' to type 'int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from '__local int *' to 'int *' converts between mismatching address spaces}}
+#endif
+ p = (__private int *)c;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting '__constant int *' to type 'int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from '__constant int *' to 'int *' converts between mismatching address spaces}}
+#endif
+ p = (__private int *)cc;
+#if !__OPENCL_CPP_VERSION__
+// expected-error at -2 {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
+#else
+// expected-error at -4 {{C-style cast from 'const __constant int *' to 'int *' converts between mismatching address spaces}}
+#endif
}
void ok_explicit_casts(__global int *g, __global int *g2, __local int *l, __local int *l2, __private int *p, __private int *p2) {
More information about the cfe-commits
mailing list