accelerator.hpp

Go to the documentation of this file.

// Copyright (c) Lawrence Livermore National Security, LLC and
// other Smith Project Developers. See the top-level LICENSE file for
// details.
//
// SPDX-License-Identifier: (BSD-3-Clause)
 
#pragma once
 
#include "smith/smith_config.hpp"
 
#if defined(SMITH_USE_CUDA) || defined(SMITH_USE_HIP)
#define SMITH_HOST_DEVICE __host__ __device__
#define SMITH_HOST __host__
#define SMITH_DEVICE __device__
#else
#define SMITH_HOST_DEVICE
#define SMITH_HOST
#define SMITH_DEVICE
#endif
 
#if defined(__CUDACC__)
#if __CUDAVER__ >= 75000
#define SMITH_SUPPRESS_NVCC_HOSTDEVICE_WARNING #pragma nv_exec_check_disable
#else
#define SMITH_SUPPRESS_NVCC_HOSTDEVICE_WARNING #pragma hd_warning_disable
#endif
#include <cuda_runtime.h>
#else
#define SMITH_SUPPRESS_NVCC_HOSTDEVICE_WARNING
#endif
 
#include <memory>
#include <cstring>
#include <tuple>
 
#include "axom/core.hpp"
 
#include "smith/infrastructure/logger.hpp"
#include "smith/infrastructure/memory.hpp"
#include "smith/infrastructure/profiling.hpp"
#include "smith/smith_config.hpp"
 
namespace smith {
 
enum class ExecutionSpace
{
  CPU,
  GPU,
  Dynamic  // Corresponds to execution that can "legally" happen on either the host or device
};
 
constexpr ExecutionSpace default_execution_space = ExecutionSpace::CPU;
 
namespace detail {
 
template <ExecutionSpace space>
struct execution_to_memory {
  static constexpr axom::MemorySpace value = axom::MemorySpace::Dynamic;
};
 
#ifdef SMITH_USE_UMPIRE
template <>
struct execution_to_memory<ExecutionSpace::CPU> {
  static constexpr axom::MemorySpace value = axom::MemorySpace::Host;
};
 
template <>
struct execution_to_memory<ExecutionSpace::GPU> {
  static constexpr axom::MemorySpace value = axom::MemorySpace::Device;
};
 
template <>
struct execution_to_memory<ExecutionSpace::Dynamic> {
  static constexpr axom::MemorySpace value = axom::MemorySpace::Unified;
};
#endif
 
template <ExecutionSpace space>
inline constexpr axom::MemorySpace execution_to_memory_v = execution_to_memory<space>::value;
 
template <typename T, int dim, axom::MemorySpace space>
void zero_out(axom::Array<T, dim, space>& arr)
{
  std::memset(arr.data(), 0, static_cast<std::size_t>(arr.size()) * sizeof(T));
}
 
template <typename T, int dim>
void zero_out(axom::ArrayView<T, dim, detail::host_memory_space>& arr)
{
  std::memset(arr.data(), 0, static_cast<std::size_t>(arr.size()) * sizeof(T));
}
#ifdef __CUDACC__
template <typename T, int dim>
void zero_out(axom::Array<T, dim, execution_to_memory_v<ExecutionSpace::GPU>>& arr)
{
  cudaMemset(arr.data(), 0, static_cast<std::size_t>(arr.size()) * sizeof(T));
}
#endif
 
}  // namespace detail
 
template <typename T, int dim, ExecutionSpace space>
using ExecArray = axom::Array<T, dim, detail::execution_to_memory_v<space>>;
 
template <typename T, int dim = 1>
using CPUArray = ExecArray<T, dim, ExecutionSpace::CPU>;
 
#ifdef SMITH_USE_CUDA
 
template <typename T, int dim = 1>
using GPUArray = ExecArray<T, dim, ExecutionSpace::GPU>;
 
template <typename T, int dim = 1>
using UnifiedArray = ExecArray<T, dim, ExecutionSpace::Dynamic>;
 
#else
// If not a CUDA build then force all arrays to be CPU
 
template <typename T, int dim = 1>
using GPUArray = ExecArray<T, dim, ExecutionSpace::CPU>;
 
template <typename T, int dim = 1>
using UnifiedArray = ExecArray<T, dim, ExecutionSpace::CPU>;
 
#endif
 
template <typename T, int dim, ExecutionSpace space>
using ExecArrayView = axom::ArrayView<T, dim, detail::execution_to_memory_v<space>>;
 
template <typename T, int dim = 1>
using CPUArrayView = ExecArrayView<T, dim, ExecutionSpace::CPU>;
 
#ifdef SMITH_USE_CUDA
template <typename T, int dim = 1>
using GPUArrayView = ExecArrayView<T, dim, ExecutionSpace::GPU>;
#endif
 
template <typename T, int dim, axom::MemorySpace space>
auto view(axom::Array<T, dim, space>& arr)
{
  return axom::ArrayView<T, dim, space>(arr);
}
 
namespace accelerator {
 
void initializeDevice(ExecutionSpace exec_space);
 
void terminateDevice();
 
#if defined(__CUDACC__)
 
inline void displayLastCUDAMessage(const char* success_string = "", bool exit_on_error = false)
{
  auto error = cudaGetLastError();
  if (error != cudaError::cudaSuccess) {
    if (exit_on_error) {
      SLIC_ERROR_ROOT(smith::profiling::concat("Last CUDA Error Message :", cudaGetErrorString(error)));
    } else {
      SLIC_WARNING_ROOT(smith::profiling::concat("Last CUDA Error Message :", cudaGetErrorString(error)));
    }
  } else if (strlen(success_string) > 0) {
    SLIC_INFO_ROOT(success_string);
  }
}
 
inline std::tuple<std::size_t, std::size_t> getCUDAMemInfo()
{
  std::size_t free_memory, total_memory;
  cudaMemGetInfo(&free_memory, &total_memory);
  displayLastCUDAMessage();
  return std::make_tuple(free_memory, total_memory);
}
 
inline std::string getCUDAMemInfoString()
{
  auto [free_memory, total_memory] = getCUDAMemInfo();
  return std::format("Free memory: {} Total_memory: {}", free_memory, total_memory);
}
 
#endif
 
template <ExecutionSpace exec, typename T>
std::shared_ptr<T[]> make_shared_array(std::size_t n)
{
  if constexpr (exec == ExecutionSpace::CPU) {
    return std::shared_ptr<T[]>(new T[n]);
  }
 
#if defined(__CUDACC__)
  if constexpr (exec == ExecutionSpace::GPU) {
    T* data;
    cudaMalloc(&data, sizeof(T) * n);
    auto deleter = [](T* ptr) { cudaFree(ptr); };
    return std::shared_ptr<T[]>(data, deleter);
  }
#endif
}
 
template <ExecutionSpace exec, typename... T>
auto make_shared_arrays(std::size_t n)
{
  return std::tuple{make_shared_array<exec, T>(n)...};
}
 
}  // namespace accelerator
 
}  // namespace smith