DPC++ Runtime
Runtime libraries for oneAPI DPC++
fpga_lsu.hpp
Go to the documentation of this file.
1 //==-------------- fpga_lsu.hpp --- SYCL FPGA LSU Extensions ---------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 #pragma once
9 
10 #include "fpga_utils.hpp"
11 #include <sycl/detail/defines.hpp>
12 #include <sycl/pointers.hpp>
13 
14 namespace sycl {
15 inline namespace _V1 {
16 namespace ext::intel {
17 constexpr uint8_t BURST_COALESCE = 0x1;
18 constexpr uint8_t CACHE = 0x2;
19 constexpr uint8_t STATICALLY_COALESCE = 0x4;
20 constexpr uint8_t PREFETCH = 0x8;
21 
22 template <int32_t _N> struct burst_coalesce_impl {
23  static constexpr int32_t value = _N;
24  static constexpr int32_t default_value = 0;
25 };
26 
27 template <int32_t _N> struct cache {
28  static constexpr int32_t value = _N;
29  static constexpr int32_t default_value = 0;
30 };
31 
32 template <int32_t _N> struct prefetch_impl {
33  static constexpr int32_t value = _N;
34  static constexpr int32_t default_value = 0;
35 };
36 
37 template <int32_t _N> struct statically_coalesce_impl {
38  static constexpr int32_t value = _N;
39  static constexpr int32_t default_value = 1;
40 };
41 
42 template <bool _B> using burst_coalesce = burst_coalesce_impl<_B>;
43 template <bool _B> using prefetch = prefetch_impl<_B>;
45 
46 template <class... _mem_access_params> class lsu final {
47 public:
48  lsu() = delete;
49 
50  template <typename _T, access::address_space _space,
51  access::decorated _is_decorated>
53  check_space<_space>();
54  check_load();
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,
59  _cache_val);
60 #else
61  return *Ptr;
62 #endif
63  }
64 
65  template <typename _T, access::address_space _space,
66  access::decorated _is_decorated>
68  check_space<_space>();
69  check_store();
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,
74  _cache_val) = Val;
75 #else
76  *Ptr = Val;
77 #endif
78  }
79 
80 private:
81  static constexpr int32_t _burst_coalesce_val =
82  _GetValue<burst_coalesce_impl, _mem_access_params...>::value;
83  static constexpr uint8_t _burst_coalesce =
84  _burst_coalesce_val == 1 ? BURST_COALESCE : 0;
85 
86  static constexpr int32_t _cache_val =
87  _GetValue<cache, _mem_access_params...>::value;
88  static constexpr uint8_t _cache = (_cache_val > 0) ? CACHE : 0;
89 
90  static constexpr int32_t _statically_coalesce_val =
91  _GetValue<statically_coalesce_impl, _mem_access_params...>::value;
92  static constexpr uint8_t _dont_statically_coalesce =
93  _statically_coalesce_val == 0 ? STATICALLY_COALESCE : 0;
94 
95  static constexpr int32_t _prefetch_val =
96  _GetValue<prefetch_impl, _mem_access_params...>::value;
97  static constexpr uint8_t _prefetch = _prefetch_val ? PREFETCH : 0;
98 
99  static_assert(_cache_val >= 0, "cache size parameter must be non-negative");
100 
101  template <access::address_space _space> static void check_space() {
102  static_assert(
106  "lsu controls are only supported for global_ptr, "
107  "device_ptr, and host_ptr objects");
108  }
109 
110  static void check_load() {
111  static_assert(_cache == 0 || _burst_coalesce == BURST_COALESCE,
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 "
115  "simulataneously");
116  static_assert(
117  _prefetch == 0 || _cache == 0,
118  "unable to implement a prefetcher and a cache simulataneously");
119  }
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.");
124  }
125 };
126 } // namespace ext::intel
127 
128 } // namespace _V1
129 } // namespace sycl
static void store(sycl::multi_ptr< _T, _space, _is_decorated > Ptr, _T Val)
Definition: fpga_lsu.hpp:67
static _T load(sycl::multi_ptr< _T, _space, _is_decorated > Ptr)
Definition: fpga_lsu.hpp:52
constexpr uint8_t CACHE
Definition: fpga_lsu.hpp:18
constexpr uint8_t PREFETCH
Definition: fpga_lsu.hpp:20
constexpr uint8_t BURST_COALESCE
Definition: fpga_lsu.hpp:17
constexpr uint8_t STATICALLY_COALESCE
Definition: fpga_lsu.hpp:19
void prefetch_impl(T *ptr, size_t bytes, Properties properties)
Definition: prefetch.hpp:71
Definition: access.hpp:18
static constexpr int32_t default_value
Definition: fpga_lsu.hpp:24
static constexpr int32_t default_value
Definition: fpga_lsu.hpp:29
static constexpr int32_t value
Definition: fpga_lsu.hpp:28
static constexpr int32_t value
Definition: fpga_lsu.hpp:33
static constexpr int32_t default_value
Definition: fpga_lsu.hpp:34