[FPGA] Update matmul sample to use annotated_arg

DirectProgramming/C++SYCL_FPGA/ReferenceDesigns/matmul/src/memory_transfers.hpp

@@ -4,6 +4,9 @@
 #include "constexpr_math.hpp"
 #include "tuple.hpp"
 #include "unrolled_loop.hpp"
+#include <sycl/sycl.hpp>
+
+using namespace sycl::ext::oneapi::experimental;
 
 /**
  * Feeder A Kernel.
@@ -31,21 +34,24 @@ template <typename TT,              // Datatype of the elements of the matrix
           typename PipeA,           // Input pipe for matrix
           typename PipeDone, // Pipe to notify compute kernel when to stop
                              // reading inputs
-          int dwidth = elems_per_ddr_access * sizeof(TT) * 8>
+          int data_width = elems_per_ddr_access * sizeof(TT) * 8>
 class MatrixReadFromDDRToPipeA {
 public:
 #if !defined(IS_BSP)
   // Customizing mmhost only supported when targetting an FPGA part/family
-  mmhost(aspace, // buffer_location or aspace
-         28,     // address width
-         dwidth, // data width
-         0,      // latency
-         1,      // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
-         1,      // maxburst
-         0,      // align, 0 defaults to alignment of the type
-         1)      // waitrequest, 0: false, 1: true
+  annotated_arg<TT*,
+    decltype(properties{
+      buffer_location<aspace>
+    , awidth<28>
+    , dwidth<data_width>
+    , latency<0>
+    , read_write_mode_read
+    , maxburst<1>
+    , wait_request_requested
+  })> a_ptr;    // Annotated input matrix pointer
+#else
+  TT *a_ptr;   // Input matrix pointer
 #endif
-      TT *a_ptr;   // Input matrix pointer
   int repetitions; // Number of times to write the same matrix to the pipe
 
   void operator()() const {
@@ -184,21 +190,24 @@ template <typename TT,              // Datatype of the elements of the matrix
           int elems_per_ddr_access, // Number of elements per DDR access
           int num_matrices,         // Number of pairs of matrices to multiply
           typename PipeB,           // Input pipe for matrix
-          int dwidth = elems_per_ddr_access * sizeof(TT) * 8>
+          int data_width = elems_per_ddr_access * sizeof(TT) * 8>
 class MatrixReadFromDDRToPipeB {
 public:
 #if !defined(IS_BSP)
   // Customizing mmhost only supported when targetting an FPGA part/family
-  mmhost(aspace, // buffer_location or aspace
-         28,     // address width
-         dwidth, // data width
-         0,      // latency
-         1,      // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
-         1,      // maxburst
-         0,      // align, 0 defaults to alignment of the type
-         1)      // waitrequest, 0: false, 1: true
-#endif
+  annotated_arg<TT*,
+    decltype(properties{
+      buffer_location<aspace>
+    , awidth<28>
+    , dwidth<data_width>
+    , latency<0>
+    , read_write_mode_read
+    , maxburst<1>
+    , wait_request_requested
+  })> b_ptr;    // Annotated input matrix pointer
+#else
       TT *b_ptr;   // Input matrix pointer
+#endif
   int repetitions; // Number of times to write the same matrix to the pipe
 
   void operator()() const {
@@ -329,21 +338,24 @@ template <typename TT,              // Datatype of the elements of the matrix
           int elems_per_ddr_access, // Number of elements per DDR access
           int num_matrices,         // Number of pairs of matrices to multiply
           typename PipeC,           // Output pipe for matrix
-          int dwidth = elems_per_ddr_access * sizeof(TT) * 8>
+          int data_width = elems_per_ddr_access * sizeof(TT) * 8>
 class MatrixReadPipeToDDR {
 public:
 #if !defined(IS_BSP)
   // Customizing mmhost only supported when targetting an FPGA part/family
-  mmhost(aspace, // buffer_location or aspace
-         28,     // address width
-         dwidth, // data width
-         0,      // latency
-         2,      // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
-         1,      // maxburst
-         0,      // align, 0 defaults to alignment of the type
-         1)      // waitrequest, 0: false, 1: true
-#endif
+  annotated_arg<TT*,
+    decltype(properties{
+      buffer_location<aspace>
+    , awidth<28>
+    , dwidth<data_width>
+    , latency<0>
+    , read_write_mode_write
+    , maxburst<1>
+    , wait_request_requested
+  })> c_ptr;    // Annotated output matrix pointer
+#else
       TT *c_ptr;   // Output matrix pointer
+#endif
   int repetitions; // Number of time to read the same matrix to the pipe
 
   void operator()() const {

DirectProgramming/C++SYCL_FPGA/Tutorials/Features/experimental/invocation_interfaces/src/register_map_functor_model.cpp

@@ -4,6 +4,8 @@
 
 #include "exception_handler.hpp"
 
+using namespace sycl::ext::oneapi::experimental;
+
 using ValueT = int;
 
 // offloaded computation
@@ -14,14 +16,15 @@ ValueT SomethingComplicated(ValueT val) { return (ValueT)(val * (val + 1)); }
 struct FunctorRegisterMapIP {
   // Use the 'register_map' annotation on a kernel argument to specify it to be
   // a register map kernel argument.
-  register_map ValueT *input;
+  annotated_arg<ValueT, decltype(properties{register_map})> input;
   // Without the annotations, kernel arguments will be inferred to be register
   // map kernel arguments if the kernel invocation interface is register mapped,
   // and vise-versa.
   ValueT *output;
   // A kernel with a register map invocation interface can also independently
   // have streaming kernel arguments, when annotated by 'conduit'.
-  conduit size_t n;
+  annotated_arg<size_t, decltype(properties{conduit})> n;
+
   register_map_interface void operator()() const {
     for (int i = 0; i < n; i++) {
       output[i] = SomethingComplicated(input[i]);

DirectProgramming/C++SYCL_FPGA/Tutorials/Features/experimental/invocation_interfaces/src/streaming_functor_model.cpp

@@ -4,6 +4,8 @@
 
 #include "exception_handler.hpp"
 
+using namespace sycl::ext::oneapi::experimental;
+
 using ValueT = int;
 
 // offloaded computation
@@ -14,10 +16,10 @@ ValueT SomethingComplicated(ValueT val) { return (ValueT)(val * (val + 1)); }
 struct FunctorStreamingIP {
   // Use the 'conduit' annotation on a kernel argument to specify it to be
   // a streaming kernel argument.
-  conduit ValueT *input;
+  annotated_arg<ValueT *, decltype(properties{conduit})> input;
   // A kernel with a streaming invocation interface can also independently
   // have register map kernel arguments, when annotated by 'register_map'.
-  register_map ValueT *output;
+  annotated_arg<ValueT *, decltype(properties{register_map})> output;
   // Without the annotations, kernel arguments will be inferred to be streaming
   // kernel arguments if the kernel invocation interface is streaming, and
   // vise-versa.

DirectProgramming/C++SYCL_FPGA/Tutorials/Features/experimental/pipelined_kernels/src/pipelined_kernels.cpp

@@ -4,13 +4,16 @@
 
 #include "exception_handler.hpp"
 
+using namespace sycl::ext::oneapi::experimental;
+
 using ValueT = int;
 
 // offloaded computation
 ValueT SomethingComplicated(ValueT val) { return (ValueT)(val * (val + 1)); }
 
 struct MyIP {
-  conduit ValueT *input;
+  annotated_arg<ValueT *, decltype(properties{conduit})> input;
+
   streaming_pipelined_interface void operator()() const {
     ValueT temp = *input;
     *input = SomethingComplicated(temp);