Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

add sycl backend #771

Draft
wants to merge 1 commit into
base: main
Choose a base branch
from
Draft

add sycl backend #771

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 9 additions & 3 deletions include/ddc/detail/dual_discretization.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,20 +26,22 @@ class DualDiscretization
using DDimImplDevice = typename DDim::template Impl<DDim, Kokkos::CudaSpace>;
#elif defined(__HIPCC__)
using DDimImplDevice = typename DDim::template Impl<DDim, Kokkos::HIPSpace>;
#elif defined(KOKKOS_ENABLE_SYCL)
using DDimImplDevice = typename DDim::template Impl<DDim, Kokkos::SYCLDeviceUSMSpace>;
#else
using DDimImplDevice = DDimImplHost;
#endif

DDimImplHost m_host;
#if defined(__CUDACC__) || defined(__HIPCC__)
#if defined(__CUDACC__) || defined(__HIPCC__) || defined(KOKKOS_ENABLE_SYCL)
DDimImplDevice m_device_on_host;
#endif

public:
template <class... Args>
explicit DualDiscretization(Args&&... args)
: m_host(std::forward<Args>(args)...)
#if defined(__CUDACC__) || defined(__HIPCC__)
#if defined(__CUDACC__) || defined(__HIPCC__) || defined(KOKKOS_ENABLE_SYCL)
, m_device_on_host(m_host)
#endif
{
Expand All @@ -59,6 +61,10 @@ class DualDiscretization
else if constexpr (std::is_same_v<MemorySpace, Kokkos::HIPSpace>) {
return m_device_on_host;
}
#elif defined(KOKKOS_ENABLE_SYCL)
else if constexpr (std::is_same_v<MemorySpace, Kokkos::SYCLDeviceUSMSpace>) {
return m_device_on_host;
}
#endif
else {
static_assert(!std::is_same_v<MemorySpace, MemorySpace>);
Copy link
Member Author

@yasahi-hpc yasahi-hpc Feb 10, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this line meaningful? @tpadioleau

static_assert(!std::is_same_v<MemorySpace, MemorySpace>);

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should be, ideally i would like to put false but the compiler will then stop compiling even if the branch is not chosen

Expand All @@ -72,7 +78,7 @@ class DualDiscretization

KOKKOS_FUNCTION DDimImplDevice const& get_device()
{
#if defined(__CUDACC__) || defined(__HIPCC__)
#if defined(__CUDACC__) || defined(__HIPCC__) || defined(KOKKOS_ENABLE_SYCL)
return m_device_on_host;
#else
return m_host;
Expand Down
2 changes: 2 additions & 0 deletions include/ddc/detail/macros.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,6 +40,8 @@
#define DDC_CURRENT_KOKKOS_SPACE Kokkos::HIPSpace
#elif defined(__CUDA_ARCH__)
#define DDC_CURRENT_KOKKOS_SPACE Kokkos::CudaSpace
#elif defined(__SYCL_DEVICE_ONLY__)
#define DDC_CURRENT_KOKKOS_SPACE Kokkos::SYCLDeviceUSMSpace
#else
#define DDC_CURRENT_KOKKOS_SPACE Kokkos::HostSpace
#endif
Expand Down
19 changes: 19 additions & 0 deletions include/ddc/discrete_space.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,6 +53,12 @@ static_assert(false, "DDC requires option -DKokkos_ENABLE_HIP_RELOCATABLE_DEVICE
#endif
#endif

#if defined(KOKKOS_ENABLE_SYCL)
#if !defined(KOKKOS_ENABLE_SYCL_RELOCATABLE_DEVICE_CODE)
static_assert(false, "DDC requires option -DKokkos_ENABLE_SYCL_RELOCATABLE_DEVICE_CODE=ON");
#endif
#endif

namespace ddc {

namespace detail {
Expand Down Expand Up @@ -140,6 +146,10 @@ __constant__ gpu_proxy<ddim_impl_t<DDim, Kokkos::CudaSpace>> g_discrete_space_de
// WARNING: do not put the `inline` keyword, seems to fail on MI100 rocm/4.5.0
template <class DDim>
__constant__ gpu_proxy<ddim_impl_t<DDim, Kokkos::HIPSpace>> g_discrete_space_device;
#elif defined(KOKKOS_ENABLE_SYCL)
// Global GPU variable viewing data owned by the CPU
template <class DDim>
SYCL_EXTERNAL inline sycl::ext::oneapi::experimental::device_global<gpu_proxy<ddim_impl_t<DDim, Kokkos::SYCLDeviceUSMSpace>>> g_discrete_space_device;
#endif

inline void display_discretization_store(std::ostream& os)
Expand Down Expand Up @@ -189,6 +199,11 @@ void init_discrete_space(Args&&... args)
detail::g_discrete_space_device<DDim>,
&detail::g_discrete_space_dual<DDim>->get_device(),
sizeof(detail::g_discrete_space_dual<DDim>->get_device())));
#elif defined(KOKKOS_ENABLE_SYCL)
Kokkos::DefaultExecutionSpace exec;
sycl::queue q = exec.sycl_queue();
q.memcpy(detail::g_discrete_space_device<DDim>,
&detail::g_discrete_space_dual<DDim>->get_device()).wait();
#endif
}

Expand Down Expand Up @@ -236,6 +251,10 @@ KOKKOS_FUNCTION detail::ddim_impl_t<DDim, MemorySpace> const& discrete_space()
else if constexpr (std::is_same_v<MemorySpace, Kokkos::HIPSpace>) {
return *detail::g_discrete_space_device<DDim>;
}
#elif defined(KOKKOS_ENABLE_SYCL)
else if constexpr (std::is_same_v<MemorySpace, Kokkos::SYCLDeviceUSMSpace>) {
return *detail::g_discrete_space_device<DDim>.get();
}
#endif
else {
static_assert(std::is_same_v<MemorySpace, MemorySpace>, "Memory space not handled");
Expand Down
Loading