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

This is a convenience for users who may template their functions/classes on SYCL pipe types.
This allows them to do something like:

```c++
template<typename MyPipe>
void foo(/* ... */) {
    using PipeT = MyPipe::value_type;
    // ...
}
```

Test update: [intel/llvm-test-suite#800](intel/llvm-test-suite#800)

Author: tyoungsc

sycl/doc/extensions/supported/sycl_ext_intel_dataflow_pipes.asciidoc

@@ -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;
 }
 ----
 
@@ -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
 

sycl/include/sycl/ext/intel/pipes.hpp

@@ -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.
@@ -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__
 };
 
@@ -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.
@@ -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.
@@ -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__
 };
 

sycl/test/extensions/fpga.cpp

@@ -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;
@@ -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() {