ebevhan created this revision.
ebevhan added reviewers: Anastasia, rjmccall.
Herald added subscribers: cfe-commits, jdoerfert.
Herald added a project: clang.
The semantics for converting nested pointers between address
spaces are not very well defined. Some conversions which do not
really carry any meaning only produce warnings, and in some cases
warnings hide invalid conversions, such as 'global int*' to
'local float*'!
This patch changes the logic in checkPointerTypesForAssignment
and checkAddressSpaceCast to fail properly on conversions that
should definitely not be permitted. We also dig deeper into the
pointer types and fail on conversions where the address space
in a nested pointer changes, regardless of whether the address
space is compatible with the corresponding pointer nesting level
on the destination type.
See https://bugs.llvm.org/show_bug.cgi?id=39674
Repository:
rC Clang
https://reviews.llvm.org/D58236
Files:
lib/Sema/SemaCast.cpp
lib/Sema/SemaExpr.cpp
test/CodeGenOpenCL/numbered-address-space.cl
test/SemaOpenCL/address-spaces.cl
test/SemaOpenCL/event_t_overload.cl
test/SemaOpenCL/numbered-address-space.cl
test/SemaOpenCL/queue_t_overload.cl
Index: test/SemaOpenCL/queue_t_overload.cl
===================================================================
--- test/SemaOpenCL/queue_t_overload.cl
+++ test/SemaOpenCL/queue_t_overload.cl
@@ -7,6 +7,6 @@
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'}}
}
Index: test/SemaOpenCL/numbered-address-space.cl
===================================================================
--- test/SemaOpenCL/numbered-address-space.cl
+++ test/SemaOpenCL/numbered-address-space.cl
@@ -26,6 +26,6 @@
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}}
}
Index: test/SemaOpenCL/event_t_overload.cl
===================================================================
--- test/SemaOpenCL/event_t_overload.cl
+++ test/SemaOpenCL/event_t_overload.cl
@@ -7,5 +7,5 @@
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'}}
}
Index: test/SemaOpenCL/address-spaces.cl
===================================================================
--- test/SemaOpenCL/address-spaces.cl
+++ test/SemaOpenCL/address-spaces.cl
@@ -52,6 +52,76 @@
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 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 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 pointer}}
+ ll = (__local int * __private *)gg_f; // expected-error {{casting '__global float **' to type '__local int **' changes address space of pointer}}
+
+ 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 pointer}}
+ gg_f = (__global float * __private *)gg;
+
+ // FIXME: This doesn't seem right.
+ __local int * __global * __private * lll;
+ lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global **' from '__global int **'}}
+}
+#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 {{casting '__global float **' to type '__global int *' changes address space of pointer}}
+
+ 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 {{casting '__global float **' to type '__local int *' changes address space of pointer}}
+
+ 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 **'}}
+ ll = (__local int * __private *)gg_f; // expected-error {{casting '__global float **' to type '__local int **' changes address space of pointer}}
+
+ 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;
+}
+#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}}
Index: test/CodeGenOpenCL/numbered-address-space.cl
===================================================================
--- test/CodeGenOpenCL/numbered-address-space.cl
+++ test/CodeGenOpenCL/numbered-address-space.cl
@@ -11,12 +11,6 @@
*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 @@
}
// 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);
}
-
Index: lib/Sema/SemaExpr.cpp
===================================================================
--- lib/Sema/SemaExpr.cpp
+++ lib/Sema/SemaExpr.cpp
@@ -7705,7 +7705,7 @@
if (!lhq.compatiblyIncludes(rhq)) {
// Treat address-space mismatches as fatal. TODO: address subspaces
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*.
@@ -7778,6 +7778,12 @@
// level of indirection, this must be the issue.
if (isa<PointerType>(lhptee) && isa<PointerType>(rhptee)) {
do {
+ // Inconsistent address spaces at this point is invalid, even if the
+ // address spaces would be compatible.
+ if (cast<PointerType>(lhptee)->getPointeeType().getAddressSpace() !=
+ cast<PointerType>(rhptee)->getPointeeType().getAddressSpace())
+ return Sema::IncompatibleNestedPointerQualifiers;
+
lhptee = cast<PointerType>(lhptee)->getPointeeType().getTypePtr();
rhptee = cast<PointerType>(rhptee)->getPointeeType().getTypePtr();
} while (isa<PointerType>(lhptee) && isa<PointerType>(rhptee));
@@ -14179,9 +14185,34 @@
return false;
DiagKind = diag::ext_typecheck_convert_discards_qualifiers;
break;
- case IncompatibleNestedPointerQualifiers:
+ case IncompatibleNestedPointerQualifiers: {
+ // Perform array-to-pointer decay if necessary.
+ if (SrcType->isArrayType()) SrcType = Context.getArrayDecayedType(SrcType);
+
+ // Address space qualifier mismatch in nested pointers is always an error.
+ // Strip the pointers until we either hit a mismatch or hit the end of the
+ // nesting.
+ // Use pointee type here since this should only care about nested
+ // qualifiers.
+ // XXX: Canonical types?
+ const Type *SrcPtr = SrcType->getPointeeType().getTypePtr();
+ const Type *DstPtr = DstType->getPointeeType().getTypePtr();
+
DiagKind = diag::ext_nested_pointer_qualifier_mismatch;
+ while (SrcPtr->isPointerType() && DstPtr->isPointerType()) {
+ if (SrcPtr->getPointeeType().getAddressSpace() !=
+ DstPtr->getPointeeType().getAddressSpace()) {
+ // XXX: Should this be a different variation of the error, like
+ // 'changes address space in nested pointer qualifiers'?
+ DiagKind = diag::err_typecheck_incompatible_address_space;
+ isInvalid = true;
+ break;
+ }
+ SrcPtr = SrcPtr->getPointeeType().getTypePtr();
+ DstPtr = DstPtr->getPointeeType().getTypePtr();
+ }
break;
+ }
case IntToBlockPointer:
DiagKind = diag::err_int_to_block_pointer;
break;
Index: lib/Sema/SemaCast.cpp
===================================================================
--- lib/Sema/SemaCast.cpp
+++ lib/Sema/SemaCast.cpp
@@ -2282,19 +2282,39 @@
// 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 consider it illegal
+ // regardless of addrspace compatibility. This makes
+ // local int ** p;
+ // return (generic int **) p;
+ // illegal 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;
+ 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(),
+ diag::err_typecheck_incompatible_address_space)
+ << SrcType << DestType << Sema::AA_Casting
+ << SrcExpr.get()->getSourceRange();
+ SrcExpr = ExprError();
+ return;
+ }
+
+ DestPtr = DestPPtr->getPointeeType().getTypePtr();
+ SrcPtr = SrcPPtr->getPointeeType().getTypePtr();
+ Nested = true;
}
}
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
