[SYCLLowerIR][SemaSYCL] Support indirect hierarchical parallelism

[asudarsa]

This PR adds a missing feature in SYCL hierarchical parallelism support. Specifically, this PR adds support for the case when there are functions between parallel_for_work_group and parallel_for_work_item in the call stack. For example: void foo(sycl::group<1> group, ...) { group.parallel_for_work_item(range<1>(), [&](h_item<1> i) { ... }); } // ... cgh.parallel_for_work_group( range<1>(...), range<1>(...), [=](group<1> g) { foo(g, ...); });

[asudarsa]

@premanandrao

Can you please take a look at the SemaSYCL changes and the test case?

Thanks

[bader]

Hierarchical parallelism is technically deprecated - https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_hierarchical_data_parallel_kernels. Is it truly necessary to invest in its repair?

[asudarsa]

There is some fundamental issue with existing LIT tests in llvm/test/SYCLLowerIR (related to LowerWGScope pass) which I am trying to address. Thanks

[asudarsa]

Hi @bader

I wanted to get your take on this. I have vetted my change to a certain level. Please let me know if it is agreeable to open this up for review. This task has been pending for a while and I would like to close this one way or another.

Thanks Sincerely

[bader]

Hi @bader

I wanted to get your take on this. I have vetted my change to a certain level. Please let me know if it is agreeable to open this up for review. This task has been pending for a while and I would like to close this one way or another.

Thanks Sincerely

@asudarsa, if I get it right, you already did all the work and the patch is ready for review. If so, feel free to request the review. My previous comment was about work prioritization in general.

[asudarsa]

FE changes are okay with me, modulo that one change in SemaSYCL. We agreed not to retrigger testing over an NFC change.

There are couple of other review comments to be resolved. So, I will handle your comment also. Thanks much