Skip to content

Commit 12bf4aa

Browse files
aditikumaditikum
authored
[SYCL][FPGA][NFC] Minor update to fpga_lsu header templates (#2375)
Updating fpga_lsu.hpp templates to limit parameter scope and prevent conflicts with any defines.
1 parent 1202ccd commit 12bf4aa

File tree

4 files changed

+76
-76
lines changed

4 files changed

+76
-76
lines changed

‎sycl/include/CL/sycl/INTEL/fpga_lsu.hpp

Lines changed: 20 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -19,38 +19,38 @@ constexpr uint8_t CACHE = 0x2;
1919
constexpr uint8_t STATICALLY_COALESCE = 0x4;
2020
constexpr uint8_t PREFETCH = 0x8;
2121

22-
template <int32_t N> struct burst_coalesce_impl {
23-
static constexpr int32_t value = N;
22+
template <int32_t _N> struct burst_coalesce_impl {
23+
static constexpr int32_t value = _N;
2424
static constexpr int32_t default_value = 0;
2525
};
2626

27-
template <int32_t N> struct cache {
28-
static constexpr int32_t value = N;
27+
template <int32_t _N> struct cache {
28+
static constexpr int32_t value = _N;
2929
static constexpr int32_t default_value = 0;
3030
};
3131

32-
template <int32_t N> struct prefetch_impl {
33-
static constexpr int32_t value = N;
32+
template <int32_t _N> struct prefetch_impl {
33+
static constexpr int32_t value = _N;
3434
static constexpr int32_t default_value = 0;
3535
};
3636

37-
template <int32_t N> struct statically_coalesce_impl {
38-
static constexpr int32_t value = N;
37+
template <int32_t _N> struct statically_coalesce_impl {
38+
static constexpr int32_t value = _N;
3939
static constexpr int32_t default_value = 1;
4040
};
4141

42-
template <bool B> using burst_coalesce = burst_coalesce_impl<B>;
43-
template <bool B> using prefetch = prefetch_impl<B>;
44-
template <bool B> using statically_coalesce = statically_coalesce_impl<B>;
42+
template <bool _B> using burst_coalesce = burst_coalesce_impl<_B>;
43+
template <bool _B> using prefetch = prefetch_impl<_B>;
44+
template <bool _B> using statically_coalesce = statically_coalesce_impl<_B>;
4545

46-
template <class... mem_access_params> class lsu final {
46+
template <class... _mem_access_params> class lsu final {
4747
public:
4848
lsu() = delete;
4949

50-
template <typename T> static T load(sycl::global_ptr<T> Ptr) {
50+
template <typename _T> static _T load(sycl::global_ptr<_T> Ptr) {
5151
check_load();
5252
#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
53-
return *__builtin_intel_fpga_mem((T *)Ptr,
53+
return *__builtin_intel_fpga_mem((_T *)Ptr,
5454
_burst_coalesce | _cache |
5555
_dont_statically_coalesce | _prefetch,
5656
_cache_val);
@@ -59,10 +59,10 @@ template <class... mem_access_params> class lsu final {
5959
#endif
6060
}
6161

62-
template <typename T> static void store(sycl::global_ptr<T> Ptr, T Val) {
62+
template <typename _T> static void store(sycl::global_ptr<_T> Ptr, _T Val) {
6363
check_store();
6464
#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
65-
*__builtin_intel_fpga_mem((T *)Ptr,
65+
*__builtin_intel_fpga_mem((_T *)Ptr,
6666
_burst_coalesce | _cache |
6767
_dont_statically_coalesce | _prefetch,
6868
_cache_val) = Val;
@@ -73,21 +73,21 @@ template <class... mem_access_params> class lsu final {
7373

7474
private:
7575
static constexpr int32_t _burst_coalesce_val =
76-
GetValue<burst_coalesce_impl, mem_access_params...>::value;
76+
_GetValue<burst_coalesce_impl, _mem_access_params...>::value;
7777
static constexpr uint8_t _burst_coalesce =
7878
_burst_coalesce_val == 1 ? BURST_COALESCE : 0;
7979

8080
static constexpr int32_t _cache_val =
81-
GetValue<cache, mem_access_params...>::value;
81+
_GetValue<cache, _mem_access_params...>::value;
8282
static constexpr uint8_t _cache = (_cache_val > 0) ? CACHE : 0;
8383

8484
static constexpr int32_t _statically_coalesce_val =
85-
GetValue<statically_coalesce_impl, mem_access_params...>::value;
85+
_GetValue<statically_coalesce_impl, _mem_access_params...>::value;
8686
static constexpr uint8_t _dont_statically_coalesce =
8787
_statically_coalesce_val == 0 ? STATICALLY_COALESCE : 0;
8888

8989
static constexpr int32_t _prefetch_val =
90-
GetValue<prefetch_impl, mem_access_params...>::value;
90+
_GetValue<prefetch_impl, _mem_access_params...>::value;
9191
static constexpr uint8_t _prefetch = _prefetch_val ? PREFETCH : 0;
9292

9393
static_assert(_cache_val >= 0, "cache size parameter must be non-negative");

‎sycl/include/CL/sycl/INTEL/fpga_reg.hpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,7 @@ __SYCL_INLINE_NAMESPACE(cl) {
1414
namespace sycl {
1515
namespace INTEL {
1616

17-
template <typename T> T fpga_reg(const T &t) {
17+
template <typename _T> _T fpga_reg(const _T &t) {
1818
#if __has_builtin(__builtin_intel_fpga_reg)
1919
return __builtin_intel_fpga_reg(t);
2020
#else
@@ -29,7 +29,7 @@ template <typename T> T fpga_reg(const T &t) {
2929
// Keep it consistent with FPGA attributes like intelfpga::memory()
3030
// Currently clang does not support nested namespace for attributes
3131
namespace intelfpga {
32-
template <typename T> T fpga_reg(const T &t) {
32+
template <typename _T> _T fpga_reg(const _T &t) {
3333
return cl::sycl::INTEL::fpga_reg(t);
3434
}
3535
} // namespace intelfpga

‎sycl/include/CL/sycl/INTEL/fpga_utils.hpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -15,18 +15,18 @@ __SYCL_INLINE_NAMESPACE(cl) {
1515
namespace sycl {
1616
namespace INTEL {
1717

18-
template <template <int32_t> class Type, class T>
19-
struct MatchType : std::is_same<Type<T::value>, T> {};
18+
template <template <int32_t> class _Type, class _T>
19+
struct _MatchType : std::is_same<_Type<_T::value>, _T> {};
2020

21-
template <template <int32_t> class Type, class... T> struct GetValue {
22-
static constexpr auto value = Type<0>::default_value;
21+
template <template <int32_t> class _Type, class... _T> struct _GetValue {
22+
static constexpr auto value = _Type<0>::default_value;
2323
};
2424

25-
template <template <int32_t> class Type, class T1, class... T>
26-
struct GetValue<Type, T1, T...> {
25+
template <template <int32_t> class _Type, class _T1, class... _T>
26+
struct _GetValue<_Type, _T1, _T...> {
2727
static constexpr auto value =
28-
std::conditional<MatchType<Type, T1>::value, T1,
29-
GetValue<Type, T...>>::type::value;
28+
std::conditional<_MatchType<_Type, _T1>::value, _T1,
29+
_GetValue<_Type, _T...>>::type::value;
3030
};
3131
} // namespace INTEL
3232
} // namespace sycl

‎sycl/include/CL/sycl/INTEL/pipes.hpp

Lines changed: 46 additions & 46 deletions
Original file line numberDiff line numberDiff line change
@@ -16,16 +16,16 @@ __SYCL_INLINE_NAMESPACE(cl) {
1616
namespace sycl {
1717
namespace INTEL {
1818

19-
template <class name, class dataT, int32_t min_capacity = 0> class pipe {
19+
template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
2020
public:
2121
// Non-blocking pipes
2222
// Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
2323
// friendly LLVM IR.
24-
static dataT read(bool &Success) {
24+
static _dataT read(bool &Success) {
2525
#ifdef __SYCL_DEVICE_ONLY__
26-
RPipeTy<dataT> RPipe =
27-
__spirv_CreatePipeFromPipeStorage_read<dataT>(&m_Storage);
28-
dataT TempData;
26+
RPipeTy<_dataT> RPipe =
27+
__spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
28+
_dataT TempData;
2929
Success = !static_cast<bool>(
3030
__spirv_ReadPipe(RPipe, &TempData, m_Size, m_Alignment));
3131
return TempData;
@@ -37,10 +37,10 @@ template <class name, class dataT, int32_t min_capacity = 0> class pipe {
3737

3838
// Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
3939
// friendly LLVM IR.
40-
static void write(const dataT &Data, bool &Success) {
40+
static void write(const _dataT &Data, bool &Success) {
4141
#ifdef __SYCL_DEVICE_ONLY__
42-
WPipeTy<dataT> WPipe =
43-
__spirv_CreatePipeFromPipeStorage_write<dataT>(&m_Storage);
42+
WPipeTy<_dataT> WPipe =
43+
__spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
4444
Success = !static_cast<bool>(
4545
__spirv_WritePipe(WPipe, &Data, m_Size, m_Alignment));
4646
#else
@@ -53,11 +53,11 @@ template <class name, class dataT, int32_t min_capacity = 0> class pipe {
5353
// Blocking pipes
5454
// Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
5555
// friendly LLVM IR.
56-
static dataT read() {
56+
static _dataT read() {
5757
#ifdef __SYCL_DEVICE_ONLY__
58-
RPipeTy<dataT> RPipe =
59-
__spirv_CreatePipeFromPipeStorage_read<dataT>(&m_Storage);
60-
dataT TempData;
58+
RPipeTy<_dataT> RPipe =
59+
__spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
60+
_dataT TempData;
6161
__spirv_ReadPipeBlockingINTEL(RPipe, &TempData, m_Size, m_Alignment);
6262
return TempData;
6363
#else
@@ -67,10 +67,10 @@ template <class name, class dataT, int32_t min_capacity = 0> class pipe {
6767

6868
// Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
6969
// friendly LLVM IR.
70-
static void write(const dataT &Data) {
70+
static void write(const _dataT &Data) {
7171
#ifdef __SYCL_DEVICE_ONLY__
72-
WPipeTy<dataT> WPipe =
73-
__spirv_CreatePipeFromPipeStorage_write<dataT>(&m_Storage);
72+
WPipeTy<_dataT> WPipe =
73+
__spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
7474
__spirv_WritePipeBlockingINTEL(WPipe, &Data, m_Size, m_Alignment);
7575
#else
7676
(void)Data;
@@ -79,9 +79,9 @@ template <class name, class dataT, int32_t min_capacity = 0> class pipe {
7979
}
8080

8181
private:
82-
static constexpr int32_t m_Size = sizeof(dataT);
83-
static constexpr int32_t m_Alignment = alignof(dataT);
84-
static constexpr int32_t m_Capacity = min_capacity;
82+
static constexpr int32_t m_Size = sizeof(_dataT);
83+
static constexpr int32_t m_Alignment = alignof(_dataT);
84+
static constexpr int32_t m_Capacity = _min_capacity;
8585
#ifdef __SYCL_DEVICE_ONLY__
8686
static constexpr struct ConstantPipeStorage m_Storage = {m_Size, m_Alignment,
8787
m_Capacity};
@@ -99,26 +99,26 @@ struct ethernet_pipe_id {
9999
static constexpr int32_t id = ID;
100100
};
101101
102-
template <class dataT, size_t min_capacity>
102+
template <class _dataT, size_t _min_capacity>
103103
using ethernet_read_pipe =
104-
kernel_readable_io_pipe<ethernet_pipe_id<0>, dataT, min_capacity>;
104+
kernel_readable_io_pipe<ethernet_pipe_id<0>, _dataT, _min_capacity>;
105105
106-
template <class dataT, size_t min_capacity>
106+
template <class _dataT, size_t _min_capacity>
107107
using ethernet_write_pipe =
108-
kernel_writeable_io_pipe<ethernet_pipe_id<1>, dataT, min_capacity>;
108+
kernel_writeable_io_pipe<ethernet_pipe_id<1>, _dataT, _min_capacity>;
109109
} // namespace intelfpga */
110110

111-
template <class name, class dataT, size_t min_capacity = 0>
111+
template <class _name, class _dataT, size_t _min_capacity = 0>
112112
class kernel_readable_io_pipe {
113113
public:
114114
// Non-blocking pipes
115115
// Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
116116
// friendly LLVM IR.
117-
static dataT read(bool &Success) {
117+
static _dataT read(bool &Success) {
118118
#ifdef __SYCL_DEVICE_ONLY__
119-
RPipeTy<dataT> RPipe =
120-
__spirv_CreatePipeFromPipeStorage_read<dataT>(&m_Storage);
121-
dataT TempData;
119+
RPipeTy<_dataT> RPipe =
120+
__spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
121+
_dataT TempData;
122122
Success = !static_cast<bool>(
123123
__spirv_ReadPipe(RPipe, &TempData, m_Size, m_Alignment));
124124
return TempData;
@@ -131,11 +131,11 @@ class kernel_readable_io_pipe {
131131
// Blocking pipes
132132
// Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
133133
// friendly LLVM IR.
134-
static dataT read() {
134+
static _dataT read() {
135135
#ifdef __SYCL_DEVICE_ONLY__
136-
RPipeTy<dataT> RPipe =
137-
__spirv_CreatePipeFromPipeStorage_read<dataT>(&m_Storage);
138-
dataT TempData;
136+
RPipeTy<_dataT> RPipe =
137+
__spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
138+
_dataT TempData;
139139
__spirv_ReadPipeBlockingINTEL(RPipe, &TempData, m_Size, m_Alignment);
140140
return TempData;
141141
#else
@@ -144,26 +144,26 @@ class kernel_readable_io_pipe {
144144
}
145145

146146
private:
147-
static constexpr int32_t m_Size = sizeof(dataT);
148-
static constexpr int32_t m_Alignment = alignof(dataT);
149-
static constexpr int32_t m_Capacity = min_capacity;
150-
static constexpr int32_t ID = name::id;
147+
static constexpr int32_t m_Size = sizeof(_dataT);
148+
static constexpr int32_t m_Alignment = alignof(_dataT);
149+
static constexpr int32_t m_Capacity = _min_capacity;
150+
static constexpr int32_t ID = _name::id;
151151
#ifdef __SYCL_DEVICE_ONLY__
152152
static constexpr struct ConstantPipeStorage m_Storage
153153
__attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity};
154154
#endif // __SYCL_DEVICE_ONLY__
155155
};
156156

157-
template <class name, class dataT, size_t min_capacity = 0>
157+
template <class _name, class _dataT, size_t _min_capacity = 0>
158158
class kernel_writeable_io_pipe {
159159
public:
160160
// Non-blocking pipes
161161
// Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
162162
// friendly LLVM IR.
163-
static void write(const dataT &Data, bool &Success) {
163+
static void write(const _dataT &Data, bool &Success) {
164164
#ifdef __SYCL_DEVICE_ONLY__
165-
WPipeTy<dataT> WPipe =
166-
__spirv_CreatePipeFromPipeStorage_write<dataT>(&m_Storage);
165+
WPipeTy<_dataT> WPipe =
166+
__spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
167167
Success = !static_cast<bool>(
168168
__spirv_WritePipe(WPipe, &Data, m_Size, m_Alignment));
169169
#else
@@ -176,10 +176,10 @@ class kernel_writeable_io_pipe {
176176
// Blocking pipes
177177
// Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
178178
// friendly LLVM IR.
179-
static void write(const dataT &Data) {
179+
static void write(const _dataT &Data) {
180180
#ifdef __SYCL_DEVICE_ONLY__
181-
WPipeTy<dataT> WPipe =
182-
__spirv_CreatePipeFromPipeStorage_write<dataT>(&m_Storage);
181+
WPipeTy<_dataT> WPipe =
182+
__spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
183183
__spirv_WritePipeBlockingINTEL(WPipe, &Data, m_Size, m_Alignment);
184184
#else
185185
(void)Data;
@@ -188,10 +188,10 @@ class kernel_writeable_io_pipe {
188188
}
189189

190190
private:
191-
static constexpr int32_t m_Size = sizeof(dataT);
192-
static constexpr int32_t m_Alignment = alignof(dataT);
193-
static constexpr int32_t m_Capacity = min_capacity;
194-
static constexpr int32_t ID = name::id;
191+
static constexpr int32_t m_Size = sizeof(_dataT);
192+
static constexpr int32_t m_Alignment = alignof(_dataT);
193+
static constexpr int32_t m_Capacity = _min_capacity;
194+
static constexpr int32_t ID = _name::id;
195195
#ifdef __SYCL_DEVICE_ONLY__
196196
static constexpr struct ConstantPipeStorage m_Storage
197197
__attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity};

0 commit comments

Comments
 (0)