[OpenCL] Add LangAS::opencl_private to represent private address space in AST

[yaxunl]

Summary

Currently Clang uses default address space (0) to represent private address space for OpenCL
in AST. There are two issues with this:

1. Multiple address spaces including private address space cannot be diagnosed.

2. There is no mangling for default address space. For example, if private int* is emitted as

i32 addrspace(5)* in IR. It is supposed to be mangled as PUAS5i but it is mangled as
Pi instead.

This patch attempts to represent OpenCL private address space explicitly in AST. It adds
a new enum LangAS::opencl_private and adds it to the variable types which are implicitly
private:

automatic variables without address space qualifier

function parameter

pointee type without address space qualifier (OpenCL 1.2 and below)

For printing and semantic check of implicit addr space, there is another patch https://reviews.llvm.org/D38857

[yaxunl]

Add private address space qualifier to automatic variable and function parameter. Update sema tests.

ToDo: drop address space when converting l-value to r-value.

[yaxunl]

Update lit tests.

It is essentially done. Clang already drops all qualifiers when converting l-value to r-value.

The only missing part is to let Clang treat (void*)0 as a generic null pointer for OpenCL 1.2 and below since now it is a pointer to private address space.

[yaxunl]

Treat (void*)0 as null pointer for OpenCL 1.2. Rebase to ToT.

[bader]

LGTM, thanks.

[Anastasia]

lib/AST/Expr.cpp

3235

Would we still be accepting the following:

global int * ptr = (global void*)0;

lib/Sema/SemaDecl.cpp

7964

Could we use LangAS::Default instead?

11846

Could we use LangAS::Default here too?

11851

Would it be better to lift this into if condition of line 11846?

lib/Sema/SemaType.cpp

6969

Would it be nicer to not append any address space at all neither here nor down at the end of this function and keep it default instead until the Codegen? If it's doable I would very much prefer that. It seems like it would make the implementation potentially a bit cleaner to understand and easier to improve semantical analysis. See one example of improving original type printing in my comment to the test below.

Also there are at least these 3 related bugs open currently:
https://bugs.llvm.org//show_bug.cgi?id=33418
https://bugs.llvm.org//show_bug.cgi?id=33419
https://bugs.llvm.org//show_bug.cgi?id=33420

Does your change address any of those?

test/SemaOpenCL/access-qualifier.cl

0–1

Ok, I think that here it would be less confusing not to add any address space since it's missing in the original source.

[yaxunl]

lib/AST/Expr.cpp

3235

Yes. There is a test SemaOpenCL/null_literal.cl

lib/Sema/SemaDecl.cpp

7964

done

11846

done

11851

will do.

lib/Sema/SemaType.cpp

6969

On the contrary, I think using default address space for automatic variable and function parameter will cause ambiguity and inconsistency in AST, making it more difficult to understand and process, and making some bug (e.g. https://bugs.llvm.org//show_bug.cgi?id=33419) unfixable. For example, private int f(void) and int f(void) will be identical in AST, therefore we cannot diagnose private int f(void).

With current representation I am able to fix the 3 bugs. I will update the diff.

test/SemaOpenCL/access-qualifier.cl

0–1

Although it looks strange at first sight, __private __read_write image1d_t is the true type of the argument. The function argument is allocated from stack and is an l-value. If we take its address, we get a private pointer.

[yaxunl]

Revised by Anastasia's comments.

[Anastasia]

lib/Sema/SemaDecl.cpp

11846

Sorry, I wasn't clear. I think we could have:

if (T.getAddressSpace() != LangAS::Default && T.getAddressSpace() != LangAS::opencl_private)

and then original condition. It is a bit clearer I think.

lib/Sema/SemaType.cpp

6969

I don't see why?

private int f(void) -> will have an address space attribute in AST as it is provided explicitly.

int f(void) -> will have no address space attribute because it's not provided explicitly and not attached implicitly either.

All I was asking is not to deduce the address space here if it's not specified explicitly until later step when we need to put it in the IR.

test/SemaOpenCL/access-qualifier.cl

0–1

Yes, it is a true type in the Clang internal representation indeed, but not in the original source though. Here we are giving the feedback about the source code so it's nicer to keep it as close to the original as possible. But we are doing similar "magic" for blocks, images and other places, so it's not that critical at the end. Just if it can be avoided it would be better.

[yaxunl]

lib/Sema/SemaDecl.cpp

11846

No. For OpenCL, the condition is on line 11847 and 11848. An array type in other address spaces is allowed.

lib/Sema/SemaType.cpp

6969

Clang already deduce global and generic address spaces and use them in the diagnostic messages. I don't see why we can use deduced global and generic address space in diagnostics whereas cannot use deduced private address space in diagnostics. Why users can accept deduced global and generic address spaces but cannot accept deduced private address space?

Automatic variables and function parameters have private address space. This is the reality and as true as a global variable has global or constant address spaces. Not using private address space in diagnostics gives user illusion that automatic variables and function parameters do not have address space, which is not true.

Besides, allowing default address space to represent private address space in AST causes ambiguity in AST. Instead of just check if a type has private address space, now we need to check if a type has private or default address spaces. Also if an expression has default address space, it is not clear if it is an l-value or r-value. This will complicate semantic checking unnecessarily. Also I am not sure if it is acceptable to modify AST between Sema and CodeGen since it seems to change the paradigm of how clang does Sema/CodeGen now.

[Anastasia]

lib/Sema/SemaDecl.cpp

11846

I think the initial condition was T.getAddressSpace() != 0 i.e. if not private address space.

So now since we added private this should be T.getAddressSpace() !=LangAS::opencl_private but we can have the default case as well hence 'T.getAddressSpace() !=LangAS::opencl_private || T.getAddressSpace() !=LangAS::Default'.

This entire case seems to be for OpenCL anyways. So you could also move out the OpenCL check if you prefer. I am just trying to see if we can make this easier to understand.

lib/Sema/SemaType.cpp

6969

Clang already deduce global and generic address spaces and use them in the diagnostic messages. I don't see why we can use deduced global and generic address space in diagnostics whereas cannot use deduced private address space in diagnostics. Why users can accept deduced global and generic address spaces but cannot accept deduced private address space?

Yes, we did this initially as a workaround because there was no way to distinguish the private and default address space. So now since we are adding the private AS explicitly, it would be nice to remove the workaround if that's possible. It's just a messy code which is hard to understand and we had places in which we need to know if the address space was specified explicitly or now. NULL pointer is one of them. There were also bugs associated to this because it wasn't expected that the address space is 'magically' attached during parsing. Even though most of the things are solved now... I would prefer not to deduce any attributes at this place not just private but global and default as well...

Besides, allowing default address space to represent private address space in AST causes ambiguity in AST. Instead of just check if a type has private address space, now we need to check if a type has private or default address spaces.

I don't see any ambiguity, this is a difference in OpenCL wrt C that we have implicit address space which we have to handle in the compiler too. The question is what's the best place to loose the information about default address space... so we first went for removing it early in parsing but there were a number of issues for example NULL couldn't be handled properly and some other bugs. The original source couldn't be printed back the same way. I personally find the current solution a bit hacky because this function was just supposed to simply parse the attribute and not to add bits that don't exist in the original source code. Therefore, the code is overcomplicated because we are trying to get some extra information about Declarator properties which this function shouldn't really touch.

Once again I have never said it is wrong to add private address space to those objects, just it will be cleaner in my view if we don't do it at this stage. If it's too difficult to do it differently (e.g too big of the change) I am OK to accept this solution. But if it's possible to do it later, I would prefer to keep default address space with no address space attribute until the last point where we generate the IR.

Also if an expression has default address space, it is not clear if it is an l-value or r-value. This will complicate semantic checking unnecessarily. Also I am not sure if it is acceptable to modify AST between Sema and CodeGen since it seems to change the paradigm of how clang does Sema/CodeGen now.

Could you please elaborate what is the difficulty here?

[yaxunl]

lib/Sema/SemaType.cpp

6969

Address space is a type qualifier applies to all expressions, not just variables. Using default address space in AST means that the real address space of any expression will depend on the context and needs to be evaluated whenever it is needed for semantic checking, which may result in lots of changes in Sema.

Since most semantic checking uses the real address space, it is more efficient to use real address space in AST.

The only semantic checking that requires to know the original default address space seems to be (void*)0 for OpenCL 1.2 since by OpenCL spec it should be a private pointer but by convention it is treated as null pointer. It does not make sense to use default address space in AST just for this one special case. I think we should be able to handle it through some other approach.

John, do you have any comments? Thanks.

[rjmccall]

The meaning we've agreed on for LangAS::Default is to be the address space of local declarations, which corresponds quite well to __private in OpenCL. I think your concern about diagnostics is better addressed by changing the pretty-printer than by changing Sema to give all local declarations qualified type.

[yaxunl]

In D35082#844620, @rjmccall wrote:

The meaning we've agreed on for LangAS::Default is to be the address space of local declarations, which corresponds quite well to __private in OpenCL. I think your concern about diagnostics is better addressed by changing the pretty-printer than by changing Sema to give all local declarations qualified type.

How about adding a ImplicitAddrSpace bit to Qualifier to indicate address space is implicit, then does not print the address space qualifier in AST printer if the bit is set.

[rjmccall]

Well, we don't currently have a concept of a non-canonical qualifier, but I think it probably wouldn't be difficult to support; you would just need ASTContext::getQualifiedType to be aware of it.

[yaxunl]

Add a flag to indicate whether address space qualifier is implicit and only print explicit address space in diagnostics.

[yaxunl]

lib/Sema/SemaDecl.cpp

11846

This is not just for OpenCL.

The condition expresses the following logic:

For non-OpenCL languages, only default addr space is allowed on function argument.

For OpenCL, for array type argument, any addr space is allowed; for non-array type argument, only private addr space is allowed.

[rjmccall]

lib/Sema/SemaDecl.cpp

11846

Parameters declared with type 'foo T[]' will be canonicalized to 'foo T*' when forming the function type, so that's not actually a different rule in the end.

[yaxunl]

ping

[Anastasia]

include/clang/AST/Type.h

332 ↗

(On Diff #112288)

Could we add a bit of comment somewhere explaining why we added implicit AS flag please.

lib/AST/Expr.cpp

3254

Is this comment right wrt to named address spaces. So is (private void*)0 not a nullptr?

3257

Can we now simplify this by checking implicit AS bit instead?

lib/Sema/SemaType.cpp

6974

I am not very convinced we need to do all this check really... I think since we have to live with this code further and maintain it let's try to simplify it a bit. This can also help us avoid redundant checks. Perhaps we could even move it out in a separate function?

How about some sort of hierarchical check structure like:

	CL version?
<2.0		>=2.0
private		Type?
	pointer		non-pointer
	generic		Scope?
		program		function
		global		Qualifiers?
			none		static/extern
			private		global

6994

The generic address space -> The default address space

[yaxunl]

include/clang/AST/Type.h

332 ↗

(On Diff #112288)

Will do.

lib/AST/Expr.cpp

3254

Yes this will be true.

Neither (private void*)0 nor (__generic void*)0 should be treated nullptr since they cannot be assigned to a constant pointer.

This need some change below to check implicity of the address space. I will do that. Also add more tests about this.

lib/Sema/SemaType.cpp

6994

'The generic address space name for arguments' is literally from the OpenCL 1.2 spec. Note it refers 'generic address space name', which is not the 'generic address space' defined by OpenCL 2.0 spec.

[yaxunl]

Add comments for getImplicitAddressSpaceFlag and fix checking of null pointer.

[Anastasia]

lib/Sema/SemaType.cpp

6994

True, but this spec was written before v2.0 was released. And I feel now it makes things confusing considering that we have v2.0 implementation too.

[yaxunl]

lib/Sema/SemaType.cpp

6994

I can make the change. I just feel a little bit uneasy since this looks like a citation but actually is rephrased.

[Anastasia]

lib/Sema/SemaType.cpp

6994

Cool! Thanks!

[yaxunl]

Rebase to ToT and clean up logic.

[yaxunl]

lib/Sema/SemaType.cpp

6974

fixed.

[Anastasia]

lib/Sema/SemaType.cpp

6808

Great! This looks so clear now!

6810

I think this could be checked before calling the function.

6863

Do we cover extern too?

6872

I think we can't have this case for CL1.2 see s6.8. But I think it could happen for extern though.

[yaxunl]

lib/Sema/SemaType.cpp

6810

will do.

6863

will add.

6872

Right. I will remove setting implicit addr space for static var for CL1.2.

For extern var, for CL2.0 I will set implicit addr space to global.

However, for CL1.2 since only constant addr space is supported for file-scope var, I can only set the implicit addr space of an extern var to be constant. However I feel this may cause more confusion than convenience, therefore I will not set implicit addr space for extern var for CL1.2. If user does not use constant addr space with extern var explicitly, they will see diagnostics that extern var must have constant addr space. This is also the current behavior before my change.

[yaxunl]

Revised by Anastasia's comments.

[rjmccall]

Why is most of this patch necessary under the design of adding a non-canonical __private address space?

include/clang/AST/Type.h

336 ↗

(On Diff #117020)

isAddressSpaceImplicit()

337 ↗

(On Diff #117020)

setAddressSpaceImplicit

lib/AST/ItaniumMangle.cpp

2232

In what situation is this mangled? I thought we agree this was non-canonical.

[Anastasia]

lib/Sema/SemaType.cpp

6872

Makes sense!

test/SemaOpenCL/storageclass.cl

63 ↗

(On Diff #117020)

Does it make sense to put different address spaces here since this code is rejected earlier anyway?

[yaxunl]

Revised by John's comments.

[yaxunl]

In D35082#887855, @rjmccall wrote:

Why is most of this patch necessary under the design of adding a non-canonical __private address space?

There are two reasons that we need a flag to indicate an address space is simplicit:

1. We need a consistent way to print the address space qualifier depending on whether it is implicit or not.

We only print address space qualifier when it is explicit. This is not just for private address space. It is for all
address spaces.

2. In some rare situations we need to know whether an address space is implicit when doing the semantic

checking.

Since the implicit property is per address space qualifier, we need this flag to be on the qualifier.

include/clang/AST/Type.h

336 ↗

(On Diff #117020)

Will change.

337 ↗

(On Diff #117020)

Will change.

lib/AST/ItaniumMangle.cpp

2232

OpenCL has overloaded builtin functions, e.g. __attribute__((overloadable)) void f(private int*) and __attribute__((overloadable)) void f(global int*). These functions need to be mangled so that the mangled names are different.

test/SemaOpenCL/storageclass.cl

63 ↗

(On Diff #117020)

In Sema I saw code handling different address spaces in various places. I want to make sure that all address spaces are handled correctly.

[rjmccall]

Are you sure it's a good idea to not print the address space when it's implicit? Won't that often lead to really confusing diagnostics?

Also, we do already have a way of expressing that an extended qualifier was explicit: AttributedType. We have very similar sorts of superficial well-formedness checks to what I think you're trying to do in ObjC ARC.

[yaxunl]

In D35082#889053, @rjmccall wrote:

Are you sure it's a good idea to not print the address space when it's implicit? Won't that often lead to really confusing diagnostics?

Also, we do already have a way of expressing that an extended qualifier was explicit: AttributedType. We have very similar sorts of superficial well-formedness checks to what I think you're trying to do in ObjC ARC.

Based on my observation, in most cases it is OK not to print the implicit address space, and printing implicit address space could cause quite annoying cluttering. In some cases printing implicit address space may be desired. I can improve the printing in some future patch, e.g. only hide the implicit address space in situations which causes cluttering but not providing much useful information.

I think AttributedType sounds an interesting idea and worth exploring.

I just felt this review dragged too long (~ 3 months) already. We have an important backend change depending on this feature. Since the current solution achieves its goals already, can we leave re-implementing it by AttributedType for the future? Thanks.

[rjmccall]

You have an important backend change relying on being able to preserve type sugar better in diagnostics? The only apparent semantic change in this patch is that you're changing the mangling, which frankly seems incorrect.

[yaxunl]

In D35082#890143, @rjmccall wrote:

You have an important backend change relying on being able to preserve type sugar better in diagnostics? The only apparent semantic change in this patch is that you're changing the mangling, which frankly seems incorrect.

Can you elaborate on why the mangling is incorrect?

[rjmccall]

Non-canonical information is not supposed to be mangled.

It's not clear to me what the language rule you're really trying to implement is. Maybe you really do need a canonical __private address space, in which case you are going to have to change a lot of code in Sema to add the address space qualifier to every local and temporary allocation. But that's not obvious to me yet, because you haven't really explained why it needs to be explicit in the representation.

If you just want pointers to private to be mangled with an address-space qualifier — meaning I guess that the mangling e.g. Pi will be completely unused — that should be easy enough to handle in the mangler. But if you need to distinguish between private-qualified types and unqualified types, and that distinction isn't purely to implement some syntactic restriction about not writing e.g. private global, then that's not good enough and you do need a canonical __private.

Telling me that you're in a hurry isn't helpful; preserving a reasonable representation and not allowing corner cases to become maintenance problems is far more important to the project than landing patches in trunk on some particular schedule.

[yaxunl]

In D35082#890150, @rjmccall wrote:

Non-canonical information is not supposed to be mangled.

It's not clear to me what the language rule you're really trying to implement is. Maybe you really do need a canonical __private address space, in which case you are going to have to change a lot of code in Sema to add the address space qualifier to every local and temporary allocation. But that's not obvious to me yet, because you haven't really explained why it needs to be explicit in the representation.

In OpenCL all memory is in certain address space, therefore every l-value and pointee in a pointer should have an address space. Private address space has equal status as global or local address space. There are language rules about pointers to what address space can be assigned to pointers to what address space. Therefore all address space needs to be canonical. This patch already puts every local variable and function parameter in private address space, as is done in deduceOpenCLImplicitAddrSpace(). We need private address space to be explicit because we need to display explicit private address space and implicit address space differently, also because in certain semantic checks we need to know if a private address space is explicit or not.

If you just want pointers to private to be mangled with an address-space qualifier — meaning I guess that the mangling e.g. Pi will be completely unused — that should be easy enough to handle in the mangler. But if you need to distinguish between private-qualified types and unqualified types, and that distinction isn't purely to implement some syntactic restriction about not writing e.g. private global, then that's not good enough and you do need a canonical __private.

The mangler does not mangle an empty type qualifier, therefore if private address space is represented as the default address space (i.e., no address space), it will not be mangled at all. This is also related to the substitution and cannot be easily changed without breaking lots of stuff in the mangler.

Telling me that you're in a hurry isn't helpful; preserving a reasonable representation and not allowing corner cases to become maintenance problems is far more important to the project than landing patches in trunk on some particular schedule.

The introduction of explicit private address space and the implicit address space flag in AST is precisely for handling those corner cases in the OpenCL language rules so that they won't become maintenance problems.

[rjmccall]

Okay. I think I see your point about why it would be better to have a canonical __private address space that is different from the implicit address space 0. I assume this means that there should basically never be pointers, references, or l-values in address space 0 in OpenCL. You will lose a significant amount of representational efficiency by doing this, but it's probably not overwhelming.

I know you aren't implementing OpenCL C++ yet, so most of the cases where temporaries are introduced aren't meaningful to you, but you will at least need to consider compound literals. I suspect the right rule is that file-scope literals should be inferred as being in global or constant memory, depending on whether they're const, and local-scope literals should be inferred as __private.

I'll try to review your patch tomorrow.

[Anastasia]

In D35082#890162, @rjmccall wrote:

Okay. I think I see your point about why it would be better to have a canonical __private address space that is different from the implicit address space 0. I assume this means that there should basically never be pointers, references, or l-values in address space 0 in OpenCL.

If you mean 0 is private address space then no. There can be private AS pointers. But if you mean 0 is no address space then this doesn't exist in OpenCL. I think the right design in the first place would have been to keep address spaces in AST just as they are in the source code and add explicit address spaces to all types in IR instead. In this case absence of any address spaces in AST would signal implicit AS to be used. This would make some Sema checks a bit more complicated though, but the design would become a lot cleaner. Unfortunately I think it would not be worth doing this big change now.

You will lose a significant amount of representational efficiency by doing this, but it's probably not overwhelming.

I know you aren't implementing OpenCL C++ yet, so most of the cases where temporaries are introduced aren't meaningful to you, but you will at least need to consider compound literals. I suspect the right rule is that file-scope literals should be inferred as being in global or constant memory, depending on whether they're const, and local-scope literals should be inferred as __private.

I'll try to review your patch tomorrow.

[rjmccall]

In D35082#890359, @Anastasia wrote:

In D35082#890162, @rjmccall wrote:

Okay. I think I see your point about why it would be better to have a canonical __private address space that is different from the implicit address space 0. I assume this means that there should basically never be pointers, references, or l-values in address space 0 in OpenCL.

If you mean 0 is private address space then no. There can be private AS pointers. But if you mean 0 is no address space then this doesn't exist in OpenCL. I think the right design in the first place would have been to keep address spaces in AST just as they are in the source code and add explicit address spaces to all types in IR instead. In this case absence of any address spaces in AST would signal implicit AS to be used. This would make some Sema checks a bit more complicated though, but the design would become a lot cleaner. Unfortunately I think it would not be worth doing this big change now.

My understanding is that that is exactly what the majority of this patch is trying to achieve: making private its own address space, separate from address space 0. The patch is certainly introducing a new address-space enumerator for private; if that's meant to be treated as canonically equal to address space 0, well, (1) this patch doesn't seem to add any of the APIs that I'd expect would be required to make that work, like a way to canonicalize an address space or ask if two address spaces are canonically the same, and (2) that would be a major new complexity in the representation that we wouldn't welcome anyway.

The patch is also tracking explicitness of address spaces in the non-canonical type system, but the majority of the patch is spent dealing with the consequences of splitting __private away from address space 0.

Yaxun has convinced me that it's important to represent a private-qualified type differently from an unqualified type.(*) It seems that you agree, but you're concerned about how that will affect your schedule. I can't speak to your schedule; it seems to me that you and Yaxun need to figure that out privately. If you decide not to separate them yet, then the majority of this patch needs to be put on hold. If you do decide to separate them, it seems to me that we should probably split the change to private out from the introduction of implicit address spaces. Just let me know what you want to do.

(*) I know that you aren't considering OpenCL C++ yet, but often these representation/model questions are easier to answer when thinking about C++ instead of C because type differences are so much more directly important in C++. In OpenCL C++, I assume it will be important to distinguish between <int> and <private int> as template arguments. That tells us straight up that private needs to be represented differently from a lack of qualifier. Note that the result, where certain type representations (e.g. an unqualified int*) become basically meaningless and should never be constructed by Sema, is already precedented in Clang: ObjC ARC does the same thing to unqualified ObjC pointer types, like id*.

[Anastasia]

(*) I know that you aren't considering OpenCL C++ yet, but often these representation/model questions are easier to answer when thinking about C++ instead of C because type differences are so much more directly important in C++. In OpenCL C++, I assume it will be important to distinguish between <int> and <private int> as template arguments. That tells us straight up that private needs to be represented differently from a lack of qualifier. Note that the result, where certain type representations (e.g. an unqualified int*) become basically meaningless and should never be constructed by Sema, is already precedented in Clang: ObjC ARC does the same thing to unqualified ObjC pointer types, like id*.

It's pretty complicated in OpenCL as int and private int won't be different, but int* and private int* would be different only in OpenCL2.0. The rules are pretty convoluted and I have summarized them in the table earlier: https://reviews.llvm.org/D35082#inline-327303. The issue is that there is not anything like this in other languages and therefore we have this issue trying to fit this concept with something similar but not exactly the same. I thought you and Sam decided to use implicit AS flag to make printing messages consistent with the original source code. I am fine with this approach, however, I would prefer to just keep no AS qualifier in the default AS mode instead of deducing it during parcing in processTypeAttrs, and only to add the AS to the IR at the end of the CodeGen phase. I think it would be a lot easier to understand. However, as Sam has pointed out a lot of code in Sema has been written assuming the deduction of AS and is using AS qualifiers explicitly and therefore it seems like it would be a bigger change to go for. At the same type we have refactored the deduction of default AS now in the parcing phase and it looks a lot better than I thought it would be. So I don't mind if we continue this way.

[yaxunl]

If there is no other issues. May I commit this patch now? Thanks.

[rjmccall]

It sounds like there's agreement about the basic technical direction of introducing LangAS::opencl_private. Please introduce isAddressSpaceImplicit() in a different patch and make this patch just about the introduction of LangAS::opencl_private. You can have the pretty-printer just ignore __private for now, which should avoid gratuitous diagnostic changes.

I would like you to investigate using AttributedType for the pretty-printing and address-space semantic checks before adding isAddressSpaceImplicit().

include/clang/AST/Type.h

562 ↗

(On Diff #117770)

"I" is not an appropriate abbreviation for "AddressSpaceImplicit".

include/clang/Basic/AddressSpaces.h

1

I think you need a real comment about the design of OpenCL address spaces here. Specifically, it is important to note that OpenCL no longer uses LangAS::Default for anything except r-values.

[yaxunl]

In D35082#895230, @rjmccall wrote:

It sounds like there's agreement about the basic technical direction of introducing LangAS::opencl_private. Please introduce isAddressSpaceImplicit() in a different patch and make this patch just about the introduction of LangAS::opencl_private. You can have the pretty-printer just ignore __private for now, which should avoid gratuitous diagnostic changes.

I would like you to investigate using AttributedType for the pretty-printing and address-space semantic checks before adding isAddressSpaceImplicit().

Thanks. I will separate the implicit addr space flag to another patch.

include/clang/AST/Type.h

562 ↗

(On Diff #117770)

Will change it to ImplictAddrSpace.

include/clang/Basic/AddressSpaces.h

1

Will do.

[rjmccall]

In D35082#895312, @yaxunl wrote:

Thanks. I will separate the implicit addr space flag to another patch.

Thanks, appreciated.

include/clang/AST/Type.h

562 ↗

(On Diff #117770)

Sounds good.

[yaxunl]

Separate implicit addr space flag to another patch as John suggested.

This patch only introduces the private addr space but does not print it.

[rjmccall]

A few minor comments; feel free to commit after addressing them.

include/clang/Basic/AddressSpaces.h

1

Thanks, WFM.

lib/AST/TypePrinter.cpp

1703

Please add the case here, even if it's unreachable for now.

lib/Sema/SemaType.cpp

5754

Please just make this a case, and then the default can be llvm_unreachable.

include/clang/Basic/AddressSpaces.h

 //===--- AddressSpaces.h - Language-specific address spaces -----*- C++ -*-===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 ///
enum ID {
   // the source code, except for OpenCL, where the address space value 0 in
   // QualType represents private address space in OpenCL source code.
   Default = 0,
 
   // OpenCL specific address spaces.
   opencl_global,
   opencl_local,
   opencl_constant,
+  opencl_private,
   opencl_generic,
 
   // CUDA specific address spaces.
   cuda_device,
   cuda_constant,
   cuda_shared,
 
   // This denotes the count of language-specific address spaces and also

lib/AST/ASTContext.cpp

static const LangAS::Map *getAddressSpaceMap(const TargetInfo &T,
   if (LOpts.FakeAddressSpaceMap) {
     // The fake address space map must have a distinct entry for each
     // language-specific address space.
     static const unsigned FakeAddrSpaceMap[] = {
       0, // Default
       1, // opencl_global
       3, // opencl_local
       2, // opencl_constant
+      0, // opencl_private
       4, // opencl_generic
       5, // cuda_device
       6, // cuda_constant
       7  // cuda_shared
     };
     return &FakeAddrSpaceMap;
   } else {
     return &T.getAddressSpaceMap();

lib/AST/Expr.cpp

   }
 
   // Strip off a cast to void*, if it exists. Except in C++.
   if (const ExplicitCastExpr *CE = dyn_cast<ExplicitCastExpr>(this)) {
     if (!Ctx.getLangOpts().CPlusPlus) {
       // Check that it is a cast to void*.
       if (const PointerType *PT = CE->getType()->getAs<PointerType>()) {
         QualType Pointee = PT->getPointeeType();
-        Qualifiers Q = Pointee.getQualifiers();
-        // In OpenCL v2.0 generic address space acts as a placeholder
-        // and should be ignored.
-        bool IsASValid = true;
-        if (Ctx.getLangOpts().OpenCLVersion >= 200) {
-          if (Pointee.getAddressSpace() == LangAS::opencl_generic)
-            Q.removeAddressSpace();
-          else
-            IsASValid = false;
-        }
+        // Only (void*)0 or equivalent are treated as nullptr. If pointee type
+        // has non-default address space it is not treated as nullptr.
+        bool PointeeHasDefaultAS = false;
+        if (Ctx.getLangOpts().OpenCL)
+          PointeeHasDefaultAS =
+              (Ctx.getLangOpts().OpenCLVersion >= 200 &&
+               Pointee.getAddressSpace() == LangAS::opencl_generic) ||
+              (Ctx.getLangOpts().OpenCLVersion < 200 &&
+               Pointee.getAddressSpace() == LangAS::opencl_private);
+        else
+          PointeeHasDefaultAS = Pointee.getAddressSpace() == LangAS::Default;
 
-        if (IsASValid && !Q.hasQualifiers() &&
-            Pointee->isVoidType() &&                      // to void*
-            CE->getSubExpr()->getType()->isIntegerType()) // from int.
+        if (PointeeHasDefaultAS && Pointee->isVoidType() && // to void*
+            CE->getSubExpr()->getType()->isIntegerType())   // from int.
           return CE->getSubExpr()->isNullPointerConstant(Ctx, NPC);
       }
     }
   } else if (const ImplicitCastExpr *ICE = dyn_cast<ImplicitCastExpr>(this)) {
     // Ignore the ImplicitCastExpr type entirely.
     return ICE->getSubExpr()->isNullPointerConstant(Ctx, NPC);
   } else if (const ParenExpr *PE = dyn_cast<ParenExpr>(this)) {
     // Accept ((void*)0) as a null pointer constant, as many other
     // implementations do.
     return PE->getSubExpr()->isNullPointerConstant(Ctx, NPC);
   } else if (const GenericSelectionExpr *GE =
                dyn_cast<GenericSelectionExpr>(this)) {
     if (GE->isResultDependent())
       return NPCK_NotNull;
     return GE->getResultExpr()->isNullPointerConstant(Ctx, NPC);
   } else if (const ChooseExpr *CE = dyn_cast<ChooseExpr>(this)) {
     if (CE->isConditionDependent())
       return NPCK_NotNull;
     return CE->getChosenSubExpr()->isNullPointerConstant(Ctx, NPC);

lib/AST/ItaniumMangle.cpp

     //   <type> ::= U <CUDA-addrspace>
 
     SmallString<64> ASString;
     unsigned AS = Quals.getAddressSpace();
 
     if (Context.getASTContext().addressSpaceMapManglingFor(AS)) {
       //  <target-addrspace> ::= "AS" <address-space-number>
       unsigned TargetAS = Context.getASTContext().getTargetAddressSpace(AS);
-      ASString = "AS" + llvm::utostr(TargetAS);
+      if (TargetAS != 0)
+        ASString = "AS" + llvm::utostr(TargetAS);
     } else {
       switch (AS) {
       default: llvm_unreachable("Not a language specific address space");
-      //  <OpenCL-addrspace> ::= "CL" [ "global" | "local" | "constant |
-      //                                "generic" ]
+      //  <OpenCL-addrspace> ::= "CL" [ "global" | "local" | "constant" |
+      //                                "private"| "generic" ]
       case LangAS::opencl_global:   ASString = "CLglobal";   break;
       case LangAS::opencl_local:    ASString = "CLlocal";    break;
       case LangAS::opencl_constant: ASString = "CLconstant"; break;
+      case LangAS::opencl_private:  ASString = "CLprivate";  break;
       case LangAS::opencl_generic:  ASString = "CLgeneric";  break;
       //  <CUDA-addrspace> ::= "CU" [ "device" | "constant" | "shared" ]
       case LangAS::cuda_device:     ASString = "CUdevice";   break;
       case LangAS::cuda_constant:   ASString = "CUconstant"; break;
       case LangAS::cuda_shared:     ASString = "CUshared";   break;
       }
     }
-    mangleVendorQualifier(ASString);
+    if (!ASString.empty())
+      mangleVendorQualifier(ASString);
   }
 
   // The ARC ownership qualifiers start with underscores.
   // Objective-C ARC Extension:
   //
   //   <type> ::= U "__strong"
   //   <type> ::= U "__weak"
   //   <type> ::= U "__autoreleasing"

   // <CV-qualifiers> ::= [r] [V] [K]    # restrict (C99), volatile, const
   if (Quals.hasRestrict())
     Out << 'r';
   if (Quals.hasVolatile())
     Out << 'V';
   if (Quals.hasConst())
     Out << 'K';
 }
 
 void CXXNameMangler::mangleVendorQualifier(StringRef name) {
   Out << 'U' << name.size() << name;
 }
 
 void CXXNameMangler::mangleRefQualifier(RefQualifierKind RefQualifier) {
   // <ref-qualifier> ::= R                # lvalue reference
   //                 ::= O                # rvalue-reference
   switch (RefQualifier) {

lib/AST/TypePrinter.cpp

         break;
       case LangAS::opencl_local:
         OS << "__local";
         break;
       case LangAS::opencl_constant:
       case LangAS::cuda_constant:
         OS << "__constant";
         break;
+      case LangAS::opencl_private:
+        OS << "__private";
+        break;
       case LangAS::opencl_generic:
         OS << "__generic";
         break;
       case LangAS::cuda_device:
         OS << "__device";
         break;
       case LangAS::cuda_shared:
         OS << "__shared";

 
     switch (lifetime) {
     case Qualifiers::OCL_None: llvm_unreachable("none but true");
     case Qualifiers::OCL_ExplicitNone: OS << "__unsafe_unretained"; break;
     case Qualifiers::OCL_Strong: 
       if (!Policy.SuppressStrongLifetime)
         OS << "__strong"; 
       break;
         
     case Qualifiers::OCL_Weak: OS << "__weak"; break;
     case Qualifiers::OCL_Autoreleasing: OS << "__autoreleasing"; break;
     }
   }
 
   if (appendSpaceIfNonEmpty && addSpace)
     OS << ' ';
 }

lib/Basic/Targets.cpp

   }
 };
 
 static const unsigned NVPTXAddrSpaceMap[] = {
     0, // Default
     1, // opencl_global
     3, // opencl_local
     4, // opencl_constant
+    0, // opencl_private
     // FIXME: generic has to be added to the target
     0, // opencl_generic
     1, // cuda_device
     4, // cuda_constant
     3, // cuda_shared
 };
 
 class NVPTXTargetInfo : public TargetInfo {

   return llvm::makeArrayRef(GCCRegNames);
 }
 
 static const LangAS::Map AMDGPUPrivIsZeroDefIsGenMap = {
     4, // Default
     1, // opencl_global
     3, // opencl_local
     2, // opencl_constant
+    0, // opencl_private
     4, // opencl_generic
     1, // cuda_device
     2, // cuda_constant
     3  // cuda_shared
 };
 static const LangAS::Map AMDGPUGenIsZeroDefIsGenMap = {
     0, // Default
     1, // opencl_global
     3, // opencl_local
     2, // opencl_constant
+    5, // opencl_private
     0, // opencl_generic
     1, // cuda_device
     2, // cuda_constant
     3  // cuda_shared
 };
 static const LangAS::Map AMDGPUPrivIsZeroDefIsPrivMap = {
     0, // Default
     1, // opencl_global
     3, // opencl_local
     2, // opencl_constant
+    0, // opencl_private
     4, // opencl_generic
     1, // cuda_device
     2, // cuda_constant
     3  // cuda_shared
 };
 static const LangAS::Map AMDGPUGenIsZeroDefIsPrivMap = {
     5, // Default
     1, // opencl_global
     3, // opencl_local
     2, // opencl_constant
+    5, // opencl_private
     0, // opencl_generic
     1, // cuda_device
     2, // cuda_constant
     3  // cuda_shared
 };
 
 // If you edit the description strings, make sure you update
 // getPointerWidthV().

 // target processor and program binary. TCE co-design environment is
 // publicly available in http://tce.cs.tut.fi
 
 static const unsigned TCEOpenCLAddrSpaceMap[] = {
     0, // Default
     3, // opencl_global
     4, // opencl_local
     5, // opencl_constant
+    0, // opencl_private
     // FIXME: generic has to be added to the target
     0, // opencl_generic
     0, // cuda_device
     0, // cuda_constant
     0  // cuda_shared
 };
 
 class TCETargetInfo : public TargetInfo {

 #include "clang/Basic/BuiltinsLe64.def"
 };
 
 static const unsigned SPIRAddrSpaceMap[] = {
     0, // Default
     1, // opencl_global
     3, // opencl_local
     2, // opencl_constant
+    0, // opencl_private
     4, // opencl_generic
     0, // cuda_device
     0, // cuda_constant
     0  // cuda_shared
 };
 class SPIRTargetInfo : public TargetInfo {
 public:
   SPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &)

lib/CodeGen/CGDecl.cpp

void CodeGenFunction::EmitLifetimeEnd(llvm::Value *Size, llvm::Value *Addr) {
   C->setDoesNotThrow();
 }
 
 /// EmitAutoVarAlloca - Emit the alloca and debug information for a
 /// local variable.  Does not emit initialization or destruction.
 CodeGenFunction::AutoVarEmission
 CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
   QualType Ty = D.getType();
-  assert(Ty.getAddressSpace() == LangAS::Default);
+  assert(
+      Ty.getAddressSpace() == LangAS::Default ||
+      (Ty.getAddressSpace() == LangAS::opencl_private && getLangOpts().OpenCL));
 
   AutoVarEmission emission(D);
 
   bool isByRef = D.hasAttr<BlocksAttr>();
   emission.IsByRef = isByRef;
 
   CharUnits alignment = getContext().getDeclAlign(&D);
 

lib/Sema/SemaChecking.cpp

static bool SemaOpenCLBuiltinToAddr(Sema &S, unsigned BuiltinID,
   auto Qual = RT.getQualifiers();
   switch (BuiltinID) {
   case Builtin::BIto_global:
     Qual.setAddressSpace(LangAS::opencl_global);
     break;
   case Builtin::BIto_local:
     Qual.setAddressSpace(LangAS::opencl_local);
     break;
+  case Builtin::BIto_private:
+    Qual.setAddressSpace(LangAS::opencl_private);
+    break;
   default:
     Qual.removeAddressSpace();
   }
   Call->setType(S.Context.getPointerType(S.Context.getQualifiedType(
       RT.getUnqualifiedType(), Qual)));
 
   return false;
 }

lib/Sema/SemaDecl.cpp

         D.setInvalidType();
       }
     }
 
     // OpenCL v1.2 s6.9.r:
     // The event type cannot be used with the __local, __constant and __global
     // address space qualifiers.
     if (R->isEventT()) {
-      if (R.getAddressSpace()) {
+      if (R.getAddressSpace() != LangAS::opencl_private) {
         Diag(D.getLocStart(), diag::err_event_t_addr_space_qual);
         D.setInvalidType();
       }
     }
   }
 
   DeclSpec::SCS SCSpec = D.getDeclSpec().getStorageClassSpec();
   StorageClass SC = StorageClassSpecToVarDeclStorageClass(D.getDeclSpec());

                   << "constant";
             else
               Diag(NewVD->getLocation(), diag::err_opencl_addrspace_scope)
                   << "local";
             NewVD->setInvalidDecl();
             return;
           }
         }
-      } else if (T.getAddressSpace() != LangAS::Default) {
+      } else if (T.getAddressSpace() != LangAS::opencl_private) {
         // Do not allow other address spaces on automatic variable.
         Diag(NewVD->getLocation(), diag::err_as_qualified_auto_decl) << 1;
         NewVD->setInvalidDecl();
         return;
       }
     }
   }
 

 };
 
 static OpenCLParamType getOpenCLKernelParameterType(Sema &S, QualType PT) {
   if (PT->isPointerType()) {
     QualType PointeeType = PT->getPointeeType();
     if (PointeeType->isPointerType())
       return PtrPtrKernelParam;
     if (PointeeType.getAddressSpace() == LangAS::opencl_generic ||
-        PointeeType.getAddressSpace() == 0)
+        PointeeType.getAddressSpace() == LangAS::opencl_private ||
+        PointeeType.getAddressSpace() == LangAS::Default)
       return InvalidAddrSpacePtrKernelParam;
     return PtrKernelParam;
   }
 
   // TODO: Forbid the other integer types (size_t, ptrdiff_t...) when they can
   // be used as builtin types.
 
   if (PT->isImageType())

 
   // Handle attributes.
   ProcessDeclAttributes(S, NewFD, D);
 
   if (getLangOpts().OpenCL) {
     // OpenCL v1.1 s6.5: Using an address space qualifier in a function return
     // type declaration will generate a compilation error.
     unsigned AddressSpace = NewFD->getReturnType().getAddressSpace();
-    if (AddressSpace == LangAS::opencl_local ||
-        AddressSpace == LangAS::opencl_global ||
-        AddressSpace == LangAS::opencl_constant) {
+    if (AddressSpace != LangAS::Default) {
       Diag(NewFD->getLocation(),
            diag::err_opencl_return_value_with_address_space);
       NewFD->setInvalidDecl();
     }
   }
 
   if (!getLangOpts().CPlusPlus) {
     // Perform semantic checking on the function declaration.

     T = Context.getObjCObjectPointerType(T);
     New->setType(T);
   }
 
   // ISO/IEC TR 18037 S6.7.3: "The type of an object with automatic storage
   // duration shall not be qualified by an address-space qualifier."
   // Since all parameters have automatic store duration, they can not have
   // an address space.
-  if (T.getAddressSpace() != 0) {
+  if (T.getAddressSpace() != LangAS::Default &&
       // OpenCL allows function arguments declared to be an array of a type
       // to be qualified with an address space.
-    if (!(getLangOpts().OpenCL && T->isArrayType())) {
-      Diag(NameLoc, diag::err_arg_with_address_space);
-      New->setInvalidDecl();
-    }
+      !(getLangOpts().OpenCL &&
+        (T->isArrayType() || T.getAddressSpace() == LangAS::opencl_private))) {
+    Diag(NameLoc, diag::err_arg_with_address_space);
+    New->setInvalidDecl();
   }
 
   return New;
 }
 
 void Sema::ActOnFinishKNRParamDeclarations(Scope *S, Declarator &D,
                                            SourceLocation LocAfterDecls) {
   DeclaratorChunk::FunctionTypeInfo &FTI = D.getFunctionTypeInfo();
 

lib/Sema/SemaType.cpp

     case AttributeList::AT_OpenCLLocalAddressSpace:
       ASIdx = LangAS::opencl_local; break;
     case AttributeList::AT_OpenCLConstantAddressSpace:
       ASIdx = LangAS::opencl_constant; break;
     case AttributeList::AT_OpenCLGenericAddressSpace:
       ASIdx = LangAS::opencl_generic; break;
     default:
       assert(Attr.getKind() == AttributeList::AT_OpenCLPrivateAddressSpace);
-      ASIdx = 0; break;
+      ASIdx = LangAS::opencl_private; break;
     }
   }
   
   Type = S.Context.getAddrSpaceQualType(Type, ASIdx);
 }
 
 /// Does this type have a "direct" ownership qualifier?  That is,
 /// is it written like "__strong id", as opposed to something like

   auto diagnoseOrDelay = [](Sema &S, SourceLocation loc,
                             unsigned diagnostic, QualType type) {
     if (S.DelayedDiagnostics.shouldDelayDiagnostics()) {
       S.DelayedDiagnostics.add(
           sema::DelayedDiagnostic::makeForbiddenType(
               S.getSourceManager().getExpansionLoc(loc),
               diagnostic, type, /*ignored*/ 0));
     } else {
       S.Diag(loc, diagnostic);
     }
   };
 
   // Sometimes, __weak isn't allowed.
   if (lifetime == Qualifiers::OCL_Weak &&
       !S.getLangOpts().ObjCWeak && !NonObjCPointer) {
 
     // Use a specialized diagnostic if the runtime just doesn't support them.

   while (attrs) {
     AttributeList &attr = *attrs;
     attrs = attr.getNext(); // reset to the next here due to early loop continue
                             // stmts
 
     // Skip attributes that were marked to be invalid.
     if (attr.isInvalid())
       continue;
 
     if (attr.isCXX11Attribute()) {
       // [[gnu::...]] attributes are treated as declaration attributes, so may
       // not appertain to a DeclaratorChunk, even if we handle them as type
       // attributes.
       if (attr.getScopeName() && attr.getScopeName()->isStr("gnu")) {
         if (TAL == TAL_DeclChunk) {
           state.getSema().Diag(attr.getLoc(),
                                diag::warn_cxx11_gnu_attribute_on_type)
               << attr.getName();
           continue;

     case AttributeList::AT_OpenCLGlobalAddressSpace:
     case AttributeList::AT_OpenCLLocalAddressSpace:
     case AttributeList::AT_OpenCLConstantAddressSpace:
     case AttributeList::AT_OpenCLGenericAddressSpace:
     case AttributeList::AT_AddressSpace:
       HandleAddressSpaceTypeAttribute(type, attr, state.getSema());
       attr.setUsedAsTypeAttr();
       hasOpenCLAddressSpace = true;
       break;
     OBJC_POINTER_TYPE_ATTRS_CASELIST:
       if (!handleObjCPointerTypeAttr(state, attr, type))
         distributeObjCPointerTypeAttr(state, attr, type);
       attr.setUsedAsTypeAttr();
       break;
     case AttributeList::AT_VectorSize:
       HandleVectorSizeAttr(type, attr, state.getSema());
       attr.setUsedAsTypeAttr();
       break;
     case AttributeList::AT_ExtVectorType:
       HandleExtVectorTypeAttr(type, attr, state.getSema());
       attr.setUsedAsTypeAttr();
       break;
     case AttributeList::AT_NeonVectorType:
       HandleNeonVectorTypeAttr(type, attr, state.getSema(),
                                VectorType::NeonVector);
       attr.setUsedAsTypeAttr();

 
       // Otherwise, handle the possible delays.
       else if (!handleFunctionTypeAttr(state, attr, type))
         distributeFunctionTypeAttr(state, attr, type);
       break;
     }
   }
 
+  if (state.getSema().getLangOpts().OpenCL &&
+      !hasOpenCLAddressSpace && type.getAddressSpace() == 0 &&
+      (TAL == TAL_DeclSpec || TAL == TAL_DeclChunk)) {
+    Declarator &D = state.getDeclarator();
+    // Put OpenCL automatic variable in private address space.
+    if (state.getCurrentChunkIndex() == 0 && !D.isFunctionDeclarator() &&
+        !D.isFunctionDefinition() &&
+        D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_typedef &&
+        !type->isVoidType()) {
+      if (D.getContext() == Declarator::BlockContext ||
+          D.getContext() == Declarator::ForContext ||
+          D.getContext() == Declarator::InitStmtContext ||
+          D.getContext() == Declarator::ConditionContext ||
+          D.getContext() == Declarator::PrototypeContext) {
+        if (D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_static) {
+          type = state.getSema().Context.getAddrSpaceQualType(
+              type, LangAS::opencl_private);
+        }
+      }
+      // OpenCL v1.2 s6.5:
+      // The generic address space name for arguments to a function in a
+      // program, or local variables of a function is __private. All function
+      // arguments shall be in the __private address space.
+    } else if (state.getCurrentChunkIndex() > 0 &&
+               state.getSema().getLangOpts().OpenCLVersion <= 120 &&
+               (D.getTypeObject(state.getCurrentChunkIndex() - 1).Kind ==
+                DeclaratorChunk::Pointer)) {
+      type = state.getSema().Context.getAddrSpaceQualType(
+          type, LangAS::opencl_private);
+    }
+  }
+
   // If address space is not set, OpenCL 2.0 defines non private default
   // address spaces for some cases:
   // OpenCL 2.0, section 6.5:
   // The address space for a variable at program scope or a static variable
   // inside a function can either be __global or __constant, but defaults to
   // __global if not specified.
   // (...)
   // Pointers that are declared without pointing to a named address space point
   // to the generic address space.
   if (state.getSema().getLangOpts().OpenCLVersion >= 200 &&
       !hasOpenCLAddressSpace && type.getAddressSpace() == 0 &&
       (TAL == TAL_DeclSpec || TAL == TAL_DeclChunk)) {
     Declarator &D = state.getDeclarator();
     if (state.getCurrentChunkIndex() > 0 &&
         (D.getTypeObject(state.getCurrentChunkIndex() - 1).Kind ==
              DeclaratorChunk::Pointer ||
          D.getTypeObject(state.getCurrentChunkIndex() - 1).Kind ==
              DeclaratorChunk::BlockPointer)) {
       type = state.getSema().Context.getAddrSpaceQualType(
           type, LangAS::opencl_generic);
     } else if (state.getCurrentChunkIndex() == 0 &&
-               D.getContext() == Declarator::FileContext &&
                !D.isFunctionDeclarator() && !D.isFunctionDefinition() &&
                D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_typedef &&
-               !type->isSamplerT())
-      type = state.getSema().Context.getAddrSpaceQualType(
+               !type->isSamplerT()) {
+      if (D.getContext() == Declarator::FileContext) {
+        type = state.getSema().Context.getAddrSpaceQualType(
           type, LangAS::opencl_global);
-    else if (state.getCurrentChunkIndex() == 0 &&
-             D.getContext() == Declarator::BlockContext &&
-             D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_static)
+      } else if (D.getContext() == Declarator::BlockContext) {
+        if (D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_static) {
+          type = state.getSema().Context.getAddrSpaceQualType(
+              type, LangAS::opencl_private);
+        } else {
+          type = state.getSema().Context.getAddrSpaceQualType(
+              type, LangAS::opencl_global);
+        }
+      }
+    } else if (state.getCurrentChunkIndex() == 0 &&
+               D.getContext() == Declarator::BlockContext &&
+               D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_static)
       type = state.getSema().Context.getAddrSpaceQualType(
           type, LangAS::opencl_global);
   }
 }
 
 void Sema::completeExprArrayBound(Expr *E) {
   if (DeclRefExpr *DRE = dyn_cast<DeclRefExpr>(E->IgnoreParens())) {
     if (VarDecl *Var = dyn_cast<VarDecl>(DRE->getDecl())) {

test/CodeGenOpenCL/address-spaces-mangling.cl

-// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefix=ASMANG %s
-// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefix=NOASMANG %s
+// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=ASMANG,ASMAN10 %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=ASMANG,ASMAN20 %s
+// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=NOASMANG,NOASMAN10 %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=NOASMANG,NOASMAN20 %s
 
 // We check that the address spaces are mangled the same in both version of OpenCL
 // RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CL2.0 -emit-llvm -o - | FileCheck -check-prefix=OCL-20 %s
 // RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CL1.2 -emit-llvm -o - | FileCheck -check-prefix=OCL-12 %s
 
 // We can't name this f as private is equivalent to default
 // no specifier given address space so we get multiple definition
 // warnings, but we do want it for comparison purposes.
 __attribute__((overloadable))
 void ff(int *arg) { }
-// ASMANG: @_Z2ffPi
-// NOASMANG: @_Z2ffPi
+// ASMANG10: @_Z2ffPi
+// ASMANG20: @_Z2ffPU3AS4i
+// NOASMANG10: @_Z2ffPi
+// NOASMANG20: @_Z2ffPU9CLgenerici
 // OCL-20-DAG: @_Z2ffPU3AS4i
 // OCL-12-DAG: @_Z2ffPi
 
 __attribute__((overloadable))
 void f(private int *arg) { }
 // ASMANG: @_Z1fPi
-// NOASMANG: @_Z1fPi
+// NOASMANG: @_Z1fPU9CLprivatei
 // OCL-20-DAG: @_Z1fPi
 // OCL-12-DAG: @_Z1fPi
 
 __attribute__((overloadable))
 void f(global int *arg) { }
 // ASMANG: @_Z1fPU3AS1i
 // NOASMANG: @_Z1fPU8CLglobali
 // OCL-20-DAG: @_Z1fPU3AS1i
 // OCL-12-DAG: @_Z1fPU3AS1i
 
 __attribute__((overloadable))
 void f(local int *arg) { }
 // ASMANG: @_Z1fPU3AS3i
 // NOASMANG: @_Z1fPU7CLlocali
 // OCL-20-DAG: @_Z1fPU3AS3i
 // OCL-12-DAG: @_Z1fPU3AS3i
 
 __attribute__((overloadable))
 void f(constant int *arg) { }
 // ASMANG: @_Z1fPU3AS2i
 // NOASMANG: @_Z1fPU10CLconstanti
 // OCL-20-DAG: @_Z1fPU3AS2i
 // OCL-12-DAG: @_Z1fPU3AS2i
+
+#if __OPENCL_C_VERSION__ >= 200
+__attribute__((overloadable))
+void f(generic int *arg) { }
+// ASMANG20: @_Z1fPU3AS4i
+// NOASMANG20: @_Z1fPU9CLgenerici
+// OCL-20-DAG: @_Z1fPU3AS4i
+#endif

test/Index/opencl-types.cl

void kernel testFloatTypes() {
   half scalarHalf;
   half4 vectorHalf;
   float scalarFloat;
   float4 vectorFloat;
   double scalarDouble;
   double4 vectorDouble;
 }
 
-// CHECK: VarDecl=scalarHalf:11:8 (Definition) [type=half] [typekind=Half] [isPOD=1]
-// CHECK: VarDecl=vectorHalf:12:9 (Definition) [type=half4] [typekind=Typedef] [canonicaltype=half __attribute__((ext_vector_type(4)))] [canonicaltypekind=Unexposed] [isPOD=1]
-// CHECK: VarDecl=scalarFloat:13:9 (Definition) [type=float] [typekind=Float] [isPOD=1]
-// CHECK: VarDecl=vectorFloat:14:10 (Definition) [type=float4] [typekind=Typedef] [canonicaltype=float __attribute__((ext_vector_type(4)))] [canonicaltypekind=Unexposed] [isPOD=1]
-// CHECK: VarDecl=scalarDouble:15:10 (Definition) [type=double] [typekind=Double] [isPOD=1]
-// CHECK: VarDecl=vectorDouble:16:11 (Definition) [type=double4] [typekind=Typedef] [canonicaltype=double __attribute__((ext_vector_type(4)))] [canonicaltypekind=Unexposed] [isPOD=1]
+// CHECK: VarDecl=scalarHalf:11:8 (Definition) [type=__private half] [typekind=Half] [isPOD=1]
+// CHECK: VarDecl=vectorHalf:12:9 (Definition) [type=__private half4] [typekind=Typedef] [canonicaltype=half __private __attribute__((ext_vector_type(4)))] [canonicaltypekind=Unexposed] [isPOD=1]
+// CHECK: VarDecl=scalarFloat:13:9 (Definition) [type=__private float] [typekind=Float] [isPOD=1]
+// CHECK: VarDecl=vectorFloat:14:10 (Definition) [type=__private float4] [typekind=Typedef] [canonicaltype=float __private __attribute__((ext_vector_type(4)))] [canonicaltypekind=Unexposed] [isPOD=1]
+// CHECK: VarDecl=scalarDouble:15:10 (Definition) [type=__private double] [typekind=Double] [isPOD=1]
+// CHECK: VarDecl=vectorDouble:16:11 (Definition) [type=__private double4] [typekind=Typedef] [canonicaltype=double __private __attribute__((ext_vector_type(4)))] [canonicaltypekind=Unexposed] [isPOD=1]
 
 #pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable
 
 void kernel OCLImage1dROTest(read_only image1d_t scalarOCLImage1dRO);
 void kernel OCLImage1dArrayROTest(read_only image1d_array_t scalarOCLImage1dArrayRO);
 void kernel OCLImage1dBufferROTest(read_only image1d_buffer_t scalarOCLImage1dBufferRO);
 void kernel OCLImage2dROTest(read_only image2d_t scalarOCLImage2dRO);
 void kernel OCLImage2dArrayROTest(read_only image2d_array_t scalarOCLImage2dArrayRO);
 void kernel OCLImage2dDepthROTest(read_only image2d_depth_t scalarOCLImage2dDepthRO);
 void kernel OCLImage2dArrayDepthROTest(read_only image2d_array_depth_t scalarOCLImage2dArrayDepthRO);
 void kernel OCLImage2dMSAAROTest(read_only image2d_msaa_t scalarOCLImage2dMSAARO);
 void kernel OCLImage2dArrayMSAAROTest(read_only image2d_array_msaa_t scalarOCLImage2dArrayMSAARO);
 void kernel OCLImage2dMSAADepthROTest(read_only image2d_msaa_depth_t scalarOCLImage2dMSAADepthRO);
 void kernel OCLImage2dArrayMSAADepthROTest(read_only image2d_array_msaa_depth_t scalarOCLImage2dArrayMSAADepthRO);
 void kernel OCLImage3dROTest(read_only image3d_t scalarOCLImage3dRO);
 
-// CHECK: ParmDecl=scalarOCLImage1dRO:28:50 (Definition) [type=__read_only image1d_t] [typekind=OCLImage1dRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dArrayRO:29:61 (Definition) [type=__read_only image1d_array_t] [typekind=OCLImage1dArrayRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dBufferRO:30:63 (Definition) [type=__read_only image1d_buffer_t] [typekind=OCLImage1dBufferRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dRO:31:50 (Definition) [type=__read_only image2d_t] [typekind=OCLImage2dRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayRO:32:61 (Definition) [type=__read_only image2d_array_t] [typekind=OCLImage2dArrayRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dDepthRO:33:61 (Definition) [type=__read_only image2d_depth_t] [typekind=OCLImage2dDepthRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRO:34:72 (Definition) [type=__read_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAARO:35:59 (Definition) [type=__read_only image2d_msaa_t] [typekind=OCLImage2dMSAARO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARO:36:70 (Definition) [type=__read_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRO:37:70 (Definition) [type=__read_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRO:38:81 (Definition) [type=__read_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage3dRO:39:50 (Definition) [type=__read_only image3d_t] [typekind=OCLImage3dRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dRO:28:50 (Definition) [type=__private __read_only image1d_t] [typekind=OCLImage1dRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dArrayRO:29:61 (Definition) [type=__private __read_only image1d_array_t] [typekind=OCLImage1dArrayRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dBufferRO:30:63 (Definition) [type=__private __read_only image1d_buffer_t] [typekind=OCLImage1dBufferRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dRO:31:50 (Definition) [type=__private __read_only image2d_t] [typekind=OCLImage2dRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayRO:32:61 (Definition) [type=__private __read_only image2d_array_t] [typekind=OCLImage2dArrayRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dDepthRO:33:61 (Definition) [type=__private __read_only image2d_depth_t] [typekind=OCLImage2dDepthRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRO:34:72 (Definition) [type=__private __read_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAARO:35:59 (Definition) [type=__private __read_only image2d_msaa_t] [typekind=OCLImage2dMSAARO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARO:36:70 (Definition) [type=__private __read_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRO:37:70 (Definition) [type=__private __read_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRO:38:81 (Definition) [type=__private __read_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage3dRO:39:50 (Definition) [type=__private __read_only image3d_t] [typekind=OCLImage3dRO] [isPOD=1]
 
 void kernel OCLImage1dWOTest(write_only image1d_t scalarOCLImage1dWO);
 void kernel OCLImage1dArrayWOTest(write_only image1d_array_t scalarOCLImage1dArrayWO);
 void kernel OCLImage1dBufferWOTest(write_only image1d_buffer_t scalarOCLImage1dBufferWO);
 void kernel OCLImage2dWOTest(write_only image2d_t scalarOCLImage2dWO);
 void kernel OCLImage2dArrayWOTest(write_only image2d_array_t scalarOCLImage2dArrayWO);
 void kernel OCLImage2dDepthWOTest(write_only image2d_depth_t scalarOCLImage2dDepthWO);
 void kernel OCLImage2dArrayDepthWOTest(write_only image2d_array_depth_t scalarOCLImage2dArrayDepthWO);
 void kernel OCLImage2dMSAAWOTest(write_only image2d_msaa_t scalarOCLImage2dMSAAWO);
 void kernel OCLImage2dArrayMSAAWOTest(write_only image2d_array_msaa_t scalarOCLImage2dArrayMSAAWO);
 void kernel OCLImage2dMSAADepthWOTest(write_only image2d_msaa_depth_t scalarOCLImage2dMSAADepthWO);
 void kernel OCLImage2dArrayMSAADepthWOTest(write_only image2d_array_msaa_depth_t scalarOCLImage2dArrayMSAADepthWO);
 void kernel OCLImage3dWOTest(write_only image3d_t scalarOCLImage3dWO);
 
-// CHECK: ParmDecl=scalarOCLImage1dWO:54:51 (Definition) [type=__write_only image1d_t] [typekind=OCLImage1dWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dArrayWO:55:62 (Definition) [type=__write_only image1d_array_t] [typekind=OCLImage1dArrayWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dBufferWO:56:64 (Definition) [type=__write_only image1d_buffer_t] [typekind=OCLImage1dBufferWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dWO:57:51 (Definition) [type=__write_only image2d_t] [typekind=OCLImage2dWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayWO:58:62 (Definition) [type=__write_only image2d_array_t] [typekind=OCLImage2dArrayWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dDepthWO:59:62 (Definition) [type=__write_only image2d_depth_t] [typekind=OCLImage2dDepthWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayDepthWO:60:73 (Definition) [type=__write_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAAWO:61:60 (Definition) [type=__write_only image2d_msaa_t] [typekind=OCLImage2dMSAAWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAAWO:62:71 (Definition) [type=__write_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAAWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAADepthWO:63:71 (Definition) [type=__write_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthWO:64:82 (Definition) [type=__write_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage3dWO:65:51 (Definition) [type=__write_only image3d_t] [typekind=OCLImage3dWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dWO:54:51 (Definition) [type=__private __write_only image1d_t] [typekind=OCLImage1dWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dArrayWO:55:62 (Definition) [type=__private __write_only image1d_array_t] [typekind=OCLImage1dArrayWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dBufferWO:56:64 (Definition) [type=__private __write_only image1d_buffer_t] [typekind=OCLImage1dBufferWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dWO:57:51 (Definition) [type=__private __write_only image2d_t] [typekind=OCLImage2dWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayWO:58:62 (Definition) [type=__private __write_only image2d_array_t] [typekind=OCLImage2dArrayWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dDepthWO:59:62 (Definition) [type=__private __write_only image2d_depth_t] [typekind=OCLImage2dDepthWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayDepthWO:60:73 (Definition) [type=__private __write_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAAWO:61:60 (Definition) [type=__private __write_only image2d_msaa_t] [typekind=OCLImage2dMSAAWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAAWO:62:71 (Definition) [type=__private __write_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAAWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAADepthWO:63:71 (Definition) [type=__private __write_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthWO:64:82 (Definition) [type=__private __write_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage3dWO:65:51 (Definition) [type=__private __write_only image3d_t] [typekind=OCLImage3dWO] [isPOD=1]
 
 void kernel OCLImage1dRWTest(read_write image1d_t scalarOCLImage1dRW);
 void kernel OCLImage1dArrayRWTest(read_write image1d_array_t scalarOCLImage1dArrayRW);
 void kernel OCLImage1dBufferRWTest(read_write image1d_buffer_t scalarOCLImage1dBufferRW);
 void kernel OCLImage2dRWTest(read_write image2d_t scalarOCLImage2dRW);
 void kernel OCLImage2dArrayRWTest(read_write image2d_array_t scalarOCLImage2dArrayRW);
 void kernel OCLImage2dDepthRWTest(read_write image2d_depth_t scalarOCLImage2dDepthRW);
 void kernel OCLImage2dArrayDepthRWTest(read_write image2d_array_depth_t scalarOCLImage2dArrayDepthRW);
 void kernel OCLImage2dMSAARWTest(read_write image2d_msaa_t scalarOCLImage2dMSAARW);
 void kernel OCLImage2dArrayMSAARWTest(read_write image2d_array_msaa_t scalarOCLImage2dArrayMSAARW);
 void kernel OCLImage2dMSAADepthRWTest(read_write image2d_msaa_depth_t scalarOCLImage2dMSAADepthRW);
 void kernel OCLImage2dArrayMSAADepthRWTest(read_write image2d_array_msaa_depth_t scalarOCLImage2dArrayMSAADepthRW);
 void kernel OCLImage3dRWTest(read_write image3d_t scalarOCLImage3dRW);
 
-// CHECK: ParmDecl=scalarOCLImage1dRW:80:51 (Definition) [type=__read_write image1d_t] [typekind=OCLImage1dRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dArrayRW:81:62 (Definition) [type=__read_write image1d_array_t] [typekind=OCLImage1dArrayRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dBufferRW:82:64 (Definition) [type=__read_write image1d_buffer_t] [typekind=OCLImage1dBufferRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dRW:83:51 (Definition) [type=__read_write image2d_t] [typekind=OCLImage2dRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayRW:84:62 (Definition) [type=__read_write image2d_array_t] [typekind=OCLImage2dArrayRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dDepthRW:85:62 (Definition) [type=__read_write image2d_depth_t] [typekind=OCLImage2dDepthRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRW:86:73 (Definition) [type=__read_write image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAARW:87:60 (Definition) [type=__read_write image2d_msaa_t] [typekind=OCLImage2dMSAARW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARW:88:71 (Definition) [type=__read_write image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRW:89:71 (Definition) [type=__read_write image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRW:90:82 (Definition) [type=__read_write image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage3dRW:91:51 (Definition) [type=__read_write image3d_t] [typekind=OCLImage3dRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dRW:80:51 (Definition) [type=__private __read_write image1d_t] [typekind=OCLImage1dRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dArrayRW:81:62 (Definition) [type=__private __read_write image1d_array_t] [typekind=OCLImage1dArrayRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dBufferRW:82:64 (Definition) [type=__private __read_write image1d_buffer_t] [typekind=OCLImage1dBufferRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dRW:83:51 (Definition) [type=__private __read_write image2d_t] [typekind=OCLImage2dRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayRW:84:62 (Definition) [type=__private __read_write image2d_array_t] [typekind=OCLImage2dArrayRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dDepthRW:85:62 (Definition) [type=__private __read_write image2d_depth_t] [typekind=OCLImage2dDepthRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRW:86:73 (Definition) [type=__private __read_write image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAARW:87:60 (Definition) [type=__private __read_write image2d_msaa_t] [typekind=OCLImage2dMSAARW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARW:88:71 (Definition) [type=__private __read_write image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRW:89:71 (Definition) [type=__private __read_write image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRW:90:82 (Definition) [type=__private __read_write image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage3dRW:91:51 (Definition) [type=__private __read_write image3d_t] [typekind=OCLImage3dRW] [isPOD=1]
 
 void kernel intPipeTestRO(read_only pipe int scalarPipe);
 void kernel intPipeTestWO(write_only pipe int scalarPipe);
 
-// CHECK: ParmDecl=scalarPipe:106:46 (Definition) [type=read_only pipe int] [typekind=Pipe] [isPOD=0]
-// CHECK: ParmDecl=scalarPipe:107:47 (Definition) [type=write_only pipe int] [typekind=Pipe] [isPOD=0]
+// CHECK: ParmDecl=scalarPipe:106:46 (Definition) [type=__private read_only pipe int] [typekind=Pipe] [isPOD=0]
+// CHECK: ParmDecl=scalarPipe:107:47 (Definition) [type=__private write_only pipe int] [typekind=Pipe] [isPOD=0]
 
 #define CLK_ADDRESS_CLAMP_TO_EDGE       2
 #define CLK_NORMALIZED_COORDS_TRUE      1
 #define CLK_FILTER_NEAREST              0x10
 
 void kernel testMiscOpenCLTypes() {
   const sampler_t scalarOCLSampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
   clk_event_t scalarOCLEvent;
   queue_t scalarOCLQueue;
   reserve_id_t scalarOCLReserveID;
 }
 
-// CHECK: VarDecl=scalarOCLSampler:117:19 (Definition) [type=const sampler_t] [typekind=Typedef] const [canonicaltype=const sampler_t] [canonicaltypekind=OCLSampler] [isPOD=1]
-// CHECK: VarDecl=scalarOCLEvent:118:15 (Definition) [type=clk_event_t] [typekind=Typedef] [canonicaltype=clk_event_t] [canonicaltypekind=Unexposed] [isPOD=1]
-// CHECK: VarDecl=scalarOCLQueue:119:11 (Definition) [type=queue_t] [typekind=Typedef] [canonicaltype=queue_t] [canonicaltypekind=OCLQueue] [isPOD=1]
-// CHECK: VarDecl=scalarOCLReserveID:120:16 (Definition) [type=reserve_id_t] [typekind=Typedef] [canonicaltype=reserve_id_t] [canonicaltypekind=OCLReserveID] [isPOD=1]
+// CHECK: VarDecl=scalarOCLSampler:117:19 (Definition) [type=const __private sampler_t] [typekind=Typedef] const [canonicaltype=const __private sampler_t] [canonicaltypekind=OCLSampler] [isPOD=1]
+// CHECK: VarDecl=scalarOCLEvent:118:15 (Definition) [type=__private clk_event_t] [typekind=Typedef] [canonicaltype=__private clk_event_t] [canonicaltypekind=Unexposed] [isPOD=1]
+// CHECK: VarDecl=scalarOCLQueue:119:11 (Definition) [type=__private queue_t] [typekind=Typedef] [canonicaltype=__private queue_t] [canonicaltypekind=OCLQueue] [isPOD=1]
+// CHECK: VarDecl=scalarOCLReserveID:120:16 (Definition) [type=__private reserve_id_t] [typekind=Typedef] [canonicaltype=__private reserve_id_t] [canonicaltypekind=OCLReserveID] [isPOD=1]

test/Parser/opencl-astype.cl

 // RUN: %clang_cc1 -fsyntax-only -verify %s -triple spir-unknown-unknown
 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
 
 void test_astype() {
   float f = 1.0f;
   unsigned int i = __builtin_astype(f, unsigned int); 
   
   typedef __attribute__(( ext_vector_type(4) ))  int int4;
   typedef __attribute__(( ext_vector_type(3) ))  float float3;
   typedef __attribute__(( ext_vector_type(4) ))  float float4;
   typedef __attribute__(( ext_vector_type(4) ))  double double4;
   
   float4 f4;
-  double4 d4 = __builtin_astype(f4, double4); // expected-error{{invalid reinterpretation: sizes of 'double4' (vector of 4 'double' values) and 'float4' (vector of 4 'float' values) must match}}
+  double4 d4 = __builtin_astype(f4, double4); // expected-error{{invalid reinterpretation: sizes of 'double4' (vector of 4 'double' values) and '__private float4' (vector of 4 'float' values) must match}}
   
   // Verify int4->float3, float3->int4 works.
   int4 i4;
   float3 f3 = __builtin_astype(i4, float3);
   i4 = __builtin_astype(f3, int4);
 }

test/Parser/opencl-atomics-cl20.cl

 #endif
 
 #ifdef CL20
 void foo(atomic_int * ptr) {}
 void atomic_ops_test() {
   atomic_int i;
   foo(&i);
 // OpenCL v2.0 s6.13.11.8, arithemtic operations are not permitted on atomic types.
-  i++; // expected-error {{invalid argument type 'atomic_int' (aka '_Atomic(int)') to unary expression}}
+  i++; // expected-error {{invalid argument type '__private atomic_int' (aka '__private _Atomic(int)') to unary expression}}
   i = 1; // expected-error {{atomic variable can be assigned to a variable only in global address space}}
-  i += 1; // expected-error {{invalid operands to binary expression ('atomic_int' (aka '_Atomic(int)') and 'int')}}
-  i = i + i; // expected-error {{invalid operands to binary expression ('atomic_int' (aka '_Atomic(int)') and 'atomic_int')}}
+  i += 1; // expected-error {{invalid operands to binary expression ('__private atomic_int' (aka '__private _Atomic(int)') and 'int')}}
+  i = i + i; // expected-error {{invalid operands to binary expression ('__private atomic_int' (aka '__private _Atomic(int)') and '__private atomic_int')}}
 }
 #endif

test/SemaOpenCL/access-qualifier.cl

 // RUN: %clang_cc1 -verify -pedantic -fsyntax-only -cl-std=CL1.2 %s
 // RUN: %clang_cc1 -verify -pedantic -fsyntax-only -cl-std=CL2.0 %s
 
 typedef image1d_t img1d_ro_default; // expected-note {{previously declared 'read_only' here}}
 
 typedef write_only image1d_t img1d_wo; // expected-note {{previously declared 'write_only' here}}
 typedef read_only image1d_t img1d_ro;
 
 #if __OPENCL_C_VERSION__ >= 200
 typedef read_write image1d_t img1d_rw;
 #endif
 
 typedef int Int;
 typedef read_only int IntRO; // expected-error {{access qualifier can only be used for pipe and image type}}
 
 
 void myWrite(write_only image1d_t); // expected-note {{passing argument to parameter here}} expected-note {{passing argument to parameter here}}
 void myRead(read_only image1d_t); // expected-note {{passing argument to parameter here}}
 
 #if __OPENCL_C_VERSION__ >= 200
 void myReadWrite(read_write image1d_t);
 #else
-void myReadWrite(read_write image1d_t); // expected-error {{access qualifier 'read_write' can not be used for '__read_write image1d_t' prior to OpenCL version 2.0}}
+void myReadWrite(read_write image1d_t); // expected-error {{access qualifier 'read_write' can not be used for '__private __read_write image1d_t' prior to OpenCL version 2.0}}
 #endif
 
 
 kernel void k1(img1d_wo img) {
-  myRead(img); // expected-error {{passing 'img1d_wo' (aka '__write_only image1d_t') to parameter of incompatible type '__read_only image1d_t'}}
+  myRead(img); // expected-error {{passing '__private img1d_wo' (aka '__private __write_only image1d_t') to parameter of incompatible type '__read_only image1d_t'}}
 }
 
 kernel void k2(img1d_ro img) {
-  myWrite(img); // expected-error {{passing 'img1d_ro' (aka '__read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}}
+  myWrite(img); // expected-error {{passing '__private img1d_ro' (aka '__private __read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}}
 }
 
 kernel void k3(img1d_wo img) {
   myWrite(img);
 }
 
 #if __OPENCL_C_VERSION__ >= 200
 kernel void k4(img1d_rw img) {
   myReadWrite(img);
 }
 #endif
 
 kernel void k5(img1d_ro_default img) {
-  myWrite(img); // expected-error {{passing 'img1d_ro_default' (aka '__read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}}
+  myWrite(img); // expected-error {{passing '__private img1d_ro_default' (aka '__private __read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}}
 }
 
 kernel void k6(img1d_ro img) {
   myRead(img);
 }
 
 kernel void k7(read_only img1d_wo img){} // expected-error {{multiple access qualifiers}}
 
 kernel void k8(write_only img1d_ro_default img){} // expected-error {{multiple access qualifiers}}
 
 kernel void k9(read_only int i){} // expected-error{{access qualifier can only be used for pipe and image type}}
 
 kernel void k10(read_only Int img){} // expected-error {{access qualifier can only be used for pipe and image type}}
 
 kernel void k11(read_only write_only image1d_t i){} // expected-error{{multiple access qualifiers}}
 
 kernel void k12(read_only read_only image1d_t i){} // expected-error{{multiple access qualifiers}}
 
 #if __OPENCL_C_VERSION__ >= 200
-kernel void k13(read_write pipe int i){} // expected-error{{access qualifier 'read_write' can not be used for 'read_only pipe int'}}
+kernel void k13(read_write pipe int i){} // expected-error{{access qualifier 'read_write' can not be used for '__private read_only pipe int'}}
 #else
-kernel void k13(__read_write image1d_t i){} // expected-error{{access qualifier '__read_write' can not be used for '__read_write image1d_t' prior to OpenCL version 2.0}}
+kernel void k13(__read_write image1d_t i){} // expected-error{{access qualifier '__read_write' can not be used for '__private __read_write image1d_t' prior to OpenCL version 2.0}}
 #endif
 
 #if __OPENCL_C_VERSION__ >= 200
 void myPipeWrite(write_only pipe int); // expected-note {{passing argument to parameter here}}
 kernel void k14(read_only pipe int p) {
-  myPipeWrite(p); // expected-error {{passing 'read_only pipe int' to parameter of incompatible type 'write_only pipe int'}}
+  myPipeWrite(p); // expected-error {{passing '__private read_only pipe int' to parameter of incompatible type 'write_only pipe int'}}
 }
 #endif
 
 #if __OPENCL_C_VERSION__ < 200
 kernel void test_image3d_wo(write_only image3d_t img) {} // expected-error {{use of type '__write_only image3d_t' requires cl_khr_3d_image_writes extension to be enabled}}
 #endif

test/SemaOpenCL/address-spaces-conversions-cl2.0.cl

 #endif
 
 void test_conversion(global int *arg_glob, local int *arg_loc,
                      constant int *arg_const, private int *arg_priv,
                      generic int *arg_gen) {
 
   AS int *var_init1 = arg_glob;
 #ifdef CONSTANT
-// expected-error@-2{{initializing '__constant int *' with an expression of type '__global int *' changes address space of pointer}}
+// expected-error@-2{{initializing '__constant int *__private' with an expression of type '__global int *__private' changes address space of pointer}}
 #endif
 
   AS int *var_init2 = arg_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type '__local int *' changes address space of pointer}}
+// expected-error-re@-2{{initializing '__{{global|constant}} int *__private' with an expression of type '__local int *__private' changes address space of pointer}}
 #endif
 
   AS int *var_init3 = arg_const;
 #ifndef CONSTANT
-// expected-error-re@-2{{initializing '__{{global|generic}} int *' with an expression of type '__constant int *' changes address space of pointer}}
+// expected-error-re@-2{{initializing '__{{global|generic}} int *__private' with an expression of type '__constant int *__private' changes address space of pointer}}
 #endif
 
   AS int *var_init4 = arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type 'int *' changes address space of pointer}}
+// expected-error-re@-2{{initializing '__{{global|constant}} int *__private' with an expression of type '__private int *__private' changes address space of pointer}}
 #endif
 
   AS int *var_init5 = arg_gen;
 #ifndef GENERIC
-// expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type '__generic int *' changes address space of pointer}}
+// expected-error-re@-2{{initializing '__{{global|constant}} int *__private' with an expression of type '__generic int *__private' changes address space of pointer}}
 #endif
 
   AS int *var_cast1 = (AS int *)arg_glob;
 #ifdef CONSTANT
 // expected-error@-2{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
 #endif
 
   AS int *var_cast2 = (AS int *)arg_loc;
 #ifndef GENERIC
 // expected-error-re@-2{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #endif
 
   AS int *var_cast3 = (AS int *)arg_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
 #endif
 
   AS int *var_cast4 = (AS int *)arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-2{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #endif
 
   AS int *var_cast5 = (AS int *)arg_gen;
 #ifdef CONSTANT
 // expected-error@-2{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
 #endif
 
   AS int *var_impl;
   var_impl = arg_glob;
 #ifdef CONSTANT
-// expected-error@-2{{assigning '__global int *' to '__constant int *' changes address space of pointer}}
+// expected-error@-2{{assigning '__global int *__private' to '__constant int *__private' changes address space of pointer}}
 #endif
 
   var_impl = arg_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{assigning '__local int *' to '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-2{{assigning '__local int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}}
 #endif
 
   var_impl = arg_const;
 #ifndef CONSTANT
-// expected-error-re@-2{{assigning '__constant int *' to '__{{global|generic}} int *' changes address space of pointer}}
+// expected-error-re@-2{{assigning '__constant int *__private' to '__{{global|generic}} int *__private' changes address space of pointer}}
 #endif
 
   var_impl = arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{assigning 'int *' to '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-2{{assigning '__private int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}}
 #endif
 
   var_impl = arg_gen;
 #ifndef GENERIC
-// expected-error-re@-2{{assigning '__generic int *' to '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-2{{assigning '__generic int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}}
 #endif
 
   var_cast1 = (AS int *)arg_glob;
 #ifdef CONSTANT
 // expected-error@-2{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
 #endif
 
   var_cast2 = (AS int *)arg_loc;
 #ifndef GENERIC
 // expected-error-re@-2{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #endif
 
   var_cast3 = (AS int *)arg_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
 #endif
 
   var_cast4 = (AS int *)arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-2{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #endif
 
   var_cast5 = (AS int *)arg_gen;
 #ifdef CONSTANT
 // expected-error@-2{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
 #endif
 
   AS int *var_cmp;
#endif
 
   b = var_cmp == arg_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{comparison between  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   b = var_cmp <= arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{comparison between  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{comparison between  ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   b = var_cmp >= arg_gen;
 #ifdef CONSTANT
 // expected-error@-2{{comparison between  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   AS int *var_sub;
#endif
 
   b = var_sub - arg_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{arithmetic operation with operands of type  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   b = var_sub - arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{arithmetic operation with operands of type  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{arithmetic operation with operands of type  ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   b = var_sub - arg_gen;
 #ifdef CONSTANT
 // expected-error@-2{{arithmetic operation with operands of type  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   f_glob(var_sub);
 #ifndef GLOBAL
-// expected-error-re@-2{{passing '__{{constant|generic}} int *' to parameter of type '__global int *' changes address space of pointer}}
+// expected-error-re@-2{{passing '__{{constant|generic}} int *__private' to parameter of type '__global int *' changes address space of pointer}}
 #endif
 
-  f_loc(var_sub); // expected-error-re{{passing '__{{global|constant|generic}} int *' to parameter of type '__local int *' changes address space of pointer}}
+  f_loc(var_sub); // expected-error-re{{passing '__{{global|constant|generic}} int *__private' to parameter of type '__local int *' changes address space of pointer}}
 
   f_const(var_sub);
 #ifndef CONSTANT
-// expected-error-re@-2{{passing '__{{global|generic}} int *' to parameter of type '__constant int *' changes address space of pointer}}
+// expected-error-re@-2{{passing '__{{global|generic}} int *__private' to parameter of type '__constant int *' changes address space of pointer}}
 #endif
 
-  f_priv(var_sub); // expected-error-re{{passing '__{{global|constant|generic}} int *' to parameter of type 'int *' changes address space of pointer}}
+  f_priv(var_sub); // expected-error-re{{passing '__{{global|constant|generic}} int *__private' to parameter of type '__private int *' changes address space of pointer}}
 
   f_gen(var_sub);
 #ifdef CONSTANT
-// expected-error@-2{{passing '__constant int *' to parameter of type '__generic int *' changes address space of pointer}}
+// expected-error@-2{{passing '__constant int *__private' to parameter of type '__generic int *' changes address space of pointer}}
 #endif
 }
 
 void test_ternary() {
   AS int *var_cond;
   generic int *var_gen;
   global int *var_glob;
   var_gen = 0 ? var_cond : var_glob;
#endif
   var_cond = 0 ? var_cond : var_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   private int *var_priv;
   var_gen = 0 ? var_cond : var_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   var_gen = 0 ? var_cond : var_gen;
 #ifdef CONSTANT
 // expected-error@-2{{conditional operator with the second and third operands of type  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   void *var_void_gen;

 // expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|generic}} int *' and '__constant char *') which are pointers to non-overlapping address spaces}}
 #else
 // expected-warning@-4{{pointer type mismatch ('__constant int *' and '__constant char *')}}
 #endif
 
   private char *var_priv_ch;
   var_void_gen = 0 ? var_cond : var_priv_ch;
 #ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and 'char *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__private char *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-warning@-4{{pointer type mismatch ('__generic int *' and 'char *')}}
+// expected-warning@-4{{pointer type mismatch ('__generic int *' and '__private char *')}}
 #endif
 
   generic char *var_gen_ch;
   var_void_gen = 0 ? var_cond : var_gen_ch;
 #ifdef CONSTANT
 // expected-error@-2{{conditional operator with the second and third operands of type  ('__constant int *' and '__generic char *') which are pointers to non-overlapping address spaces}}
 #else
 // expected-warning-re@-4{{pointer type mismatch ('__{{global|generic}} int *' and '__generic char *')}}

test/SemaOpenCL/address-spaces.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 
 __constant int ci = 1;
 
 __kernel void foo(__global int *gip) {
   __local int li;
   __local int lj = 2; // expected-error {{'__local' variable cannot have an initializer}}
 
   int *ip;
-  ip = gip; // expected-error {{assigning '__global int *' to 'int *' changes address space of pointer}}
-  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}}
+#if __OPENCL_C_VERSION__ < 200
+  ip = gip; // expected-error {{assigning '__global int *__private' to '__private int *__private' changes address space of pointer}}
+  ip = &li; // expected-error {{assigning '__local int *' to '__private int *__private' changes address space of pointer}}
+  ip = &ci; // expected-error {{assigning '__constant int *' to '__private int *__private' changes address space of pointer}}
+#else
+  ip = gip;
+  ip = &li;
+  ip = &ci; // expected-error {{assigning '__constant int *' to '__generic int *__private' changes address space of pointer}}
+#endif
 }
 
 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}}
+  g = (global int*) p;    // expected-error {{casting '__private 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}}
+  l = (local int*) p;     // expected-error {{casting '__private 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}}
+  c = (constant int*) p;  // expected-error {{casting '__private 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}}
+  p = (private int*) g;   // expected-error {{casting '__global int *' to type '__private int *' changes address space of pointer}}
+  p = (private int*) l;   // expected-error {{casting '__local int *' to type '__private int *' changes address space of pointer}}
+  p = (private int*) c;   // expected-error {{casting '__constant int *' to type '__private int *' changes address space of pointer}}
+  p = (private int*) cc;  // expected-error {{casting 'const __constant int *' to type '__private 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;
 }
+
+__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}}
+__constant int func_return_constant(void);  //expected-error {{return value cannot be qualified with address space}}
+#if __OPENCL_C_VERSION__ >= 200
+__generic int func_return_generic(void);    //expected-error {{return value cannot be qualified with address space}}
+#endif
+
+void func_multiple_addr(void) {
+  typedef __private int private_int_t;
+  __local __private int var1;   // expected-error {{multiple address spaces specified for type}}
+  __local __private int *var2;  // expected-error {{multiple address spaces specified for type}}
+  __local private_int_t var3;   // expected-error {{multiple address spaces specified for type}}
+  __local private_int_t *var4;  // expected-error {{multiple address spaces specified for type}}
+}

test/SemaOpenCL/arithmetic-conversions.cl

 // RUN: %clang_cc1 %s -triple spir-unknown-unknown -verify -pedantic -fsyntax-only -cl-std=CL1.2
 
 typedef float float2 __attribute__((ext_vector_type(2)));
 typedef long long2 __attribute__((ext_vector_type(2)));
 typedef int int2 __attribute__((ext_vector_type(2)));
 
 kernel void foo1(float2 in, global float2 *out) { *out = in + 0.5;} // expected-error {{scalar operand type has greater rank than the type of the vector element. ('float2' (vector of 2 'float' values) and 'double')}}
 
 kernel void foo2(float2 in, global float2 *out) { *out = 0.5 + in;} // expected-error {{scalar operand type has greater rank than the type of the vector element. ('double' and 'float2' (vector of 2 'float' values))}}
 
 kernel void foo3(float2 in, global float2 *out) { *out = 0.5f + in;}
 
 kernel void foo4(long2 in, global long2 *out) { *out = 5 + in;}
 
 kernel void foo5(float2 in, global float2 *out) {
     float* f;
-    *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('float *' and 'float2' (vector of 2 'float' values))}}
+    *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('__private float *' and 'float2' (vector of 2 'float' values))}}
 }
 
 kernel void foo6(int2 in, global int2 *out) {
     int* f;
-    *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('int *' and 'int2' (vector of 2 'int' values))}}
+    *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('__private int *' and 'int2' (vector of 2 'int' values))}}
 }

test/SemaOpenCL/as_type.cl

 // RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -finclude-default-header -o - -verify -fsyntax-only
 
 char3 f1(char16 x) {
-  return  __builtin_astype(x, char3); // expected-error{{invalid reinterpretation: sizes of 'char3' (vector of 3 'char' values) and 'char16' (vector of 16 'char' values) must match}}
+  return  __builtin_astype(x, char3); // expected-error{{invalid reinterpretation: sizes of 'char3' (vector of 3 'char' values) and '__private char16' (vector of 16 'char' values) must match}}
 }
 
 char16 f3(int x) {
-  return __builtin_astype(x, char16); // expected-error{{invalid reinterpretation: sizes of 'char16' (vector of 16 'char' values) and 'int' must match}}
+  return __builtin_astype(x, char16); // expected-error{{invalid reinterpretation: sizes of 'char16' (vector of 16 'char' values) and '__private int' must match}}
 }
 
 void foo() {
     char src = 1;
-    int dst = as_int(src); // expected-error{{invalid reinterpretation: sizes of 'int' and 'char' must match}}
+    int dst = as_int(src); // expected-error{{invalid reinterpretation: sizes of 'int' and '__private char' must match}}
 }

test/SemaOpenCL/cl20-device-side-enqueue.cl

#endif
 
   char c;
   enqueue_kernel(default_queue, flags, ndrange,
                  ^(local void *a, local void *b) {
                    return 0;
                  },
                  c, 1024L);
 #ifdef WCONV
-// expected-warning-re@-2{{implicit conversion changes signedness: 'char' to 'unsigned {{int|long}}'}}
+// expected-warning-re@-2{{implicit conversion changes signedness: '__private char' to 'unsigned {{int|long}}'}}
 #endif
 #define UINT_MAX 4294967295
 
   enqueue_kernel(default_queue, flags, ndrange,
                  ^(local void *a, local void *b) {
                    return 0;
                  },
                  sizeof(int), sizeof(int) * UINT_MAX);

test/SemaOpenCL/event_t.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
 
 event_t glb_evt; // expected-error {{the 'event_t' type cannot be used to declare a program scope variable}}
 
 constant struct evt_s {
   event_t evt; // expected-error {{the 'event_t' type cannot be used to declare a structure or union field}}
 } evt_str = {0};
 
 void foo(event_t evt); // expected-note {{passing argument to parameter 'evt' here}}
 
-void kernel ker(event_t argevt) { // expected-error {{'event_t' cannot be used as the type of a kernel parameter}}
+void kernel ker(event_t argevt) { // expected-error {{'__private event_t' cannot be used as the type of a kernel parameter}}
   event_t e;
   constant event_t const_evt; // expected-error {{the event_t type can only be used with __private address space qualifier}}
   foo(e);
   foo(0);
   foo(5); // expected-error {{passing 'int' to parameter of incompatible type 'event_t'}}
   foo((event_t)1); // expected-error {{cannot cast non-zero value '1' to 'event_t'}}
 }
 

test/SemaOpenCL/extension-begin.cl

void test_f1(void) {
   g(0);
 }
 
 #pragma OPENCL EXTENSION my_ext : disable 
 void test_f2(void) {
   struct A test_A2; // expected-error {{use of type 'struct A' requires my_ext extension to be enabled}}
   const struct A test_A_local; // expected-error {{use of type 'struct A' requires my_ext extension to be enabled}}
   TypedefOfA test_typedef_A; // expected-error {{use of type 'TypedefOfA' (aka 'struct A') requires my_ext extension to be enabled}}
-  PointerOfA test_A_pointer; // expected-error {{use of type 'PointerOfA' (aka 'const struct A *') requires my_ext extension to be enabled}}
+  PointerOfA test_A_pointer; // expected-error {{use of type 'PointerOfA' (aka 'const __private struct A *') requires my_ext extension to be enabled}}
   f(); // expected-error {{use of declaration 'f' requires my_ext extension to be enabled}}
   g(0); // expected-error {{no matching function for call to 'g'}}
         // expected-note@-26 {{candidate disabled due to OpenCL extension}}
         // expected-note@-22 {{candidate function not viable: requires 0 arguments, but 1 was provided}}
 }
 
 #endif // HEADER_USER
 

test/SemaOpenCL/half.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -Wno-unused-value -triple spir-unknown-unknown
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : disable
 constant float f = 1.0h; // expected-error{{half precision constant requires cl_khr_fp16}}
 
 half half_disabled(half *p, // expected-error{{declaring function return value of type 'half' is not allowed}}
-                   half h)  // expected-error{{declaring function parameter of type 'half' is not allowed}}
+                   half h)  // expected-error{{declaring function parameter of type '__private half' is not allowed}}
 {
-  half a[2]; // expected-error{{declaring variable of type 'half [2]' is not allowed}}
-  half b;    // expected-error{{declaring variable of type 'half' is not allowed}}
-  *p; // expected-error{{loading directly from pointer to type 'half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
+  half a[2]; // expected-error{{declaring variable of type 'half __private[2]' is not allowed}}
+  half b;    // expected-error{{declaring variable of type '__private half' is not allowed}}
+  *p; // expected-error{{loading directly from pointer to type '__private half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
   *p = 0; // expected-error{{assigning directly to pointer to type 'half' requires cl_khr_fp16. Use vector data store builtin functions instead}}
-  p[1]; // expected-error{{loading directly from pointer to type 'half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
+  p[1]; // expected-error{{loading directly from pointer to type '__private half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
   p[1] = 0; // expected-error{{assigning directly to pointer to type 'half' requires cl_khr_fp16. Use vector data store builtin functions instead}}
 
   float c = 1.0f;
   b = (half) c;  // expected-error{{casting to type 'half' is not allowed}}
   c = (float) 1.0h;  // expected-error{{half precision constant requires cl_khr_fp16}}
   b = 1.0h; // expected-error{{half precision constant requires cl_khr_fp16}}
 
   half *allowed = &p[1];
   half *allowed2 = &*p;
   half *allowed3 = p + 1;
 
   return h;
 }
 
 kernel void half_disabled_kernel(global half *p,
-                                 half h);  // expected-error{{declaring function parameter of type 'half' is not allowed; did you forget * ?}}
+                                 half h);  // expected-error{{declaring function parameter of type '__private half' is not allowed; did you forget * ?}}
 
 // Exactly the same as above but with the cl_khr_fp16 extension enabled.
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 constant half a = 1.0h;
 half half_enabled(half *p, half h)
 {
   half a[2];
   half b;

test/SemaOpenCL/images.cl

 // RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 
 void img2d_ro(read_only image2d_t); // expected-note 3{{passing argument to parameter here}}
 void img2d_wo(write_only image2d_t); // expected-note 2{{passing argument to parameter here}}
 void img2d_rw(read_write image2d_t); // expected-note 2{{passing argument to parameter here}}
 void img2d_default(image2d_t); // expected-note 2{{passing argument to parameter here}}
 
 void imgage_access_test(image2d_t img2dro, image3d_t img3dro) {
   img2d_ro(img2dro); // read_only = read_only
-  img2d_ro(img3dro); // expected-error{{passing '__read_only image3d_t' to parameter of incompatible type '__read_only image2d_t'}}
+  img2d_ro(img3dro); // expected-error{{passing '__private __read_only image3d_t' to parameter of incompatible type '__read_only image2d_t'}}
 }
 
 kernel void read_only_access_test(read_only image2d_t img) {
   img2d_ro(img); // read_only = read_only
-  img2d_wo(img); // expected-error {{passing '__read_only image2d_t' to parameter of incompatible type '__write_only image2d_t'}}
-  img2d_rw(img); // expected-error {{passing '__read_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}}
+  img2d_wo(img); // expected-error {{passing '__private __read_only image2d_t' to parameter of incompatible type '__write_only image2d_t'}}
+  img2d_rw(img); // expected-error {{passing '__private __read_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}}
   img2d_default(img); // read_only = read_only
 }
 
 kernel void write_only_access_test(write_only image2d_t img) {
-  img2d_ro(img); // expected-error {{passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
+  img2d_ro(img); // expected-error {{passing '__private __write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
   img2d_wo(img); // write_only = write_only
-  img2d_rw(img); // expected-error {{passing '__write_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}}
-  img2d_default(img); // expected-error {{passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
+  img2d_rw(img); // expected-error {{passing '__private __write_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}}
+  img2d_default(img); // expected-error {{passing '__private __write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
 }
 
 kernel void read_write_access_test(read_write image2d_t img) {
-  img2d_ro(img);  // expected-error {{passing '__read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
-  img2d_wo(img); // expected-error {{passing '__read_write image2d_t' to parameter of incompatible type '__write_only image2d_t'}}
+  img2d_ro(img);  // expected-error {{passing '__private __read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
+  img2d_wo(img); // expected-error {{passing '__private __read_write image2d_t' to parameter of incompatible type '__write_only image2d_t'}}
   img2d_rw(img); //read_write = read_write
-  img2d_default(img); // expected-error {{passing '__read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
+  img2d_default(img); // expected-error {{passing '__private __read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
 }

test/SemaOpenCL/invalid-block.cl

 // RUN: %clang_cc1 -verify -fblocks -cl-std=CL2.0 %s
 
 // OpenCL v2.0 s6.12.5
 void f0(int (^const bl)());
 // All blocks declarations must be const qualified and initialized.
 void f1() {
   int (^bl1)(void) = ^() {
     return 1;
   };
   int (^const bl2)(void) = ^() {
     return 1;
   };
   f0(bl1);
   f0(bl2);
-  bl1 = bl2;          // expected-error{{invalid operands to binary expression ('int (__generic ^const)(void)' and 'int (__generic ^const)(void)')}}
+  bl1 = bl2;          // expected-error{{invalid operands to binary expression ('int (__generic ^const __private)(void)' and 'int (__generic ^const __private)(void)')}}
   int (^const bl3)(); // expected-error{{invalid block variable declaration - must be initialized}}
 }
 
 // A block with extern storage class is not allowed.
 extern int (^bl)(void) = ^() { // expected-error{{invalid block variable declaration - using 'extern' storage class is disallowed}}
   return 1;
 };
 void f2() {

 typedef int (^bl2_t)(int);
 void f5(int i) {
   bl2_t bl1 = ^(int i) {
     return 1;
   };
   bl2_t bl2 = ^(int i) {
     return 2;
   };
-  bl2_t arr[] = {bl1, bl2}; // expected-error {{array of 'bl2_t' (aka 'int (__generic ^const)(int)') type is invalid in OpenCL}}
+  bl2_t arr[] = {bl1, bl2}; // expected-error {{array of 'bl2_t' (aka 'int (__generic ^const)(__private int)') type is invalid in OpenCL}}
   int tmp = i ? bl1(i)      // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}}
               : bl2(i);     // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}}
 }
 // A block pointer type and all pointer operations are disallowed
-void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
+void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(__private int)') is invalid in OpenCL}}
   bl2_t bl = ^(int i) {
     return 1;
   };
-  bl2_t *p; // expected-error {{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
-  *bl;      // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
-  &bl;      // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
+  bl2_t *p; // expected-error {{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(__private int)') is invalid in OpenCL}}
+  *bl;      // expected-error {{invalid argument type '__private bl2_t' (aka 'int (__generic ^const __private)(__private int)') to unary expression}}
+  &bl;      // expected-error {{invalid argument type '__private bl2_t' (aka 'int (__generic ^const __private)(__private int)') to unary expression}}
 }
 // A block can't reference another block
 kernel void f7() {
   bl2_t bl1 = ^(int i) {
     return 1;
   };
   void (^bl2)(void) = ^{
     int i = bl1(1); // expected-error {{cannot refer to a block inside block}}
   };
   void (^bl3)(void) = ^{
   };
   void (^bl4)(void) = ^{
     bl3(); // expected-error {{cannot refer to a block inside block}}
   };
   return;
 }

test/SemaOpenCL/invalid-image.cl

 // RUN: %clang_cc1 -verify %s
 // RUN: %clang_cc1 -verify -D=ATTR_TEST -fms-compatibility %s
 
-void test1(image1d_t *i) {} // expected-error{{pointer to type '__read_only image1d_t' is invalid in OpenCL}}
+void test1(image1d_t *i) {} // expected-error{{pointer to type '__private __read_only image1d_t' is invalid in OpenCL}}
 
 void test2(image1d_t i) {
-  image1d_t ti;            // expected-error{{type '__read_only image1d_t' can only be used as a function parameter}}
+  image1d_t ti;            // expected-error{{type '__private __read_only image1d_t' can only be used as a function parameter}}
   image1d_t ai[] = {i, i}; // expected-error{{array of '__read_only image1d_t' type is invalid in OpenCL}}
-  i=i; // expected-error{{invalid operands to binary expression ('__read_only image1d_t' and '__read_only image1d_t')}}
-  i+1; // expected-error{{invalid operands to binary expression ('__read_only image1d_t' and 'int')}}
-  &i; // expected-error{{invalid argument type '__read_only image1d_t' to unary expression}}
-  *i; // expected-error{{invalid argument type '__read_only image1d_t' to unary expression}}
+  i=i; // expected-error{{invalid operands to binary expression ('__private __read_only image1d_t' and '__private __read_only image1d_t')}}
+  i+1; // expected-error{{invalid operands to binary expression ('__private __read_only image1d_t' and 'int')}}
+  &i; // expected-error{{invalid argument type '__private __read_only image1d_t' to unary expression}}
+  *i; // expected-error{{invalid argument type '__private __read_only image1d_t' to unary expression}}
 }
 
 image1d_t test3() {} // expected-error{{declaring function return value of type '__read_only image1d_t' is not allowed}}
 
 #ifdef ATTR_TEST
 // Test case for an infinite loop bug.
 kernel void foob(read_only __ptr32  image2d_t i) { } // expected-error{{'__ptr32' attribute only applies to pointer arguments}}
 #endif

test/SemaOpenCL/invalid-kernel-parameters.cl

 // RUN: %clang_cc1 -fsyntax-only -verify %s -triple spir-unknown-unknown
 
-kernel void half_arg(half x) { } // expected-error{{declaring function parameter of type 'half' is not allowed; did you forget * ?}}
+kernel void half_arg(half x) { } // expected-error{{declaring function parameter of type '__private half' is not allowed; did you forget * ?}}
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 
 // Disallowed: parameters with type
 // bool, half, size_t, ptrdiff_t, intptr_t, and uintptr_t
 // or a struct / union with any of these types in them
 
 // TODO: Ban int types, size_t, ptrdiff_t ...
 
-kernel void bool_arg(bool x) { } // expected-error{{'bool' cannot be used as the type of a kernel parameter}}
+kernel void bool_arg(bool x) { } // expected-error{{'__private bool' cannot be used as the type of a kernel parameter}}
 
 // half kernel argument is allowed when cl_khr_fp16 is enabled.
 kernel void half_arg(half x) { }
 
 typedef struct ContainsBool // expected-note{{within field of type 'ContainsBool' declared here}}
 {
   bool x; // expected-note{{field of illegal type 'bool' declared here}}
 } ContainsBool;
 
-kernel void bool_in_struct_arg(ContainsBool x) { } // expected-error{{'ContainsBool' (aka 'struct ContainsBool') cannot be used as the type of a kernel parameter}}
+kernel void bool_in_struct_arg(ContainsBool x) { } // expected-error{{'__private ContainsBool' (aka '__private struct ContainsBool') cannot be used as the type of a kernel parameter}}
 
 
 
 typedef struct FooImage2D // expected-note{{within field of type 'FooImage2D' declared here}}
 {
   // TODO: Clean up needed - we don't really need to check for image, event, etc
   // as a note here any longer.
   // They are diagnosed as an error for all struct fields (OpenCL v1.2 s6.9b,r).
   image2d_t imageField; // expected-note{{field of illegal type '__read_only image2d_t' declared here}} expected-error{{the '__read_only image2d_t' type cannot be used to declare a structure or union field}}
 } FooImage2D;
 
 kernel void image_in_struct_arg(FooImage2D arg) { } // expected-error{{struct kernel parameters may not contain pointers}}
 
 typedef struct Foo // expected-note{{within field of type 'Foo' declared here}}
 {
-  int* ptrField; // expected-note{{field of illegal pointer type 'int *' declared here}}
+  int* ptrField; // expected-note{{field of illegal pointer type '__private int *' declared here}}
 } Foo;
 
 kernel void pointer_in_struct_arg(Foo arg) { } // expected-error{{struct kernel parameters may not contain pointers}}
 
 typedef union FooUnion // expected-note{{within field of type 'FooUnion' declared here}}
 {
-  int* ptrField; // expected-note{{field of illegal pointer type 'int *' declared here}}
+  int* ptrField; // expected-note{{field of illegal pointer type '__private int *' declared here}}
 } FooUnion;
 
 kernel void pointer_in_union_arg(FooUnion arg) { }// expected-error{{union kernel parameters may not contain pointers}}
 
 typedef struct NestedPointer // expected-note 2 {{within field of type 'NestedPointer' declared here}}
 {
   int x;
   struct InnerNestedPointer
   {
-    int* ptrField; // expected-note 3 {{field of illegal pointer type 'int *' declared here}}
+    int* ptrField; // expected-note 3 {{field of illegal pointer type '__private int *' declared here}}
   } inner; // expected-note 3 {{within field of type 'struct InnerNestedPointer' declared here}}
 } NestedPointer;
 
 kernel void pointer_in_nested_struct_arg(NestedPointer arg) { }// expected-error{{struct kernel parameters may not contain pointers}}
 
 struct NestedPointerComplex // expected-note{{within field of type 'NestedPointerComplex' declared here}}
 {
   int foo;
   float bar;
 
   struct InnerNestedPointerComplex
   {
     int innerFoo;
-    int* innerPtrField; // expected-note{{field of illegal pointer type 'int *' declared here}}
+    int* innerPtrField; // expected-note{{field of illegal pointer type '__private int *' declared here}}
   } inner; // expected-note{{within field of type 'struct InnerNestedPointerComplex' declared here}}
 
   float y;
   float z[4];
 };
 
 kernel void pointer_in_nested_struct_arg_complex(struct NestedPointerComplex arg) { }// expected-error{{struct kernel parameters may not contain pointers}}
 
 typedef struct NestedBool // expected-note 2 {{within field of type 'NestedBool' declared here}}
 {
   int x;
   struct InnerNestedBool
   {
     bool boolField; // expected-note 2 {{field of illegal type 'bool' declared here}}
   } inner; // expected-note 2 {{within field of type 'struct InnerNestedBool' declared here}}
 } NestedBool;
 
-kernel void bool_in_nested_struct_arg(NestedBool arg) { } // expected-error{{'NestedBool' (aka 'struct NestedBool') cannot be used as the type of a kernel parameter}}
+kernel void bool_in_nested_struct_arg(NestedBool arg) { } // expected-error{{'__private NestedBool' (aka '__private struct NestedBool') cannot be used as the type of a kernel parameter}}
 
 // Warning emitted again for argument used in other kernel
-kernel void bool_in_nested_struct_arg_again(NestedBool arg) { } // expected-error{{'NestedBool' (aka 'struct NestedBool') cannot be used as the type of a kernel parameter}}
+kernel void bool_in_nested_struct_arg_again(NestedBool arg) { } // expected-error{{'__private NestedBool' (aka '__private struct NestedBool') cannot be used as the type of a kernel parameter}}
 
 
 // Check for note with a struct not defined inside the struct
 typedef struct NestedBool2Inner
 {
   bool boolField; // expected-note{{field of illegal type 'bool' declared here}}
 } NestedBool2Inner;
 
 typedef struct NestedBool2 // expected-note{{within field of type 'NestedBool2' declared here}}
 {
   int x;
   NestedBool2Inner inner; // expected-note{{within field of type 'NestedBool2Inner' (aka 'struct NestedBool2Inner') declared here}}
 } NestedBool2;
 
-kernel void bool_in_nested_struct_2_arg(NestedBool2 arg) { } // expected-error{{'NestedBool2' (aka 'struct NestedBool2') cannot be used as the type of a kernel parameter}}
+kernel void bool_in_nested_struct_2_arg(NestedBool2 arg) { } // expected-error{{'__private NestedBool2' (aka '__private struct NestedBool2') cannot be used as the type of a kernel parameter}}
 
 
 struct InnerInner
 {
   int* foo;
   bool x;
 };
 

test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
 
 void test1(read_only pipe int p, global int* ptr){
   int tmp;
   reserve_id_t rid;
 
   // read/write_pipe
   read_pipe(p, &tmp);
   read_pipe(p, ptr);
   read_pipe(tmp, p);    // expected-error {{first argument to 'read_pipe' must be a pipe type}}
   read_pipe(p);   // expected-error {{invalid number of arguments to function: 'read_pipe'}}
   read_pipe(p, rid, tmp, ptr);
-  read_pipe(p, tmp, tmp, ptr);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'reserve_id_t' having 'int')}}
-  read_pipe(p, rid, rid, ptr);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'unsigned int' having 'reserve_id_t')}}
-  read_pipe(p, tmp);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'int *' having 'int')}}
+  read_pipe(p, tmp, tmp, ptr);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'reserve_id_t' having '__private int')}}
+  read_pipe(p, rid, rid, ptr);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'unsigned int' having '__private reserve_id_t')}}
+  read_pipe(p, tmp);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'int *' having '__private int')}}
   write_pipe(p, ptr);    // expected-error {{invalid pipe access modifier (expecting write_only)}}
   write_pipe(p, rid, tmp, ptr);    // expected-error {{invalid pipe access modifier (expecting write_only)}}
 
   // reserve_read/write_pipe
   reserve_read_pipe(p, tmp);
-  reserve_read_pipe(p, ptr);    // expected-error{{invalid argument type to function 'reserve_read_pipe' (expecting 'unsigned int' having '__global int *')}}
+  reserve_read_pipe(p, ptr);    // expected-error{{invalid argument type to function 'reserve_read_pipe' (expecting 'unsigned int' having '__global int *__private')}}
   work_group_reserve_read_pipe(tmp, tmp);    // expected-error{{first argument to 'work_group_reserve_read_pipe' must be a pipe type}}
   sub_group_reserve_write_pipe(p, tmp);    // expected-error{{invalid pipe access modifier (expecting write_only)}}
 
   // commit_read/write_pipe
   commit_read_pipe(p, rid);
   commit_read_pipe(tmp, rid);    // expected-error{{first argument to 'commit_read_pipe' must be a pipe type}}
-  work_group_commit_read_pipe(p, tmp);    // expected-error{{invalid argument type to function 'work_group_commit_read_pipe' (expecting 'reserve_id_t' having 'int')}}
+  work_group_commit_read_pipe(p, tmp);    // expected-error{{invalid argument type to function 'work_group_commit_read_pipe' (expecting 'reserve_id_t' having '__private int')}}
   sub_group_commit_write_pipe(p, tmp);    // expected-error{{invalid pipe access modifier (expecting write_only)}}
 }
 
 void test2(write_only pipe int p, global int* ptr){
   int tmp;
   reserve_id_t rid;
 
   // read/write_pipe
   write_pipe(p, &tmp);
   write_pipe(p, ptr);
   write_pipe(tmp, p);    // expected-error {{first argument to 'write_pipe' must be a pipe type}}
   write_pipe(p);   // expected-error {{invalid number of arguments to function: 'write_pipe'}}
   write_pipe(p, rid, tmp, ptr);
-  write_pipe(p, tmp, tmp, ptr);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'reserve_id_t' having 'int')}}
-  write_pipe(p, rid, rid, ptr);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'unsigned int' having 'reserve_id_t')}}
-  write_pipe(p, tmp);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'int *' having 'int')}}
+  write_pipe(p, tmp, tmp, ptr);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'reserve_id_t' having '__private int')}}
+  write_pipe(p, rid, rid, ptr);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'unsigned int' having '__private reserve_id_t')}}
+  write_pipe(p, tmp);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'int *' having '__private int')}}
   read_pipe(p, ptr);    // expected-error {{invalid pipe access modifier (expecting read_only)}}
   read_pipe(p, rid, tmp, ptr);    // expected-error {{invalid pipe access modifier (expecting read_only)}}
 
   // reserve_read/write_pipe
   reserve_write_pipe(p, tmp);
-  reserve_write_pipe(p, ptr);    // expected-error{{invalid argument type to function 'reserve_write_pipe' (expecting 'unsigned int' having '__global int *')}}
+  reserve_write_pipe(p, ptr);    // expected-error{{invalid argument type to function 'reserve_write_pipe' (expecting 'unsigned int' having '__global int *__private')}}
   work_group_reserve_write_pipe(tmp, tmp);    // expected-error{{first argument to 'work_group_reserve_write_pipe' must be a pipe type}}
   sub_group_reserve_read_pipe(p, tmp);    // expected-error{{invalid pipe access modifier (expecting read_only)}}
 
   // commit_read/write_pipe
   commit_write_pipe(p, rid);
   commit_write_pipe(tmp, rid);    // expected-error{{first argument to 'commit_write_pipe' must be a pipe type}}
-  work_group_commit_write_pipe(p, tmp);    // expected-error{{invalid argument type to function 'work_group_commit_write_pipe' (expecting 'reserve_id_t' having 'int')}}
+  work_group_commit_write_pipe(p, tmp);    // expected-error{{invalid argument type to function 'work_group_commit_write_pipe' (expecting 'reserve_id_t' having '__private int')}}
   sub_group_commit_read_pipe(p, tmp);    // expected-error{{invalid pipe access modifier (expecting read_only)}}
 }
 
 void test3(){
   int tmp;
   get_pipe_num_packets(tmp);    // expected-error {{first argument to 'get_pipe_num_packets' must be a pipe type}}
   get_pipe_max_packets(tmp);    // expected-error {{first argument to 'get_pipe_max_packets' must be a pipe type}}
 }

test/SemaOpenCL/invalid-pipes-cl2.0.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
 
 global pipe int gp;            // expected-error {{type '__global read_only pipe int' can only be used as a function parameter in OpenCL}}
 global reserve_id_t rid;          // expected-error {{the '__global reserve_id_t' type cannot be used to declare a program scope variable}}
 
 extern pipe write_only int get_pipe(); // expected-error {{type '__global write_only pipe int ()' can only be used as a function parameter in OpenCL}}
 
-kernel void test_invalid_reserved_id(reserve_id_t ID) { // expected-error {{'reserve_id_t' cannot be used as the type of a kernel parameter}}
+kernel void test_invalid_reserved_id(reserve_id_t ID) { // expected-error {{'__private reserve_id_t' cannot be used as the type of a kernel parameter}}
 }
 
 void test1(pipe int *p) {// expected-error {{pipes packet types cannot be of reference type}}
 }
 void test2(pipe p) {// expected-error {{missing actual type specifier for pipe}}
 }
 void test3(int pipe p) {// expected-error {{cannot combine with previous 'int' declaration specifier}}
 }
 void test4() {
-  pipe int p; // expected-error {{type 'read_only pipe int' can only be used as a function parameter}}
+  pipe int p; // expected-error {{type '__private read_only pipe int' can only be used as a function parameter}}
   //TODO: fix parsing of this pipe int (*p);
 }
 
 void test5(pipe int p) {
-  p+p; // expected-error{{invalid operands to binary expression ('read_only pipe int' and 'read_only pipe int')}}
-  p=p; // expected-error{{invalid operands to binary expression ('read_only pipe int' and 'read_only pipe int')}}
-  &p; // expected-error{{invalid argument type 'read_only pipe int' to unary expression}}
-  *p; // expected-error{{invalid argument type 'read_only pipe int' to unary expression}}
+  p+p; // expected-error{{invalid operands to binary expression ('__private read_only pipe int' and '__private read_only pipe int')}}
+  p=p; // expected-error{{invalid operands to binary expression ('__private read_only pipe int' and '__private read_only pipe int')}}
+  &p; // expected-error{{invalid argument type '__private read_only pipe int' to unary expression}}
+  *p; // expected-error{{invalid argument type '__private read_only pipe int' to unary expression}}
 }
 
 typedef pipe int pipe_int_t;
 pipe_int_t test6() {} // expected-error{{declaring function return value of type 'pipe_int_t' (aka 'read_only pipe int') is not allowed}}
 
 bool test_id_comprision(void) {
   reserve_id_t id1, id2;
   return (id1 == id2);          // expected-error {{invalid operands to binary expression ('reserve_id_t' and 'reserve_id_t')}}

test/SemaOpenCL/null_literal.cl

 // RUN: %clang_cc1 -verify %s
-// RUN: %clang_cc1 -cl-std=CL2.0 -DCL20 -verify %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -verify %s
 
 #define NULL ((void*)0)
 
 void foo(){
 
 global int* ptr1 = NULL;
 
 global int* ptr2 = (global void*)0;
 
 constant int* ptr3 = NULL;
 
-constant int* ptr4 = (global void*)0; // expected-error{{initializing '__constant int *' with an expression of type '__global void *' changes address space of pointer}}
+constant int* ptr4 = (global void*)0; // expected-error{{initializing '__constant int *__private' with an expression of type '__global void *' changes address space of pointer}}
+
+#if __OPENCL_C_VERSION__ >= 200
 
-#ifdef CL20
 // Accept explicitly pointer to generic address space in OpenCL v2.0.
 global int* ptr5 = (generic void*)0;
+
+global int *ptr = (private void *)0; // expected-error{{initializing '__global int *__private' with an expression of type '__private void *' changes address space of pointer}}
+
+#else
+
+global int *ptr = (private void *)0;
+
 #endif
 
-global int* ptr6 = (local void*)0; // expected-error{{initializing '__global int *' with an expression of type '__local void *' changes address space of pointer}}
+global int* ptr6 = (local void*)0; // expected-error{{initializing '__global int *__private' with an expression of type '__local void *' changes address space of pointer}}
 
 bool cmp = ptr1 == NULL;
 
 cmp = ptr1 == (local void*)0; // expected-error{{comparison between  ('__global int *' and '__local void *') which are pointers to non-overlapping address spaces}}
 
 cmp = ptr3 == NULL;
 
 }

test/SemaOpenCL/null_queue.cl

 // RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 extern queue_t get_default_queue();
 
 bool compare() {
   return 1 == get_default_queue() && // expected-error{{invalid operands to binary expression ('int' and 'queue_t')}}
          get_default_queue() == 1; // expected-error{{invalid operands to binary expression ('queue_t' and 'int')}}
 }
 
 void init() {
-  queue_t q1 = 1; // expected-error{{initializing 'queue_t' with an expression of incompatible type 'int'}}
+  queue_t q1 = 1; // expected-error{{initializing '__private queue_t' with an expression of incompatible type 'int'}}
   queue_t q = 0;
 }

test/SemaOpenCL/queue_t_overload.cl

 // RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 
-void __attribute__((overloadable)) foo(queue_t, __local char *); // expected-note {{candidate function not viable: no known conversion from 'int' to 'queue_t' for 1st argument}} // expected-note {{candidate function}}
-void __attribute__((overloadable)) foo(queue_t, __local float *); // expected-note {{candidate function not viable: no known conversion from 'int' to 'queue_t' for 1st argument}} // expected-note {{candidate function}}
+void __attribute__((overloadable)) foo(queue_t, __local char *); // expected-note {{candidate function not viable: no known conversion from 'int' to '__private queue_t' for 1st argument}} // expected-note {{candidate function}}
+void __attribute__((overloadable)) foo(queue_t, __local float *); // expected-note {{candidate function not viable: no known conversion from 'int' to '__private queue_t' for 1st argument}} // expected-note {{candidate function}}
 
 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(1, src3); // expected-error {{no matching function for call to 'foo'}}
 }

test/SemaOpenCL/sampler_t.cl

 const constant sampler_t glb_smp11 = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
 
 void kernel ker(sampler_t argsmp) {
   local sampler_t smp; // expected-error{{sampler type cannot be used with the __local and __global address space qualifiers}}
   const sampler_t const_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
   const sampler_t const_smp2;
   const sampler_t const_smp3 = const_smp;
   const sampler_t const_smp4 = f();
-  const sampler_t const_smp5 = 1.0f; // expected-error{{initializing 'const sampler_t' with an expression of incompatible type 'float'}}
+  const sampler_t const_smp5 = 1.0f; // expected-error{{initializing 'const __private sampler_t' with an expression of incompatible type 'float'}}
   const sampler_t const_smp6 = 0x100000000LL; // expected-error{{sampler_t initialization requires 32-bit integer, not 'long long'}}
 
   foo(glb_smp);
   foo(glb_smp2);
   foo(glb_smp3);
   foo(glb_smp4);
   foo(glb_smp5);
   foo(glb_smp6);
void kernel ker(sampler_t argsmp) {
   foo(argsmp);
   foo(5);
   foo(5.0f); // expected-error {{passing 'float' to parameter of incompatible type 'sampler_t'}}
   sampler_t sa[] = {argsmp, const_smp}; // expected-error {{array of 'sampler_t' type is invalid in OpenCL}}
   foo(sa[0]);
   foo(bad());
 }
 
-void bad(sampler_t*); // expected-error{{pointer to type 'sampler_t' is invalid in OpenCL}}
+void bad(sampler_t*); // expected-error{{pointer to type '__private sampler_t' is invalid in OpenCL}}
 
 void bar() {
   sampler_t smp1 = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
   sampler_t smp2 = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_NEAREST;
-  smp1=smp2; //expected-error{{invalid operands to binary expression ('sampler_t' and 'sampler_t')}}
-  smp1+1; //expected-error{{invalid operands to binary expression ('sampler_t' and 'int')}}
-  &smp1; //expected-error{{invalid argument type 'sampler_t' to unary expression}}
-  *smp2; //expected-error{{invalid argument type 'sampler_t' to unary expression}}
-  foo(smp1+1); //expected-error{{invalid operands to binary expression ('sampler_t' and 'int')}}
+  smp1=smp2; //expected-error{{invalid operands to binary expression ('__private sampler_t' and '__private sampler_t')}}
+  smp1+1; //expected-error{{invalid operands to binary expression ('__private sampler_t' and 'int')}}
+  &smp1; //expected-error{{invalid argument type '__private sampler_t' to unary expression}}
+  *smp2; //expected-error{{invalid argument type '__private sampler_t' to unary expression}}
+  foo(smp1+1); //expected-error{{invalid operands to binary expression ('__private sampler_t' and 'int')}}
 }
 

test/SemaOpenCL/shifts.cl

void ptest07() {
   char3 w = {1,2,3};
 
   v <<= w;
 }
 
 // ** Negative tests **
 
 char2 ntest01(char c, char2 s) {
-  return c << s; // expected-error {{requested shift is a vector of type 'char2' (vector of 2 'char' values) but the first operand is not a vector ('char')}}
+  return c << s; // expected-error {{requested shift is a vector of type '__private char2' (vector of 2 'char' values) but the first operand is not a vector ('__private char')}}
 }
 
 char3 ntest02(char3 c, char2 s) {
   return c << s; // expected-error {{vector operands do not have the same number of elements ('char3' (vector of 3 'char' values) and 'char2' (vector of 2 'char' values))}}
 }
 
 float2 ntest03(float2 c, char s) {
   return c << s; // expected-error {{used type 'float2' (vector of 2 'float' values) where integer is required}}
 }
 
 int2 ntest04(int2 c, float s) {
   return c << s; // expected-error {{used type 'float' where integer is required}}
 }

test/SemaOpenCL/to_addr_builtin.cl

 // RUN: %clang_cc1 -verify -fsyntax-only %s
 // RUN: %clang_cc1 -verify -fsyntax-only -cl-std=CL2.0 %s
 
 void test(void) {
   global int *glob;
   local int *loc;
   constant int *con;
   typedef constant int const_int_ty;
   const_int_ty *con_typedef;
 
   glob = to_global(glob, loc);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
   // expected-error@-2{{implicit declaration of function 'to_global' is invalid in OpenCL}}
-  // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error@-5{{invalid number of arguments to function: 'to_global'}}
 #endif
 
   int x;
   glob = to_global(x);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error@-4{{invalid argument x to function: 'to_global', expecting a generic pointer argument}}
 #endif
 
   glob = to_global(con);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error@-4{{invalid argument con to function: 'to_global', expecting a generic pointer argument}}
 #endif
 
   glob = to_global(con_typedef);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error@-4{{invalid argument con_typedef to function: 'to_global', expecting a generic pointer argument}}
 #endif
 
   loc = to_global(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}}
 #else
-  // expected-error@-4{{assigning '__global int *' to '__local int *' changes address space of pointer}}
+  // expected-error@-4{{assigning '__global int *' to '__local int *__private' changes address space of pointer}}
 #endif
 
   global char *glob_c = to_global(loc);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion initializing '__global char *' with an expression of type 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion initializing '__global char *__private' with an expression of type 'int'}}
 #else
-  // expected-warning@-4{{incompatible pointer types initializing '__global char *' with an expression of type '__global int *'}}
+  // expected-warning@-4{{incompatible pointer types initializing '__global char *__private' with an expression of type '__global int *'}}
 #endif
 
 }

test/SemaOpenCL/vec_step.cl

void foo(int3 arg1, int8 arg2) {
   int res6[vec_step(auto4) == 2 ? 1 : -1];
   int res7[vec_step(int2) == 2 ? 1 : -1];
   int res8[vec_step(int3) == 4 ? 1 : -1];
   int res9[vec_step(int4) == 4 ? 1 : -1];
   int res10[vec_step(int8) == 8 ? 1 : -1];
   int res11[vec_step(int16) == 16 ? 1 : -1];
   int res12[vec_step(void) == 1 ? 1 : -1];
 
-  int res13 = vec_step(*incomplete1); // expected-error {{'vec_step' requires built-in scalar or vector type, 'struct S' invalid}}
-  int res14 = vec_step(int16*); // expected-error {{'vec_step' requires built-in scalar or vector type, 'int16 *' invalid}}
+  int res13 = vec_step(*incomplete1); // expected-error {{'vec_step' requires built-in scalar or vector type, '__private struct S' invalid}}
+  int res14 = vec_step(int16 *); // expected-error {{'vec_step' requires built-in scalar or vector type, '__private int16 *' invalid}}
   int res15 = vec_step(void(void)); // expected-error {{'vec_step' requires built-in scalar or vector type, 'void (void)' invalid}}
 }

test/SemaOpenCL/vector_conv_invalid.cl

 // RUN: %clang_cc1 -verify %s
 
 typedef unsigned int uint4 __attribute((ext_vector_type(4)));
 typedef int int4 __attribute((ext_vector_type(4)));
 typedef int int3 __attribute((ext_vector_type(3)));
 typedef unsigned uint3 __attribute((ext_vector_type(3)));
 
 void vector_conv_invalid() {
   uint4 u = (uint4)(1);
-  int4 i = u; // expected-error{{initializing 'int4' (vector of 4 'int' values) with an expression of incompatible type 'uint4' (vector of 4 'unsigned int' values)}}
+  int4 i = u; // expected-error{{initializing '__private int4' (vector of 4 'int' values) with an expression of incompatible type '__private uint4' (vector of 4 'unsigned int' values)}}
   int4 e = (int4)u; // expected-error{{invalid conversion between ext-vector type 'int4' (vector of 4 'int' values) and 'uint4' (vector of 4 'unsigned int' values)}}
 
   uint3 u4 = (uint3)u; // expected-error{{invalid conversion between ext-vector type 'uint3' (vector of 3 'unsigned int' values) and 'uint4' (vector of 4 'unsigned int' values)}}
 }