mirror of
https://gitlab.com/libeigen/eigen.git
synced 2026-04-10 11:34:33 +08:00
9ea490c82c
* Modifying TensorDeviceSYCL to use `EIGEN_THROW_X`. * Modifying TensorMacro to use `EIGEN_TRY/CATCH(X)` macro. * Modifying TensorReverse.h to use `EIGEN_DEVICE_REF` instead of `&`. * Fixing the SYCL device macro in SpecialFunctionsImpl.h.
950 lines
35 KiB
C++
950 lines
35 KiB
C++
// This file is part of Eigen, a lightweight C++ template library
|
|
// for linear algebra.
|
|
//
|
|
// Mehdi Goli Codeplay Software Ltd.
|
|
// Ralph Potter Codeplay Software Ltd.
|
|
// Luke Iwanski Codeplay Software Ltd.
|
|
// Contact: <eigen@codeplay.com>
|
|
// Copyright (C) 2016 Benoit Steiner <benoit.steiner.goog@gmail.com>
|
|
|
|
//
|
|
// This Source Code Form is subject to the terms of the Mozilla
|
|
// Public License v. 2.0. If a copy of the MPL was not distributed
|
|
// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
|
|
|
|
#if defined(EIGEN_USE_SYCL) && !defined(EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H)
|
|
#define EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H
|
|
#include <unordered_set>
|
|
|
|
|
|
namespace Eigen {
|
|
|
|
namespace TensorSycl {
|
|
namespace internal {
|
|
|
|
/// Cache all the device information needed
|
|
struct SyclDeviceInfo {
|
|
SyclDeviceInfo(cl::sycl::queue queue)
|
|
: local_mem_type(
|
|
queue.get_device()
|
|
.template get_info<cl::sycl::info::device::local_mem_type>()),
|
|
max_work_item_sizes(
|
|
queue.get_device()
|
|
.template get_info<
|
|
cl::sycl::info::device::max_work_item_sizes>()),
|
|
max_mem_alloc_size(
|
|
queue.get_device()
|
|
.template get_info<
|
|
cl::sycl::info::device::max_mem_alloc_size>()),
|
|
max_compute_units(queue.get_device()
|
|
.template get_info<
|
|
cl::sycl::info::device::max_compute_units>()),
|
|
max_work_group_size(
|
|
queue.get_device()
|
|
.template get_info<
|
|
cl::sycl::info::device::max_work_group_size>()),
|
|
local_mem_size(
|
|
queue.get_device()
|
|
.template get_info<cl::sycl::info::device::local_mem_size>()),
|
|
platform_name(queue.get_device()
|
|
.get_platform()
|
|
.template get_info<cl::sycl::info::platform::name>()),
|
|
device_name(queue.get_device()
|
|
.template get_info<cl::sycl::info::device::name>()),
|
|
device_vendor(
|
|
queue.get_device()
|
|
.template get_info<cl::sycl::info::device::vendor>()) {}
|
|
|
|
cl::sycl::info::local_mem_type local_mem_type;
|
|
cl::sycl::id<3> max_work_item_sizes;
|
|
unsigned long max_mem_alloc_size;
|
|
unsigned long max_compute_units;
|
|
unsigned long max_work_group_size;
|
|
size_t local_mem_size;
|
|
std::string platform_name;
|
|
std::string device_name;
|
|
std::string device_vendor;
|
|
};
|
|
|
|
} // end namespace internal
|
|
} // end namespace TensorSycl
|
|
|
|
typedef TensorSycl::internal::buffer_data_type_t buffer_scalar_t;
|
|
// All devices (even AMD CPU with intel OpenCL runtime) that support OpenCL and
|
|
// can consume SPIR or SPIRV can use the Eigen SYCL backend and consequently
|
|
// TensorFlow via the Eigen SYCL Backend.
|
|
EIGEN_STRONG_INLINE auto get_sycl_supported_devices()
|
|
-> decltype(cl::sycl::device::get_devices()) {
|
|
#ifdef EIGEN_SYCL_USE_DEFAULT_SELECTOR
|
|
return {cl::sycl::device(cl::sycl::default_selector())};
|
|
#else
|
|
std::vector<cl::sycl::device> supported_devices;
|
|
auto platform_list = cl::sycl::platform::get_platforms();
|
|
for (const auto &platform : platform_list) {
|
|
auto device_list = platform.get_devices();
|
|
auto platform_name =
|
|
platform.template get_info<cl::sycl::info::platform::name>();
|
|
std::transform(platform_name.begin(), platform_name.end(),
|
|
platform_name.begin(), ::tolower);
|
|
for (const auto &device : device_list) {
|
|
auto vendor = device.template get_info<cl::sycl::info::device::vendor>();
|
|
std::transform(vendor.begin(), vendor.end(), vendor.begin(), ::tolower);
|
|
bool unsupported_condition =
|
|
(device.is_cpu() && platform_name.find("amd") != std::string::npos &&
|
|
vendor.find("apu") == std::string::npos) ||
|
|
(platform_name.find("experimental") != std::string::npos) ||
|
|
device.is_host();
|
|
if (!unsupported_condition) {
|
|
supported_devices.push_back(device);
|
|
}
|
|
}
|
|
}
|
|
return supported_devices;
|
|
#endif
|
|
}
|
|
|
|
class QueueInterface {
|
|
public:
|
|
/// Creating device by using cl::sycl::selector or cl::sycl::device.
|
|
template <typename DeviceOrSelector>
|
|
explicit QueueInterface(
|
|
const DeviceOrSelector &dev_or_sel, cl::sycl::async_handler handler,
|
|
unsigned num_threads = std::thread::hardware_concurrency())
|
|
: m_queue(dev_or_sel, handler),
|
|
#ifdef EIGEN_SYCL_USE_PROGRAM_CLASS
|
|
m_prog(m_queue.get_context(), get_sycl_supported_devices()),
|
|
#endif
|
|
m_thread_pool(num_threads),
|
|
m_device_info(m_queue) {
|
|
#ifdef EIGEN_SYCL_USE_PROGRAM_CLASS
|
|
m_prog.build_with_kernel_type<DeviceOrSelector>();
|
|
auto f = [&](cl::sycl::handler &cgh) {
|
|
cgh.single_task<DeviceOrSelector>(m_prog.get_kernel<DeviceOrSelector>(),
|
|
[=]() {})
|
|
};
|
|
EIGEN_SYCL_TRY_CATCH(m_queue.submit(f));
|
|
#endif
|
|
}
|
|
|
|
template <typename DeviceOrSelector>
|
|
explicit QueueInterface(
|
|
const DeviceOrSelector &dev_or_sel,
|
|
unsigned num_threads = std::thread::hardware_concurrency())
|
|
: QueueInterface(dev_or_sel,
|
|
[this](cl::sycl::exception_list l) {
|
|
this->exception_caught_ = this->sycl_async_handler(l);
|
|
},
|
|
num_threads) {}
|
|
|
|
#ifdef EIGEN_SYCL_USE_PROGRAM_CLASS
|
|
EIGEN_STRONG_INLINE cl::sycl::program &program() const { return m_prog; }
|
|
#endif
|
|
|
|
/// Attach an existing buffer to the pointer map, Eigen will not reuse it
|
|
EIGEN_STRONG_INLINE void *attach_buffer(
|
|
cl::sycl::buffer<buffer_scalar_t, 1> &buf) const {
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
return static_cast<void *>(pMapper.add_pointer(buf));
|
|
}
|
|
|
|
/// Detach previously attached buffer
|
|
EIGEN_STRONG_INLINE void detach_buffer(void *p) const {
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
TensorSycl::internal::SYCLfree<false>(p, pMapper);
|
|
}
|
|
|
|
/// Allocating device pointer. This pointer is actually an 8 bytes host
|
|
/// pointer used as key to access the sycl device buffer. The reason is that
|
|
/// we cannot use device buffer as a pointer as a m_data in Eigen leafNode
|
|
/// expressions. So we create a key pointer to be used in Eigen expression
|
|
/// construction. When we convert the Eigen construction into the sycl
|
|
/// construction we use this pointer as a key in our buffer_map and we make
|
|
/// sure that we dedicate only one buffer only for this pointer. The device
|
|
/// pointer would be deleted by calling deallocate function.
|
|
EIGEN_STRONG_INLINE void *allocate(size_t num_bytes) const {
|
|
#if EIGEN_MAX_ALIGN_BYTES > 0
|
|
size_t align = num_bytes % EIGEN_MAX_ALIGN_BYTES;
|
|
if (align > 0) {
|
|
num_bytes += EIGEN_MAX_ALIGN_BYTES - align;
|
|
}
|
|
#endif
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
return TensorSycl::internal::SYCLmalloc(num_bytes, pMapper);
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE void *allocate_temp(size_t num_bytes) const {
|
|
#if EIGEN_MAX_ALIGN_BYTES > 0
|
|
size_t align = num_bytes % EIGEN_MAX_ALIGN_BYTES;
|
|
if (align > 0) {
|
|
num_bytes += EIGEN_MAX_ALIGN_BYTES - align;
|
|
}
|
|
#endif
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
#ifndef EIGEN_SYCL_NO_REUSE_BUFFERS
|
|
if (scratch_buffers.empty()) {
|
|
return TensorSycl::internal::SYCLmalloc(num_bytes, pMapper);
|
|
;
|
|
} else {
|
|
for (auto it = scratch_buffers.begin(); it != scratch_buffers.end();) {
|
|
auto buff = pMapper.get_buffer(*it);
|
|
if (buff.get_size() >= num_bytes) {
|
|
auto ptr = *it;
|
|
scratch_buffers.erase(it);
|
|
return ptr;
|
|
} else {
|
|
++it;
|
|
}
|
|
}
|
|
return TensorSycl::internal::SYCLmalloc(num_bytes, pMapper);
|
|
}
|
|
#else
|
|
return TensorSycl::internal::SYCLmalloc(num_bytes, pMapper);
|
|
#endif
|
|
}
|
|
template <typename data_t>
|
|
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorSycl::internal::RangeAccess<
|
|
cl::sycl::access::mode::read_write, data_t>
|
|
get(data_t *data) const {
|
|
return get_range_accessor<cl::sycl::access::mode::read_write, data_t>(data);
|
|
}
|
|
template <typename data_t>
|
|
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE data_t *get(
|
|
TensorSycl::internal::RangeAccess<cl::sycl::access::mode::read_write,
|
|
data_t>
|
|
data) const {
|
|
return static_cast<data_t *>(data.get_virtual_pointer());
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE void deallocate_temp(void *p) const {
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
#ifndef EIGEN_SYCL_NO_REUSE_BUFFERS
|
|
scratch_buffers.insert(p);
|
|
#else
|
|
TensorSycl::internal::SYCLfree(p, pMapper);
|
|
#endif
|
|
}
|
|
template <cl::sycl::access::mode AcMd, typename T>
|
|
EIGEN_STRONG_INLINE void deallocate_temp(
|
|
const TensorSycl::internal::RangeAccess<AcMd, T> &p) const {
|
|
deallocate_temp(p.get_virtual_pointer());
|
|
}
|
|
|
|
/// This is used to deallocate the device pointer. p is used as a key inside
|
|
/// the map to find the device buffer and delete it.
|
|
EIGEN_STRONG_INLINE void deallocate(void *p) const {
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
TensorSycl::internal::SYCLfree(p, pMapper);
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE void deallocate_all() const {
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
TensorSycl::internal::SYCLfreeAll(pMapper);
|
|
#ifndef EIGEN_SYCL_NO_REUSE_BUFFERS
|
|
scratch_buffers.clear();
|
|
#endif
|
|
}
|
|
|
|
/// The memcpyHostToDevice is used to copy the data from host to device
|
|
/// The destination pointer could be deleted before the copy happend which is
|
|
/// why a callback function is needed. By default if none is provided, the
|
|
/// function is blocking.
|
|
EIGEN_STRONG_INLINE void memcpyHostToDevice(
|
|
void *dst, const void *src, size_t n,
|
|
std::function<void()> callback) const {
|
|
static const auto write_mode = cl::sycl::access::mode::discard_write;
|
|
static const auto global_access = cl::sycl::access::target::global_buffer;
|
|
typedef cl::sycl::accessor<buffer_scalar_t, 1, write_mode, global_access>
|
|
write_accessor;
|
|
if (n == 0) {
|
|
if (callback) callback();
|
|
return;
|
|
}
|
|
n /= sizeof(buffer_scalar_t);
|
|
auto f = [&](cl::sycl::handler &cgh) {
|
|
write_accessor dst_acc = get_range_accessor<write_mode>(cgh, dst, n);
|
|
buffer_scalar_t const *ptr = static_cast<buffer_scalar_t const *>(src);
|
|
auto non_deleter = [](buffer_scalar_t const *) {};
|
|
std::shared_ptr<const buffer_scalar_t> s_ptr(ptr, non_deleter);
|
|
cgh.copy(s_ptr, dst_acc);
|
|
};
|
|
cl::sycl::event e;
|
|
EIGEN_SYCL_TRY_CATCH(e = m_queue.submit(f));
|
|
synchronize_and_callback(e, callback);
|
|
}
|
|
|
|
/// The memcpyDeviceToHost is used to copy the data from device to host.
|
|
/// The source pointer could be deleted before the copy happend which is
|
|
/// why a callback function is needed. By default if none is provided, the
|
|
/// function is blocking.
|
|
EIGEN_STRONG_INLINE void memcpyDeviceToHost(
|
|
void *dst, const void *src, size_t n,
|
|
std::function<void()> callback) const {
|
|
static const auto read_mode = cl::sycl::access::mode::read;
|
|
static const auto global_access = cl::sycl::access::target::global_buffer;
|
|
typedef cl::sycl::accessor<buffer_scalar_t, 1, read_mode, global_access>
|
|
read_accessor;
|
|
if (n == 0) {
|
|
if (callback) callback();
|
|
return;
|
|
}
|
|
n /= sizeof(buffer_scalar_t);
|
|
auto f = [&](cl::sycl::handler &cgh) {
|
|
read_accessor src_acc = get_range_accessor<read_mode>(cgh, src, n);
|
|
buffer_scalar_t *ptr = static_cast<buffer_scalar_t *>(dst);
|
|
auto non_deleter = [](buffer_scalar_t *) {};
|
|
std::shared_ptr<buffer_scalar_t> s_ptr(ptr, non_deleter);
|
|
cgh.copy(src_acc, s_ptr);
|
|
};
|
|
cl::sycl::event e;
|
|
EIGEN_SYCL_TRY_CATCH(e = m_queue.submit(f));
|
|
synchronize_and_callback(e, callback);
|
|
}
|
|
|
|
/// The memcpy function.
|
|
/// No callback is required here as both arguments are on the device
|
|
/// and SYCL can handle the dependency.
|
|
EIGEN_STRONG_INLINE void memcpy(void *dst, const void *src, size_t n) const {
|
|
static const auto read_mode = cl::sycl::access::mode::read;
|
|
static const auto write_mode = cl::sycl::access::mode::discard_write;
|
|
if (n == 0) {
|
|
return;
|
|
}
|
|
n /= sizeof(buffer_scalar_t);
|
|
auto f = [&](cl::sycl::handler &cgh) {
|
|
auto src_acc = get_range_accessor<read_mode>(cgh, src, n);
|
|
auto dst_acc = get_range_accessor<write_mode>(cgh, dst, n);
|
|
cgh.copy(src_acc, dst_acc);
|
|
};
|
|
cl::sycl::event e;
|
|
EIGEN_SYCL_TRY_CATCH(e = m_queue.submit(f));
|
|
async_synchronize(e);
|
|
}
|
|
|
|
/// the memset function.
|
|
/// No callback is required here as both arguments are on the device
|
|
/// and SYCL can handle the dependency.
|
|
EIGEN_STRONG_INLINE void memset(void *data, int c, size_t n) const {
|
|
static const auto write_mode = cl::sycl::access::mode::discard_write;
|
|
if (n == 0) {
|
|
return;
|
|
}
|
|
n /= sizeof(buffer_scalar_t);
|
|
auto f = [&](cl::sycl::handler &cgh) {
|
|
auto dst_acc = get_range_accessor<write_mode>(cgh, data, n);
|
|
// The cast to uint8_t is here to match the behaviour of the standard
|
|
// memset. The cast to buffer_scalar_t is needed to match the type of the
|
|
// accessor (in case buffer_scalar_t is not uint8_t)
|
|
cgh.fill(dst_acc, static_cast<buffer_scalar_t>(static_cast<uint8_t>(c)));
|
|
};
|
|
cl::sycl::event e;
|
|
EIGEN_SYCL_TRY_CATCH(e = m_queue.submit(f));
|
|
async_synchronize(e);
|
|
}
|
|
|
|
/// Get a range accessor to the virtual pointer's device memory. This range
|
|
/// accessor will allow access to the memory from the pointer to the end of
|
|
/// the buffer.
|
|
///
|
|
/// NOTE: Inside a kernel the range accessor will always be indexed from the
|
|
/// start of the buffer, so the offset in the accessor is only used by
|
|
/// methods like handler::copy and will not be available inside a kernel.
|
|
template <cl::sycl::access::mode AcMd, typename T>
|
|
EIGEN_STRONG_INLINE TensorSycl::internal::RangeAccess<AcMd, T>
|
|
get_range_accessor(const void *ptr) const {
|
|
static const auto global_access = cl::sycl::access::target::global_buffer;
|
|
static const auto is_place_holder = cl::sycl::access::placeholder::true_t;
|
|
typedef TensorSycl::internal::RangeAccess<AcMd, T> ret_type;
|
|
typedef const TensorSycl::internal::buffer_data_type_t *internal_ptr_t;
|
|
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
|
|
auto original_buffer = pMapper.get_buffer(ptr);
|
|
const ptrdiff_t offset = pMapper.get_offset(ptr);
|
|
const ptrdiff_t typed_offset = offset / sizeof(T);
|
|
eigen_assert(typed_offset >= 0);
|
|
const auto typed_size = original_buffer.get_size() / sizeof(T);
|
|
auto buffer = original_buffer.template reinterpret<
|
|
typename Eigen::internal::remove_const<T>::type>(
|
|
cl::sycl::range<1>(typed_size));
|
|
const ptrdiff_t size = buffer.get_count() - typed_offset;
|
|
eigen_assert(size >= 0);
|
|
typedef cl::sycl::accessor<typename Eigen::internal::remove_const<T>::type,
|
|
1, AcMd, global_access, is_place_holder>
|
|
placeholder_accessor_t;
|
|
const auto start_ptr = static_cast<internal_ptr_t>(ptr) - offset;
|
|
return ret_type(placeholder_accessor_t(buffer, cl::sycl::range<1>(size),
|
|
cl::sycl::id<1>(typed_offset)),
|
|
static_cast<size_t>(typed_offset),
|
|
reinterpret_cast<std::intptr_t>(start_ptr));
|
|
}
|
|
|
|
/// Get a range accessor to the virtual pointer's device memory with a
|
|
/// specified size.
|
|
template <cl::sycl::access::mode AcMd, typename Index>
|
|
EIGEN_STRONG_INLINE cl::sycl::accessor<
|
|
buffer_scalar_t, 1, AcMd, cl::sycl::access::target::global_buffer>
|
|
get_range_accessor(cl::sycl::handler &cgh, const void *ptr,
|
|
const Index n_bytes) const {
|
|
static const auto global_access = cl::sycl::access::target::global_buffer;
|
|
eigen_assert(n_bytes >= 0);
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
auto buffer = pMapper.get_buffer(ptr);
|
|
const ptrdiff_t offset = pMapper.get_offset(ptr);
|
|
eigen_assert(offset >= 0);
|
|
eigen_assert(offset + n_bytes <= buffer.get_size());
|
|
return buffer.template get_access<AcMd, global_access>(
|
|
cgh, cl::sycl::range<1>(n_bytes), cl::sycl::id<1>(offset));
|
|
}
|
|
|
|
/// Creation of sycl accessor for a buffer. This function first tries to find
|
|
/// the buffer in the buffer_map. If found it gets the accessor from it, if
|
|
/// not, the function then adds an entry by creating a sycl buffer for that
|
|
/// particular pointer.
|
|
template <cl::sycl::access::mode AcMd>
|
|
EIGEN_STRONG_INLINE cl::sycl::accessor<
|
|
buffer_scalar_t, 1, AcMd, cl::sycl::access::target::global_buffer>
|
|
get_sycl_accessor(cl::sycl::handler &cgh, const void *ptr) const {
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
return pMapper.get_buffer(ptr)
|
|
.template get_access<AcMd, cl::sycl::access::target::global_buffer>(
|
|
cgh);
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE cl::sycl::buffer<buffer_scalar_t, 1> get_sycl_buffer(
|
|
const void *ptr) const {
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
return pMapper.get_buffer(ptr);
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE ptrdiff_t get_offset(const void *ptr) const {
|
|
std::lock_guard<std::mutex> lock(pmapper_mutex_);
|
|
return pMapper.get_offset(ptr);
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE void synchronize() const {
|
|
#ifdef EIGEN_EXCEPTIONS
|
|
m_queue.wait_and_throw();
|
|
#else
|
|
m_queue.wait();
|
|
#endif
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE void async_synchronize(cl::sycl::event e) const {
|
|
set_latest_event(e);
|
|
#ifndef EIGEN_SYCL_ASYNC_EXECUTION
|
|
synchronize();
|
|
#endif
|
|
}
|
|
|
|
template <typename Index>
|
|
EIGEN_STRONG_INLINE void parallel_for_setup(Index n, Index &tileSize,
|
|
Index &rng, Index &GRange) const {
|
|
tileSize = static_cast<Index>(getNearestPowerOfTwoWorkGroupSize());
|
|
tileSize = std::min(static_cast<Index>(EIGEN_SYCL_LOCAL_THREAD_DIM0 *
|
|
EIGEN_SYCL_LOCAL_THREAD_DIM1),
|
|
static_cast<Index>(tileSize));
|
|
rng = n;
|
|
if (rng == 0) rng = static_cast<Index>(1);
|
|
GRange = rng;
|
|
if (tileSize > GRange)
|
|
tileSize = GRange;
|
|
else if (GRange > tileSize) {
|
|
Index xMode = static_cast<Index>(GRange % tileSize);
|
|
if (xMode != 0) GRange += static_cast<Index>(tileSize - xMode);
|
|
}
|
|
}
|
|
|
|
/// This is used to prepare the number of threads and also the number of
|
|
/// threads per block for sycl kernels
|
|
template <typename Index>
|
|
EIGEN_STRONG_INLINE void parallel_for_setup(Index dim0, Index dim1,
|
|
Index &tileSize0,
|
|
Index &tileSize1, Index &rng0,
|
|
Index &rng1, Index &GRange0,
|
|
Index &GRange1) const {
|
|
Index max_workgroup_Size =
|
|
static_cast<Index>(getNearestPowerOfTwoWorkGroupSize());
|
|
max_workgroup_Size =
|
|
std::min(static_cast<Index>(EIGEN_SYCL_LOCAL_THREAD_DIM0 *
|
|
EIGEN_SYCL_LOCAL_THREAD_DIM1),
|
|
static_cast<Index>(max_workgroup_Size));
|
|
Index pow_of_2 = static_cast<Index>(std::log2(max_workgroup_Size));
|
|
tileSize1 =
|
|
static_cast<Index>(std::pow(2, static_cast<Index>(pow_of_2 / 2)));
|
|
rng1 = dim1;
|
|
if (rng1 == 0) rng1 = static_cast<Index>(1);
|
|
GRange1 = rng1;
|
|
if (tileSize1 > GRange1)
|
|
tileSize1 = GRange1;
|
|
else if (GRange1 > tileSize1) {
|
|
Index xMode = static_cast<Index>(GRange1 % tileSize1);
|
|
if (xMode != 0) GRange1 += static_cast<Index>(tileSize1 - xMode);
|
|
}
|
|
tileSize0 = static_cast<Index>(max_workgroup_Size / tileSize1);
|
|
rng0 = dim0;
|
|
if (rng0 == 0) rng0 = static_cast<Index>(1);
|
|
GRange0 = rng0;
|
|
if (tileSize0 > GRange0)
|
|
tileSize0 = GRange0;
|
|
else if (GRange0 > tileSize0) {
|
|
Index xMode = static_cast<Index>(GRange0 % tileSize0);
|
|
if (xMode != 0) GRange0 += static_cast<Index>(tileSize0 - xMode);
|
|
}
|
|
}
|
|
|
|
/// This is used to prepare the number of threads and also the number of
|
|
/// threads per block for sycl kernels
|
|
template <typename Index>
|
|
EIGEN_STRONG_INLINE void parallel_for_setup(
|
|
Index dim0, Index dim1, Index dim2, Index &tileSize0, Index &tileSize1,
|
|
Index &tileSize2, Index &rng0, Index &rng1, Index &rng2, Index &GRange0,
|
|
Index &GRange1, Index &GRange2) const {
|
|
Index max_workgroup_Size =
|
|
static_cast<Index>(getNearestPowerOfTwoWorkGroupSize());
|
|
max_workgroup_Size =
|
|
std::min(static_cast<Index>(EIGEN_SYCL_LOCAL_THREAD_DIM0 *
|
|
EIGEN_SYCL_LOCAL_THREAD_DIM1),
|
|
static_cast<Index>(max_workgroup_Size));
|
|
Index pow_of_2 = static_cast<Index>(std::log2(max_workgroup_Size));
|
|
tileSize2 =
|
|
static_cast<Index>(std::pow(2, static_cast<Index>(pow_of_2 / 3)));
|
|
rng2 = dim2;
|
|
if (rng2 == 0) rng1 = static_cast<Index>(1);
|
|
GRange2 = rng2;
|
|
if (tileSize2 > GRange2)
|
|
tileSize2 = GRange2;
|
|
else if (GRange2 > tileSize2) {
|
|
Index xMode = static_cast<Index>(GRange2 % tileSize2);
|
|
if (xMode != 0) GRange2 += static_cast<Index>(tileSize2 - xMode);
|
|
}
|
|
pow_of_2 = static_cast<Index>(
|
|
std::log2(static_cast<Index>(max_workgroup_Size / tileSize2)));
|
|
tileSize1 =
|
|
static_cast<Index>(std::pow(2, static_cast<Index>(pow_of_2 / 2)));
|
|
rng1 = dim1;
|
|
if (rng1 == 0) rng1 = static_cast<Index>(1);
|
|
GRange1 = rng1;
|
|
if (tileSize1 > GRange1)
|
|
tileSize1 = GRange1;
|
|
else if (GRange1 > tileSize1) {
|
|
Index xMode = static_cast<Index>(GRange1 % tileSize1);
|
|
if (xMode != 0) GRange1 += static_cast<Index>(tileSize1 - xMode);
|
|
}
|
|
tileSize0 =
|
|
static_cast<Index>(max_workgroup_Size / (tileSize1 * tileSize2));
|
|
rng0 = dim0;
|
|
if (rng0 == 0) rng0 = static_cast<Index>(1);
|
|
GRange0 = rng0;
|
|
if (tileSize0 > GRange0)
|
|
tileSize0 = GRange0;
|
|
else if (GRange0 > tileSize0) {
|
|
Index xMode = static_cast<Index>(GRange0 % tileSize0);
|
|
if (xMode != 0) GRange0 += static_cast<Index>(tileSize0 - xMode);
|
|
}
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE bool has_local_memory() const {
|
|
#if !defined(EIGEN_SYCL_LOCA_MEM) && defined(EIGEN_SYCL_NO_LOCAL_MEM)
|
|
return false;
|
|
#elif defined(EIGEN_SYCL_LOCAL_MEM) && !defined(EIGEN_SYCL_NO_LOCAL_MEM)
|
|
return true;
|
|
#else
|
|
return m_device_info.local_mem_type ==
|
|
cl::sycl::info::local_mem_type::local;
|
|
#endif
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE unsigned long max_buffer_size() const {
|
|
return m_device_info.max_mem_alloc_size;
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE unsigned long getNumSyclMultiProcessors() const {
|
|
return m_device_info.max_compute_units;
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE unsigned long maxSyclThreadsPerBlock() const {
|
|
return m_device_info.max_work_group_size;
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE cl::sycl::id<3> maxWorkItemSizes() const {
|
|
return m_device_info.max_work_item_sizes;
|
|
}
|
|
|
|
/// No need for sycl it should act the same as CPU version
|
|
EIGEN_STRONG_INLINE int majorDeviceVersion() const { return 1; }
|
|
|
|
EIGEN_STRONG_INLINE unsigned long maxSyclThreadsPerMultiProcessor() const {
|
|
// OpenCL doesnot have such concept
|
|
return 2;
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE size_t sharedMemPerBlock() const {
|
|
return m_device_info.local_mem_size;
|
|
}
|
|
|
|
// This function returns the nearest power of 2 Work-group size which is <=
|
|
// maximum device workgroup size.
|
|
EIGEN_STRONG_INLINE size_t getNearestPowerOfTwoWorkGroupSize() const {
|
|
return getPowerOfTwo(m_device_info.max_work_group_size, false);
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE std::string getPlatformName() const {
|
|
return m_device_info.platform_name;
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE std::string getDeviceName() const {
|
|
return m_device_info.device_name;
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE std::string getDeviceVendor() const {
|
|
return m_device_info.device_vendor;
|
|
}
|
|
|
|
// This function returns the nearest power of 2
|
|
// if roundup is true returns result>=wgsize
|
|
// else it return result <= wgsize
|
|
EIGEN_STRONG_INLINE size_t getPowerOfTwo(size_t wGSize, bool roundUp) const {
|
|
if (roundUp) --wGSize;
|
|
wGSize |= (wGSize >> 1);
|
|
wGSize |= (wGSize >> 2);
|
|
wGSize |= (wGSize >> 4);
|
|
wGSize |= (wGSize >> 8);
|
|
wGSize |= (wGSize >> 16);
|
|
#if EIGEN_ARCH_x86_64 || EIGEN_ARCH_ARM64 || EIGEN_OS_WIN64
|
|
wGSize |= (wGSize >> 32);
|
|
#endif
|
|
return ((!roundUp) ? (wGSize - (wGSize >> 1)) : ++wGSize);
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE cl::sycl::queue &sycl_queue() const { return m_queue; }
|
|
|
|
// This function checks if the runtime recorded an error for the
|
|
// underlying stream device.
|
|
EIGEN_STRONG_INLINE bool ok() const {
|
|
if (!exception_caught_) {
|
|
synchronize();
|
|
}
|
|
return !exception_caught_;
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE cl::sycl::event get_latest_event() const {
|
|
#ifdef EIGEN_SYCL_STORE_LATEST_EVENT
|
|
std::lock_guard<std::mutex> lock(event_mutex_);
|
|
return latest_events_[std::this_thread::get_id()];
|
|
#else
|
|
eigen_assert(false);
|
|
return cl::sycl::event();
|
|
#endif
|
|
}
|
|
|
|
// destructor
|
|
~QueueInterface() {
|
|
pMapper.clear();
|
|
#ifndef EIGEN_SYCL_NO_REUSE_BUFFERS
|
|
scratch_buffers.clear();
|
|
#endif
|
|
}
|
|
|
|
protected:
|
|
EIGEN_STRONG_INLINE void set_latest_event(cl::sycl::event e) const {
|
|
#ifdef EIGEN_SYCL_STORE_LATEST_EVENT
|
|
std::lock_guard<std::mutex> lock(event_mutex_);
|
|
latest_events_[std::this_thread::get_id()] = e;
|
|
#else
|
|
EIGEN_UNUSED_VARIABLE(e);
|
|
#endif
|
|
}
|
|
|
|
void synchronize_and_callback(cl::sycl::event e,
|
|
const std::function<void()> &callback) const {
|
|
set_latest_event(e);
|
|
if (callback) {
|
|
auto callback_ = [=]() {
|
|
#ifdef EIGEN_EXCEPTIONS
|
|
cl::sycl::event(e).wait_and_throw();
|
|
#else
|
|
cl::sycl::event(e).wait();
|
|
#endif
|
|
callback();
|
|
};
|
|
m_thread_pool.Schedule(std::move(callback_));
|
|
} else {
|
|
#ifdef EIGEN_EXCEPTIONS
|
|
m_queue.wait_and_throw();
|
|
#else
|
|
m_queue.wait();
|
|
#endif
|
|
}
|
|
}
|
|
|
|
bool sycl_async_handler(cl::sycl::exception_list exceptions) const {
|
|
bool exception_caught = false;
|
|
for (const auto &e : exceptions) {
|
|
if (e) {
|
|
exception_caught = true;
|
|
EIGEN_THROW_X(e);
|
|
}
|
|
}
|
|
return exception_caught;
|
|
}
|
|
|
|
/// class members:
|
|
bool exception_caught_ = false;
|
|
|
|
mutable std::mutex pmapper_mutex_;
|
|
|
|
#ifdef EIGEN_SYCL_STORE_LATEST_EVENT
|
|
mutable std::mutex event_mutex_;
|
|
mutable std::unordered_map<std::thread::id, cl::sycl::event> latest_events_;
|
|
#endif
|
|
|
|
/// std::map is the container used to make sure that we create only one buffer
|
|
/// per pointer. The lifespan of the buffer now depends on the lifespan of
|
|
/// SyclDevice. If a non-read-only pointer is needed to be accessed on the
|
|
/// host we should manually deallocate it.
|
|
mutable TensorSycl::internal::PointerMapper pMapper;
|
|
#ifndef EIGEN_SYCL_NO_REUSE_BUFFERS
|
|
mutable std::unordered_set<void *> scratch_buffers;
|
|
#endif
|
|
/// sycl queue
|
|
mutable cl::sycl::queue m_queue;
|
|
#ifdef EIGEN_SYCL_USE_PROGRAM_CLASS
|
|
mutable cl::sycl::program m_prog;
|
|
#endif
|
|
|
|
/// The thread pool is used to wait on events and call callbacks
|
|
/// asynchronously
|
|
mutable Eigen::ThreadPool m_thread_pool;
|
|
|
|
const TensorSycl::internal::SyclDeviceInfo m_device_info;
|
|
};
|
|
|
|
struct SyclDeviceBase {
|
|
/// QueueInterface is not owned. it is the caller's responsibility to destroy
|
|
/// it
|
|
const QueueInterface *m_queue_stream;
|
|
explicit SyclDeviceBase(const QueueInterface *queue_stream)
|
|
: m_queue_stream(queue_stream) {}
|
|
EIGEN_STRONG_INLINE const QueueInterface *queue_stream() const {
|
|
return m_queue_stream;
|
|
}
|
|
};
|
|
|
|
// Here is a sycl device struct which accept the sycl queue interface
|
|
// as an input
|
|
struct SyclDevice : public SyclDeviceBase {
|
|
explicit SyclDevice(const QueueInterface *queue_stream)
|
|
: SyclDeviceBase(queue_stream) {}
|
|
|
|
// this is the accessor used to construct the evaluator
|
|
template <cl::sycl::access::mode AcMd, typename T>
|
|
EIGEN_STRONG_INLINE TensorSycl::internal::RangeAccess<AcMd, T>
|
|
get_range_accessor(const void *ptr) const {
|
|
return queue_stream()->template get_range_accessor<AcMd, T>(ptr);
|
|
}
|
|
|
|
// get sycl accessor
|
|
template <cl::sycl::access::mode AcMd>
|
|
EIGEN_STRONG_INLINE cl::sycl::accessor<
|
|
buffer_scalar_t, 1, AcMd, cl::sycl::access::target::global_buffer>
|
|
get_sycl_accessor(cl::sycl::handler &cgh, const void *ptr) const {
|
|
return queue_stream()->template get_sycl_accessor<AcMd>(cgh, ptr);
|
|
}
|
|
|
|
/// Accessing the created sycl device buffer for the device pointer
|
|
EIGEN_STRONG_INLINE cl::sycl::buffer<buffer_scalar_t, 1> get_sycl_buffer(
|
|
const void *ptr) const {
|
|
return queue_stream()->get_sycl_buffer(ptr);
|
|
}
|
|
|
|
/// This is used to prepare the number of threads and also the number of
|
|
/// threads per block for sycl kernels
|
|
template <typename Index>
|
|
EIGEN_STRONG_INLINE void parallel_for_setup(Index n, Index &tileSize,
|
|
Index &rng, Index &GRange) const {
|
|
queue_stream()->parallel_for_setup(n, tileSize, rng, GRange);
|
|
}
|
|
|
|
/// This is used to prepare the number of threads and also the number of
|
|
/// threads per block for sycl kernels
|
|
template <typename Index>
|
|
EIGEN_STRONG_INLINE void parallel_for_setup(Index dim0, Index dim1,
|
|
Index &tileSize0,
|
|
Index &tileSize1, Index &rng0,
|
|
Index &rng1, Index &GRange0,
|
|
Index &GRange1) const {
|
|
queue_stream()->parallel_for_setup(dim0, dim1, tileSize0, tileSize1, rng0,
|
|
rng1, GRange0, GRange1);
|
|
}
|
|
|
|
/// This is used to prepare the number of threads and also the number of
|
|
/// threads per block for sycl kernels
|
|
template <typename Index>
|
|
EIGEN_STRONG_INLINE void parallel_for_setup(
|
|
Index dim0, Index dim1, Index dim2, Index &tileSize0, Index &tileSize1,
|
|
Index &tileSize2, Index &rng0, Index &rng1, Index &rng2, Index &GRange0,
|
|
Index &GRange1, Index &GRange2) const {
|
|
queue_stream()->parallel_for_setup(dim0, dim1, dim2, tileSize0, tileSize1,
|
|
tileSize2, rng0, rng1, rng2, GRange0,
|
|
GRange1, GRange2);
|
|
}
|
|
|
|
/// allocate device memory
|
|
EIGEN_STRONG_INLINE void *allocate(size_t num_bytes) const {
|
|
return queue_stream()->allocate(num_bytes);
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE void *allocate_temp(size_t num_bytes) const {
|
|
return queue_stream()->allocate_temp(num_bytes);
|
|
}
|
|
|
|
/// deallocate device memory
|
|
EIGEN_STRONG_INLINE void deallocate(void *p) const {
|
|
queue_stream()->deallocate(p);
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE void deallocate_temp(void *buffer) const {
|
|
queue_stream()->deallocate_temp(buffer);
|
|
}
|
|
template <cl::sycl::access::mode AcMd, typename T>
|
|
EIGEN_STRONG_INLINE void deallocate_temp(
|
|
const TensorSycl::internal::RangeAccess<AcMd, T> &buffer) const {
|
|
queue_stream()->deallocate_temp(buffer);
|
|
}
|
|
EIGEN_STRONG_INLINE void deallocate_all() const {
|
|
queue_stream()->deallocate_all();
|
|
}
|
|
|
|
template <typename data_t>
|
|
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorSycl::internal::RangeAccess<
|
|
cl::sycl::access::mode::read_write, data_t>
|
|
get(data_t *data) const {
|
|
return queue_stream()->get(data);
|
|
}
|
|
template <typename data_t>
|
|
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE data_t *get(
|
|
TensorSycl::internal::RangeAccess<cl::sycl::access::mode::read_write,
|
|
data_t>
|
|
data) const {
|
|
return queue_stream()->get(data);
|
|
}
|
|
|
|
/// attach existing buffer
|
|
EIGEN_STRONG_INLINE void *attach_buffer(
|
|
cl::sycl::buffer<buffer_scalar_t, 1> &buf) const {
|
|
return queue_stream()->attach_buffer(buf);
|
|
}
|
|
/// detach buffer
|
|
EIGEN_STRONG_INLINE void detach_buffer(void *p) const {
|
|
queue_stream()->detach_buffer(p);
|
|
}
|
|
EIGEN_STRONG_INLINE ptrdiff_t get_offset(const void *ptr) const {
|
|
return queue_stream()->get_offset(ptr);
|
|
}
|
|
|
|
// some runtime conditions that can be applied here
|
|
EIGEN_STRONG_INLINE bool isDeviceSuitable() const { return true; }
|
|
|
|
/// memcpyHostToDevice
|
|
template <typename Index>
|
|
EIGEN_STRONG_INLINE void memcpyHostToDevice(
|
|
Index *dst, const Index *src, size_t n,
|
|
std::function<void()> callback = {}) const {
|
|
queue_stream()->memcpyHostToDevice(dst, src, n, callback);
|
|
}
|
|
/// memcpyDeviceToHost
|
|
template <typename Index>
|
|
EIGEN_STRONG_INLINE void memcpyDeviceToHost(
|
|
void *dst, const Index *src, size_t n,
|
|
std::function<void()> callback = {}) const {
|
|
queue_stream()->memcpyDeviceToHost(dst, src, n, callback);
|
|
}
|
|
/// the memcpy function
|
|
template <typename Index>
|
|
EIGEN_STRONG_INLINE void memcpy(void *dst, const Index *src, size_t n) const {
|
|
queue_stream()->memcpy(dst, src, n);
|
|
}
|
|
/// the memset function
|
|
EIGEN_STRONG_INLINE void memset(void *data, int c, size_t n) const {
|
|
queue_stream()->memset(data, c, n);
|
|
}
|
|
/// returning the sycl queue
|
|
EIGEN_STRONG_INLINE cl::sycl::queue &sycl_queue() const {
|
|
return queue_stream()->sycl_queue();
|
|
}
|
|
#ifdef EIGEN_SYCL_USE_PROGRAM_CLASS
|
|
EIGEN_STRONG_INLINE cl::sycl::program &program() const {
|
|
return queue_stream()->program();
|
|
}
|
|
#endif
|
|
|
|
EIGEN_STRONG_INLINE size_t firstLevelCacheSize() const { return 48 * 1024; }
|
|
|
|
EIGEN_STRONG_INLINE size_t lastLevelCacheSize() const {
|
|
// We won't try to take advantage of the l2 cache for the time being, and
|
|
// there is no l3 cache on sycl devices.
|
|
return firstLevelCacheSize();
|
|
}
|
|
EIGEN_STRONG_INLINE unsigned long getNumSyclMultiProcessors() const {
|
|
return queue_stream()->getNumSyclMultiProcessors();
|
|
}
|
|
EIGEN_STRONG_INLINE unsigned long maxSyclThreadsPerBlock() const {
|
|
return queue_stream()->maxSyclThreadsPerBlock();
|
|
}
|
|
EIGEN_STRONG_INLINE cl::sycl::id<3> maxWorkItemSizes() const {
|
|
return queue_stream()->maxWorkItemSizes();
|
|
}
|
|
EIGEN_STRONG_INLINE unsigned long maxSyclThreadsPerMultiProcessor() const {
|
|
// OpenCL doesnot have such concept
|
|
return queue_stream()->maxSyclThreadsPerMultiProcessor();
|
|
}
|
|
EIGEN_STRONG_INLINE size_t sharedMemPerBlock() const {
|
|
return queue_stream()->sharedMemPerBlock();
|
|
}
|
|
EIGEN_STRONG_INLINE size_t getNearestPowerOfTwoWorkGroupSize() const {
|
|
return queue_stream()->getNearestPowerOfTwoWorkGroupSize();
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE size_t getPowerOfTwo(size_t val, bool roundUp) const {
|
|
return queue_stream()->getPowerOfTwo(val, roundUp);
|
|
}
|
|
/// No need for sycl it should act the same as CPU version
|
|
EIGEN_STRONG_INLINE int majorDeviceVersion() const {
|
|
return queue_stream()->majorDeviceVersion();
|
|
}
|
|
|
|
EIGEN_STRONG_INLINE void synchronize() const {
|
|
queue_stream()->synchronize();
|
|
}
|
|
EIGEN_STRONG_INLINE void async_synchronize(
|
|
cl::sycl::event e = cl::sycl::event()) const {
|
|
queue_stream()->async_synchronize(e);
|
|
}
|
|
EIGEN_STRONG_INLINE cl::sycl::event get_latest_event() const {
|
|
return queue_stream()->get_latest_event();
|
|
}
|
|
|
|
// This function checks if the runtime recorded an error for the
|
|
// underlying stream device.
|
|
EIGEN_STRONG_INLINE bool ok() const { return queue_stream()->ok(); }
|
|
|
|
EIGEN_STRONG_INLINE bool has_local_memory() const {
|
|
return queue_stream()->has_local_memory();
|
|
}
|
|
EIGEN_STRONG_INLINE long max_buffer_size() const {
|
|
return queue_stream()->max_buffer_size();
|
|
}
|
|
EIGEN_STRONG_INLINE std::string getPlatformName() const {
|
|
return queue_stream()->getPlatformName();
|
|
}
|
|
EIGEN_STRONG_INLINE std::string getDeviceName() const {
|
|
return queue_stream()->getDeviceName();
|
|
}
|
|
EIGEN_STRONG_INLINE std::string getDeviceVendor() const {
|
|
return queue_stream()->getDeviceVendor();
|
|
}
|
|
};
|
|
} // end namespace Eigen
|
|
|
|
#endif // EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H
|