Skip to content

Commit

Permalink
[SYCL][ESIMD] Add more stringent compile time checks for accessor ver…
Browse files Browse the repository at this point in the history 
…sions of block_load/block_store, gather/scatter API (#11145)
  • Loading branch information
@fineg74
fineg74 authored Oct 23, 2023
1 parent 87ff465 commit 803a77f
Show file tree
Hide file tree
Showing 6 changed files with 342 additions and 164 deletions.
61 changes: 50 additions & 11 deletions sycl/include/sycl/ext/intel/esimd/detail/simd_obj_impl.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -330,13 +330,11 @@ class [[__sycl_detail__::__uses_aspects__(
/// argument.
/// @param acc The accessor to read from.
/// @param offset offset in bytes of the first element.
template <typename AccessorT, typename Flags = element_aligned_tag,
typename = std::enable_if_t<
(sycl::detail::acc_properties::is_local_accessor_v<AccessorT> ||
detail::is_sycl_accessor_with<
AccessorT, accessor_mode_cap::can_read,
sycl::access::target::device>::value) &&
is_simd_flag_type_v<Flags>>>
template <
typename AccessorT, typename Flags = element_aligned_tag,
typename = std::enable_if_t<
detail::is_accessor_with_v<AccessorT, accessor_mode_cap::can_read> &&
is_simd_flag_type_v<Flags>>>
simd_obj_impl(AccessorT acc,
#ifdef __ESIMD_FORCE_STATELESS_MEM
uint64_t offset,
Expand Down Expand Up @@ -744,8 +742,7 @@ class [[__sycl_detail__::__uses_aspects__(
template <typename AccessorT, typename Flags = element_aligned_tag,
int ChunkSize = 32,
typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read,
sycl::access::target::device, void>
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read, void>
copy_from(AccessorT acc,
#ifdef __ESIMD_FORCE_STATELESS_MEM
uint64_t offset,
Expand All @@ -754,6 +751,25 @@ class [[__sycl_detail__::__uses_aspects__(
#endif
Flags = {}) SYCL_ESIMD_FUNCTION;

/// Copy a contiguous block of data from memory into this simd_obj_impl
/// object. The amount of memory copied equals the total size of vector
/// elements in this object. Source memory location is represented via a
/// local accessor and offset.
/// None of the template parameters except documented ones can/should be
/// specified by callers.
/// @tparam AccessorT Type of the accessor (auto-deduced).
/// @tparam Flags Alignment control for the copy operation.
/// See @ref sycl_esimd_core_align for more info.
/// @param acc accessor to copy from.
/// @param offset offset to copy from (in bytes).
template <typename AccessorT, typename Flags = element_aligned_tag,
int ChunkSize = 32,
typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE std::enable_if_t<
detail::is_local_accessor_with_v<AccessorT, accessor_mode_cap::can_read>,
void>
copy_from(AccessorT acc, uint32_t offset, Flags = {}) SYCL_ESIMD_FUNCTION;

/// Copy all vector elements of this object into a contiguous block in memory.
/// None of the template parameters should be be specified by callers.
/// @tparam Flags Alignment control for the copy operation.
Expand All @@ -776,8 +792,7 @@ class [[__sycl_detail__::__uses_aspects__(
template <typename AccessorT, typename Flags = element_aligned_tag,
int ChunkSize = 32,
typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write,
sycl::access::target::device, void>
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write, void>
copy_to(AccessorT acc,
#ifdef __ESIMD_FORCE_STATELESS_MEM
uint64_t offset,
Expand All @@ -786,6 +801,23 @@ class [[__sycl_detail__::__uses_aspects__(
#endif
Flags = {}) const SYCL_ESIMD_FUNCTION;

/// Copy all vector elements of this object into a contiguous block in memory.
/// Destination memory location is represented via a local accessor and
/// offset.
/// None of the template parameters should be be specified by callers.
/// @tparam AccessorT Type of the accessor (auto-deduced).
/// @tparam Flags Alignment control for the copy operation.
/// See @ref sycl_esimd_core_align for more info.
/// @param acc accessor to copy from.
/// @param offset offset to copy from.
template <typename AccessorT, typename Flags = element_aligned_tag,
int ChunkSize = 32,
typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
ESIMD_INLINE std::enable_if_t<
detail::is_local_accessor_with_v<AccessorT, accessor_mode_cap::can_write>,
void>
copy_to(AccessorT acc, uint32_t offset, Flags = {}) const SYCL_ESIMD_FUNCTION;

// Unary operations.

/// Per-element bitwise inversion, available in all subclasses, but only for
Expand Down Expand Up @@ -916,6 +948,13 @@ class [[__sycl_detail__::__uses_aspects__(
// The underlying data for this vector.
raw_vector_type M_data;

template <int ChunkSize, typename Flags, typename AccessorT, typename TOffset>
ESIMD_INLINE void copy_to_impl(AccessorT acc,
TOffset offset) const SYCL_ESIMD_FUNCTION;
template <int ChunkSize, typename Flags, typename AccessorT, typename TOffset>
ESIMD_INLINE void copy_from_impl(AccessorT acc,
TOffset offset) SYCL_ESIMD_FUNCTION;

protected:
// The test proxy if enabled
__ESIMD_DECLARE_TEST_PROXY
Expand Down
48 changes: 38 additions & 10 deletions sycl/include/sycl/ext/intel/esimd/detail/sycl_util.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -63,23 +63,51 @@ constexpr bool accessor_mode_has_capability() {
(Mode == sycl::access::mode::read);
}

template <typename T> struct local_accessor_access_mode {
static constexpr sycl::access::mode mode =
static_cast<sycl::access::mode>(-1);
};

template <typename DataT, int Dimensions>
struct local_accessor_access_mode<local_accessor<DataT, Dimensions>> {
static constexpr sycl::access::mode mode =
sycl::detail::accessModeFromConstness<DataT>();
};

// Checks that given type is a SYCL accessor type with given capability and
// target.
template <typename T, accessor_mode_cap_val_t Capability,
sycl::access::target AccessTarget>
struct is_sycl_accessor_with
template <typename T, accessor_mode_cap_val_t Capability>
struct is_device_accessor_with
: public std::conditional_t<
accessor_mode_has_capability<is_sycl_accessor<T>::mode,
Capability>() &&
(is_sycl_accessor<T>::target == AccessTarget),
(is_sycl_accessor<T>::target == sycl::access::target::device),
std::true_type, std::false_type> {};

template <typename T, accessor_mode_cap_val_t Capability,
sycl::access::target AccessTarget, typename RetT>
using EnableIfAccessor = std::enable_if_t<
detail::is_sycl_accessor_with<T, Capability, AccessTarget>::value ||
sycl::detail::acc_properties::is_local_accessor_v<T>,
RetT>;
template <typename T, accessor_mode_cap_val_t Capability>
struct is_local_accessor_with
: public std::conditional_t<
sycl::detail::acc_properties::is_local_accessor_v<T> &&
accessor_mode_has_capability<local_accessor_access_mode<T>::mode,
Capability>(),
std::true_type, std::false_type> {};

template <typename T, accessor_mode_cap_val_t Capability>
inline constexpr bool is_local_accessor_with_v =
is_local_accessor_with<T, Capability>::value;

template <typename T, accessor_mode_cap_val_t Capability>
inline constexpr bool is_device_accessor_with_v =
is_device_accessor_with<T, Capability>::value;

template <typename T, accessor_mode_cap_val_t Capability>
inline constexpr bool is_accessor_with_v =
is_device_accessor_with_v<T, Capability> ||
is_local_accessor_with_v<T, Capability>;

template <typename T, accessor_mode_cap_val_t Capability, typename RetT>
using EnableIfAccessor =
std::enable_if_t<detail::is_device_accessor_with_v<T, Capability>, RetT>;

template <typename T, int Dimensions>
__ESIMD_API uint32_t localAccessorToOffset(local_accessor<T, Dimensions> acc) {
Expand Down
Loading

0 comments on commit 803a77f

Please sign in to comment.