Skip to content

Commit

Permalink
[Sema][OpenCL] Make address space conversions a bit stricter.
Browse files Browse the repository at this point in the history 
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 implicit conversions
that should definitely not be permitted. We also dig deeper into the
pointer types and warn on explicit 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.

Fixes PR39674!

Patch by ebevhan (Bevin Hansson)!

Differential Revision: https://reviews.llvm.org/D58236

llvm-svn: 360258
  • Loading branch information
@MrSidims
Anastasia Stulova authored and MrSidims committed May 17, 2019
1 parent bc99fc3 commit 245f992
Show file tree
Hide file tree
Showing 9 changed files with 179 additions and 30 deletions.
13 changes: 13 additions & 0 deletions clang/include/clang/Basic/DiagnosticSemaKinds.td
View file
Expand Up @@ -7013,6 +7013,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 $|"
Expand Down
6 changes: 6 additions & 0 deletions clang/include/clang/Sema/Sema.h
View file
Expand Up @@ -9850,6 +9850,12 @@ class Sema {
/// 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
Expand Down
46 changes: 34 additions & 12 deletions clang/lib/Sema/SemaCast.cpp
View file
Expand Up @@ -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;
}
}
}
Expand Down
25 changes: 21 additions & 4 deletions clang/lib/Sema/SemaExpr.cpp
View file
Expand Up @@ -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*.
Expand Down Expand Up @@ -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)
Expand Down Expand Up @@ -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;
Expand Down
13 changes: 2 additions & 11 deletions clang/test/CodeGenOpenCL/numbered-address-space.cl
View file
Expand Up @@ -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) {
Expand All @@ -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);
}

100 changes: 100 additions & 0 deletions clang/test/SemaOpenCL/address-spaces.cl
View file
Expand Up @@ -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}}
Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaOpenCL/event_t_overload.cl
View file
Expand Up @@ -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'}}
}
2 changes: 1 addition & 1 deletion clang/test/SemaOpenCL/numbered-address-space.cl
View file
Expand Up @@ -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}}
}

2 changes: 1 addition & 1 deletion clang/test/SemaOpenCL/queue_t_overload.cl
View file
Expand Up @@ -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'}}
}

0 comments on commit 245f992

Please sign in to comment.