[SYCL] Use dim_loop to unroll loops in reduce_over_group in cuda backend.

[JackAKirk]

A performance regression was reported when using reduce_over_group with sycl::vec.
This was due to a loop over calls to the scalar reduce_over_group for each of the sycl::vec components that was not unrolled and led to register spills even at -O3.
It was initially possible to fix the performance by calling #pragma unroll and declare reduce_over_group with __attribute__((always_inline)). However the SYCL_UNROLL macro that calls #pragma unroll has been removed in favour of dim_loop (#6939).
I have used dim_loop to fix the loop unrolling. However, in the cuda backend, just using dim_loop in this way actually makes the performance worse. This is because dim_loop introduces new non inlined function calls in the cuda backend that lead to register spills. The solution to this coincides with the solution of several user reports that the cuda backend is not aggressive enough with inlining. In this PR I have also therefore increased the inlining threshold multiplier value to 11.

See https://reviews.llvm.org/D142232/new/ for the corresponding upstream PR (for the inlining threshold change) that includes much more details on benchmarking dpc++ cuda with this change. In short, for dpc++ cuda, there is no other downside apart from a very small increase in compile time in some cases, but there is a massive benefit to increasing the inlining threshold across a large amount of applications.

Testing using opencl cpu backend reveals that this code change has no effect on this backend. This change is required for the cuda backend but should have no performance effect for other backends.

fixes #6583.

Signed-off-by: JackAKirk [email protected]

[JackAKirk]

Seems that dim_loop is favoured instead of #pragma unroll, which is making CI build fail here due to warnings. Marked as draft. I will try to get dim_loop to work correctly.

[JackAKirk]

Seems that dim_loop is favoured instead of #pragma unroll, which is making CI build fail here due to warnings. Marked as draft. I will try to get dim_loop to work correctly.

Done. I've updated the PR description accordingly.

[JackAKirk]

ping @intel/dpcpp-tools-reviewers

[JackAKirk]

Fixes #6583. esimd and opencl CI just didn't run this time and should be unrelated: they ran before and passed. There was an unrelated hip failure before but this time the hip ci passed.F

[JackAKirk]

Fixes #6583.

[JackAKirk]

ping @intel/dpcpp-tools-reviewers

@intel/dpcpp-tools-reviewers Would it be possible to get a review?
Thanks

[bader]

fixes #7948.

Please, fix the issue reference in the description. This is the id of this PR.

NOTE: to link PRs with issues, keywords must be used in the description. See https://docs.github.com/en/issues/tracking-your-work-with-issues/linking-a-pull-request-to-an-issue#linking-a-pull-request-to-an-issue-using-a-keyword.

[JackAKirk]

fixes #7948.

Please, fix the issue reference in the description. This is the id of this PR.

NOTE: to link PRs with issues, keywords must be used in the description. See https://docs.github.com/en/issues/tracking-your-work-with-issues/linking-a-pull-request-to-an-issue#linking-a-pull-request-to-an-issue-using-a-keyword.

I see thanks. Done.