15 inline namespace _V1 {
16 namespace ext::intel {
23 static constexpr int32_t
value = _N;
27 template <
int32_t _N>
struct cache {
28 static constexpr int32_t
value = _N;
33 static constexpr int32_t
value = _N;
38 static constexpr int32_t
value = _N;
46 template <
class... _mem_access_params>
class lsu final {
53 check_space<_space>();
55 #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
56 return *__builtin_intel_fpga_mem((_T *)Ptr,
57 _burst_coalesce | _cache |
58 _dont_statically_coalesce | _prefetch,
68 check_space<_space>();
70 #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
71 *__builtin_intel_fpga_mem((_T *)Ptr,
72 _burst_coalesce | _cache |
73 _dont_statically_coalesce | _prefetch,
81 static constexpr int32_t _burst_coalesce_val =
83 static constexpr uint8_t _burst_coalesce =
86 static constexpr int32_t _cache_val =
88 static constexpr uint8_t _cache = (_cache_val > 0) ?
CACHE : 0;
90 static constexpr int32_t _statically_coalesce_val =
92 static constexpr uint8_t _dont_statically_coalesce =
95 static constexpr int32_t _prefetch_val =
97 static constexpr uint8_t _prefetch = _prefetch_val ?
PREFETCH : 0;
99 static_assert(_cache_val >= 0,
"cache size parameter must be non-negative");
101 template <access::address_space _space>
static void check_space() {
106 "lsu controls are only supported for global_ptr, "
107 "device_ptr, and host_ptr objects");
110 static void check_load() {
112 "unable to implement a cache without a burst coalescer");
113 static_assert(_prefetch == 0 || _burst_coalesce == 0,
114 "unable to implement a prefetcher and a burst coalescer "
117 _prefetch == 0 || _cache == 0,
118 "unable to implement a prefetcher and a cache simulataneously");
120 static void check_store() {
121 static_assert(_cache == 0,
"unable to implement a store LSU with a cache.");
122 static_assert(_prefetch == 0,
123 "unable to implement a store LSU with a prefetcher.");
static void store(sycl::multi_ptr< _T, _space, _is_decorated > Ptr, _T Val)
static _T load(sycl::multi_ptr< _T, _space, _is_decorated > Ptr)
@ ext_intel_global_host_space
@ ext_intel_global_device_space
constexpr uint8_t PREFETCH
constexpr uint8_t BURST_COALESCE
constexpr uint8_t STATICALLY_COALESCE
void prefetch_impl(T *ptr, size_t bytes, Properties properties)
static constexpr int32_t value
static constexpr int32_t default_value
static constexpr int32_t default_value
static constexpr int32_t value
static constexpr int32_t value
static constexpr int32_t default_value
static constexpr int32_t value
static constexpr int32_t default_value