summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--include/clang/Basic/DiagnosticSemaKinds.td13
-rw-r--r--include/clang/Sema/Sema.h6
-rw-r--r--lib/Sema/SemaCast.cpp46
-rw-r--r--lib/Sema/SemaExpr.cpp25
-rw-r--r--test/CodeGenOpenCL/numbered-address-space.cl13
-rw-r--r--test/SemaOpenCL/address-spaces.cl100
-rw-r--r--test/SemaOpenCL/event_t_overload.cl2
-rw-r--r--test/SemaOpenCL/numbered-address-space.cl2
-rw-r--r--test/SemaOpenCL/queue_t_overload.cl2
9 files changed, 179 insertions, 30 deletions
diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index 9205405853..ba376d5f10 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -7005,6 +7005,19 @@ def err_typecheck_incompatible_address_space : Error<
"sending to parameter of different type}0,1"
"|%diff{casting $ to type $|casting between types}0,1}2"
" changes address space of pointer">;
+def err_typecheck_incompatible_nested_address_space : Error<
+ "%select{%diff{assigning $ to $|assigning to different types}1,0"
+ "|%diff{passing $ to parameter of type $|"
+ "passing to parameter of different type}0,1"
+ "|%diff{returning $ from a function with result type $|"
+ "returning from function with different return type}0,1"
+ "|%diff{converting $ to type $|converting between types}0,1"
+ "|%diff{initializing $ with an expression of type $|"
+ "initializing with expression of different type}0,1"
+ "|%diff{sending $ to parameter of type $|"
+ "sending to parameter of different type}0,1"
+ "|%diff{casting $ to type $|casting between types}0,1}2"
+ " changes address space of nested pointer">;
def err_typecheck_incompatible_ownership : Error<
"%select{%diff{assigning $ to $|assigning to different types}1,0"
"|%diff{passing $ to parameter of type $|"
diff --git a/include/clang/Sema/Sema.h b/include/clang/Sema/Sema.h
index 0ce34251e9..5f29fb8244 100644
--- a/include/clang/Sema/Sema.h
+++ b/include/clang/Sema/Sema.h
@@ -9726,6 +9726,12 @@ public:
/// like address spaces.
IncompatiblePointerDiscardsQualifiers,
+ /// IncompatibleNestedPointerAddressSpaceMismatch - The assignment
+ /// changes address spaces in nested pointer types which is not allowed.
+ /// For instance, converting __private int ** to __generic int ** is
+ /// illegal even though __private could be converted to __generic.
+ IncompatibleNestedPointerAddressSpaceMismatch,
+
/// IncompatibleNestedPointerQualifiers - The assignment is between two
/// nested pointer types, and the qualifiers other than the first two
/// levels differ e.g. char ** -> const char **, but we accept them as an
diff --git a/lib/Sema/SemaCast.cpp b/lib/Sema/SemaCast.cpp
index 21a4d107ba..0958943f9e 100644
--- a/lib/Sema/SemaCast.cpp
+++ b/lib/Sema/SemaCast.cpp
@@ -2323,19 +2323,41 @@ 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
// with any named one, except for constant.
+
+ // Converting the top level pointee addrspace is permitted for compatible
+ // addrspaces (such as 'generic int *' to 'local int *' or vice versa), but
+ // if any of the nested pointee addrspaces differ, we emit a warning
+ // regardless of addrspace compatibility. This makes
+ // local int ** p;
+ // return (generic int **) p;
+ // warn even though local -> generic is permitted.
if (Self.getLangOpts().OpenCL) {
- auto SrcPtrType = SrcType->getAs<PointerType>();
- if (!SrcPtrType)
- return;
- auto DestPtrType = DestType->getAs<PointerType>();
- if (!DestPtrType)
- return;
- if (!DestPtrType->isAddressSpaceOverlapping(*SrcPtrType)) {
- Self.Diag(OpRange.getBegin(),
- diag::err_typecheck_incompatible_address_space)
- << SrcType << DestType << Sema::AA_Casting
- << SrcExpr.get()->getSourceRange();
- SrcExpr = ExprError();
+ const Type *DestPtr, *SrcPtr;
+ bool Nested = false;
+ unsigned DiagID = diag::err_typecheck_incompatible_address_space;
+ DestPtr = Self.getASTContext().getCanonicalType(DestType.getTypePtr()),
+ SrcPtr = Self.getASTContext().getCanonicalType(SrcType.getTypePtr());
+
+ while (isa<PointerType>(DestPtr) && isa<PointerType>(SrcPtr)) {
+ const PointerType *DestPPtr = cast<PointerType>(DestPtr);
+ const PointerType *SrcPPtr = cast<PointerType>(SrcPtr);
+ QualType DestPPointee = DestPPtr->getPointeeType();
+ QualType SrcPPointee = SrcPPtr->getPointeeType();
+ if (Nested ? DestPPointee.getAddressSpace() !=
+ SrcPPointee.getAddressSpace()
+ : !DestPPtr->isAddressSpaceOverlapping(*SrcPPtr)) {
+ Self.Diag(OpRange.getBegin(), DiagID)
+ << SrcType << DestType << Sema::AA_Casting
+ << SrcExpr.get()->getSourceRange();
+ if (!Nested)
+ SrcExpr = ExprError();
+ return;
+ }
+
+ DestPtr = DestPPtr->getPointeeType().getTypePtr();
+ SrcPtr = SrcPPtr->getPointeeType().getTypePtr();
+ Nested = true;
+ DiagID = diag::ext_nested_pointer_qualifier_mismatch;
}
}
}
diff --git a/lib/Sema/SemaExpr.cpp b/lib/Sema/SemaExpr.cpp
index d8c0896147..8bb306e66d 100644
--- a/lib/Sema/SemaExpr.cpp
+++ b/lib/Sema/SemaExpr.cpp
@@ -7724,9 +7724,9 @@ checkPointerTypesForAssignment(Sema &S, QualType LHSType, QualType RHSType) {
}
if (!lhq.compatiblyIncludes(rhq)) {
- // Treat address-space mismatches as fatal. TODO: address subspaces
+ // Treat address-space mismatches as fatal.
if (!lhq.isAddressSpaceSupersetOf(rhq))
- ConvTy = Sema::IncompatiblePointerDiscardsQualifiers;
+ return Sema::IncompatiblePointerDiscardsQualifiers;
// It's okay to add or remove GC or lifetime qualifiers when converting to
// and from void*.
@@ -7799,8 +7799,22 @@ checkPointerTypesForAssignment(Sema &S, QualType LHSType, QualType RHSType) {
// level of indirection, this must be the issue.
if (isa<PointerType>(lhptee) && isa<PointerType>(rhptee)) {
do {
- lhptee = cast<PointerType>(lhptee)->getPointeeType().getTypePtr();
- rhptee = cast<PointerType>(rhptee)->getPointeeType().getTypePtr();
+ std::tie(lhptee, lhq) =
+ cast<PointerType>(lhptee)->getPointeeType().split().asPair();
+ std::tie(rhptee, rhq) =
+ cast<PointerType>(rhptee)->getPointeeType().split().asPair();
+
+ // Inconsistent address spaces at this point is invalid, even if the
+ // address spaces would be compatible.
+ // FIXME: This doesn't catch address space mismatches for pointers of
+ // different nesting levels, like:
+ // __local int *** a;
+ // int ** b = a;
+ // It's not clear how to actually determine when such pointers are
+ // invalidly incompatible.
+ if (lhq.getAddressSpace() != rhq.getAddressSpace())
+ return Sema::IncompatibleNestedPointerAddressSpaceMismatch;
+
} while (isa<PointerType>(lhptee) && isa<PointerType>(rhptee));
if (lhptee == rhptee)
@@ -14213,6 +14227,9 @@ bool Sema::DiagnoseAssignmentResult(AssignConvertType ConvTy,
case IncompatibleNestedPointerQualifiers:
DiagKind = diag::ext_nested_pointer_qualifier_mismatch;
break;
+ case IncompatibleNestedPointerAddressSpaceMismatch:
+ DiagKind = diag::err_typecheck_incompatible_nested_address_space;
+ break;
case IntToBlockPointer:
DiagKind = diag::err_int_to_block_pointer;
break;
diff --git a/test/CodeGenOpenCL/numbered-address-space.cl b/test/CodeGenOpenCL/numbered-address-space.cl
index dbaba87476..296923ee50 100644
--- a/test/CodeGenOpenCL/numbered-address-space.cl
+++ b/test/CodeGenOpenCL/numbered-address-space.cl
@@ -11,12 +11,6 @@ void test_numbered_as_to_generic(__attribute__((address_space(42))) int *arbitar
*generic_ptr = 4;
}
-// CHECK-LABEL: @test_numbered_as_to_builtin(
-// CHECK: addrspacecast i32 addrspace(42)* %0 to float addrspace(3)*
-void test_numbered_as_to_builtin(__attribute__((address_space(42))) int *arbitary_numbered_ptr, float src) {
- volatile float result = __builtin_amdgcn_ds_fmaxf(arbitary_numbered_ptr, src, 0, 0, false);
-}
-
// CHECK-LABEL: @test_generic_as_to_builtin_parameter_explicit_cast(
// CHECK: addrspacecast i32 addrspace(3)* %0 to i32*
void test_generic_as_to_builtin_parameter_explicit_cast(__local int *local_ptr, float src) {
@@ -25,10 +19,7 @@ void test_generic_as_to_builtin_parameter_explicit_cast(__local int *local_ptr,
}
// CHECK-LABEL: @test_generic_as_to_builtin_parameter_implicit_cast(
-// CHECK: addrspacecast i32* %2 to float addrspace(3)*
+// CHECK: bitcast i32 addrspace(3)* %0 to float addrspace(3)*
void test_generic_as_to_builtin_parameter_implicit_cast(__local int *local_ptr, float src) {
- generic int* generic_ptr = local_ptr;
-
- volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false);
+ volatile float result = __builtin_amdgcn_ds_fmaxf(local_ptr, src, 0, 0, false);
}
-
diff --git a/test/SemaOpenCL/address-spaces.cl b/test/SemaOpenCL/address-spaces.cl
index f5fa43d922..13e9442b95 100644
--- a/test/SemaOpenCL/address-spaces.cl
+++ b/test/SemaOpenCL/address-spaces.cl
@@ -124,6 +124,106 @@ void ok_explicit_casts(__global int *g, __global int *g2, __local int *l, __loca
p = (__private int *)p2;
}
+#if !__OPENCL_CPP_VERSION__
+void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {
+ g = gg; // expected-error {{assigning '__global int **' to '__global int *' changes address space of pointer}}
+ g = l; // expected-error {{assigning '__local int *' to '__global int *' changes address space of pointer}}
+ g = ll; // expected-error {{assigning '__local int **' to '__global int *' changes address space of pointer}}
+ g = gg_f; // expected-error {{assigning '__global float **' to '__global int *' changes address space of pointer}}
+ g = (__global int *)gg_f; // expected-error {{casting '__global float **' to type '__global int *' changes address space of pointer}}
+
+ gg = g; // expected-error {{assigning '__global int *' to '__global int **' changes address space of pointer}}
+ gg = l; // expected-error {{assigning '__local int *' to '__global int **' changes address space of pointer}}
+ gg = ll; // expected-error {{assigning '__local int **' to '__global int **' changes address space of nested pointer}}
+ gg = gg_f; // expected-warning {{incompatible pointer types assigning to '__global int **' from '__global float **'}}
+ gg = (__global int * __private *)gg_f;
+
+ l = g; // expected-error {{assigning '__global int *' to '__local int *' changes address space of pointer}}
+ l = gg; // expected-error {{assigning '__global int **' to '__local int *' changes address space of pointer}}
+ l = ll; // expected-error {{assigning '__local int **' to '__local int *' changes address space of pointer}}
+ l = gg_f; // expected-error {{assigning '__global float **' to '__local int *' changes address space of pointer}}
+ l = (__local int *)gg_f; // expected-error {{casting '__global float **' to type '__local int *' changes address space of pointer}}
+
+ ll = g; // expected-error {{assigning '__global int *' to '__local int **' changes address space of pointer}}
+ ll = gg; // expected-error {{assigning '__global int **' to '__local int **' changes address space of nested pointer}}
+ ll = l; // expected-error {{assigning '__local int *' to '__local int **' changes address space of pointer}}
+ ll = gg_f; // expected-error {{assigning '__global float **' to '__local int **' changes address space of nested pointer}}
+ ll = (__local int * __private *)gg_f; // expected-warning {{casting '__global float **' to type '__local int **' discards qualifiers in nested pointer types}}
+
+ gg_f = g; // expected-error {{assigning '__global int *' to '__global float **' changes address space of pointer}}
+ gg_f = gg; // expected-warning {{incompatible pointer types assigning to '__global float **' from '__global int **'}}
+ gg_f = l; // expected-error {{assigning '__local int *' to '__global float **' changes address space of pointer}}
+ gg_f = ll; // expected-error {{assigning '__local int **' to '__global float **' changes address space of nested pointer}}
+ gg_f = (__global float * __private *)gg;
+
+ // FIXME: This doesn't seem right. This should be an error, not a warning.
+ __local int * __global * __private * lll;
+ lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global **' from '__global int **'}}
+
+ typedef __local int * l_t;
+ typedef __global int * g_t;
+ __private l_t * pl;
+ __private g_t * pg;
+ gg = pl; // expected-error {{assigning 'l_t *' (aka '__local int **') to '__global int **' changes address space of nested pointer}}
+ pl = gg; // expected-error {{assigning '__global int **' to 'l_t *' (aka '__local int **') changes address space of nested pointer}}
+ gg = pg;
+ pg = gg;
+ pg = pl; // expected-error {{assigning 'l_t *' (aka '__local int **') to 'g_t *' (aka '__global int **') changes address space of nested pointer}}
+ pl = pg; // expected-error {{assigning 'g_t *' (aka '__global int **') to 'l_t *' (aka '__local int **') changes address space of nested pointer}}
+
+ ll = (__local int * __private *)(void *)gg;
+ void *vp = ll;
+}
+#else
+void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {
+ g = gg; // expected-error {{assigning to '__global int *' from incompatible type '__global int **'}}
+ g = l; // expected-error {{assigning to '__global int *' from incompatible type '__local int *'}}
+ g = ll; // expected-error {{assigning to '__global int *' from incompatible type '__local int **'}}
+ g = gg_f; // expected-error {{assigning to '__global int *' from incompatible type '__global float **'}}
+ g = (__global int *)gg_f; // expected-error {{C-style cast from '__global float **' to '__global int *' converts between mismatching address spaces}}
+
+ gg = g; // expected-error {{assigning to '__global int **' from incompatible type '__global int *'; take the address with &}}
+ gg = l; // expected-error {{assigning to '__global int **' from incompatible type '__local int *'}}
+ gg = ll; // expected-error {{assigning to '__global int **' from incompatible type '__local int **'}}
+ gg = gg_f; // expected-error {{assigning to '__global int **' from incompatible type '__global float **'}}
+ gg = (__global int * __private *)gg_f;
+
+ l = g; // expected-error {{assigning to '__local int *' from incompatible type '__global int *'}}
+ l = gg; // expected-error {{assigning to '__local int *' from incompatible type '__global int **'}}
+ l = ll; // expected-error {{assigning to '__local int *' from incompatible type '__local int **'}}
+ l = gg_f; // expected-error {{assigning to '__local int *' from incompatible type '__global float **'}}
+ l = (__local int *)gg_f; // expected-error {{C-style cast from '__global float **' to '__local int *' converts between mismatching address spaces}}
+
+ ll = g; // expected-error {{assigning to '__local int **' from incompatible type '__global int *'}}
+ ll = gg; // expected-error {{assigning to '__local int **' from incompatible type '__global int **'}}
+ ll = l; // expected-error {{assigning to '__local int **' from incompatible type '__local int *'; take the address with &}}
+ ll = gg_f; // expected-error {{assigning to '__local int **' from incompatible type '__global float **'}}
+ // 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;
+
+ gg_f = g; // expected-error {{assigning to '__global float **' from incompatible type '__global int *'}}
+ gg_f = gg; // expected-error {{assigning to '__global float **' from incompatible type '__global int **'}}
+ gg_f = l; // expected-error {{assigning to '__global float **' from incompatible type '__local int *'}}
+ gg_f = ll; // expected-error {{assigning to '__global float **' from incompatible type '__local int **'}}
+ gg_f = (__global float * __private *)gg;
+
+ typedef __local int * l_t;
+ typedef __global int * g_t;
+ __private l_t * pl;
+ __private g_t * pg;
+ gg = pl; // expected-error {{assigning to '__global int **' from incompatible type 'l_t *' (aka '__local int **')}}
+ pl = gg; // expected-error {{assigning to 'l_t *' (aka '__local int **') from incompatible type '__global int **'}}
+ gg = pg;
+ pg = gg;
+ pg = pl; // expected-error {{assigning to 'g_t *' (aka '__global int **') from incompatible type 'l_t *' (aka '__local int **')}}
+ pl = pg; // expected-error {{assigning to 'l_t *' (aka '__local int **') from incompatible type 'g_t *' (aka '__global int **')}}
+
+ ll = (__local int * __private *)(void *)gg;
+ void *vp = ll;
+}
+#endif
+
__private int func_return_priv(void); //expected-error {{return value cannot be qualified with address space}}
__global int func_return_global(void); //expected-error {{return value cannot be qualified with address space}}
__local int func_return_local(void); //expected-error {{return value cannot be qualified with address space}}
diff --git a/test/SemaOpenCL/event_t_overload.cl b/test/SemaOpenCL/event_t_overload.cl
index 0bc67c13c7..eaf05a0489 100644
--- a/test/SemaOpenCL/event_t_overload.cl
+++ b/test/SemaOpenCL/event_t_overload.cl
@@ -7,5 +7,5 @@ void kernel ker(__local char *src1, __local float *src2, __global int *src3) {
event_t evt;
foo(evt, src1);
foo(0, src2);
- foo(evt, src3); // expected-error {{call to 'foo' is ambiguous}}
+ foo(evt, src3); // expected-error {{no matching function for call to 'foo'}}
}
diff --git a/test/SemaOpenCL/numbered-address-space.cl b/test/SemaOpenCL/numbered-address-space.cl
index 423d03274c..8a24b3ce97 100644
--- a/test/SemaOpenCL/numbered-address-space.cl
+++ b/test/SemaOpenCL/numbered-address-space.cl
@@ -26,6 +26,6 @@ void test_generic_as_to_builtin_parameter_explicit_cast_numeric(__attribute__((a
void test_generic_as_to_builtin_parameterimplicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) {
generic int* generic_ptr = as3_ptr;
- volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__local float *'}}
+ volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-error {{passing '__generic int *' to parameter of type '__local float *' changes address space of pointer}}
}
diff --git a/test/SemaOpenCL/queue_t_overload.cl b/test/SemaOpenCL/queue_t_overload.cl
index 0048895ac8..bc3b241bbe 100644
--- a/test/SemaOpenCL/queue_t_overload.cl
+++ b/test/SemaOpenCL/queue_t_overload.cl
@@ -7,6 +7,6 @@ void kernel ker(__local char *src1, __local float *src2, __global int *src3) {
queue_t q;
foo(q, src1);
foo(0, src2);
- foo(q, src3); // expected-error {{call to 'foo' is ambiguous}}
+ foo(q, src3); // expected-error {{no matching function for call to 'foo'}}
foo(1, src3); // expected-error {{no matching function for call to 'foo'}}
}
View Lorry Controller Status