18#include <Kokkos_Core.hpp>
20#include "ddc/detail/dual_discretization.hpp"
21#include "ddc/detail/macros.hpp"
23#if defined(KOKKOS_ENABLE_CUDA)
27#elif defined(KOKKOS_ENABLE_HIP)
30# include <hip/hip_runtime.h>
33#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
34# define DDC_DETAIL_DEVICE_THROW_ON_ERROR(val)
35 ddc::detail::device_throw_on_error((val), #val, __FILE__, __LINE__)
42#if defined(KOKKOS_ENABLE_CUDA)
43inline void device_throw_on_error(
44 cudaError_t
const err,
45 const char*
const func,
46 const char*
const file,
49 if (err != cudaSuccess) {
51 ss <<
"CUDA Runtime Error at: " << file <<
":" << line <<
"\n";
52 ss << cudaGetErrorString(err) <<
" " << func <<
"\n";
53 throw std::runtime_error(ss.str());
56#elif defined(KOKKOS_ENABLE_HIP)
57inline void device_throw_on_error(
59 const char*
const func,
60 const char*
const file,
63 if (err != hipSuccess) {
65 ss <<
"HIP Runtime Error at: " << file <<
":" << line <<
"\n";
66 ss << hipGetErrorString(err) <<
" " << func <<
"\n";
67 throw std::runtime_error(ss.str());
72template <
class DDim,
class MemorySpace>
73using ddim_impl_t =
typename DDim::
template Impl<DDim, MemorySpace>;
89 alignas(T) Kokkos::Array<std::byte,
sizeof(T)> m_data;
95 return reinterpret_cast<T*>(m_data.data());
101 return *
reinterpret_cast<T*>(m_data.data());
107 return reinterpret_cast<T*>(m_data.data());
112inline std::optional<std::map<std::string, std::function<
void()>>> g_discretization_store;
116inline std::optional<DualDiscretization<DDim>> g_discrete_space_dual;
118#if defined(KOKKOS_ENABLE_CUDA)
121__constant__ gpu_proxy<ddim_impl_t<DDim, GlobalVariableDeviceSpace>> g_discrete_space_device;
122#elif defined(KOKKOS_ENABLE_HIP)
126__constant__ gpu_proxy<ddim_impl_t<DDim, GlobalVariableDeviceSpace>> g_discrete_space_device;
127#elif defined(KOKKOS_ENABLE_SYCL)
130SYCL_EXTERNAL
inline sycl::ext::oneapi::experimental::device_global<
131 gpu_proxy<ddim_impl_t<DDim, GlobalVariableDeviceSpace>>>
132 g_discrete_space_device;
135inline void display_discretization_store(std::ostream& os)
137 if (g_discretization_store) {
138 os <<
"The host discretization store is initialized:\n";
139 for (
auto const& [key, value] : *g_discretization_store) {
140 os <<
" - " << key <<
"\n";
143 os <<
"The host discretization store is not initialized:\n";
147template <
class Tuple, std::size_t... Ids>
148auto extract_after(Tuple&& t, std::index_sequence<Ids...>)
150 return std::make_tuple(std::move(std::get<Ids + 1>(t))...);
156
157
158
159template <
class DDim,
class... Args>
163 !std::is_same_v<DDim,
typename DDim::discrete_dimension_type>,
164 "Discrete dimensions should inherit from the discretization, not use an alias");
165 if (detail::g_discrete_space_dual<DDim>) {
166 throw std::runtime_error(
"Discrete space function already initialized.");
168 detail::g_discrete_space_dual<DDim>.emplace(std::forward<Args>(args)...);
169 detail::g_discretization_store->emplace(
typeid(DDim).name(), []() {
170 detail::g_discrete_space_dual<DDim>.reset();
172#if defined(KOKKOS_ENABLE_CUDA)
173 DDC_DETAIL_DEVICE_THROW_ON_ERROR(cudaMemcpyToSymbol(
174 detail::g_discrete_space_device<DDim>,
175 &detail::g_discrete_space_dual<DDim>->get_device(),
176 sizeof(detail::g_discrete_space_dual<DDim>->get_device())));
177#elif defined(KOKKOS_ENABLE_HIP)
178 DDC_DETAIL_DEVICE_THROW_ON_ERROR(hipMemcpyToSymbol(
179 detail::g_discrete_space_device<DDim>,
180 &detail::g_discrete_space_dual<DDim>->get_device(),
181 sizeof(detail::g_discrete_space_dual<DDim>->get_device())));
182#elif defined(KOKKOS_ENABLE_SYCL)
183 Kokkos::DefaultExecutionSpace exec;
184 sycl::queue q = exec.sycl_queue();
185 q.memcpy(detail::g_discrete_space_device<DDim>,
186 &detail::g_discrete_space_dual<DDim>->get_device())
192
193
194
195
196template <
class DDim,
class DDimImpl,
class Arg0>
199 init_discrete_space<DDim>(std::move(std::get<0>(a)));
200 return std::get<1>(a);
204
205
206
207
208template <
class DDim,
class DDimImpl,
class Arg0,
class Arg1,
class... Args>
209std::tuple<Arg0, Arg1, Args...>
init_discrete_space(std::tuple<DDimImpl, Arg0, Arg1, Args...>&& a)
211 init_discrete_space<DDim>(std::move(std::get<0>(a)));
212 return detail::extract_after(std::move(a), std::index_sequence_for<Arg0, Arg1, Args...>());
216
217
218
219
220
221template <
class DDim,
class MemorySpace = DDC_CURRENT_KOKKOS_SPACE>
224 if constexpr (std::is_same_v<MemorySpace, Kokkos::HostSpace>) {
225 return detail::g_discrete_space_dual<DDim>->get_host();
227#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
228 else if constexpr (std::is_same_v<MemorySpace, detail::GlobalVariableDeviceSpace>) {
229 return *detail::g_discrete_space_device<DDim>;
231#elif defined(KOKKOS_ENABLE_SYCL)
232 else if constexpr (std::is_same_v<MemorySpace, detail::GlobalVariableDeviceSpace>) {
233 return *detail::g_discrete_space_device<DDim>.get();
237 static_assert(std::is_same_v<MemorySpace, MemorySpace>,
"Memory space not handled");
244 return detail::g_discrete_space_dual<DDim>.has_value();
250 return detail::g_discrete_space_dual<DDim>->get_host();
255#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
256# undef DDC_DETAIL_DEVICE_THROW_ON_ERROR
The top-level namespace of DDC.
bool is_discrete_space_initialized() noexcept
void init_discrete_space(Args &&... args)
Initialize (emplace) a global singleton discrete space.
detail::ddim_impl_t< DDim, Kokkos::HostSpace > const & host_discrete_space()
KOKKOS_FUNCTION detail::ddim_impl_t< DDim, MemorySpace > const & discrete_space()
Arg0 init_discrete_space(std::tuple< DDimImpl, Arg0 > &&a)
Move construct a global singleton discrete space and pass through the other argument.
std::tuple< Arg0, Arg1, Args... > init_discrete_space(std::tuple< DDimImpl, Arg0, Arg1, Args... > &&a)
Move construct a global singleton discrete space and pass through remaining arguments.