Skip to content

[SYCL][FPGA] Expose value_type and min_capacity from SYCL pipes extension class #5471

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.

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 9 commits into from
Mar 3, 2022
Original file line number Diff line number Diff line change
Expand Up @@ -183,6 +183,10 @@ class pipe {
// Non-blocking
static dataT read( bool &success_code );
static void write( const dataT &data, bool &success_code );

// Static members
using value_type = dataT;
size_t min_capacity = min_capacity;
}
----

Expand All @@ -191,8 +195,8 @@ The read and write member functions may be invoked within device code, or within
The template parameters of the device type are defined as:

* `name`: Type that is the basis of pipe identification. Typically a user-defined class, in a user namespace. Forward declaration of the type is sufficient, and the type does not need to be defined.
* `dataT`: The type of data word/packet contained within a pipe. This is the data type that is read during a successful `pipe::read` operation, or written during a successful `pipe::write` operation. The type must be standard layout and trivially copyable.
* `min_capacity`: User defined minimum number of words in units of `dataT` that the pipe must be able to store without any being read out. A minimum capacity is required in some algorithms to avoid deadlock, or for performance tuning. An implementation can include more capacity than this parameter, but not less.
* `dataT`: The type of data word/packet contained within a pipe. This is the data type that is read during a successful `pipe::read` operation, or written during a successful `pipe::write` operation. The type must be standard layout and trivially copyable. This template parameter can be queried by using the `value_type` type alias.
* `min_capacity`: User defined minimum number of words in units of `dataT` that the pipe must be able to store without any being read out. A minimum capacity is required in some algorithms to avoid deadlock, or for performance tuning. An implementation can include more capacity than this parameter, but not less. This template parameter can be queried by using the `min_capacity` static member.

== Pipe types and {cpp} scope

Expand Down
15 changes: 9 additions & 6 deletions sycl/include/sycl/ext/intel/pipes.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -19,6 +19,8 @@ namespace intel {

template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
public:
using value_type = _dataT;
static constexpr int32_t min_capacity = _min_capacity;
// Non-blocking pipes
// Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
// friendly LLVM IR.
Expand Down Expand Up @@ -82,10 +84,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;
#ifdef __SYCL_DEVICE_ONLY__
static constexpr struct ConstantPipeStorage m_Storage = {m_Size, m_Alignment,
m_Capacity};
min_capacity};
#endif // __SYCL_DEVICE_ONLY__
};

Expand All @@ -112,6 +113,8 @@ using ethernet_write_pipe =
template <class _name, class _dataT, size_t _min_capacity = 0>
class kernel_readable_io_pipe {
public:
using value_type = _dataT;
static constexpr int32_t min_capacity = _min_capacity;
// Non-blocking pipes
// Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
// friendly LLVM IR.
Expand Down Expand Up @@ -147,17 +150,18 @@ 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;
#ifdef __SYCL_DEVICE_ONLY__
static constexpr struct ConstantPipeStorage m_Storage
__attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity};
__attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, min_capacity};
#endif // __SYCL_DEVICE_ONLY__
};

template <class _name, class _dataT, size_t _min_capacity = 0>
class kernel_writeable_io_pipe {
public:
using value_type = _dataT;
static constexpr int32_t min_capacity = _min_capacity;
// Non-blocking pipes
// Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
// friendly LLVM IR.
Expand Down Expand Up @@ -191,11 +195,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;
#ifdef __SYCL_DEVICE_ONLY__
static constexpr struct ConstantPipeStorage m_Storage
__attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity};
__attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, min_capacity};
#endif // __SYCL_DEVICE_ONLY__
};

Expand Down
8 changes: 8 additions & 0 deletions sycl/test/extensions/fpga.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,9 @@

#include <CL/sycl.hpp>
#include <sycl/ext/intel/fpga_extensions.hpp>

#include <type_traits>

namespace intelfpga {
template <unsigned ID> struct ethernet_pipe_id {
static constexpr unsigned id = ID;
Expand Down Expand Up @@ -36,6 +39,11 @@ using ethernet_read_pipe =
sycl::ext::intel::kernel_readable_io_pipe<ethernet_pipe_id<0>, int, 0>;
using ethernet_write_pipe =
sycl::ext::intel::kernel_writeable_io_pipe<ethernet_pipe_id<1>, int, 0>;

static_assert(std::is_same_v<ethernet_read_pipe::value_type, int>);
static_assert(std::is_same_v<ethernet_write_pipe::value_type, int>);
static_assert(ethernet_read_pipe::min_capacity == 0);
static_assert(ethernet_write_pipe::min_capacity == 0);
} // namespace intelfpga

int main() {
Expand Down