DDC 0.5.2
Loading...
Searching...
No Matches
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(KOKKOS_ENABLE_CUDA)
24# include <sstream>
25
26# include <cuda.h>
27#elif defined(KOKKOS_ENABLE_HIP)
28# include <sstream>
29
30# include <hip/hip_runtime.h>
31#endif
32
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__)
36#endif
37
38namespace ddc {
39
40namespace detail {
41
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,
47 const int line)
48{
49 if (err != cudaSuccess) {
50 std::stringstream ss;
51 ss << "CUDA Runtime Error at: " << file << ":" << line << "\n";
52 ss << cudaGetErrorString(err) << " " << func << "\n";
53 throw std::runtime_error(ss.str());
54 }
55}
56#elif defined(KOKKOS_ENABLE_HIP)
57inline void device_throw_on_error(
58 hipError_t const err,
59 const char* const func,
60 const char* const file,
61 const int line)
62{
63 if (err != hipSuccess) {
64 std::stringstream ss;
65 ss << "HIP Runtime Error at: " << file << ":" << line << "\n";
66 ss << hipGetErrorString(err) << " " << func << "\n";
67 throw std::runtime_error(ss.str());
68 }
69}
70#endif
71
72template <class DDim, class MemorySpace>
73using ddim_impl_t = typename DDim::template Impl<DDim, MemorySpace>;
74
75template <class T>
76class gpu_proxy
77{
78 // Here are some reasonable concepts that T should satisfy to avoid undefined behaviors:
79 // - copy-constructible: objects may be memcopied to the device,
80 // - standard layout: objects will be ensured to have the same, standard, representation on the host and the device,
81 // - trivially destructible: the destructor of objects located on a device may not be called.
82 // static_assert(std::is_standard_layout_v<T>, "Not standard layout");
83 // static_assert(std::is_trivially_destructible_v<T>, "Not trivially destructible");
84 // static_assert(std::is_trivially_copy_constructible_v<T>, "Not trivially copy-constructible");
85 // Currently not trivially destructible because for example of the Kokkos::View (mostly reference-counting)
86 // Currently not trivially copy-constructible because of discrete spaces that have deleted copy-constructors and Kokkos::View (mostly reference-counting)
87
88private:
89 alignas(T) Kokkos::Array<std::byte, sizeof(T)> m_data;
90
91public:
92 KOKKOS_FUNCTION
93 T* operator->()
94 {
95 return reinterpret_cast<T*>(m_data.data());
96 }
97
98 KOKKOS_FUNCTION
99 T& operator*()
100 {
101 return *reinterpret_cast<T*>(m_data.data());
102 }
103
104 KOKKOS_FUNCTION
105 T* data()
106 {
107 return reinterpret_cast<T*>(m_data.data());
108 }
109};
110
111// Global CPU variable storing resetters. Required to correctly free data.
112inline std::optional<std::map<std::string, std::function<void()>>> g_discretization_store;
113
114// Global CPU variable owning discrete spaces data for CPU and GPU
115template <class DDim>
116inline std::optional<DualDiscretization<DDim>> g_discrete_space_dual;
117
118#if defined(KOKKOS_ENABLE_CUDA)
119// Global GPU variable viewing data owned by the CPU
120template <class DDim>
121__constant__ gpu_proxy<ddim_impl_t<DDim, GlobalVariableDeviceSpace>> g_discrete_space_device;
122#elif defined(KOKKOS_ENABLE_HIP)
123// Global GPU variable viewing data owned by the CPU
124// WARNING: do not put the `inline` keyword, seems to fail on MI100 rocm/4.5.0
125template <class DDim>
126__constant__ gpu_proxy<ddim_impl_t<DDim, GlobalVariableDeviceSpace>> g_discrete_space_device;
127#elif defined(KOKKOS_ENABLE_SYCL)
128// Global GPU variable viewing data owned by the CPU
129template <class DDim>
130SYCL_EXTERNAL inline sycl::ext::oneapi::experimental::device_global<
131 gpu_proxy<ddim_impl_t<DDim, GlobalVariableDeviceSpace>>>
132 g_discrete_space_device;
133#endif
134
135inline void display_discretization_store(std::ostream& os)
136{
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";
141 }
142 } else {
143 os << "The host discretization store is not initialized:\n";
144 }
145}
146
147template <class Tuple, std::size_t... Ids>
148auto extract_after(Tuple&& t, std::index_sequence<Ids...>)
149{
150 return std::make_tuple(std::move(std::get<Ids + 1>(t))...);
151}
152
153} // namespace detail
154
155/** Initialize (emplace) a global singleton discrete space
156 *
157 * @param args the constructor arguments
158 */
159template <class DDim, class... Args>
160void init_discrete_space(Args&&... args)
161{
162 static_assert(
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.");
167 }
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();
171 });
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())
187 .wait();
188#endif
189}
190
191/** Move construct a global singleton discrete space and pass through the other argument
192 *
193 * @param a - the discrete space to move at index 0
194 * - the arguments to pass through at index 1
195 */
196template <class DDim, class DDimImpl, class Arg0>
197Arg0 init_discrete_space(std::tuple<DDimImpl, Arg0>&& a)
198{
199 init_discrete_space<DDim>(std::move(std::get<0>(a)));
200 return std::get<1>(a);
201}
202
203/** Move construct a global singleton discrete space and pass through remaining arguments
204 *
205 * @param a - the discrete space to move at index 0
206 * - the (2+) arguments to pass through in other indices
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)
210{
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...>());
213}
214
215/**
216 * @tparam DDim a discrete dimension
217 * @return the discrete space instance associated with `DDim`.
218 * This function must be called from a `KOKKOS_FUNCTION`.
219 * Call `ddc::host_discrete_space` for a host-only function instead.
220 */
221template <class DDim, class MemorySpace = DDC_CURRENT_KOKKOS_SPACE>
222KOKKOS_FUNCTION detail::ddim_impl_t<DDim, MemorySpace> const& discrete_space()
223{
224 if constexpr (std::is_same_v<MemorySpace, Kokkos::HostSpace>) {
225 return detail::g_discrete_space_dual<DDim>->get_host();
226 }
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>;
230 }
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();
234 }
235#endif
236 else {
237 static_assert(std::is_same_v<MemorySpace, MemorySpace>, "Memory space not handled");
238 }
239}
240
241template <class DDim>
242bool is_discrete_space_initialized() noexcept
243{
244 return detail::g_discrete_space_dual<DDim>.has_value();
245}
246
247template <class DDim>
248detail::ddim_impl_t<DDim, Kokkos::HostSpace> const& host_discrete_space()
249{
250 return detail::g_discrete_space_dual<DDim>->get_host();
251}
252
253} // namespace ddc
254
255#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
256# undef DDC_DETAIL_DEVICE_THROW_ON_ERROR
257#endif
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.