space_settings.hpp

Go to the documentation of this file.

#ifndef SPACE_SETTINGS_HPP
#define SPACE_SETTINGS_HPP
#include <Kokkos_Core.hpp>
#include <Kokkos_Random.hpp>
#include <vector>
// Declare the memory and execution spaces.
#ifdef USE_GPU
// 4 eventual possible GPU options: Cuda, Hip, Sycl, OpenMP+
#ifdef USE_CUDA
using MemSpace = Kokkos::CudaSpace;
using ExSpace = Kokkos::Cuda;
#elif defined(USE_HIP)
using MemSpace = Kokkos::Experimental::HIPSpace;
using ExSpace = Kokkos::Experimental::HIP;
#elif defined(USE_SYCL)
using MemSpace = Kokkos::Experimental::SYCLDeviceUSMSpace;
using ExSpace = Kokkos::Experimental::SYCL;
#elif defined(USE_OMPT)
using MemSpace = Kokkos::Experimental::OpenMPTargetSpace;
using ExSpace = Kokkos::Experimental::OpenMPTarget;
#else
#error USE_GPU is on, but no Kokkos backend specified
#endif

#else
// USE_GPU undefined
using MemSpace = Kokkos::HostSpace;
#ifdef USE_OMP
using ExSpace = Kokkos::OpenMP;
#elif defined(USE_SERIAL)
using ExSpace = Kokkos::Serial;
#else
#error No CPU backend specified (USE_OMP or USE_SERIAL)
#endif

#endif

#ifndef USE_CUDA
// Native cuda min/max is used with cuda since std library is not available
// Otherwise, use std min/max.
using std::min;
using std::max;
using std::abs;
using std::isnan;
#endif

using DeviceType = Kokkos::Device<ExSpace,MemSpace>;

// Designate a CPU-specific Memory and Execution space
using HostMemSpace = Kokkos::HostSpace;
#ifdef USE_OMP
using HostExSpace = Kokkos::OpenMP;
#elif defined(USE_SERIAL)
using HostExSpace = Kokkos::Serial;
#endif
using HostType = Kokkos::Device<HostExSpace,HostMemSpace>;

// MPI
#ifdef USE_GPU_AWARE_MPI
using MPIDeviceType = DeviceType;
#else
using MPIDeviceType = HostType;
#endif

/* Some commonly used Kokkos features that are renamed to increase legibility */
using Kokkos::View; // Now just View
using CLayout = Kokkos::LayoutRight; // 'C' Layout, i.e. the right-most index is the contiguous one
using NoInit = Kokkos::ViewAllocateWithoutInitializing; // Turns off default View behavior of initializing to zero upon allocation

// For random number generation
typedef Kokkos::Random_XorShift64_Pool<DeviceType> pool_type;
typedef typename pool_type::generator_type RandGen;

// Which level of memory team scratch is placed in. 0 is kilobytes (L1), 1 is a few gigabytes, 2 is generally node capacity
#ifndef KOKKOS_TEAM_SCRATCH_OPT
#define KOKKOS_TEAM_SCRATCH_OPT 0
#endif

// Default Kokkos behavior is that a kernel cannot be launched until the previous one is complete (essentially a fence)
// HintLightWeight removes this fence and is necessary to allow asynchronous CPU/GPU kernels.
namespace Opt = Kokkos::Experimental;
constexpr static const Kokkos::Experimental::WorkItemProperty::HintLightWeight_t Async = Kokkos::Experimental::WorkItemProperty::HintLightWeight;

#ifndef USE_SYCL
#  define DEVICE_PRINTF(...) printf(__VA_ARGS__)
#else
#  define DEVICE_PRINTF(...) sycl::ext::oneapi::experimental::printf(__VA_ARGS__)
#endif

// SPACE_SETTINGS_HPP
#endif