18#include <Kokkos_Core.hpp>
20#include "ddc/detail/dual_discretization.hpp"
21#include "ddc/detail/macros.hpp"
23#if defined(__CUDACC__)
31#include <hip/hip_runtime.h>
39#if defined(__CUDACC__)
40#define CUDA_THROW_ON_ERROR(val) ddc::detail::cuda_throw_on_error((val), #val, __FILE__, __LINE__)
42void cuda_throw_on_error(
44 const char*
const func,
45 const char*
const file,
48 if (err != cudaSuccess) {
50 ss <<
"CUDA Runtime Error at: " << file <<
":" << line << std::endl;
51 ss << cudaGetErrorString(err) <<
" " << func << std::endl;
52 throw std::runtime_error(ss.str());
55#elif defined(__HIPCC__)
56#define HIP_THROW_ON_ERROR(val) ddc::detail::hip_throw_on_error((val), #val, __FILE__, __LINE__)
58void hip_throw_on_error(T
const err,
const char*
const func,
const char*
const file,
const int line)
60 if (err != hipSuccess) {
62 ss <<
"HIP Runtime Error at: " << file <<
":" << line << std::endl;
63 ss << hipGetErrorString(err) <<
" " << func << std::endl;
64 throw std::runtime_error(ss.str());
69template <
class DDim,
class MemorySpace>
70using ddim_impl_t =
typename DDim::
template Impl<MemorySpace>;
86 alignas(T) Kokkos::Array<std::byte,
sizeof(T)> m_data;
92 return reinterpret_cast<T*>(m_data.data());
98 return *
reinterpret_cast<T*>(m_data.data());
104 return reinterpret_cast<T*>(m_data.data());
109inline std::optional<std::map<std::string, std::function<
void()>>> g_discretization_store;
113inline std::optional<DualDiscretization<DDim>> g_discrete_space_dual;
115#if defined(__CUDACC__)
118__constant__ gpu_proxy<ddim_impl_t<DDim, Kokkos::CudaSpace>> g_discrete_space_device;
119#elif defined(__HIPCC__)
123__constant__ gpu_proxy<ddim_impl_t<DDim, Kokkos::HIPSpace>> g_discrete_space_device;
126inline void display_discretization_store(std::ostream& os)
128 if (g_discretization_store) {
129 os <<
"The host discretization store is initialized:\n";
130 for (
auto const& [key, value] : *g_discretization_store) {
131 os <<
" - " << key <<
"\n";
134 os <<
"The host discretization store is not initialized:\n";
138template <
class Tuple, std::size_t... Ids>
139auto extract_after(Tuple&& t, std::index_sequence<Ids...>)
141 return std::make_tuple(std::move(std::get<Ids + 1>(t))...);
147
148
149
150template <
class DDim,
class... Args>
153 if (detail::g_discrete_space_dual<DDim>) {
154 throw std::runtime_error(
"Discrete space function already initialized.");
156 detail::g_discrete_space_dual<DDim>.emplace(std::forward<Args>(args)...);
157 detail::g_discretization_store->emplace(
typeid(DDim).name(), []() {
158 detail::g_discrete_space_dual<DDim>.reset();
160#if defined(__CUDACC__)
161 CUDA_THROW_ON_ERROR(cudaMemcpyToSymbol(
162 detail::g_discrete_space_device<DDim>,
163 &detail::g_discrete_space_dual<DDim>->get_device(),
164 sizeof(detail::g_discrete_space_dual<DDim>->get_device())));
165#elif defined(__HIPCC__)
166 HIP_THROW_ON_ERROR(hipMemcpyToSymbol(
167 detail::g_discrete_space_device<DDim>,
168 &detail::g_discrete_space_dual<DDim>->get_device(),
169 sizeof(detail::g_discrete_space_dual<DDim>->get_device())));
174
175
176
177
178template <
class DDimImpl,
class Arg>
181 using DDim =
typename DDimImpl::discrete_dimension_type;
182 init_discrete_space<DDim>(std::move(std::get<0>(a)));
183 return std::get<1>(a);
187
188
189
190
191template <
class DDimImpl,
class... Args>
192std::enable_if_t<2 <=
sizeof...(Args), std::tuple<Args...>> init_discrete_space(
193 std::tuple<DDimImpl, Args...>&& a)
195 using DDim =
typename DDimImpl::discrete_dimension_type;
196 init_discrete_space<DDim>(std::move(std::get<0>(a)));
197 return detail::extract_after(std::move(a), std::index_sequence_for<Args...>());
201
202
203
204
205
206template <
class DDim,
class MemorySpace = DDC_CURRENT_KOKKOS_SPACE>
209 if constexpr (std::is_same_v<MemorySpace, Kokkos::HostSpace>) {
210 return detail::g_discrete_space_dual<DDim>->get_host();
212#if defined(__CUDACC__)
213 else if constexpr (std::is_same_v<MemorySpace, Kokkos::CudaSpace>) {
214 return *detail::g_discrete_space_device<DDim>;
216#elif defined(__HIPCC__)
217 else if constexpr (std::is_same_v<MemorySpace, Kokkos::HIPSpace>) {
218 return *detail::g_discrete_space_device<DDim>;
222 static_assert(std::is_same_v<MemorySpace, MemorySpace>,
"Memory space not handled");
229 return detail::g_discrete_space_dual<DDim>.has_value();
235 return detail::g_discrete_space_dual<DDim>->get_host();
Definition: aligned_allocator.hpp:11
bool is_discrete_space_initialized() noexcept
Definition: discrete_space.hpp:227
void init_discrete_space(Args &&... args)
Initialize (emplace) a global singleton discrete space.
Definition: discrete_space.hpp:151
detail::ddim_impl_t< DDim, Kokkos::HostSpace > const & host_discrete_space()
Definition: discrete_space.hpp:233
std::enable_if_t< 2<=sizeof...(Args), std::tuple< Args... > > init_discrete_space(std::tuple< DDimImpl, Args... > &&a){ using DDim=typename DDimImpl::discrete_dimension_type;init_discrete_space< DDim >(std::move(std::get< 0 >(a)));return detail::extract_after(std::move(a), std::index_sequence_for< Args... >());}template< class DDim, class MemorySpace=DDC_CURRENT_KOKKOS_SPACE > KOKKOS_FUNCTION detail::ddim_impl_t< DDim, MemorySpace > const & discrete_space()
Move construct a global singleton discrete space and pass through remaining arguments.
Definition: discrete_space.hpp:207
Arg init_discrete_space(std::tuple< DDimImpl, Arg > &&a)
Move construct a global singleton discrete space and pass through the other argument.
Definition: discrete_space.hpp:179