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 
15 namespace sycl {
16 namespace ext {
17 namespace intel {
18 constexpr uint8_t BURST_COALESCE = 0x1;
19 constexpr uint8_t CACHE = 0x2;
20 constexpr uint8_t STATICALLY_COALESCE = 0x4;
21 constexpr uint8_t PREFETCH = 0x8;
22 
23 template <int32_t _N> struct burst_coalesce_impl {
24  static constexpr int32_t value = _N;
25  static constexpr int32_t default_value = 0;
26 };
27 
28 template <int32_t _N> struct cache {
29  static constexpr int32_t value = _N;
30  static constexpr int32_t default_value = 0;
31 };
32 
33 template <int32_t _N> struct prefetch_impl {
34  static constexpr int32_t value = _N;
35  static constexpr int32_t default_value = 0;
36 };
37 
38 template <int32_t _N> struct statically_coalesce_impl {
39  static constexpr int32_t value = _N;
40  static constexpr int32_t default_value = 1;
41 };
42 
43 template <bool _B> using burst_coalesce = burst_coalesce_impl<_B>;
44 template <bool _B> using prefetch = prefetch_impl<_B>;
46 
47 template <class... _mem_access_params> class lsu final {
48 public:
49  lsu() = delete;
50 
51  template <typename _T, access::address_space _space>
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  static void store(sycl::multi_ptr<_T, _space> Ptr, _T Val) {
67  check_space<_space>();
68  check_store();
69 #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
70  *__builtin_intel_fpga_mem((_T *)Ptr,
71  _burst_coalesce | _cache |
72  _dont_statically_coalesce | _prefetch,
73  _cache_val) = Val;
74 #else
75  *Ptr = Val;
76 #endif
77  }
78 
79 private:
80  static constexpr int32_t _burst_coalesce_val =
81  _GetValue<burst_coalesce_impl, _mem_access_params...>::value;
82  static constexpr uint8_t _burst_coalesce =
83  _burst_coalesce_val == 1 ? BURST_COALESCE : 0;
84 
85  static constexpr int32_t _cache_val =
86  _GetValue<cache, _mem_access_params...>::value;
87  static constexpr uint8_t _cache = (_cache_val > 0) ? CACHE : 0;
88 
89  static constexpr int32_t _statically_coalesce_val =
90  _GetValue<statically_coalesce_impl, _mem_access_params...>::value;
91  static constexpr uint8_t _dont_statically_coalesce =
92  _statically_coalesce_val == 0 ? STATICALLY_COALESCE : 0;
93 
94  static constexpr int32_t _prefetch_val =
95  _GetValue<prefetch_impl, _mem_access_params...>::value;
96  static constexpr uint8_t _prefetch = _prefetch_val ? PREFETCH : 0;
97 
98  static_assert(_cache_val >= 0, "cache size parameter must be non-negative");
99 
100  template <access::address_space _space> static void check_space() {
101  static_assert(
102  _space == access::address_space::global_space ||
103  _space == access::address_space::ext_intel_global_device_space ||
104  _space == access::address_space::ext_intel_global_host_space,
105  "lsu controls are only supported for global_ptr, "
106  "device_ptr, and host_ptr objects");
107  }
108 
109  static void check_load() {
110  static_assert(_cache == 0 || _burst_coalesce == BURST_COALESCE,
111  "unable to implement a cache without a burst coalescer");
112  static_assert(_prefetch == 0 || _burst_coalesce == 0,
113  "unable to implement a prefetcher and a burst coalescer "
114  "simulataneously");
115  static_assert(
116  _prefetch == 0 || _cache == 0,
117  "unable to implement a prefetcher and a cache simulataneously");
118  }
119  static void check_store() {
120  static_assert(_cache == 0, "unable to implement a store LSU with a cache.");
121  static_assert(_prefetch == 0,
122  "unable to implement a store LSU with a prefetcher.");
123  }
124 };
125 } // namespace intel
126 } // namespace ext
127 
128 } // namespace sycl
129 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::ext::intel::_GetValue
Definition: fpga_utils.hpp:23
cl::sycl::ext::intel::PREFETCH
constexpr uint8_t PREFETCH
Definition: fpga_lsu.hpp:21
cl::sycl::ext::intel::burst_coalesce_impl
Definition: fpga_lsu.hpp:23
cl::sycl::ext::intel::lsu::store
static void store(sycl::multi_ptr< _T, _space > Ptr, _T Val)
Definition: fpga_lsu.hpp:66
cl::sycl::ext::intel::STATICALLY_COALESCE
constexpr uint8_t STATICALLY_COALESCE
Definition: fpga_lsu.hpp:20
sycl
Definition: invoke_simd.hpp:68
cl::sycl::multi_ptr
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: atomic.hpp:33
cl::sycl::ext::intel::lsu::load
static _T load(sycl::multi_ptr< _T, _space > Ptr)
Definition: fpga_lsu.hpp:52
defines.hpp
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::ext::intel::statically_coalesce_impl
Definition: fpga_lsu.hpp:38
cl::sycl::ext::intel::prefetch_impl
Definition: fpga_lsu.hpp:33
cl::sycl::ext::intel::cache
Definition: fpga_lsu.hpp:28
fpga_utils.hpp
cl::sycl::ext::intel::BURST_COALESCE
constexpr uint8_t BURST_COALESCE
Definition: fpga_lsu.hpp:18
cl::sycl::ext::intel::lsu
Definition: fpga_lsu.hpp:47
cl::sycl::ext::intel::CACHE
constexpr uint8_t CACHE
Definition: fpga_lsu.hpp:19
pointers.hpp
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:11