[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 ↗

(On Diff #105785)

Would we still be accepting the following:

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

lib/Sema/SemaDecl.cpp

7964 ↗

(On Diff #105785)

Could we use LangAS::Default instead?

11846 ↗

(On Diff #105785)

Could we use LangAS::Default here too?

11851 ↗

(On Diff #105785)

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

lib/Sema/SemaType.cpp

6969 ↗

(On Diff #105785)

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

23 ↗

(On Diff #105785)

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 ↗

(On Diff #105785)

Yes. There is a test SemaOpenCL/null_literal.cl

lib/Sema/SemaDecl.cpp

7964 ↗

(On Diff #105785)

done

11846 ↗

(On Diff #105785)

done

11851 ↗

(On Diff #105785)

will do.

lib/Sema/SemaType.cpp

6969 ↗

(On Diff #105785)

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

23 ↗

(On Diff #105785)

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 ↗

(On Diff #105785)

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 ↗

(On Diff #105785)

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

23 ↗

(On Diff #105785)

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 ↗

(On Diff #105785)

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 ↗

(On Diff #105785)

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 ↗

(On Diff #105785)

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 ↗

(On Diff #105785)

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 ↗

(On Diff #105785)

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 ↗

(On Diff #105785)

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 ↗

(On Diff #105785)

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 ↗

(On Diff #112288)

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

3257 ↗

(On Diff #112288)

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

lib/Sema/SemaType.cpp

6974 ↗

(On Diff #112288)

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 ↗

(On Diff #112288)

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 ↗

(On Diff #112288)

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 ↗

(On Diff #112288)

'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 ↗

(On Diff #112288)

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 ↗

(On Diff #112288)

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 ↗

(On Diff #112288)

Cool! Thanks!

[yaxunl]

Rebase to ToT and clean up logic.

[yaxunl]

lib/Sema/SemaType.cpp

6974 ↗

(On Diff #112288)

fixed.

[Anastasia]

lib/Sema/SemaType.cpp

6808 ↗

(On Diff #116704)

Great! This looks so clear now!

6810 ↗

(On Diff #116704)

I think this could be checked before calling the function.

6863 ↗

(On Diff #116704)

Do we cover extern too?

6872 ↗

(On Diff #116704)

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 ↗

(On Diff #116704)

will do.

6863 ↗

(On Diff #116704)

will add.

6872 ↗

(On Diff #116704)

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 ↗

(On Diff #117020)

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

[Anastasia]

lib/Sema/SemaType.cpp

6872 ↗

(On Diff #116704)

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 ↗

(On Diff #117020)

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

34 ↗

(On Diff #117770)

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

34 ↗

(On Diff #117770)

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

34 ↗

(On Diff #117770)

Thanks, WFM.

lib/AST/TypePrinter.cpp

1703 ↗

(On Diff #118813)

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

lib/Sema/SemaType.cpp

5754 ↗

(On Diff #118813)

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

cfe/trunk/include/clang/Basic/AddressSpaces.h

 
 namespace LangAS {
 
 /// \brief Defines the address space values used by the address space qualifier
 /// of QualType.
 ///
 enum ID {
   // The default value 0 is the value used in QualType for the the situation
-  // where there is no address space qualifier. For most languages, this also
-  // corresponds to the situation where there is no address space qualifier in
-  // the source code, except for OpenCL, where the address space value 0 in
-  // QualType represents private address space in OpenCL source code.
+  // where there is no address space qualifier.
   Default = 0,
 
   // OpenCL specific address spaces.
+  // In OpenCL each l-value must have certain non-default address space, each
+  // r-value must have no address space (i.e. the default address space). The
+  // pointee of a pointer must have non-default address space.
   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

cfe/trunk/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();

cfe/trunk/lib/AST/Expr.cpp

Expr::isNullPointerConstant(ASTContext &Ctx,
   }
 
   // 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.
+        // (__generic void*)0 in OpenCL 2.0 should not be treated as nullptr
+        // since it cannot be assigned to a pointer to constant address space.
+        bool PointeeHasDefaultAS =
+            Pointee.getAddressSpace() == LangAS::Default ||
+            (Ctx.getLangOpts().OpenCLVersion >= 200 &&
+             Pointee.getAddressSpace() == LangAS::opencl_generic) ||
+            (Ctx.getLangOpts().OpenCL &&
+             Ctx.getLangOpts().OpenCLVersion < 200 &&
+             Pointee.getAddressSpace() == LangAS::opencl_private);
 
-        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

cfe/trunk/lib/AST/ItaniumMangle.cpp

if (Quals.hasAddressSpace()) {
     //   <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);
+      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;
       }
     }
+    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"

cfe/trunk/lib/AST/TypePrinter.cpp

void Qualifiers::print(raw_ostream &OS, const PrintingPolicy& Policy,
   }
   if (hasUnaligned()) {
     if (addSpace)
       OS << ' ';
     OS << "__unaligned";
     addSpace = true;
   }
   if (unsigned addrspace = getAddressSpace()) {
+    if (addrspace != LangAS::opencl_private) {
       if (addSpace)
         OS << ' ';
       addSpace = true;
       switch (addrspace) {
       case LangAS::opencl_global:
         OS << "__global";
         break;
       case LangAS::opencl_local:
         OS << "__local";
         break;
+      case LangAS::opencl_private:
+        break;
       case LangAS::opencl_constant:
       case LangAS::cuda_constant:
         OS << "__constant";
         break;
       case LangAS::opencl_generic:
         OS << "__generic";
         break;
       case LangAS::cuda_device:
         OS << "__device";
         break;
       case LangAS::cuda_shared:
         OS << "__shared";
         break;
       default:
         assert(addrspace >= LangAS::FirstTargetAddressSpace);
         OS << "__attribute__((address_space(";
         OS << addrspace - LangAS::FirstTargetAddressSpace;
         OS << ")))";
       }
     }
+  }
   if (Qualifiers::GC gc = getObjCGCAttr()) {
     if (addSpace)
       OS << ' ';
     addSpace = true;
     if (gc == Qualifiers::Weak)
       OS << "__weak";
     else
       OS << "__strong";

cfe/trunk/lib/Basic/Targets/AMDGPU.cpp

static const char *const DataLayoutStringSIGenericIsZero =
     "-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128"
     "-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5";
 
 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
 };
 } // namespace targets
 } // namespace clang
 

cfe/trunk/lib/Basic/Targets/NVPTX.h

 namespace clang {
 namespace targets {
 
 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 LLVM_LIBRARY_VISIBILITY NVPTXTargetInfo : public TargetInfo {

cfe/trunk/lib/Basic/Targets/SPIR.h

 namespace clang {
 namespace targets {
 
 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 LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public TargetInfo {
 public:

cfe/trunk/lib/Basic/Targets/TCE.h

 // 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 LLVM_LIBRARY_VISIBILITY TCETargetInfo : public TargetInfo {

cfe/trunk/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);
 

cfe/trunk/lib/Sema/SemaChecking.cpp

static bool SemaOpenCLBuiltinNDRangeAndBlock(Sema &S, CallExpr *TheCall) {
   if (checkArgCount(S, TheCall, 2))
     return true;
 
   if (checkOpenCLSubgroupExt(S, TheCall))
     return true;
 
   // First argument is an ndrange_t type.
   Expr *NDRangeArg = TheCall->getArg(0);
-  if (NDRangeArg->getType().getAsString() != "ndrange_t") {
+  if (NDRangeArg->getType().getUnqualifiedType().getAsString() != "ndrange_t") {
     S.Diag(NDRangeArg->getLocStart(),
            diag::err_opencl_builtin_expected_type)
         << TheCall->getDirectCallee() << "'ndrange_t'";
     return true;
   }
 
   Expr *BlockArg = TheCall->getArg(1);
   if (!isBlockPointer(BlockArg)) {
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();
+    llvm_unreachable("Invalid builtin function");
   }
   Call->setType(S.Context.getPointerType(S.Context.getQualifiedType(
       RT.getUnqualifiedType(), Qual)));
 
   return false;
 }
 
 ExprResult

cfe/trunk/lib/Sema/SemaDecl.cpp

if (R->isSamplerT()) {
         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());
if (NewVD->isFileVarDecl() || NewVD->isStaticLocal() ||
                   << "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())
Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC,
 
   // 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.
if (T->isObjCObjectType()) {
     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))) {
 
   return New;
 }
 
 void Sema::ActOnFinishKNRParamDeclarations(Scope *S, Declarator &D,
                                            SourceLocation LocAfterDecls) {
   DeclaratorChunk::FunctionTypeInfo &FTI = D.getFunctionTypeInfo();
 

cfe/trunk/lib/Sema/SemaType.cpp

 TypeSourceInfo *Sema::GetTypeForDeclarator(Declarator &D, Scope *S) {
   // Determine the type of the declarator. Not all forms of declarator
   // have a type.
 
   TypeProcessingState state(*this, D);
 
   TypeSourceInfo *ReturnTypeInfo = nullptr;
   QualType T = GetDeclSpecTypeForDeclarator(state, ReturnTypeInfo);
-
   if (D.isPrototypeContext() && getLangOpts().ObjCAutoRefCount)
     inferARCWriteback(state, T);
 
   return GetFullTypeForDeclarator(state, T, ReturnTypeInfo);
 }
 
 static void transferARCOwnershipToDeclSpec(Sema &S,
                                            QualType &declSpecTy,
if (Attr.getKind() == AttributeList::AT_AddressSpace) {
     case AttributeList::AT_OpenCLGlobalAddressSpace:
       ASIdx = LangAS::opencl_global; break;
     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;
+    case AttributeList::AT_OpenCLPrivateAddressSpace:
+      ASIdx = LangAS::opencl_private; break;
     default:
-      assert(Attr.getKind() == AttributeList::AT_OpenCLPrivateAddressSpace);
-      ASIdx = 0; break;
+      llvm_unreachable("Invalid address space");
     }
 
     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
static void HandleOpenCLAccessAttr(QualType &CurType, const AttributeList &Attr,
   } else if (CurType->isPipeType()) {
     if (Attr.getSemanticSpelling() == OpenCLAccessAttr::Keyword_write_only) {
       QualType ElemType = CurType->getAs<PipeType>()->getElementType();
       CurType = S.Context.getWritePipeType(ElemType);
     }
   }
 }
 
+static void deduceOpenCLImplicitAddrSpace(TypeProcessingState &State,
+                                          QualType &T, TypeAttrLocation TAL) {
+  Declarator &D = State.getDeclarator();
+
+  // Handle the cases where address space should not be deduced.
+  //
+  // The pointee type of a pointer type is alwasy deduced since a pointer always
+  // points to some memory location which should has an address space.
+  //
+  // There are situations that at the point of certain declarations, the address
+  // space may be unknown and better to be left as default. For example, when
+  // definining a typedef or struct type, they are not associated with any
+  // specific address space. Later on, they may be used with any address space
+  // to declare a variable.
+  //
+  // The return value of a function is r-value, therefore should not have
+  // address space.
+  //
+  // The void type does not occupy memory, therefore should not have address
+  // space, except when it is used as a pointee type.
+  //
+  // Since LLVM assumes function type is in default address space, it should not
+  // have address space.
+  auto ChunkIndex = State.getCurrentChunkIndex();
+  bool IsPointee =
+      ChunkIndex > 0 &&
+      (D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Pointer ||
+       D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer);
+  bool IsFuncReturnType =
+      ChunkIndex > 0 &&
+      D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Function;
+  bool IsFuncType =
+      ChunkIndex < D.getNumTypeObjects() &&
+      D.getTypeObject(ChunkIndex).Kind == DeclaratorChunk::Function;
+  if ( // Do not deduce addr space for function return type and function type,
+       // otherwise it will fail some sema check.
+      IsFuncReturnType || IsFuncType ||
+      // Do not deduce addr space for member types of struct, except the pointee
+      // type of a pointer member type.
+      (D.getContext() == Declarator::MemberContext && !IsPointee) ||
+      // Do not deduce addr space for types used to define a typedef and the
+      // typedef itself, except the pointee type of a pointer type which is used
+      // to define the typedef.
+      (D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_typedef &&
+       !IsPointee) ||
+      // Do not deduce addr space of the void type, e.g. in f(void), otherwise
+      // it will fail some sema check.
+      (T->isVoidType() && !IsPointee))
+    return;
+
+  unsigned ImpAddr;
+  // Put OpenCL automatic variable in private address space.
+  // OpenCL v1.2 s6.5:
+  // The default 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.
+  if (State.getSema().getLangOpts().OpenCLVersion <= 120) {
+      ImpAddr = LangAS::opencl_private;
+  } else {
+    // 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 (IsPointee) {
+      ImpAddr = LangAS::opencl_generic;
+    } else {
+      if (D.getContext() == Declarator::FileContext) {
+        ImpAddr = LangAS::opencl_global;
+      } else {
+        if (D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_static ||
+            D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_extern) {
+          ImpAddr = LangAS::opencl_global;
+        } else {
+          ImpAddr = LangAS::opencl_private;
+        }
+      }
+    }
+  }
+  T = State.getSema().Context.getAddrSpaceQualType(T, ImpAddr);
+}
+
 static void processTypeAttrs(TypeProcessingState &state, QualType &type,
                              TypeAttrLocation TAL, AttributeList *attrs) {
   // Scan through and apply attributes to this type where it makes sense.  Some
   // attributes (such as __address_space__, __vector_size__, etc) apply to the
   // type, but others can be present in the type specifiers even though they
   // apply to the decl.  Here we apply type attributes and ignore the rest.
 
   bool hasOpenCLAddressSpace = false;
FUNCTION_TYPE_ATTRS_CASELIST:
 
       // Otherwise, handle the possible delays.
       else if (!handleFunctionTypeAttr(state, attr, type))
         distributeFunctionTypeAttr(state, attr, type);
       break;
     }
   }
 
-  // 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, 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);
-  }
+  if (!state.getSema().getLangOpts().OpenCL ||
+      type.getAddressSpace() != LangAS::Default)
+    return;
+
+  deduceOpenCLImplicitAddrSpace(state, type, TAL);
 }
 
 void Sema::completeExprArrayBound(Expr *E) {
   if (DeclRefExpr *DRE = dyn_cast<DeclRefExpr>(E->IgnoreParens())) {
     if (VarDecl *Var = dyn_cast<VarDecl>(DRE->getDecl())) {
       if (isTemplateInstantiation(Var->getTemplateSpecializationKind())) {
         SourceLocation PointOfInstantiation = E->getExprLoc();
 

cfe/trunk/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

cfe/trunk/test/CodeGenOpenCL/address-spaces.cl

 // RUN: %clang_cc1 %s -O0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,SPIR
 // RUN: %clang_cc1 %s -O0 -DCL20 -cl-std=CL2.0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20SPIR
 // RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa-opencl -emit-llvm -o - | FileCheck --check-prefixes=CHECK,SPIR %s
 // RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa-opencl -DCL20 -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20SPIR
 // RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa-amdgizcl -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,GIZ
 // RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa-amdgizcl -DCL20 -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20GIZ
 // RUN: %clang_cc1 %s -O0 -triple amdgcn-mesa-mesa3d -emit-llvm -o - | FileCheck --check-prefixes=CHECK,SPIR %s
 // RUN: %clang_cc1 %s -O0 -triple r600-- -emit-llvm -o - | FileCheck --check-prefixes=CHECK,SPIR %s
 
+// SPIR: %struct.S = type { i32, i32, i32* }
+// CL20SPIR: %struct.S = type { i32, i32, i32 addrspace(4)* }
+struct S {
+  int x;
+  int y;
+  int *z;
+};
+
+// CL20-DAG: @g_extern_var = external addrspace(1) global float
+// CL20-DAG: @l_extern_var = external addrspace(1) global float
+// CL20-DAG: @test_static.l_static_var = internal addrspace(1) global float 0.000000e+00
+// CL20-DAG: @g_static_var = internal addrspace(1) global float 0.000000e+00
+
+#ifdef CL20
+// CL20-DAG: @g_s = common addrspace(1) global %struct.S zeroinitializer
+struct S g_s;
+#endif
+
 // SPIR: i32* %arg
 // GIZ: i32 addrspace(5)* %arg
 void f__p(__private int *arg) {}
 
 // CHECK: i32 addrspace(1)* %arg
 void f__g(__global int *arg) {}
 
 // CHECK: i32 addrspace(3)* %arg

 // CL20SPIR-DAG: %i = alloca i32,
 // CL20GIZ-DAG: %i = alloca i32{{.*}}addrspace(5)
 
 #ifdef CL20
   static int ii;
 // CL20-DAG: @f.ii = internal addrspace(1) global i32 0
 #endif
 }
+
+typedef int int_td;
+typedef int *intp_td;
+// SPIR: define void @test_typedef(i32 addrspace(1)* %x, i32 addrspace(2)* %y, i32* %z)
+void test_typedef(global int_td *x, constant int_td *y, intp_td z) {
+  *x = *y;
+  *z = 0;
+}
+
+// SPIR: define void @test_struct()
+void test_struct() {
+  // SPIR: %ps = alloca %struct.S*
+  // CL20SPIR: %ps = alloca %struct.S addrspace(4)*
+  struct S *ps;
+  // SPIR: store i32 0, i32* %x
+  // CL20SPIR: store i32 0, i32 addrspace(4)* %x
+  ps->x = 0;
+#ifdef CL20
+  // CL20SPIR: store i32 0, i32 addrspace(1)* getelementptr inbounds (%struct.S, %struct.S addrspace(1)* @g_s, i32 0, i32 0)
+  g_s.x = 0;
+#endif
+}
+
+// SPIR-LABEL: define void @test_void_par()
+void test_void_par(void) {}
+
+// SPIR-LABEL: define i32 @test_func_return_type()
+int test_func_return_type(void) {
+  return 0;
+}
+
+#ifdef CL20
+extern float g_extern_var;
+
+// CL20-LABEL: define {{.*}}void @test_extern(
+kernel void test_extern(global float *buf) {
+  extern float l_extern_var;
+  buf[0] += g_extern_var + l_extern_var;
+}
+
+static float g_static_var;
+
+// CL20-LABEL: define {{.*}}void @test_static(
+kernel void test_static(global float *buf) {
+  static float l_static_var;
+  buf[0] += g_static_var + l_static_var;
+}
+
+#endif

cfe/trunk/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;
+#if __OPENCL_C_VERSION__ < 200
   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}}
+#else
+  ip = gip;
+  ip = &li;
+  ip = &ci; // expected-error {{assigning '__constant int *' to '__generic int *' 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}}

 }
 
 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}}
+}

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

kernel void foo(global int *buf)
   ndrange_t n;
   buf[0] = get_kernel_max_sub_group_size_for_ndrange(n, ^(){});
   buf[0] = get_kernel_max_sub_group_size_for_ndrange(0, ^(){}); // expected-error{{illegal call to 'get_kernel_max_sub_group_size_for_ndrange', expected 'ndrange_t' argument type}}
   buf[0] = get_kernel_max_sub_group_size_for_ndrange(n, 1); // expected-error{{illegal call to 'get_kernel_max_sub_group_size_for_ndrange', expected block argument type}}
 }
 
 kernel void bar(global int *buf)
 {
-  ndrange_t n;
+  __private ndrange_t n;
   buf[0] = get_kernel_sub_group_count_for_ndrange(n, ^(){});
   buf[0] = get_kernel_sub_group_count_for_ndrange(0, ^(){}); // expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', expected 'ndrange_t' argument type}}
   buf[0] = get_kernel_sub_group_count_for_ndrange(n, 1); // expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', expected block argument type}}
 }
 
 #pragma OPENCL EXTENSION cl_khr_subgroups : disable
 
 kernel void foo1(global int *buf)

cfe/trunk/test/SemaOpenCL/extern.cl

-// RUN: %clang_cc1 -x cl -cl-opt-disable -cl-std=CL1.2 -emit-llvm -ffake-address-space-map %s -o - -verify | FileCheck %s
-// expected-no-diagnostics
-
-// CHECK: @foo = external addrspace(2) constant float
-extern constant float foo;
-
-kernel void test(global float* buf) {
-  buf[0] += foo;
-}

cfe/trunk/test/SemaOpenCL/storageclass-cl20.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
 
-static constant int G1 = 0;
 int G2 = 0;
 global int G3 = 0;
 local int G4 = 0;              // expected-error{{program scope variable must reside in global or constant address space}}
 
-void kernel foo() {
-  static int S1 = 5;
-  static global int S2 = 5;
-  static private int S3 = 5;   // expected-error{{static local variable must reside in global or constant address space}}
+static float g_implicit_static_var = 0;
+static constant float g_constant_static_var = 0;
+static global float g_global_static_var = 0;
+static local float g_local_static_var = 0;     // expected-error {{program scope variable must reside in global or constant address space}}
+static private float g_private_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
+static generic float g_generic_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
+
+extern float g_implicit_extern_var;
+extern constant float g_constant_extern_var;
+extern global float g_global_extern_var;
+extern local float g_local_extern_var;     // expected-error {{extern variable must reside in global or constant address space}}
+extern private float g_private_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
+extern generic float g_generic_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
 
+void kernel foo() {
   constant int L1 = 0;
   local int L2;
   global int L3;                              // expected-error{{function scope variable cannot be declared in global address space}}
   generic int L4;                             // expected-error{{automatic variable qualified with an invalid address space}}
   __attribute__((address_space(100))) int L5; // expected-error{{automatic variable qualified with an invalid address space}}
 
-  extern global int G5;
-  extern int G6; // expected-error{{extern variable must reside in global or constant address space}}
+  static float l_implicit_static_var = 0;
+  static constant float l_constant_static_var = 0;
+  static global float l_global_static_var = 0;
+  static local float l_local_static_var = 0;     // expected-error {{static local variable must reside in global or constant address space}}
+  static private float l_private_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
+  static generic float l_generic_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
+
+  extern float l_implicit_extern_var;
+  extern constant float l_constant_extern_var;
+  extern global float l_global_extern_var;
+  extern local float l_local_extern_var;     // expected-error {{extern variable must reside in global or constant address space}}
+  extern private float l_private_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
+  extern generic float l_generic_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
 }

cfe/trunk/test/SemaOpenCL/storageclass.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL1.2
 
 static constant int G1 = 0;
 constant int G2 = 0;
 int G3 = 0;        // expected-error{{program scope variable must reside in constant address space}}
 global int G4 = 0; // expected-error{{program scope variable must reside in constant address space}}
 
+static float g_implicit_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
+static constant float g_constant_static_var = 0;
+static global float g_global_static_var = 0;   // expected-error {{program scope variable must reside in constant address space}}
+static local float g_local_static_var = 0;     // expected-error {{program scope variable must reside in constant address space}}
+static private float g_private_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
+static generic float g_generic_static_var = 0; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{program scope variable must reside in constant address space}}
+
+extern float g_implicit_extern_var; // expected-error {{extern variable must reside in constant address space}}
+extern constant float g_constant_extern_var;
+extern global float g_global_extern_var;   // expected-error {{extern variable must reside in constant address space}}
+extern local float g_local_extern_var;     // expected-error {{extern variable must reside in constant address space}}
+extern private float g_private_extern_var; // expected-error {{extern variable must reside in constant address space}}
+extern generic float g_generic_extern_var; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{extern variable must reside in constant address space}}
+
 void kernel foo(int x) {
   // static is not allowed at local scope before CL2.0
   static int S1 = 5;          // expected-error{{variables in function scope cannot be declared static}}
   static constant int S2 = 5; // expected-error{{variables in function scope cannot be declared static}}
 
   constant int L1 = 0;
   local int L2;
 
void f() {
 
   {
     constant int L1 = 0;                        // expected-error{{non-kernel function variable cannot be declared in constant address space}}
     local int L2;                               // expected-error{{non-kernel function variable cannot be declared in local address space}}
     global int L3;                              // expected-error{{function scope variable cannot be declared in global address space}}
     __attribute__((address_space(100))) int L4; // expected-error{{automatic variable qualified with an invalid address space}}
   }
 
-
-  extern constant float L5;
-  extern local float L6; // expected-error{{extern variable must reside in constant address space}}
-
-  static int L7 = 0;     // expected-error{{variables in function scope cannot be declared static}}
-  static int L8;         // expected-error{{variables in function scope cannot be declared static}}
+  static float l_implicit_static_var = 0;          // expected-error {{variables in function scope cannot be declared static}}
+  static constant float l_constant_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
+  static global float l_global_static_var = 0;     // expected-error {{variables in function scope cannot be declared static}}
+  static local float l_local_static_var = 0;       // expected-error {{variables in function scope cannot be declared static}}
+  static private float l_private_static_var = 0;   // expected-error {{variables in function scope cannot be declared static}}
+  static generic float l_generic_static_var = 0;   // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{variables in function scope cannot be declared static}}
+
+  extern float l_implicit_extern_var; // expected-error {{extern variable must reside in constant address space}}
+  extern constant float l_constant_extern_var;
+  extern global float l_global_extern_var;   // expected-error {{extern variable must reside in constant address space}}
+  extern local float l_local_extern_var;     // expected-error {{extern variable must reside in constant address space}}
+  extern private float l_private_extern_var; // expected-error {{extern variable must reside in constant address space}}
+  extern generic float l_generic_extern_var; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{extern variable must reside in constant address space}}
 }

cfe/trunk/test/SemaTemplate/address_space-dependent.cpp

 
 template <int I>
 void neg() {
   __attribute__((address_space(I))) int *bounds; // expected-error {{address space is negative}}
 }
 
 template <long int I>
 void tooBig() {
-  __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388599)}}
+  __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388598)}}
 }
 
 template <long int I>
 void correct() {
   __attribute__((address_space(I))) int *bounds;
 }
 
 template <int I, int J>

 
 int main() {  
   int __attribute__((address_space(1))) * p1;
   int p = GetAddressSpaceValue(p1);
 
   car<1, 2, 3>(); // expected-note {{in instantiation of function template specialization 'car<1, 2, 3>' requested here}}
   HasASTemplateFields<1> HASTF;
   neg<-1>(); // expected-note {{in instantiation of function template specialization 'neg<-1>' requested here}}
-  correct<0x7FFFF7>();
+  correct<0x7FFFF6>();
   tooBig<8388650>(); // expected-note {{in instantiation of function template specialization 'tooBig<8388650>' requested here}}
 
   __attribute__((address_space(1))) char *x;
   __attribute__((address_space(2))) char *y;
   cmp<1, 2>(x, y); // expected-note {{in instantiation of function template specialization 'cmp<1, 2>' requested here}}
 
   fooFunction<1> ff;
   ff.get_0(); // expected-note {{in instantiation of member function 'fooFunction<1>::get_0' requested here}}
   ff.qf();
   ff.test3(); // expected-note {{in instantiation of member function 'fooFunction<1>::test3' requested here}}
   
   static_assert(partial_spec_deduce_as<int __attribute__((address_space(3))) *>::value == 3, "address space value has been incorrectly deduced"); 
 
   return 0;
 }