[SYCL] Re-use OpenCL address space attributes for SYCL

[bader]

Today we re-use OpenCL parsed attributes, but have separate SYCL address
space semantic attributes as current implementation of OpenCL semantics
breaks valid C++. This patch enables re-use of OpenCL semantic
attributes by allowing conversions between types qualified with OpenCL
address spaces and type w/o address space qualifiers. Clang compiler
(almost) always adds address space qualifiers in OpenCL mode, so it
should not affect OpenCL mode.

NOTE: this change also disables implicit conversion between the
unqualified types and types qualified with
__attribute__((address_space(N))), enabled by one of the previous SYCL
patches.

Signed-off-by: Alexey Bader alexey.bader@intel.com

[bader]

@AnastasiaStulova, I'd like to check with you if this is a bug in C++ for OpenCL compiler or not.
NOTE, this patch enables conversion to from/to unqualified types and as OpenCL qualifies all the types it should not be affected.
All OpenCL tests pass except this one test case.
Is it a bug in function pointer type inference for lambdas?

[AnastasiaStulova]

Can you explain why you think this is a bug?

In OpenCL all local variables are deduced to be in __private addr space so it is the case for function object priv3, however its call operator has __global qualifier and hence can only be used with function objects in __global addr space because no other addr space can implicitly be converted to __global. So giving an error on the call to this call operator seems correct.

[AnastasiaStulova]

I think the issue might be due to your downstream change in isAddressSpaceSupersetOf that I don't entirely understand.

[bader]

isAddressSpaceSupersetOf allow conversion between unqualified types and types qualified with opencl_global, opencl_local and opencl_private address space qualifiers.
As I mentioned above it should have zero impact on OpenCL compiler, which is supposed to qualify ALL types with address space qualifier.

In OpenCL all local variables are deduced to be in __private addr space so it is the case for function object priv3, however its call operator has __global qualifier and hence can only be used with function objects in __global addr space because no other addr space can implicitly be converted to __global. So giving an error on the call to this call operator seems correct.

It doesn't seem to work as you described if address space is deduced correctly, the mismatch will still be reported. It looks like compiler relies somehow on isAddressSpaceSupersetOf in deduction process and priv3 type is not qualified with the address space. I'll check this assumption.

[AnastasiaStulova]

It doesn't seem to work as you described if address space is deduced correctly, the mismatch will still be reported.

Sorry what doesn't work as I describe?

It looks like compiler relies somehow on isAddressSpaceSupersetOf in deduction process and priv3 type is not qualified with the address space. I'll check this assumption.

isAddressSpaceSupersetOf is not used in deduction it is used in call operator overloading resolution.

[AnastasiaStulova]

Do you know where isAddressSpaceSupersetOf called from for this line?

[AnastasiaStulova]

It's not how lambda parsing works in C++. The initializer is parsed separately and there is no danger of declaring it this way unless it's being called.

I would suggest to check the comments on the review too that might help to clarify the topic:
https://reviews.llvm.org/D70242
https://reviews.llvm.org/D69938

[bader]

https://reviews.llvm.org/D69938 - the discussion started with an assumption that deduced address space for lambda object can be changed to legalize function call operator usages - https://reviews.llvm.org/D69938#1737196, but ended with an open question if it's really possible - https://reviews.llvm.org/D69938#1755709.
Later the conclusion was to use "opencl_generic" for temporaries rather than "default" - https://reviews.llvm.org/D69938#1759241. Why does current implementation use default address space instead of opencl_generic?

It's still not clear whether it's the right approach to detect an error at the function call operator use or it should be done for lambda variable declaration. The case where address space of a lambda object is changed to make use of inconsistent declaration is not covered by tests.

[Fznamznon]

Alright, I dug in. This fail basically happens because It seems that lambdas have predefined conversions to function pointer form:

`-LambdaExpr 0x563552014f70 <col:16, col:31> '(lambda at address-space-lambda.cl:34:16)'
  |-CXXRecordDecl 0x563552014928 <col:16> col:16 implicit class definition
  | |-DefinitionData lambda pass_in_registers empty standard_layout trivially_copyable literal can_const_default_init
  | | |-DefaultConstructor defaulted_is_constexpr
  | | |-CopyConstructor simple trivial has_const_param needs_implicit implicit_has_const_param
  | | |-MoveConstructor exists simple trivial needs_implicit
  | | |-CopyAssignment trivial has_const_param needs_implicit implicit_has_const_param
  | | |-MoveAssignment
  | | `-Destructor simple irrelevant trivial
  | |-CXXMethodDecl 0x563552014a60 <col:19, col:31> col:16 used constexpr operator() 'void () const __global' inline
  | | `-CompoundStmt 0x563552014b10 <col:30, col:31>
  | |-[[THIS ONE]] CXXConversionDecl 0x563552014e08 <col:16> col:16 implicit used constexpr operator void (*)() 'void (*() const noexcept)()' inline
  | | `-CompoundStmt 0x563552044e78 <col:16>
  | |   `-ReturnStmt 0x563552044e68 <col:16>
  | |     `-ImplicitCastExpr 0x563552044e50 <col:16> 'void (*)()' <FunctionToPointerDecay>
  | |       `-DeclRefExpr 0x563552044e30 <col:16> 'void ()' lvalue CXXMethod 0x563552014eb8 '__invoke' 'void ()'
  | |-CXXMethodDecl 0x563552014eb8 <col:16> col:16 implicit used __invoke 'void ()' static inline
  | | `-CompoundStmt 0x563552044e20 <col:16>
  | `-CXXDestructorDecl 0x563552014fa0 <col:16> col:16 implicit referenced ~ 'void () __generic noexcept' inline default trivial
  `-CompoundStmt 0x563552014b10 <col:30, col:31>

C++ requres it:

The closure type for a lambda-expression with no lambda-capture has a public non-virtual non-explicit const conversion function to pointer to function having the same parameter and return types as the closure type’s function call operator.

But as you can see from the AST dump this conversion operator converts this lambda to a poiner to a function without address space qualifiers at all, i,e, everything what this function accepts can be in Default address space. Whereas calling operator of such lambda has __global qualifier.
After we allow conversion between between Default address space and others this conversion is used by compiler to make such call of this lambda valid:

`-CallExpr 0x563552044ec8 <line:35:3, col:9> 'void':'void'
  `-ImplicitCastExpr 0x563552044eb0 <col:3> 'void (*)()' <UserDefinedConversion>
    `-CXXMemberCallExpr 0x563552044e90 <col:3> 'void (*)()'
      `-MemberExpr 0x563552044df0 <col:3> '<bound member function type>' .operator void (*)() [[THE CONVERSION!!!!]] 0x563552014e08
        `-ImplicitCastExpr 0x563552044dd8 <col:3> 'const (lambda at address-space-lambda.cl:34:16)' lvalue <AddressSpaceConversion>
          `-DeclRefExpr 0x563552044d70 <col:3> '__private (lambda at address-space-lambda.cl:34:16)':'__private (lambda at address-space-lambda.cl:34:16)' lvalue Var 0x563552014820 'priv3' '__private (lambda

I'm not an expert in OpenCL and C++ for OpenCL but, if the philosophy of OpenCL address spaces is qualify ALL types with address space qualifier and if C++ for OpenCL should work like regular C++ for language features like lambdas, I think conversion from lambda object to function pointer must convert to a pointer to a function with the same address space qualifiers as operator () has. So, IMO this is a bug in C++ for OpenCL. @AnastasiaStulova , could you please provide your opinion on this?

[Fznamznon]

Submitted a bug https://bugs.llvm.org/show_bug.cgi?id=45472 .

[erichkeane]

I'm not privy to the motivation here, but this seems in general like a revert of a previous change I did. I remember there being a motivation for it, but if that motivation is no longer present, this looks fine.

[erichkeane]

This seems like an unrelated change. A good one, just not related.

[bader]

Not sure how to interpret this...
Do you want to me revert it? Add to the PR description? No actions implied?

[erichkeane]

Yeah, sorry for not being clear. I guess the answer is: It depends on where this lands. I'd rather you do a separate commit for this formatting change here, but if you're submitting to llvm community, I'd rather the ternary version just be part of the patch.

[bader]

This patch is to SYCL project.
First, I'd like to commit it here and see if it breaks anything. So far internal testing do not exposed any issues.
In addition to that I'd like to get high-level feedback on the approach before moving to the LLVM project.

WRT using different address space map for SYCL. I don't know if we break anything by mapping "default" address space to "4", which SPIR-V converter interprets as generic. It might that other programming models targeting SPIR rely on existing mapping.
In theory OpenCL is not supposed to use "default" address space, so it should be okay to use SYCLAddrSpaceMap in all cases.
As it was a separate patch, I suggest doing this experiment separately. Does it sounds okay to you?

Tagging other SYCL implementers: @Naghasan, @keryell, @illuhad,

[AnastasiaStulova]

Having lang address space in target information seems strange. Ideally the IR should be emitted as language agnostic as possible. What are you trying to achieve with this separate address space map?

[bader]

Having lang address space in target information seems strange. Ideally the IR should be emitted as language agnostic as possible. What are you trying to achieve with this separate address space map?

I'm trying to map "pointer w/o address space qualifiers" to "generic" address space. Currently it's mapped to "private", which doesn't seems right. As mentioned in the previous comment if all other languages do not rely on that behavior, I can remove "language dependent" customization.

[AnastasiaStulova]

Well OpenCL 2.0 does exactly the same. All pointer are in generic address space. Why don't you just add generic on AST while parsing the types?

[bader]

@erichkeane answered this in this comment: https://github.com/intel/llvm/pull/1039/files#r370236039.

[erichkeane]

What is happening here? The change to expected disables these tests, right? We dont want that, do we?

[bader]

What is happening here?

I'm not sure. I hope Anastasia can help to understand. #1039 (review).

The change to expected disables these tests, right?

Yes.

We dont want that, do we?

I'd like to understand why this patch affected the test and fix either the test or OpenCL mode.

[erichkeane]

Could you perhaps show a diff vs 'master' (or is there a link somewhere?)? It would be easier to see what we're changing in reference to community.

[bader]

Could you perhaps show a diff vs 'master' (or is there a link somewhere?)? It would be easier to see what we're changing in reference to community.

Full diff between the SYCL branch and LLVM master: https://github.com/intel/llvm/compare/356b33516c2e0ef241066dded16d7ecc1f7aa8cc..sycl
Is this the link you are looking for?

[erichkeane]

Could you perhaps show a diff vs 'master' (or is there a link somewhere?)? It would be easier to see what we're changing in reference to community.

Full diff between the SYCL branch and LLVM master: https://github.com/intel/llvm/compare/356b33516c2e0ef241066dded16d7ecc1f7aa8cc..sycl
Is this the link you are looking for?

No, I guess what I'd like to see is a diff of this patch going to LLVM community. I note for example that the isAddressSpaceSupersetOf is way different in community now than when we altered it initially (https://github.com/llvm/llvm-project/blob/master/clang/include/clang/AST/Type.h#L489).

[bader]

Could you perhaps show a diff vs 'master' (or is there a link somewhere?)? It would be easier to see what we're changing in reference to community.

Full diff between the SYCL branch and LLVM master: https://github.com/intel/llvm/compare/356b33516c2e0ef241066dded16d7ecc1f7aa8cc..sycl
Is this the link you are looking for?

No, I guess what I'd like to see is a diff of this patch going to LLVM community. I note for example that the isAddressSpaceSupersetOf is way different in community now than when we altered it initially (https://github.com/llvm/llvm-project/blob/master/clang/include/clang/AST/Type.h#L489).

I mentioned my plans here in #968 (comment). If there are no objections with re-using OpenCL attributes, I'm going to create a new patch instead of #968, but before we go to LLVM community, I'd like to align within SYCL working group. I know that @illuhad and Codeplay team are working on non-OpenCL based implementations, so re-using OpenCL attributes might not meet their needs.

I think the patch re-using OpenCL attributes will be must smaller and include (only?) changes in isAddressSpaceSupersetOf function. Something like this.

[AnastasiaStulova]

I mentioned my plans here in #968 (comment). If there are no objections with re-using OpenCL attributes, I'm going to create a new patch instead of #968, but before we go to LLVM community, I'd like to align within SYCL working group. I know that @illuhad and Codeplay team are working on non-OpenCL based implementations, so re-using OpenCL attributes might not meet their needs.

I guess you can implement the logic in multiple ways since SYCL spec defines behaviour at a much higher level using C++ libraries constructs and leaving lots of freedom to the implementers. However, if you plan to discuss this with WG can you please let me know when this is happening as I would like to be involved in the conversation.

[AnastasiaStulova]

This seems strange. Why do you need this downstream change here?

[AnastasiaStulova]

What are you trying to achieve?

[bader]

I'd like to avoid adding AS qualifiers to all types and treat all objects of "unqualified" types to be allocated in "generic" address space. This allows us to emit valid SPIR-V files for the SYCL code, which inter-operates regular C++ code satisfying SYCL kernel restrictions.
We qualify with address spaces only types required by the SPIR-V spec (e.g. kernel parameters, program scope variables, etc). All other objects (not qualified explicitly by user via SYCL pointer classes), are residing in "generic" address space. This approach allows us altering C++ type system.

[AnastasiaStulova]

In OpenCL (and other languages too) we do qualifiers deduction during parsing. All pointers without explicit address space are in generic in OpenCL v2.0. I guess you could just do the same.

This function determines qualifier relations regardless of exact type. Also this function should just have the same behaviour as the one above since they are to be used interchangeably. I think we just added this function to allow calling it as a member but it just called the helper. So I don't understand why these two have different logic in your downstream changes. They should return the same result for the same pair of qualifiers.

[AnastasiaStulova]

We have recently changed how deduction works:
https://reviews.llvm.org/D65744

My guess is you just need to alter deduceOpenCLPointeeAddrSpace slightly and make sure you invoke it in the same places.

[AnastasiaStulova]

All pointers in AST have the same address space (default).

I presume if pointer address space was provided explicitly you would keep it?

This way we follow standard C++ type system, so every valid C++ code can be compiled and executed.

What about C++ features that can't be easily supported by OpenCL devices like virtual functions?

In Clang CG, concrete address space pointers appear: alloca produces private pointers, kernel arguments can be global or local. Since AST assumes that all pointers have the same type, most of these concrete pointers are immediately casted to generic.

I don't see how they appear in CodeGen if you didn't have them in AST at all? I presume you have some address spaces somewhere in AST? Also how about address spaces that can't convert to generic like constant. Do you support this?

Later in the compilation flow we use llvm::InferAddressSpaces pass to optimize addrspace casts where the algorithm can identify the real address space behind a generic pointer.

That makes sense. I guess the limitation is if you compile separate translation modules or don't inline function calls then the inference will be limited.

[asavonic]

All pointers in AST have the same address space (default).

I presume if pointer address space was provided explicitly you would keep it?

Yes, but the only way to do that is through sycl::multi_ptr (and sycl::accessor, forgot about it) classes. User never writes

This way we follow standard C++ type system, so every valid C++ code can be compiled and executed.

What about C++ features that can't be easily supported by OpenCL devices like virtual functions?

We do support function pointers on CPU and GPU as an extension:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/SPIRV/SPV_INTEL_function_pointers.asciidoc

In Clang CG, concrete address space pointers appear: alloca produces private pointers, kernel arguments can be global or local. Since AST assumes that all pointers have the same type, most of these concrete pointers are immediately casted to generic.

I don't see how they appear in CodeGen if you didn't have them in AST at all? I presume you have some address spaces somewhere in AST?

Sorry, I should've say this better.
You're correct, some address spaces are present in AST, but they are "hidden" from a user inside C++ classes, so there is no legal way to get their type.
Alloca as I mentioned, is always addrspace(0) from data layout, this address space does not come from AST.

Also how about address spaces that can't convert to generic like constant. Do you support this?

Constant is not fully supported as a raw pointer. It cannot be converted to generic, so if we allow it, then we risk emitting LLVM IR that cannot be lowered to SPIR-V.
Constant can still be used with sycl::multi_ptr and sycl::accessor classes though.

Later in the compilation flow we use llvm::InferAddressSpaces pass to optimize addrspace casts where the algorithm can identify the real address space behind a generic pointer.

That makes sense. I guess the limitation is if you compile separate translation modules or don't inline function calls then the inference will be limited.

Right. More limitations come if a pointer is stored in memory and then loaded elsewhere.
Unless mem2reg eliminates this store/load, it is difficult to prove anything.

[AnastasiaStulova]

We do support function pointers on CPU and GPU as an extension:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/SPIRV/SPV_INTEL_function_pointers.asciidoc

Right this is great! But I was thinking of generic supported by multiple vendors. Does it mean that SYCL will define a new type of devices that it supports i.e. not just any OpenCL device?

Sorry, I should've say this better.
You're correct, some address spaces are present in AST, but they are "hidden" from a user inside C++ classes, so there is no legal way to get their type.
Alloca as I mentioned, is always addrspace(0) from data layout, this address space does not come from AST.

I am thinking of some language related aspect that you might miss to diagnose with this strategy. Let's say if user code creates pointers:

loc_ptr1 - pointer to local addr space
loc_ptr2 - pointer to local addr space
glob_ptr - pointer to global addr space

then are they cast to generic (no address space) straight away and then if they appear in the following statement:

loc_ptr1 = somevariable ? loc_ptr2 : glob_ptr;

if you don't have address spaces any more at this point you can't really provide an error saying that the code is likely illegal?

Are you aware of the cases like that and how severe they might be?

[asavonic]

We do support function pointers on CPU and GPU as an extension:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/SPIRV/SPV_INTEL_function_pointers.asciidoc

Right this is great! But I was thinking of generic supported by multiple vendors. Does it mean that SYCL will define a new type of devices that it supports i.e. not just any OpenCL device?

Since OpenCL doesn't have this in the core specification, it is hard (impossible?) to provide a solution that works everywhere (with reasonable performance). But the extension is fairly generic, so I guess any vendor interested in this feature can implement it.

Sorry, I should've say this better.
You're correct, some address spaces are present in AST, but they are "hidden" from a user inside C++ classes, so there is no legal way to get their type.
Alloca as I mentioned, is always addrspace(0) from data layout, this address space does not come from AST.

I am thinking of some language related aspect that you might miss to diagnose with this strategy. Let's say if user code creates pointers:

loc_ptr1 - pointer to local addr space
loc_ptr2 - pointer to local addr space
glob_ptr - pointer to global addr space

Address space qualifiers are not exposed in the SYCL language, so there is no way for a user to create a qualified raw pointer. (unless they use internal attributes directly; nobody should do this)

then are they cast to generic (no address space) straight away and then if they appear in the following statement:

loc_ptr1 = somevariable ? loc_ptr2 : glob_ptr;

if you don't have address spaces any more at this point you can't really provide an error saying that the code is likely illegal?

This is perfectly legal, if you write it without address space qualifiers.
Consider this (not real SYCL code, but close enough):

void main() {
  global_ptr<int> glob_ptr;
  local_ptr<int> loc_ptr;
  int cond = random();
  parallel_for([]() {
    // kernel code
    int *glob = glob_ptr.get();  // glob_ptr.get returns an int* in default AS (== generic)
    int *loc= loc_ptr.get();     // loc_ptr.get also returns an int* in default AS
    int *ptr = cond? glob : loc; // ptr is also default
    *ptr = 42; // this is going to be a store by a generic pointer, unless the compiler
               // figures out what `cond' is.
  }
}

In LLVM IR this looks like:

define kernel(i32 addrspace(1)* glob_ptr, i32 addrspace(3)* loc_ptr, i32 cond) {
    %glob = addrspacecast %glob_ptr to i32 addrspace(4)*
    %loc = addrspacecast %loc_ptr to i32 addrspace(4)*
    %ptr = select %cond,  i32 addrspace(4)* %glob,  i32 addrspace(4)* %loc
    store  i32 addrspace(4)* %ptr, 42
}

Again, this a valid code for a device that supports generic address space. If a device doesn't support it, then compilation terminates (or gives a diagnostic) from the device backend/middle-end compiler (it is not as pretty as the a diagnostic from Clang). Although, for less capable devices, it is probably better to use C++ classes instead of raw pointers where possible.

Are you aware of the cases like that and how severe they might be?

Control flow, stored/loaded pointers from memory, arrays of pointers in different address spaces. I suspect these cases are pretty common for generic C++ code, and there is a runtime overhead for supporting them. User can optimize them manually (by using sycl::multi_ptr class), or rely on the compiler.

[AnastasiaStulova]

In my original code I didn't use generic pointer. By adding generic you removed the ability to diagnose the behaviour that is illegal i.e. assigning local address space pointer into global.

Even in OpenCL 2.0 compatible devices address spaces are used for performance reasons so developers might want to avoid generic as much as possible. Then it is really helpful to get diagnostics from the compiler as debugging such issues can be very painful to nearly impossible on accelerators.

Ok I generally understand the reasoning behind your design. I just find it a bit unfortunate that you remove address spaces from AST early so you can't benefit from the power of semantical analysis fully. I don't know whether developers will find it valuable that the compiler compiles everything but then they get little help in what goes wrong.

[AnastasiaStulova]

Having lang address space in target information seems strange. Ideally the IR should be emitted as language agnostic as possible. What are you trying to achieve with this separate address space map?

[bader]

I mentioned my plans here in #968 (comment). If there are no objections with re-using OpenCL attributes, I'm going to create a new patch instead of #968, but before we go to LLVM community, I'd like to align within SYCL working group. I know that @illuhad and Codeplay team are working on non-OpenCL based implementations, so re-using OpenCL attributes might not meet their needs.

I guess you can implement the logic in multiple ways since SYCL spec defines behaviour at a much higher level using C++ libraries constructs and leaving lots of freedom to the implementers. However, if you plan to discuss this with WG can you please let me know when this is happening as I would like to be involved in the conversation.

In addition to GitHub issues/pull requests comments we use bi-weekly Skype calls on Mondays. I thought I had sent you an invite. Please, let me know if you didn't get it.

[AnastasiaStulova]

I mentioned my plans here in #968 (comment). If there are no objections with re-using OpenCL attributes, I'm going to create a new patch instead of #968, but before we go to LLVM community, I'd like to align within SYCL working group. I know that @illuhad and Codeplay team are working on non-OpenCL based implementations, so re-using OpenCL attributes might not meet their needs.

I guess you can implement the logic in multiple ways since SYCL spec defines behaviour at a much higher level using C++ libraries constructs and leaving lots of freedom to the implementers. However, if you plan to discuss this with WG can you please let me know when this is happening as I would like to be involved in the conversation.

In addition to GitHub issues/pull requests comments we use bi-weekly Skype calls on Mondays. I thought I had sent you an invite. Please, let me know if you didn't get it.

Yes, thanks! I generally don't have enough bandwidth to join regularly but I can join on demand if there are topics related to me. Just ping me if so. :)

[Fznamznon]

I think we need this patch. For upstream and to avoid hacks. @bader, do you have bandwidth to continue with it?

[bader]

@Fznamznon, I was going to get back to this after enabling optimizations in the device compiler, but it might take time due to multiple issues with SPIR-V representation of optimized LLVM IR.
I would appreciate, if you can help with this patch.

[keryell]

it might take time due to multiple issues with SPIR-V representation of optimized LLVM IR.
I would appreciate, if you can help with this patch.

If you come with some LLVM passes to generate correct SPIR-V after -O3 or figure out some SPIR-V extensions so it is valid SPIR-V, a lot of people will love you! :-)

Interestingly it looks like https://reviews.llvm.org/rG8c11bc0cd06 landed recently. Perhaps it can help you.

[Fznamznon]

@Fznamznon, I was going to get back to this after enabling optimizations in the device compiler, but it might take time due to multiple issues with SPIR-V representation of optimized LLVM IR.
I would appreciate, if you can help with this patch.

Alright, then I will continue with this.

[bader]

Replaced by #1581.