Skip to content

Commit

Permalink
Basic configuration for SYCL
Browse files Browse the repository at this point in the history
  • Loading branch information
masterleinad committed Oct 5, 2020
1 parent e4028ea commit 32d2187
Show file tree
Hide file tree
Showing 14 changed files with 839 additions and 20 deletions.
5 changes: 5 additions & 0 deletions core/src/CMakeLists.txt
Expand Up @@ -60,6 +60,11 @@ ELSE()
LIST(REMOVE_ITEM KOKKOS_CORE_SRCS ${CMAKE_CURRENT_SOURCE_DIR}/impl/Kokkos_Serial_task.cpp)
ENDIF()

IF (KOKKOS_ENABLE_SYCL)
APPEND_GLOB(KOKKOS_CORE_SRCS ${CMAKE_CURRENT_SOURCE_DIR}/SYCL/*.cpp)
APPEND_GLOB(KOKKOS_CORE_HEADERS ${CMAKE_CURRENT_SOURCE_DIR}/SYCL/*.hpp)
ENDIF()

KOKKOS_ADD_LIBRARY(
kokkoscore
SOURCES ${KOKKOS_CORE_SRCS}
Expand Down
4 changes: 4 additions & 0 deletions core/src/Kokkos_Core.hpp
Expand Up @@ -52,6 +52,10 @@

#include <KokkosCore_Config_DeclareBackend.hpp>

#if defined(KOKKOS_ENABLE_SYCL)
#include <Kokkos_SYCL.hpp>
#endif

#include <Kokkos_AnonymousSpace.hpp>
#include <Kokkos_Pair.hpp>
#include <Kokkos_MemoryPool.hpp>
Expand Down
9 changes: 8 additions & 1 deletion core/src/Kokkos_Core_fwd.hpp
Expand Up @@ -126,6 +126,9 @@ using DefaultExecutionSpace KOKKOS_IMPL_DEFAULT_EXEC_SPACE_ANNOTATION =
#elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_HIP)
using DefaultExecutionSpace KOKKOS_IMPL_DEFAULT_EXEC_SPACE_ANNOTATION =
Experimental::HIP;
#elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_SYCL)
using DefaultExecutionSpace KOKKOS_IMPL_DEFAULT_EXEC_SPACE_ANNOTATION =
Experimental::SYCL;
#elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMP)
using DefaultExecutionSpace KOKKOS_IMPL_DEFAULT_EXEC_SPACE_ANNOTATION = OpenMP;
#elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_THREADS)
Expand All @@ -137,7 +140,7 @@ using DefaultExecutionSpace KOKKOS_IMPL_DEFAULT_EXEC_SPACE_ANNOTATION =
using DefaultExecutionSpace KOKKOS_IMPL_DEFAULT_EXEC_SPACE_ANNOTATION = Serial;
#else
#error \
"At least one of the following execution spaces must be defined in order to use Kokkos: Kokkos::Cuda, Kokkos::Experimental::HIP, Kokkos::Experimental::OpenMPTarget, Kokkos::OpenMP, Kokkos::Threads, Kokkos::Experimental::HPX, or Kokkos::Serial."
"At least one of the following execution spaces must be defined in order to use Kokkos: Kokkos::Cuda, Kokkos::Experimental::HIP, Kokkos::Experimental::SYCL, Kokkos::Experimental::OpenMPTarget, Kokkos::OpenMP, Kokkos::Threads, Kokkos::Experimental::HPX, or Kokkos::Serial."
#endif
#if defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMP)
Expand Down Expand Up @@ -183,6 +186,10 @@ namespace Impl {
#if defined(KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_CUDA) && \
defined(KOKKOS_ENABLE_CUDA)
using ActiveExecutionMemorySpace = Kokkos::CudaSpace;
#elif defined(KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_SYCL)
using ActiveExecutionMemorySpace = Kokkos::Experimental::SYCLDeviceUSMSpace;
// FIXME_SYCL
// using ActiveExecutionMemorySpace = Kokkos::Experimental::SYCLHostUSMSpace;
#elif defined(KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_HIP_GPU)
using ActiveExecutionMemorySpace = Kokkos::Experimental::HIPSpace;
#elif defined(KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_HOST)
Expand Down
29 changes: 12 additions & 17 deletions core/src/Kokkos_Macros.hpp
Expand Up @@ -97,9 +97,10 @@

//----------------------------------------------------------------------------

#if !defined(KOKKOS_ENABLE_THREADS) && !defined(KOKKOS_ENABLE_CUDA) && \
!defined(KOKKOS_ENABLE_OPENMP) && !defined(KOKKOS_ENABLE_HPX) && \
!defined(KOKKOS_ENABLE_OPENMPTARGET) && !defined(KOKKOS_ENABLE_HIP)
#if !defined(KOKKOS_ENABLE_THREADS) && !defined(KOKKOS_ENABLE_CUDA) && \
!defined(KOKKOS_ENABLE_OPENMP) && !defined(KOKKOS_ENABLE_HPX) && \
!defined(KOKKOS_ENABLE_OPENMPTARGET) && !defined(KOKKOS_ENABLE_HIP) && \
!defined(KOKKOS_ENABLE_SYCL)
#define KOKKOS_INTERNAL_NOT_PARALLEL
#endif

Expand Down Expand Up @@ -186,16 +187,6 @@
#define KOKKOS_COMPILER_MSVC _MSC_VER
#endif

#if defined(KOKKOS_ENABLE_HIP)
#define KOKKOS_IMPL_FORCEINLINE_FUNCTION __device__ __host__ __forceinline__
#define KOKKOS_IMPL_INLINE_FUNCTION __device__ __host__ inline
#define KOKKOS_DEFAULTED_FUNCTION __device__ __host__ inline
#define KOKKOS_INLINE_FUNCTION_DELETED __device__ __host__ inline
#define KOKKOS_IMPL_FUNCTION __device__ __host__
#define KOKKOS_IMPL_HOST_FUNCTION __host__
#define KOKKOS_IMPL_DEVICE_FUNCTION __device__
#endif // #if defined( KOKKOS_ENABLE_HIP )

#if defined(_OPENMP)
// Compiling with OpenMP.
// The value of _OPENMP is an integer value YYYYMM
Expand Down Expand Up @@ -428,6 +419,7 @@

#if 1 < ((defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_CUDA) ? 1 : 0) + \
(defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_HIP) ? 1 : 0) + \
(defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_SYCL) ? 1 : 0) + \
(defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMPTARGET) ? 1 : 0) + \
(defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMP) ? 1 : 0) + \
(defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_THREADS) ? 1 : 0) + \
Expand All @@ -437,9 +429,10 @@
#endif

// If default is not specified then chose from enabled execution spaces.
// Priority: CUDA, HIP, OPENMPTARGET, OPENMP, THREADS, HPX, SERIAL
// Priority: CUDA, HIP, SYCL, OPENMPTARGET, OPENMP, THREADS, HPX, SERIAL
#if defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_CUDA)
#elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_HIP)
#elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_SYCL)
#elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMPTARGET)
#elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMP)
#elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_THREADS)
Expand All @@ -454,6 +447,8 @@
// as valid overload criteria
#define KOKKOS_IMPL_ENABLE_OVERLOAD_HOST_DEVICE
#endif
#elif defined(KOKKOS_ENABLE_SYCL)
#define KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_SYCL
#elif defined(KOKKOS_ENABLE_OPENMPTARGET)
#define KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMPTARGET
#elif defined(KOKKOS_ENABLE_OPENMP)
Expand All @@ -471,6 +466,8 @@

#if defined(__CUDACC__) && defined(__CUDA_ARCH__) && defined(KOKKOS_ENABLE_CUDA)
#define KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_CUDA
#elif defined(__SYCL_DEVICE_ONLY__) && defined(KOKKOS_ENABLE_SYCL)
#define KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_SYCL
#elif defined(__HIPCC__) && defined(__HIP_DEVICE_COMPILE__) && \
defined(KOKKOS_ENABLE_HIP)
#define KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_HIP_GPU
Expand All @@ -495,11 +492,9 @@
#if defined(KOKKOS_ENABLE_CUDA_RELOCATABLE_DEVICE_CODE)
#define KOKKOS_ENABLE_TASKDAG
#endif
#else
#ifndef KOKKOS_ENABLE_HIP
#elif !defined(KOKKOS_ENABLE_HIP) && !defined(KOKKOS_ENABLE_SYCL)
#define KOKKOS_ENABLE_TASKDAG
#endif
#endif

#if defined(KOKKOS_ENABLE_CUDA)
#define KOKKOS_IMPL_CUDA_VERSION_9_WORKAROUND
Expand Down
178 changes: 178 additions & 0 deletions core/src/Kokkos_SYCL.hpp
@@ -0,0 +1,178 @@
/*
//@HEADER
// ************************************************************************
//
// Kokkos v. 3.0
// Copyright (2020) National Technology & Engineering
// Solutions of Sandia, LLC (NTESS).
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are
// met:
//
// 1. Redistributions of source code must retain the above copyright
// notice, this list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimer in the
// documentation and/or other materials provided with the distribution.
//
// 3. Neither the name of the Corporation nor the names of the
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY NTESS "AS IS" AND ANY
// EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
// PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NTESS OR THE
// CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
// EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
// PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
// LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER
*/

#ifndef KOKKOS_SYCL_HPP
#define KOKKOS_SYCL_HPP

#include <Kokkos_Macros.hpp>

#ifdef KOKKOS_ENABLE_SYCL
#include <CL/sycl.hpp>
#include <Kokkos_SYCL_Space.hpp>
#include <Kokkos_Layout.hpp>
#include <Kokkos_ScratchSpace.hpp>
#include <impl/Kokkos_ExecSpaceInitializer.hpp>

namespace Kokkos {
namespace Experimental {
namespace Impl {
class SYCLInternal;
}

/// \class SYCL
/// \brief Kokkos device for multicore processors in the host memory space.
class SYCL {
public:
//------------------------------------
//! \name Type declarations that all Kokkos devices must provide.
//@{

//! Tag this class as a kokkos execution space
using execution_space = SYCL;
using memory_space = SYCLDeviceUSMSpace;
using device_type = Kokkos::Device<execution_space, memory_space>;

using array_layout = LayoutLeft;
using size_type = memory_space::size_type;

using scratch_memory_space = ScratchMemorySpace<SYCL>;

~SYCL() = default;
SYCL();

SYCL(SYCL&&) = default;
SYCL(const SYCL&) = default;
SYCL& operator=(SYCL&&) = default;
SYCL& operator=(const SYCL&) = default;

uint32_t impl_instance_id() const noexcept { return 0; }

//@}
//------------------------------------
//! \name Functions that all Kokkos devices must implement.
//@{

KOKKOS_INLINE_FUNCTION static int in_parallel() {
#if defined(__SYCL_ARCH__)
return true;
#else
return false;
#endif
}

/** \brief Set the device in a "sleep" state. */
static bool sleep();

/** \brief Wake the device from the 'sleep' state. A noop for OpenMP. */
static bool wake();

/** \brief Wait until all dispatched functors complete. A noop for OpenMP. */
static void impl_static_fence();
void fence() const;

/// \brief Print configuration information to the given output stream.
static void print_configuration(std::ostream&, const bool detail = false);

/// \brief Free any resources being consumed by the device.
static void impl_finalize();

/** \brief Initialize the device.
*
*/

struct SYCLDevice {
SYCLDevice();
explicit SYCLDevice(cl::sycl::device d);
explicit SYCLDevice(const cl::sycl::device_selector& selector);
explicit SYCLDevice(size_t id);
explicit SYCLDevice(const std::function<bool(const sycl::device&)>& pred);

cl::sycl::device get_device() const;

friend std::ostream& operator<<(std::ostream& os, const SYCLDevice& that) {
return that.info(os);
}

static std::ostream& list_devices(std::ostream& os);
static void list_devices();

private:
std::ostream& info(std::ostream& os) const;

cl::sycl::device m_device;
};

static void impl_initialize(SYCLDevice = SYCLDevice());

int sycl_device() const;

static bool impl_is_initialized();

static int concurrency();
static const char* name();

inline Impl::SYCLInternal* impl_internal_space_instance() const {
return m_space_instance;
}

private:
Impl::SYCLInternal* m_space_instance;
};

namespace Impl {

class SYCLSpaceInitializer : public Kokkos::Impl::ExecSpaceInitializerBase {
public:
void initialize(const InitArguments& args) final;
void finalize(const bool) final;
void fence() final;
void print_configuration(std::ostream& msg, const bool detail) final;
};

} // namespace Impl
} // namespace Experimental
} // namespace Kokkos

#endif
#endif
77 changes: 77 additions & 0 deletions core/src/Kokkos_SYCL_Space.hpp
@@ -0,0 +1,77 @@
/*
//@HEADER
// ************************************************************************
//
// Kokkos v. 3.0
// Copyright (2020) National Technology & Engineering
// Solutions of Sandia, LLC (NTESS).
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are
// met:
//
// 1. Redistributions of source code must retain the above copyright
// notice, this list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimer in the
// documentation and/or other materials provided with the distribution.
//
// 3. Neither the name of the Corporation nor the names of the
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY NTESS "AS IS" AND ANY
// EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
// PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NTESS OR THE
// CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
// EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
// PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
// LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER
*/

#ifndef KOKKOS_SYCLSPACE_HPP
#define KOKKOS_SYCLSPACE_HPP

#include <Kokkos_Core_fwd.hpp>

#ifdef KOKKOS_ENABLE_SYCL

namespace Kokkos {
namespace Experimental {

class SYCLDeviceUSMSpace {
public:
using execution_space = SYCL;
using memory_space = SYCLDeviceUSMSpace;
using device_type = Kokkos::Device<execution_space, memory_space>;
using size_type = unsigned int;

SYCLDeviceUSMSpace();

void* allocate(const std::size_t arg_alloc_size) const;
void deallocate(void* const arg_alloc_ptr,
const std::size_t arg_alloc_size) const;

static constexpr const char* name() { return "SYCLDeviceUSM"; };

private:
int m_device;
};
} // namespace Experimental
} // namespace Kokkos

#endif
#endif

0 comments on commit 32d2187

Please sign in to comment.