DDC 0.0.0

a discrete domain computation library

discrete_space.hpp
1// Copyright (C) The DDC development team, see COPYRIGHT.md file
2//
3// SPDX-License-Identifier: MIT
4
5#pragma once
6
7#include <cstddef>
8#include <functional>
9#include <map>
10#include <optional>
11#include <ostream>
12#include <stdexcept>
13#include <string>
14#include <tuple>
15#include <type_traits>
16#include <utility>
17
18#include <Kokkos_Core.hpp>
19
20#include "ddc/detail/dual_discretization.hpp"
21#include "ddc/detail/macros.hpp"
22
23#if defined(__CUDACC__)
24#include <sstream>
25
26#include <cuda.h>
27#endif
28#if defined(__HIPCC__)
29#include <sstream>
30
31#include <hip/hip_runtime.h>
32#endif
33
34
35namespace ddc {
36
37namespace detail {
38
39#if defined(__CUDACC__)
40#define CUDA_THROW_ON_ERROR(val) ddc::detail::cuda_throw_on_error((val), #val, __FILE__, __LINE__)
41template <class T>
42void cuda_throw_on_error(
43 T const err,
44 const char* const func,
45 const char* const file,
46 const int line)
47{
48 if (err != cudaSuccess) {
49 std::stringstream ss;
50 ss << "CUDA Runtime Error at: " << file << ":" << line << std::endl;
51 ss << cudaGetErrorString(err) << " " << func << std::endl;
52 throw std::runtime_error(ss.str());
53 }
54}
55#elif defined(__HIPCC__)
56#define HIP_THROW_ON_ERROR(val) ddc::detail::hip_throw_on_error((val), #val, __FILE__, __LINE__)
57template <class T>
58void hip_throw_on_error(T const err, const char* const func, const char* const file, const int line)
59{
60 if (err != hipSuccess) {
61 std::stringstream ss;
62 ss << "HIP Runtime Error at: " << file << ":" << line << std::endl;
63 ss << hipGetErrorString(err) << " " << func << std::endl;
64 throw std::runtime_error(ss.str());
65 }
66}
67#endif
68
69template <class DDim, class MemorySpace>
70using ddim_impl_t = typename DDim::template Impl<MemorySpace>;
71
72template <class T>
73class gpu_proxy
74{
75 // Here are some reasonable concepts that T should satisfy to avoid undefined behaviors:
76 // - copy-constructible: objects may be memcopied to the device,
77 // - standard layout: objects will be ensured to have the same, standard, representation on the host and the device,
78 // - trivially destructible: the destructor of objects located on a device may not be called.
79 // static_assert(std::is_standard_layout_v<T>, "Not standard layout");
80 // static_assert(std::is_trivially_destructible_v<T>, "Not trivially destructible");
81 // static_assert(std::is_trivially_copy_constructible_v<T>, "Not trivially copy-constructible");
82 // Currently not trivially destructible because for example of the Kokkos::View (mostly reference-counting)
83 // Currently not trivially copy-constructible because of discrete spaces that have deleted copy-constructors and Kokkos::View (mostly reference-counting)
84
85private:
86 alignas(T) Kokkos::Array<std::byte, sizeof(T)> m_data;
87
88public:
89 KOKKOS_FUNCTION
90 T* operator->()
91 {
92 return reinterpret_cast<T*>(m_data.data());
93 }
94
95 KOKKOS_FUNCTION
96 T& operator*()
97 {
98 return *reinterpret_cast<T*>(m_data.data());
99 }
100
101 KOKKOS_FUNCTION
102 T* data()
103 {
104 return reinterpret_cast<T*>(m_data.data());
105 }
106};
107
108// Global CPU variable storing resetters. Required to correctly free data.
109inline std::optional<std::map<std::string, std::function<void()>>> g_discretization_store;
110
111// Global CPU variable owning discrete spaces data for CPU and GPU
112template <class DDim>
113inline std::optional<DualDiscretization<DDim>> g_discrete_space_dual;
114
115#if defined(__CUDACC__)
116// Global GPU variable viewing data owned by the CPU
117template <class DDim>
118__constant__ gpu_proxy<ddim_impl_t<DDim, Kokkos::CudaSpace>> g_discrete_space_device;
119#elif defined(__HIPCC__)
120// Global GPU variable viewing data owned by the CPU
121// WARNING: do not put the `inline` keyword, seems to fail on MI100 rocm/4.5.0
122template <class DDim>
123__constant__ gpu_proxy<ddim_impl_t<DDim, Kokkos::HIPSpace>> g_discrete_space_device;
124#endif
125
126inline void display_discretization_store(std::ostream& os)
127{
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";
132 }
133 } else {
134 os << "The host discretization store is not initialized:\n";
135 }
136}
137
138template <class Tuple, std::size_t... Ids>
139auto extract_after(Tuple&& t, std::index_sequence<Ids...>)
140{
141 return std::make_tuple(std::move(std::get<Ids + 1>(t))...);
142}
143
144} // namespace detail
145
146/** Initialize (emplace) a global singleton discrete space
147 *
148 * @param args the constructor arguments
149 */
150template <class DDim, class... Args>
151void init_discrete_space(Args&&... args)
152{
153 if (detail::g_discrete_space_dual<DDim>) {
154 throw std::runtime_error("Discrete space function already initialized.");
155 }
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();
159 });
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())));
170#endif
171}
172
173/** Move construct a global singleton discrete space and pass through the other argument
174 *
175 * @param a - the discrete space to move at index 0
176 * - the arguments to pass through at index 1
177 */
178template <class DDimImpl, class Arg>
179Arg init_discrete_space(std::tuple<DDimImpl, Arg>&& a)
180{
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);
184}
185
186/** Move construct a global singleton discrete space and pass through remaining arguments
187 *
188 * @param a - the discrete space to move at index 0
189 * - the (2+) arguments to pass through in other indices
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)
194{
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...>());
198}
199
200/**
201 * @tparam DDim a discrete dimension
202 * @return the discrete space instance associated with `DDim`.
203 * This function must be called from a `KOKKOS_FUNCTION`.
204 * Call `ddc::host_discrete_space` for a host-only function instead.
205 */
206template <class DDim, class MemorySpace = DDC_CURRENT_KOKKOS_SPACE>
207KOKKOS_FUNCTION detail::ddim_impl_t<DDim, MemorySpace> const& discrete_space()
208{
209 if constexpr (std::is_same_v<MemorySpace, Kokkos::HostSpace>) {
210 return detail::g_discrete_space_dual<DDim>->get_host();
211 }
212#if defined(__CUDACC__)
213 else if constexpr (std::is_same_v<MemorySpace, Kokkos::CudaSpace>) {
214 return *detail::g_discrete_space_device<DDim>;
215 }
216#elif defined(__HIPCC__)
217 else if constexpr (std::is_same_v<MemorySpace, Kokkos::HIPSpace>) {
218 return *detail::g_discrete_space_device<DDim>;
219 }
220#endif
221 else {
222 static_assert(std::is_same_v<MemorySpace, MemorySpace>, "Memory space not handled");
223 }
224}
225
226template <class DDim>
227bool is_discrete_space_initialized() noexcept
228{
229 return detail::g_discrete_space_dual<DDim>.has_value();
230}
231
232template <class DDim>
233detail::ddim_impl_t<DDim, Kokkos::HostSpace> const& host_discrete_space()
234{
235 return detail::g_discrete_space_dual<DDim>->get_host();
236}
237
238} // namespace ddc
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