Skip to content

[SYCL][FPGA][NFC] Minor update to fpga_lsu header templates #2375

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Sep 1, 2020
Merged

[SYCL][FPGA][NFC] Minor update to fpga_lsu header templates #2375

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
7553285
[SYCL][FPGA] Minor update to fpga_lsu header templates
aditikum Aug 26, 2020
b6a2af8
Updating fpga headers scope
aditikum Aug 26, 2020
fe56227
Update fpga_lsu.hpp GetValue calls
aditikum Aug 27, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
40 changes: 20 additions & 20 deletions sycl/include/CL/sycl/INTEL/fpga_lsu.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,38 +19,38 @@ constexpr uint8_t CACHE = 0x2;
constexpr uint8_t STATICALLY_COALESCE = 0x4;
constexpr uint8_t PREFETCH = 0x8;

template <int32_t N> struct burst_coalesce_impl {
static constexpr int32_t value = N;
template <int32_t _N> struct burst_coalesce_impl {
static constexpr int32_t value = _N;
static constexpr int32_t default_value = 0;
};

template <int32_t N> struct cache {
static constexpr int32_t value = N;
template <int32_t _N> struct cache {
static constexpr int32_t value = _N;
static constexpr int32_t default_value = 0;
};

template <int32_t N> struct prefetch_impl {
static constexpr int32_t value = N;
template <int32_t _N> struct prefetch_impl {
static constexpr int32_t value = _N;
static constexpr int32_t default_value = 0;
};

template <int32_t N> struct statically_coalesce_impl {
static constexpr int32_t value = N;
template <int32_t _N> struct statically_coalesce_impl {
static constexpr int32_t value = _N;
static constexpr int32_t default_value = 1;
};

template <bool B> using burst_coalesce = burst_coalesce_impl<B>;
template <bool B> using prefetch = prefetch_impl<B>;
template <bool B> using statically_coalesce = statically_coalesce_impl<B>;
template <bool _B> using burst_coalesce = burst_coalesce_impl<_B>;
template <bool _B> using prefetch = prefetch_impl<_B>;
template <bool _B> using statically_coalesce = statically_coalesce_impl<_B>;

template <class... mem_access_params> class lsu final {
template <class... _mem_access_params> class lsu final {
public:
lsu() = delete;

template <typename T> static T load(sycl::global_ptr<T> Ptr) {
template <typename _T> static _T load(sycl::global_ptr<_T> Ptr) {
check_load();
#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
return *__builtin_intel_fpga_mem((T *)Ptr,
return *__builtin_intel_fpga_mem((_T *)Ptr,
_burst_coalesce | _cache |
_dont_statically_coalesce | _prefetch,
_cache_val);
Expand All @@ -59,10 +59,10 @@ template <class... mem_access_params> class lsu final {
#endif
}

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

private:
static constexpr int32_t _burst_coalesce_val =
GetValue<burst_coalesce_impl, mem_access_params...>::value;
_GetValue<burst_coalesce_impl, _mem_access_params...>::value;
static constexpr uint8_t _burst_coalesce =
_burst_coalesce_val == 1 ? BURST_COALESCE : 0;

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

static constexpr int32_t _statically_coalesce_val =
GetValue<statically_coalesce_impl, mem_access_params...>::value;
_GetValue<statically_coalesce_impl, _mem_access_params...>::value;
static constexpr uint8_t _dont_statically_coalesce =
_statically_coalesce_val == 0 ? STATICALLY_COALESCE : 0;

static constexpr int32_t _prefetch_val =
GetValue<prefetch_impl, mem_access_params...>::value;
_GetValue<prefetch_impl, _mem_access_params...>::value;
static constexpr uint8_t _prefetch = _prefetch_val ? PREFETCH : 0;

static_assert(_cache_val >= 0, "cache size parameter must be non-negative");
Expand Down
4 changes: 2 additions & 2 deletions sycl/include/CL/sycl/INTEL/fpga_reg.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@ __SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
namespace INTEL {

template <typename T> T fpga_reg(const T &t) {
template <typename _T> _T fpga_reg(const _T &t) {
#if __has_builtin(__builtin_intel_fpga_reg)
return __builtin_intel_fpga_reg(t);
#else
Expand All @@ -29,7 +29,7 @@ template <typename T> T fpga_reg(const T &t) {
// Keep it consistent with FPGA attributes like intelfpga::memory()
// Currently clang does not support nested namespace for attributes
namespace intelfpga {
template <typename T> T fpga_reg(const T &t) {
template <typename _T> _T fpga_reg(const _T &t) {
return cl::sycl::INTEL::fpga_reg(t);
}
} // namespace intelfpga
16 changes: 8 additions & 8 deletions sycl/include/CL/sycl/INTEL/fpga_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,18 +15,18 @@ __SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
namespace INTEL {

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

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

template <template <int32_t> class Type, class T1, class... T>
struct GetValue<Type, T1, T...> {
template <template <int32_t> class _Type, class _T1, class... _T>
struct _GetValue<_Type, _T1, _T...> {
static constexpr auto value =
std::conditional<MatchType<Type, T1>::value, T1,
GetValue<Type, T...>>::type::value;
std::conditional<_MatchType<_Type, _T1>::value, _T1,
_GetValue<_Type, _T...>>::type::value;
};
} // namespace INTEL
} // namespace sycl
Expand Down
92 changes: 46 additions & 46 deletions sycl/include/CL/sycl/INTEL/pipes.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,16 +16,16 @@ __SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
namespace INTEL {

template <class name, class dataT, int32_t min_capacity = 0> class pipe {
template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
public:
// Non-blocking pipes
// Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
// friendly LLVM IR.
static dataT read(bool &Success) {
static _dataT read(bool &Success) {
#ifdef __SYCL_DEVICE_ONLY__
RPipeTy<dataT> RPipe =
__spirv_CreatePipeFromPipeStorage_read<dataT>(&m_Storage);
dataT TempData;
RPipeTy<_dataT> RPipe =
__spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
_dataT TempData;
Success = !static_cast<bool>(
__spirv_ReadPipe(RPipe, &TempData, m_Size, m_Alignment));
return TempData;
Expand All @@ -37,10 +37,10 @@ template <class name, class dataT, int32_t min_capacity = 0> class pipe {

// Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
// friendly LLVM IR.
static void write(const dataT &Data, bool &Success) {
static void write(const _dataT &Data, bool &Success) {
#ifdef __SYCL_DEVICE_ONLY__
WPipeTy<dataT> WPipe =
__spirv_CreatePipeFromPipeStorage_write<dataT>(&m_Storage);
WPipeTy<_dataT> WPipe =
__spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
Success = !static_cast<bool>(
__spirv_WritePipe(WPipe, &Data, m_Size, m_Alignment));
#else
Expand All @@ -53,11 +53,11 @@ template <class name, class dataT, int32_t min_capacity = 0> class pipe {
// Blocking pipes
// Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
// friendly LLVM IR.
static dataT read() {
static _dataT read() {
#ifdef __SYCL_DEVICE_ONLY__
RPipeTy<dataT> RPipe =
__spirv_CreatePipeFromPipeStorage_read<dataT>(&m_Storage);
dataT TempData;
RPipeTy<_dataT> RPipe =
__spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
_dataT TempData;
__spirv_ReadPipeBlockingINTEL(RPipe, &TempData, m_Size, m_Alignment);
return TempData;
#else
Expand All @@ -67,10 +67,10 @@ template <class name, class dataT, int32_t min_capacity = 0> class pipe {

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

private:
static constexpr int32_t m_Size = sizeof(dataT);
static constexpr int32_t m_Alignment = alignof(dataT);
static constexpr int32_t m_Capacity = min_capacity;
static constexpr int32_t m_Size = sizeof(_dataT);
static constexpr int32_t m_Alignment = alignof(_dataT);
static constexpr int32_t m_Capacity = _min_capacity;
#ifdef __SYCL_DEVICE_ONLY__
static constexpr struct ConstantPipeStorage m_Storage = {m_Size, m_Alignment,
m_Capacity};
Expand All @@ -99,26 +99,26 @@ struct ethernet_pipe_id {
static constexpr int32_t id = ID;
};

template <class dataT, size_t min_capacity>
template <class _dataT, size_t _min_capacity>
using ethernet_read_pipe =
kernel_readable_io_pipe<ethernet_pipe_id<0>, dataT, min_capacity>;
kernel_readable_io_pipe<ethernet_pipe_id<0>, _dataT, _min_capacity>;

template <class dataT, size_t min_capacity>
template <class _dataT, size_t _min_capacity>
using ethernet_write_pipe =
kernel_writeable_io_pipe<ethernet_pipe_id<1>, dataT, min_capacity>;
kernel_writeable_io_pipe<ethernet_pipe_id<1>, _dataT, _min_capacity>;
} // namespace intelfpga */

template <class name, class dataT, size_t min_capacity = 0>
template <class _name, class _dataT, size_t _min_capacity = 0>
class kernel_readable_io_pipe {
public:
// Non-blocking pipes
// Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
// friendly LLVM IR.
static dataT read(bool &Success) {
static _dataT read(bool &Success) {
#ifdef __SYCL_DEVICE_ONLY__
RPipeTy<dataT> RPipe =
__spirv_CreatePipeFromPipeStorage_read<dataT>(&m_Storage);
dataT TempData;
RPipeTy<_dataT> RPipe =
__spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
_dataT TempData;
Success = !static_cast<bool>(
__spirv_ReadPipe(RPipe, &TempData, m_Size, m_Alignment));
return TempData;
Expand All @@ -131,11 +131,11 @@ class kernel_readable_io_pipe {
// Blocking pipes
// Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
// friendly LLVM IR.
static dataT read() {
static _dataT read() {
#ifdef __SYCL_DEVICE_ONLY__
RPipeTy<dataT> RPipe =
__spirv_CreatePipeFromPipeStorage_read<dataT>(&m_Storage);
dataT TempData;
RPipeTy<_dataT> RPipe =
__spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
_dataT TempData;
__spirv_ReadPipeBlockingINTEL(RPipe, &TempData, m_Size, m_Alignment);
return TempData;
#else
Expand All @@ -144,26 +144,26 @@ class kernel_readable_io_pipe {
}

private:
static constexpr int32_t m_Size = sizeof(dataT);
static constexpr int32_t m_Alignment = alignof(dataT);
static constexpr int32_t m_Capacity = min_capacity;
static constexpr int32_t ID = name::id;
static constexpr int32_t m_Size = sizeof(_dataT);
static constexpr int32_t m_Alignment = alignof(_dataT);
static constexpr int32_t m_Capacity = _min_capacity;
static constexpr int32_t ID = _name::id;
#ifdef __SYCL_DEVICE_ONLY__
static constexpr struct ConstantPipeStorage m_Storage
__attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity};
#endif // __SYCL_DEVICE_ONLY__
};

template <class name, class dataT, size_t min_capacity = 0>
template <class _name, class _dataT, size_t _min_capacity = 0>
class kernel_writeable_io_pipe {
public:
// Non-blocking pipes
// Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
// friendly LLVM IR.
static void write(const dataT &Data, bool &Success) {
static void write(const _dataT &Data, bool &Success) {
#ifdef __SYCL_DEVICE_ONLY__
WPipeTy<dataT> WPipe =
__spirv_CreatePipeFromPipeStorage_write<dataT>(&m_Storage);
WPipeTy<_dataT> WPipe =
__spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
Success = !static_cast<bool>(
__spirv_WritePipe(WPipe, &Data, m_Size, m_Alignment));
#else
Expand All @@ -176,10 +176,10 @@ class kernel_writeable_io_pipe {
// Blocking pipes
// Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
// friendly LLVM IR.
static void write(const dataT &Data) {
static void write(const _dataT &Data) {
#ifdef __SYCL_DEVICE_ONLY__
WPipeTy<dataT> WPipe =
__spirv_CreatePipeFromPipeStorage_write<dataT>(&m_Storage);
WPipeTy<_dataT> WPipe =
__spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
__spirv_WritePipeBlockingINTEL(WPipe, &Data, m_Size, m_Alignment);
#else
(void)Data;
Expand All @@ -188,10 +188,10 @@ class kernel_writeable_io_pipe {
}

private:
static constexpr int32_t m_Size = sizeof(dataT);
static constexpr int32_t m_Alignment = alignof(dataT);
static constexpr int32_t m_Capacity = min_capacity;
static constexpr int32_t ID = name::id;
static constexpr int32_t m_Size = sizeof(_dataT);
static constexpr int32_t m_Alignment = alignof(_dataT);
static constexpr int32_t m_Capacity = _min_capacity;
static constexpr int32_t ID = _name::id;
#ifdef __SYCL_DEVICE_ONLY__
static constexpr struct ConstantPipeStorage m_Storage
__attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity};
Expand Down