sycl_ext_oneapi_prefetch

Notice

Khronos® is a registered trademark and SYCL™ and SPIR™ are trademarks of The Khronos Group Inc. OpenCL™ is a trademark of Apple Inc. used by permission by Khronos.

Contact

To report problems with this extension, please open a new issue at:

https://github.com/intel/llvm/issues

Dependencies

This extension is written against the SYCL 2020 revision 7 specification. All references below to the "core SYCL specification" or to section numbers in the SYCL specification refer to that revision.

This extension also depends on the following other SYCL extensions:

sycl_ext_oneapi_properties

Status

This is an experimental extension specification, intended to provide early access to features and gather community feedback. Interfaces defined in this specification are implemented in DPC, but they are not finalized and may change incompatibly in future versions of DPC without prior notice. Shipping software products should not rely on APIs defined in this specification.

Overview

Many devices targeted by SYCL support software prefetch operations, which act as asynchronous memory reads intended to populate cache(s). Prefetches are generally used to hide memory latency, and their deployment may be a critical component for software tuning on some hardware.

The multi_ptr::prefetch function provided by SYCL 2020 is insufficient to cover many real use-cases — it does not acknowledge the existence of a cache hierarchy, and does not provide a mechanism for cooperative (i.e. group) prefetches.

This proposal addresses these shortcomings by introducing a set of free functions that groups of work-items can use to prefetch data into specific levels of cache (controlled via compile-time properties).

Specification

Feature test macro

This extension provides a feature-test macro as described in the core SYCL specification. An implementation supporting this extension must predefine the macro SYCL_EXT_ONEAPI_PREFETCH to one of the values defined in the table below. Applications can test for the existence of this macro to determine if the implementation supports this feature, or applications can test the macro’s value to determine which of the extension’s features the implementation supports.

Value	Description
1	The APIs of this experimental extension are not versioned, so the feature-test macro always has this value.

Specifying cache levels

The level of cache targeted by a prefetch is specified using compile-time properties. This extension defines hints for four levels of cache, corresponding to the four levels of cache currently defined in SYCL (e.g. as used by info::partition_affinity_domain).

Note	Not all devices targeted by SYCL will have four levels of cache. Some of these hints may have no meaning on some devices. However, implementations are encouraged to map these constants to the closest level of cache available.

When multiple cache levels are specified, the lowest level takes precedence (e.g. a request to prefetch into L1 and L4 is treated as a request to prefetch into L1). When no cache levels are specified, the default behavior is to prefetch into the lowest level cache (i.e. L1).

Note	Future hints may alter the default behavior of prefetches with respect to cache levels. Any such alterations are expected to be documented in the definition of those new hints.

namespace sycl::ext::oneapi::experimental {

enum class cache_level {
  L1,
  L2,
  L3,
  L4,
};

struct nontemporal;

struct prefetch_hint_key {
  template <cache_level Level, typename Hint>
  using value_t = property_value<prefetch_hint_key,
                                 std::integral_constant<cache_level, Level>,
                                 Hint>;
};

template <cache_level Level, typename Hint>
inline constexpr prefetch_hint_key::value_t<Level, Hint> prefetch_hint;

inline constexpr prefetch_hint_key::value_t<cache_level::L1, void> prefetch_hint_L1;
inline constexpr prefetch_hint_key::value_t<cache_level::L2, void> prefetch_hint_L2;
inline constexpr prefetch_hint_key::value_t<cache_level::L3, void> prefetch_hint_L3;
inline constexpr prefetch_hint_key::value_t<cache_level::L4, void> prefetch_hint_L4;

inline constexpr prefetch_hint_key::value_t<cache_level::L1, nontemporal> prefetch_hint_L1_nt;
inline constexpr prefetch_hint_key::value_t<cache_level::L2, nontemporal> prefetch_hint_L2_nt;
inline constexpr prefetch_hint_key::value_t<cache_level::L3, nontemporal> prefetch_hint_L3_nt;
inline constexpr prefetch_hint_key::value_t<cache_level::L4, nontemporal> prefetch_hint_L4_nt;

} // namespace sycl::ext::oneapi::experimental

If the nontemporal type is used as the value of the Hint template parameter, this acts as a hint that the data being prefetched will not be reused.

Note	The implementation of `nontemporal` is device-specific, owing to differences in cache hierarchies and replacement policies. On some devices non-temporal data may be evicted earlier than temporal data, while on other devices non-temporal data may not be stored in certain levels of cache at all.

Work-item prefetches

The functions in this section allow individual work-items to prefetch data.

namespace sycl::ext::oneapi::experimental {

template <typename Properties = empty_properties_t>
void prefetch(void* ptr, Properties properties = {});

template <typename Properties = empty_properties_t>
void prefetch(void* ptr, size_t bytes, Properties properties = {});

template <typename T, typename Properties = empty_properties_t>
void prefetch(T* ptr, Properties properties = {});

template <typename T, typename Properties = empty_properties_t>
void prefetch(T* ptr, size_t count, Properties properties = {});

// Only available if AddressSpace == global_space || AddressSpace == generic_space
template <access::address_space AddressSpace, access::decorated IsDecorated,
          typename Properties = empty_properties_t>
void prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, Properties properties = {});

// Only available if AddressSpace == global_space || AddressSpace == generic_space
template <access::address_space AddressSpace, access::decorated IsDecorated,
          typename Properties = empty_properties_t>
void prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes, Properties properties = {});

// Only available if AddressSpace == global_space || AddressSpace == generic_space
template <typename T, access::address_space AddressSpace, access::decorated IsDecorated,
          typename Properties = empty_properties_t>
void prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, Properties properties = {});

// Only available if AddressSpace == global_space || AddressSpace == generic_space
template <typename T, access::address_space AddressSpace, access::decorated IsDecorated,
          typename Properties = empty_properties_t>
void prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count, Properties properties = {});

// Only available if Dimensions > 0 && (AccessMode == read || AccessMode == read_write)
template <typename DataT, int Dimensions,
          access_mode AccessMode, access::placeholder IsPlaceholder,
          typename Properties = empty_properties_t>
void prefetch(accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
              id<Dimensions> offset, Properties properties = {});

// Only available if Dimensions > 0 && (AccessMode == read || AccessMode == read_write)
template <typename DataT, int Dimensions,
          access_mode AccessMode, access::placeholder IsPlaceholder,
          typename Properties = empty_properties_t>
void prefetch(accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
              id<Dimensions> offset, size_t count, Properties properties = {});

} // namespace sycl::ext::oneapi::experimental

template <typename Properties = empty_properties_t>
void prefetch(void* ptr, Properties properties = {});

Constraints: Available only if is_property_list_v<std::decay_t<Properties>> is true.