[SYCL][Graph] Implement dynamic_work_group_memory for SYCL-Graphs

[keyradical]

Implements dynamic_work_group_memory for SYCL-Graphs: #16712

With this PR we're able to update work_group_memory size on subsequent graph executions. We're also now able to use dynamic_work_group_memory with both lambdas and free function kernels.

[KseniyaTikhomirova]

It looks like you are adding new extension exp_ext::dynamic_work_group_memory.
Could you please clarify what spec/doc covers it?

update: I see, #16712. Should Spec PR be blocking for this one then?

[keyradical]

It looks like you are adding new extension exp_ext::dynamic_work_group_memory. Could you please clarify what spec/doc covers it?

update: I see, #16712. Should Spec PR be blocking for this one then?

This is actually not a new extension, it is just an addition to the experimental sycl::ext::oneapi::graph extension. So in this case we'd want to merge the implementation first, before the PR with the spec changes to avoid having something in the spec which is not supported (yet).

If this were a new specification/distinct feature it might be different (since that would be merged as "proposed" first before the implementation is done), but since we are updating an existing feature/specification it's the other way around.

[keyradical]

I had to fix some L0 leaks but this is now done and this PR is ready for final review.

[EwanC]

Nice work, LGTM. I'd like @Bensuo and @fabiomestre to approve this as well.

[Bensuo]

LGTM just a few nitpicks

[sommerlukas]

JIT compiler changes LGTM.

[EwanC]

pinging @intel/dpcpp-cfe-reviewers for review

[Fznamznon]

Are more kinds of work_group_memory coming?
The copy-paste in clang's code is somewhat unfortunate. Any chance the runtime library may not need to know which kind of work_group_memory is that?

[keyradical]

Are more kinds of work_group_memory coming? The copy-paste in clang's code is somewhat unfortunate. Any chance the runtime library may not need to know which kind of work_group_memory is that?

Short answer is yes. We're extending sycl-graph specification with dynamic_work_group_memory, dynamic_local_accessor and dynamic_accessor. These are outlined in the spec changes PR: #16712. Similarly to this PR, the plan is for all these new types to follow the similar pattern which wraps around the existing classes to add the update functionality in graphs. Altough actually they would be related to local_accessor and accessor and not work_group_memory. Specifically for work_group_memory, no, there are no more new kinds of it coming.

That's true, there has been a bit of a copy-paste in clang's code but AFAIK, this is required if we want to pass dynamic_work_group_memory type to a kernel. I can't think of a way in which the frontend doesn't need to know the exact kind of work_group_memory. I'll follow up with a refactor PR if I manage to find a solution.

[keyradical]

@intel/llvm-gatekeepers this is ready to be merged, thank you.

[EwanC]

@intel/llvm-gatekeepers this is ready to be merged, thank you.

I think this still needs a @intel/llvm-reviewers-runtime approval

[keyradical]

@intel/llvm-gatekeepers this is ready to be merged, thank you.

I think this still needs a @intel/llvm-reviewers-runtime approval

That's true, gatekeepers, sorry for the oversight.

[keyradical]

@KseniyaTikhomirova could I get a review on this, please?

[KseniyaTikhomirova]

sycl part LGTM

[keyradical]

@intel/llvm-gatekeepers this is now ready to merge