maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[OpenCL] Disallow casts between address spaces. 

llvm-svn: 199208
This commit is contained in:
Joey Gouly 2014-01-14 12:47:29 +00:00
parent 9d2e0df049
commit 8fc32f01dc
2 changed files with 41 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
clang/lib/Sema/SemaCast.cpp
View File

@ -2170,6 +2170,21 @@ void CastOperation::CheckCStyleCast() {
assert(!SrcType->isPlaceholderType());
// OpenCL v1 s6.5: Casting a pointer to address space A to a pointer to
// address space B is illegal.
if (Self.getLangOpts().OpenCL && DestType->isPointerType() &&
SrcType->isPointerType()) {
if (DestType->getPointeeType().getAddressSpace() !=
SrcType->getPointeeType().getAddressSpace()) {
Self.Diag(OpRange.getBegin(),
diag::err_typecheck_incompatible_address_space)
<< SrcType << DestType << Sema::AA_Casting
<< SrcExpr.get()->getSourceRange();
SrcExpr = ExprError();
return;
}
}
if (Self.RequireCompleteType(OpRange.getBegin(), DestType,
diag::err_typecheck_cast_to_incomplete)) {
SrcExpr = ExprError();

26
clang/test/SemaOpenCL/address-spaces.cl
View File

@ -11,3 +11,29 @@ __kernel void foo(__global int *gip) {
ip = &li; // expected-error {{assigning '__local int *' to 'int *' changes address space of pointer}}
ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' changes address space of pointer}}
}
void explicit_cast(global int* g, local int* l, constant int* c, private int* p)
{
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*) 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*) 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}}
}
void ok_explicit_casts(global int *g, global int* g2, local int* l, local int* l2, private int* p, private int* p2)
{
g = (global int*) g2;
l = (local int*) l2;
p = (private int*) p2;
}