DDC 0.0.0

a discrete domain computation library

discrete_space.hpp
1 // SPDX-License-Identifier: MIT
2 
3 #pragma once
4 
5 #include <map>
6 #include <memory>
7 #include <optional>
8 #include <ostream>
9 #include <stdexcept>
10 
11 #include <Kokkos_Core.hpp>
12 #if defined(__CUDACC__)
13 #include <cuda.h>
14 #endif
15 #if defined(__HIPCC__)
16 #include <hip/hip_runtime.h>
17 #endif
18 
19 #include "ddc/discrete_domain.hpp"
20 #include "ddc/discrete_space.hpp"
21 #include "ddc/dual_discretization.hpp"
22 
23 namespace ddc {
24 
25 namespace ddc_detail {
26 
27 template <class DDim, class MemorySpace>
28 using ddim_impl_t = typename DDim::template Impl<MemorySpace>;
29 
30 template <class T>
31 class gpu_proxy
32 {
33  // Here are some reasonable concepts that T should satisfy to avoid undefined behaviors:
34  // - copy-constructible: objects may be memcopied to the device,
35  // - standard layout: objects will be ensured to have the same, standard, representation on the host and the device,
36  // - trivially destructible: the destructor of objects located on a device may not be called.
37  static_assert(std::is_standard_layout_v<T>, "Not standard layout");
38  // static_assert(std::is_trivially_destructible_v<T>, "Not trivially destructible");
39  // static_assert(std::is_trivially_copy_constructible_v<T>, "Not trivially copy-constructible");
40  // Currently not trivially destructible because for example of the Kokkos::View (mostly reference-counting)
41  // Currently not trivially copy-constructible because of discrete spaces that have deleted copy-constructors and Kokkos::View (mostly reference-counting)
42 
43 private:
44  alignas(T) std::byte m_data[sizeof(T)];
45 
46 public:
47  DDC_INLINE_FUNCTION
49  {
50  return reinterpret_cast<T*>(m_data);
51  }
52 
53  DDC_INLINE_FUNCTION
54  T& operator*()
55  {
56  return *reinterpret_cast<T*>(m_data);
57  }
58 
59  DDC_INLINE_FUNCTION
60  T* data()
61  {
62  return reinterpret_cast<T*>(m_data);
63  }
64 };
65 
66 // Global CPU variable storing resetters. Required to correctly free data.
67 inline std::optional<std::map<std::string, std::function<void()>>> g_discretization_store;
68 
69 // Global CPU variable owning discrete spaces data for CPU and GPU
70 template <class DDim>
71 inline std::optional<DualDiscretization<DDim>> g_discrete_space_dual;
72 
73 #if defined(__CUDACC__)
74 // Global GPU variable viewing data owned by the CPU
75 template <class DDim>
76 __constant__ gpu_proxy<ddim_impl_t<DDim, Kokkos::CudaSpace>> g_discrete_space_device;
77 #elif defined(__HIPCC__)
78 // Global GPU variable viewing data owned by the CPU
79 // WARNING: do not put the `inline` keyword, seems to fail on MI100 rocm/4.5.0
80 template <class DDim>
81 __constant__ gpu_proxy<ddim_impl_t<DDim, Kokkos::Experimental::HIPSpace>> g_discrete_space_device;
82 #endif
83 
84 inline void display_discretization_store(std::ostream& os)
85 {
87  os << "The host discretization store is initialized:\n";
88  for (auto const& [key, value] : *g_discretization_store) {
89  os << " - " << key << "\n";
90  }
91  } else {
92  os << "The host discretization store is not initialized:\n";
93  }
94 }
95 
96 template <class Tuple, std::size_t... Ids>
97 auto extract_after(Tuple&& t, std::index_sequence<Ids...>)
98 {
99  return std::make_tuple(std::move(std::get<Ids + 1>(t))...);
100 }
101 
102 } // namespace ddc_detail
103 
108 template <class DDim, class... Args>
109 void init_discrete_space(Args&&... args)
110 {
111  if (ddc_detail::g_discrete_space_dual<DDim>) {
112  throw std::runtime_error("Discrete space function already initialized.");
113  }
114  ddc_detail::g_discrete_space_dual<DDim>.emplace(std::forward<Args>(args)...);
115  ddc_detail::g_discretization_store->emplace(typeid(DDim).name(), []() {
116  ddc_detail::g_discrete_space_dual<DDim>.reset();
117  });
118 #if defined(__CUDACC__)
119  cudaMemcpyToSymbol(
120  ddc_detail::g_discrete_space_device<DDim>,
121  &ddc_detail::g_discrete_space_dual<DDim>->get_device(),
122  sizeof(ddc_detail::g_discrete_space_dual<DDim>->get_device()));
123 #elif defined(__HIPCC__)
124  hipMemcpyToSymbol(
125  ddc_detail::g_discrete_space_device<DDim>,
126  &ddc_detail::g_discrete_space_dual<DDim>->get_device(),
127  sizeof(ddc_detail::g_discrete_space_dual<DDim>->get_device()));
128 #endif
129 }
130 
136 template <class DDimImpl, class Arg>
137 Arg init_discrete_space(std::tuple<DDimImpl, Arg>&& a)
138 {
139  using DDim = typename DDimImpl::discrete_dimension_type;
140  init_discrete_space<DDim>(std::move(std::get<0>(a)));
141  return std::get<1>(a);
142 }
143 
149 template <class DDimImpl, class... Args>
150 std::enable_if_t<2 <= sizeof...(Args), std::tuple<Args...>> init_discrete_space(
151  std::tuple<DDimImpl, Args...>&& a)
152 {
153  using DDim = typename DDimImpl::discrete_dimension_type;
154  init_discrete_space<DDim>(std::move(std::get<0>(a)));
155  return ddc_detail::extract_after(std::move(a), std::index_sequence_for<Args...>());
156 }
157 
158 template <class DDim, class MemorySpace = DDC_CURRENT_KOKKOS_SPACE>
159 DDC_INLINE_FUNCTION ddc_detail::ddim_impl_t<DDim, MemorySpace> const& discrete_space()
160 {
161  if constexpr (std::is_same_v<MemorySpace, Kokkos::HostSpace>) {
162  return ddc_detail::g_discrete_space_dual<DDim>->get_host();
163  }
164 #if defined(__CUDACC__)
165  else if constexpr (std::is_same_v<MemorySpace, Kokkos::CudaSpace>) {
166  return *ddc_detail::g_discrete_space_device<DDim>;
167  }
168 #elif defined(__HIPCC__)
169  else if constexpr (std::is_same_v<MemorySpace, Kokkos::Experimental::HIPSpace>) {
170  return *ddc_detail::g_discrete_space_device<DDim>;
171  }
172 #endif
173  else {
174  static_assert(std::is_same_v<MemorySpace, MemorySpace>, "Memory space not handled");
175  }
176 }
177 
178 } // namespace ddc
Definition: discrete_space.hpp:32
DDC_INLINE_FUNCTION T & operator*()
Definition: discrete_space.hpp:54
DDC_INLINE_FUNCTION T * data()
Definition: discrete_space.hpp:60
DDC_INLINE_FUNCTION T * operator->()
Definition: discrete_space.hpp:48
std::optional< std::map< std::string, std::function< void()> > > g_discretization_store
Definition: discrete_space.hpp:67
std::optional< DualDiscretization< DDim > > g_discrete_space_dual
Definition: discrete_space.hpp:71
void display_discretization_store(std::ostream &os)
Definition: discrete_space.hpp:84
auto extract_after(Tuple &&t, std::index_sequence< Ids... >)
Definition: discrete_space.hpp:97
typename DDim::template Impl< MemorySpace > ddim_impl_t
Definition: discrete_space.hpp:28
Definition: aligned_allocator.hpp:9
void init_discrete_space(Args &&... args)
Initialize (emplace) a global singleton discrete space.
Definition: discrete_space.hpp:109