[SYCL] Refactor stream class handing implementation

[zahiraam]

Currently the stream class when passed from host to device requires additional handling.
It is handled as a wrapper struct for several accessors. In order to create the AST, the stream class is processed as a whole and each accessor is processed individually.
This patch is to simplify the process. An __init and __finalize methods are introduced that eliminate the need to process the field individually. The steps are as follows:

1. Look up for the member function sycl::stream::__init
2. Copy the sycl::stream::__init arguments into the kernel function argument list
3. Allocate a sycl::stream object and call sycl::stream::__init with the kernel argument
4. Call sycl::stream::__finalize with the kernel argument at the end of the kernel

For this simple stream class:
sycl::handler H;
struct HasStreams {
// stream(totalBufferSize, workItemBufferSize, handler)
sycl::stream s1{0, 0, H};

};
int main() {
sycl::stream in_lambda{0, 0, H};

myQueue.submit([&](sycl::handler &h) {
h.single_task(= {
in_lambda.use();
});
});
return 0;
}

This is the resulting AST:

// Function Declaration
// CHECK: FunctionDecl {{.}}stream_test{{.}}

// CHECK: InitListExpr {{.}} '(lambda at
// CHECK-NEXT: CXXConstructExpr {{.}} 'sycl::stream':'sycl::stream' 'void () noexcept'

// CHECK: CXXMemberCallExpr {{.}} 'void'
// CHECK-NEXT: MemberExpr {{.}} 'void (__global char , range<1>, range<1>, id<1>, int)' lvalue .__init
// CHECK-NEXT: MemberExpr {{.}} 'sycl::stream':'sycl::stream' lvalue .
// CHECK-NEXT: DeclRefExpr {{.}} '(lambda at
// CHECK-NEXT: ImplicitCastExpr {{.}} '__global char '
// CHECK-NEXT: DeclRefExpr {{.}} '__global char *' lvalue ParmVar

// CHECK: CXXMemberCallExpr {{.}} 'void'
// CHECK-NEXT: MemberExpr {{.}} 'void ()' lvalue .__finalize
// CHECK-NEXT: MemberExpr {{.}} 'sycl::stream':'sycl::stream' lvalue .
// CHECK-NEXT: DeclRefExpr {{.}} '(lambda at

[AaronBallman]

FWIW, I don't feel particularly qualified to review this, so I'd recommend adding @elizabethandrews and @premanandrao as reviewers as well.

[bader]

@erichkeane, @Fznamznon, FYI.

[zahiraam]

This is not quite working yet (draft)! I am still in the process of understanding the code.
There stream.cpp on the sema side is failing and on the codegen side the __finalize function is not generated. If any of you can take a look, that would be great.

[elizabethandrews]

From a quick look, FE changes look like it is heading in right direction. I haven't looked at the new stream implementation in headers. I suspect you also need to change AST tests with streams in them, since the InitListExpr generated now will be different. If these tests don't exist please add one to ensure AST generated for OpenCL kernel with streams is correct.

[elizabethandrews]

FWIW, I don't feel particularly qualified to review this, so I'd recommend adding @elizabethandrews and @premanandrao as reviewers as well.

@AaronBallman the gist of this change is - stream was defined as a class holding accessors (@zahiraam please correct me if my memory serves me wrong). In SemaSYCL, we used to 'process' the stream class as a whole, and then step into this class to process each accessor individually. This change is introducing a __init( ) method in stream definition which makes it unnecessary to step into stream class and process its individual fields. We can now just call the __init( ) method to handle stream arguments like we do with other special types.

[keryell]

Always nice to simplify the code.
Perhaps the PR title should summarize more what is done in the PR?
Just curious: is it possible to have a stream to be hidden in some class, say a std::tuple in some application using meta-programming and the stream work if the object is capture by the kernel lambda?

[zahiraam]

@Fznamznon, @elizabethandrews This is not quite complete. But before going further (stream class LIT test needs editing) I want to make sure that this is moving in the right direction. Previously, the __init and __finalize of the fields were generated together, now each field has an __init and __finalize back to back. Is that correct? Did I put the createSpecialMethodCall with the finalize argument in the right place?
Thanks.

[zahiraam]

Also should I add documentation only of a description of the init/finalize function creation in the AST or should it be a complete description as it is in #2091 ? Not sure if the work in this PR is already up-streamed?

[zahiraam]

Always nice to simplify the code.
Perhaps the PR title should summarize more what is done in the PR?
Just curious: is it possible to have a stream to be hidden in some class, say a std::tuple in some application using meta-programming and the stream work if the object is capture by the kernel lambda?

@keryell Thanks for the feedback. Let me know if the description is up to your expectations.
Not sure about your question. I can make some inquiries. Would you mind illustrating this with a test case?
Thanks.

[keryell]

Just curious: is it possible to have a stream to be hidden in some class, say a std::tuple in some application using meta-programming and the stream work if the object is capture by the kernel lambda?

Not sure about your question. I can make some inquiries. Would you mind illustrating this with a test case?

Sure!
Can the following

int main() {
  std::tuple<sycl::stream> in_lambda { {0, 0, H} };

  myQueue.submit([&](sycl::handler &h) {
    h.single_task([=] {
      std::get<0>(in_lambda).use();
    });
  });
  return 0;
}

just work?
There were a lot of discussions to have accessor working even when hidden in some classes so I was wondering whether we cannot have also this generalized to other SYCL magical types...

[zahiraam]

@premanandrao, @againull Review please. Thanks.

[elizabethandrews]

In handleSyclStreamType for SyclOptReportCreator, there is a comment which can be removed with this change. Can you do that in this PR?

[elizabethandrews]

%5 here is ACC_DATA right? Since the parameters to __init is important here, I think it will be useful to check it in this test (as opposed to just checking the call to __init).

[Fznamznon]

/summary:run

[elizabethandrews]

LGTM! Thanks!

[bader]

@againull, @intel/llvm-reviewers-runtime, ping.

[againull]

Thank you! LGTM