1 #ifndef TPETRA_DETAILS_EXECUTIONSPACES_HPP
2 #define TPETRA_DETAILS_EXECUTIONSPACES_HPP
8 #include <Kokkos_Core.hpp>
10 #include <Teuchos_RCP.hpp>
31 #define TPETRA_DETAILS_SPACES_THROW(x) \
33 std::stringstream ss; \
34 ss << __FILE__ << ":" << __LINE__ << ": " << x; \
35 throw std::runtime_error(ss.str()); \
54 #if defined(KOKKOS_ENABLE_CUDA)
55 inline void success_or_throw(cudaError_t err,
const char *file,
57 if (err != cudaSuccess) {
59 ss << file <<
":" << line <<
": ";
60 ss << cudaGetErrorString(err);
61 throw std::runtime_error(ss.str());
64 #define TPETRA_DETAILS_SPACES_CUDA_RUNTIME(x) \
65 Tpetra::Details::Spaces::success_or_throw((x), __FILE__, __LINE__)
66 #endif // KOKKOS_ENABLE_CUDA
76 #if defined(KOKKOS_ENABLE_CUDA)
82 cudaEvent_t execSpaceWaitEvent_;
85 ~CudaInfo() =
default;
86 CudaInfo(
const CudaInfo &other) =
delete;
87 CudaInfo(CudaInfo &&other) =
delete;
89 extern CudaInfo cudaInfo;
90 #endif // KOKKOS_ENABLE_CUDA
93 #if defined(KOKKOS_ENABLE_CUDA)
94 template <
typename Space>
95 using IsCuda = std::enable_if_t<std::is_same_v<Space, Kokkos::Cuda>,
bool>;
96 template <
typename Space>
97 using NotCuda = std::enable_if_t<!std::is_same_v<Space, Kokkos::Cuda>,
bool>;
98 template <
typename S1,
typename S2>
99 using BothCuda = std::enable_if_t<
100 std::is_same_v<S1, Kokkos::Cuda> && std::is_same_v<S2, Kokkos::Cuda>,
bool>;
101 template <
typename S1,
typename S2>
102 using NotBothCuda = std::enable_if_t<!std::is_same_v<S1, Kokkos::Cuda> ||
103 !std::is_same_v<S2, Kokkos::Cuda>,
105 #endif // KOKKOS_ENABLE_CUDA
107 #if defined(KOKKOS_ENABLE_SERIAL)
108 template <
typename Space>
110 using IsSerial = std::enable_if_t<std::is_same_v<Space, Kokkos::Serial>,
bool>;
111 #endif // KOKKOS_ENABLE_SERIAL
113 #if defined(KOKKOS_ENABLE_OPENMP)
114 template <
typename Space>
116 using IsOpenMP = std::enable_if_t<std::is_same_v<Space, Kokkos::OpenMP>,
bool>;
117 #endif // KOKKOS_ENABLE_OPENMP
119 #if defined(KOKKOS_ENABLE_HIP)
120 template <
typename Space>
122 using IsHIP = std::enable_if_t<std::is_same_v<Space, Kokkos::HIP>,
bool>;
123 #endif // KOKKOS_ENABLE_HIP
125 #if defined(KOKKOS_ENABLE_SYCL)
126 template <
typename Space>
128 using IsSYCL = std::enable_if_t<std::is_same_v<Space, Kokkos::Experimental::SYCL>,
bool>;
129 #endif // KOKKOS_ENABLE_SYCL
136 template <
typename ExecSpace, Priority priority = Priority::medium
137 #if defined(KOKKOS_ENABLE_CUDA)
139 NotCuda<ExecSpace> =
true
140 #endif // KOKKOS_ENABLE_CUDA
142 ExecSpace make_instance() {
152 #if defined(KOKKOS_ENABLE_CUDA)
153 template <
typename ExecSpace, Priority priority = Priority::medium,
154 IsCuda<ExecSpace> =
true>
155 Kokkos::Cuda make_instance() {
161 prio = cudaInfo.highPrio_;
163 case Priority::medium:
164 prio = cudaInfo.mediumPrio_;
167 prio = cudaInfo.lowPrio_;
170 throw std::runtime_error(
"unexpected static Tpetra Space priority");
172 TPETRA_DETAILS_SPACES_CUDA_RUNTIME(
173 cudaStreamCreateWithPriority(&stream, cudaStreamNonBlocking, prio));
174 return Kokkos::Cuda(stream,
true );
176 #endif // KOKKOS_ENABLE_CUDA
183 template <
typename ExecSpace> ExecSpace make_instance(
const Priority &prio) {
186 return make_instance<ExecSpace, Priority::high>();
187 case Priority::medium:
188 return make_instance<ExecSpace, Priority::medium>();
190 return make_instance<ExecSpace, Priority::low>();
192 throw std::runtime_error(
"unexpected dynamic Tpetra Space priority");
209 using execution_space = ExecSpace;
210 using rcp_type = Teuchos::RCP<const execution_space>;
218 template <Priority priority = Priority::medium>
221 "Tpetra::Details::Spaces::space_instance");
223 constexpr
int p =
static_cast<int>(priority);
224 static_assert(p <
sizeof(instances) /
sizeof(instances[0]),
225 "Spaces::Priority enum error");
228 TPETRA_DETAILS_SPACES_THROW(
"requested instance id " << i <<
" (< 0)");
231 TPETRA_DETAILS_SPACES_THROW(
232 "requested instance id "
234 <<
") set by TPETRA_SPACES_ID_WARN_LIMIT");
239 while (
size_t(i) >= instances[p].size()) {
240 instances[p].push_back(Teuchos::ENull());
252 if (instances[p][i].is_null() || !instances[p][i].is_valid_ptr()) {
255 rcp_type r = Teuchos::RCP<const execution_space>(
256 new ExecSpace(make_instance<ExecSpace, priority>()));
259 instances[p][i] = r.create_weak();
265 auto r = instances[p][i].create_strong();
273 for (
int i = 0; i < static_cast<int>(Spaces::Priority::NUM_LEVELS); ++i) {
274 for (
const rcp_type &rcp : instances[i]) {
275 if (rcp.is_valid_ptr() && !rcp.is_null()) {
277 std::cerr << __FILE__ <<
":" << __LINE__
278 <<
" execution space instance survived to "
279 "~InstanceLifetimeManager. strong_count() = "
280 << rcp.strong_count()
281 <<
". Did a Tpetra object live past Kokkos::finalize()?"
290 std::vector<rcp_type>
291 instances[
static_cast<int>(Spaces::Priority::NUM_LEVELS)];
294 #if defined(KOKKOS_ENABLE_CUDA)
295 extern InstanceLifetimeManager<Kokkos::Cuda> cudaSpaces;
297 #if defined(KOKKOS_ENABLE_SERIAL)
298 extern InstanceLifetimeManager<Kokkos::Serial> serialSpaces;
300 #if defined(KOKKOS_ENABLE_OPENMP)
301 extern InstanceLifetimeManager<Kokkos::OpenMP> openMPSpaces;
303 #if defined(KOKKOS_ENABLE_HIP)
304 extern InstanceLifetimeManager<Kokkos::HIP> HIPSpaces;
306 #if defined(KOKKOS_ENABLE_SYCL)
307 extern InstanceLifetimeManager<Kokkos::Experimental::SYCL> SYCLSpaces;
310 #if defined(KOKKOS_ENABLE_CUDA)
315 template <
typename ExecSpace, Priority priority = Priority::medium,
316 IsCuda<ExecSpace> =
true>
317 Teuchos::RCP<const ExecSpace> space_instance(
int i = 0) {
318 return cudaSpaces.space_instance<priority>(i);
322 #if defined(KOKKOS_ENABLE_SERIAL)
326 template <
typename ExecSpace, Priority priority = Priority::medium,
327 IsSerial<ExecSpace> =
true>
328 Teuchos::RCP<const ExecSpace> space_instance(
int i = 0) {
329 return serialSpaces.space_instance<priority>(i);
333 #if defined(KOKKOS_ENABLE_OPENMP)
337 template <
typename ExecSpace, Priority priority = Priority::medium,
338 IsOpenMP<ExecSpace> =
true>
339 Teuchos::RCP<const ExecSpace> space_instance(
int i = 0) {
340 return openMPSpaces.space_instance<priority>(i);
344 #if defined(KOKKOS_ENABLE_HIP)
347 template <
typename ExecSpace, Priority priority = Priority::medium,
348 IsHIP<ExecSpace> =
true>
349 Teuchos::RCP<const ExecSpace> space_instance(
int i = 0) {
350 return HIPSpaces.space_instance<priority>(i);
353 #if defined(KOKKOS_ENABLE_SYCL)
357 template <
typename ExecSpace, Priority priority = Priority::medium,
358 IsSYCL<ExecSpace> =
true>
359 Teuchos::RCP<const ExecSpace> space_instance(
int i = 0) {
360 return SYCLSpaces.space_instance<priority>(i);
369 template <
typename ExecSpace>
370 Teuchos::RCP<const ExecSpace> space_instance(
const Priority &priority,
374 return space_instance<ExecSpace, Priority::high>(i);
375 case Priority::medium:
376 return space_instance<ExecSpace, Priority::medium>(i);
378 return space_instance<ExecSpace, Priority::low>(i);
380 throw std::runtime_error(
381 "unexpected dynamic Tpetra Space priority in space_instance");
398 template <
typename S1,
typename S2
399 #if defined(KOKKOS_ENABLE_CUDA)
401 NotBothCuda<S1, S2> =
true
404 void exec_space_wait(
const char *msg,
const S1 &waitee,
const S2 & ) {
406 "Tpetra::Details::Spaces::exec_space_wait");
411 #if defined(KOKKOS_ENABLE_CUDA)
412 template <
typename S1,
typename S2, BothCuda<S1, S2> = true>
413 void exec_space_wait(
const char *msg,
const S1 &waitee,
const S2 &waiter) {
415 "Tpetra::Details::Spaces::exec_space_wait");
419 if (waitee.impl_instance_id() !=
421 .impl_instance_id()) {
426 TPETRA_DETAILS_SPACES_CUDA_RUNTIME(
427 cudaEventRecord(cudaInfo.execSpaceWaitEvent_, waitee.cuda_stream()));
428 TPETRA_DETAILS_SPACES_CUDA_RUNTIME(cudaStreamWaitEvent(
429 waiter.cuda_stream(), cudaInfo.execSpaceWaitEvent_, 0 ));
434 template <
typename S1,
typename S2>
435 void exec_space_wait(
const S1 &waitee,
const S2 &waiter) {
437 "Tpetra::Details::Spaces::exec_space_wait");
439 exec_space_wait(
"anonymous", waitee, waiter);
442 template <
typename ExecutionSpace>
443 constexpr KOKKOS_INLINE_FUNCTION
bool is_gpu_exec_space() {
447 #if defined(KOKKOS_ENABLE_CUDA)
449 constexpr KOKKOS_INLINE_FUNCTION
bool is_gpu_exec_space<Kokkos::Cuda>() {
454 #if defined(KOKKOS_ENABLE_HIP)
456 constexpr KOKKOS_INLINE_FUNCTION
bool
457 is_gpu_exec_space<Kokkos::HIP>() {
462 #if defined(KOKKOS_ENABLE_SYCL)
464 constexpr KOKKOS_INLINE_FUNCTION
bool
465 is_gpu_exec_space<Kokkos::Experimental::SYCL>() {
474 #undef TPETRA_DETAILS_SPACES_THROW
476 #endif // TPETRA_DETAILS_EXECUTIONSPACES_HPP
Declaration of Tpetra::Details::Profiling, a scope guard for Kokkos Profiling.
static size_t spacesIdWarnLimit()
Warn if more than this many Kokkos spaces are accessed.
rcp_type space_instance(int i=0)
Retrieve a strong Teuchos::RCP<const ExecSpace> to instance i
~InstanceLifetimeManager()
Issue a warning if any Tpetra-managed execution space instances survive to the end of static lifetime...
Provides reusable Kokkos execution space instances.
Declaration of Tpetra::Details::Behavior, a class that describes Tpetra's behavior.