G-071
commented
1 year ago @diehlpk I think the Codacy warning is one we can safely ignore! Do you know how to silence Codacy about it, or how to make it skip that test?
Closed G-071 closed 1 year ago
G-071
commented
1 year ago @diehlpk I think the Codacy warning is one we can safely ignore! Do you know how to silence Codacy about it, or how to make it skip that test?
This PR adds support for running the Kokkos kernels on the SYCL execution space (effectively adding support for running Octo-Tiger on Intel GPUs).
In order to this, the utilized executors and allocators are changed to their respective SYCL versions when Kokkos SYCL support is detected during compile-time. Unfortunately, to allow the Kokkos kernels to run with SYCL, I also had to change the math functions to use the sycl versions for GPU support (for example use
sycl::sqrtfor SYCL instantiation of a kernel). The PR further adds tests for the SYCL versions of the kernels, as well as CLI arguments for selecting those kernels.The PR was tested both on a A100 GPU and on a MI100 GPU using the Intel LLVM toolchain (though the MI100 run required some slight Kokkos changes to propagate the architecture flags).
Of course, we need the actual allocators and executors for this to work, so this PR comes with multiple dependencies: For asynchronous launches with HPX futures, we need to HPX SYCL integration/executors added in STEllAR-GROUP/hpx#6085 as well as the respective hpx-kokkos integration in STEllAR-GROUP/hpx-kokkos#13. We further require the SYCL allocators to instantiate the memory pools (with SYCL buffers) which were added in sycl SC-SGS/cppuddle#15. Lastly, to make Octo-Tiger work with recent (current master) version of HPX, we also require the changes in #428 (most of the commits showing up in this PR are due to 428 not being in master yet -- hence I added this PR only as a draft until then).
For convenience, Octo-Tiger buildscripts with the SYCL toolchain are available here.