Swap x- and z-indices in SYCL MDRangePolicy

[masterleinad]

Compared to CUDA, SYCL reverses the order of blocks and threads when mapping to the hardware, see https://www.codeplay.com/portal/blogs/2019/11/18/computecpp-v1-1-6-changes-to-work-item-mapping-optimization.html. We need to take this into account for multi-dimensional sycl::nd_range ranges. This pull request deals with the MDRangePolicy part.

[DavidPoliakoff]

So, we've talked in Slack about the wisdom of this X-Z swap. Do we have confidence that all the implementations we'll care about will have good performance from this? Do we think this swap is something they'll be consistent on, or something they'll change in time?

[joeatodd]

I've tested this locally and it does the trick for the MDRange. Alongside my local changes to Kokkos_SYCL_Parallel_Team.hpp and Kokkos_SYCL_Team.hpp, it fixes the crashes we've seen in large LAMMPS runs. It's also a lot tidier than the implementation I wrote!

I suspect that, so long as other applications are using the default LayoutLeft, this should improve performance. @masterleinad I think you mentioned that you'd seen improved performance in your tests, right?

As to the permanence of this change, my understanding is that this is how the SYCL spec was always defined, and the change in 2019 simply brings the implementations in line with the standard. As such, I am quite confident that SYCL will remain row-major.

[masterleinad]

I suspect that, so long as other applications are using the default LayoutLeft, this should improve performance. @masterleinad I think you mentioned that you'd seen improved performance in your tests, right?

I checked that the performance results generated by the tests in https://github.com/kokkos/kokkos/tree/master/core/perf_test are more in line with the CUDA results, i.e. copying LayoutRight to LayoutLeft is faster than copying LayoutLeft to LayoutRight. For these tests, I didn't really see a general performance benefit, though.