Address Ronan's comments

Author: jessicadavies-intel

DirectProgramming/DPC++FPGA/Tutorials/DesignPatterns/compute_units/README.md

@@ -42,22 +42,25 @@ The `Source` kernel reads data from host memory and writes it to the first pipe:
 void SourceKernel(queue &q, float data) {
 
   q.submit([&](handler &h) {
-    h.single_task<Source>([=]() { Pipes::PipeAt<0>::write(data); });
+    h.single_task<Source>([=] { Pipes::PipeAt<0>::write(data); });
   });
 }
 ```
 
 At the end of the chain, the `Sink` kernel reads data from the last pipe and returns it to the host:
 
 ``` c++
-void SinkKernel(queue &q, std::array<float, 1> &out_data) {
+void SinkKernel(queue &q, float *out_data) {
 
-  buffer<float, 1> out_buf(out_data.data(), 1);
+  // The verbose buffer syntax is necessary here,
+  // since out_data is just a single scalar value
+  // and its size can not be inferred automatically
+  buffer<float, 1> out_buf(out_data, 1);
 
   q.submit([&](handler &h) {
     auto out_accessor = out_buf.get_access<access::mode::write>(h);
     h.single_task<Sink>(
-        [=]() { out_accessor[0] = Pipes::PipeAt<kEngines>::read(); });
+        [=] { out_accessor[0] = Pipes::PipeAt<kEngines>::read(); });
   });
 }
 ```
@@ -66,7 +69,7 @@ void SinkKernel(queue &q, std::array<float, 1> &out_data) {
 
 ``` c++
 intelfpga::submit_compute_units<kEngines, ChainComputeUnit>(q, [=](auto ID) {
-    float f = Pipes::PipeAt<ID>::read();
+    auto f = Pipes::PipeAt<ID>::read();
     Pipes::PipeAt<ID + 1>::write(f);
   });
 ```

DirectProgramming/DPC++FPGA/Tutorials/DesignPatterns/compute_units/src/compute_units.cpp

@@ -45,11 +45,12 @@ void SourceKernel(queue &q, float data) {
 }
 
 // Get the data out of the chain and return it to the host
-void SinkKernel(queue &q, std::array<float, 1> &out_data) {
+void SinkKernel(queue &q, float *out_data) {
 
-  // Use verbose SYCL 1.2 syntax for the output buffer.
-  // (This will become unnecessary in a future compiler version.)
-  buffer<float, 1> out_buf(out_data.data(), 1);
+  // The verbose buffer syntax is necessary here,
+  // since out_data is just a single scalar value
+  // and its size can not be inferred automatically
+  buffer<float, 1> out_buf(out_data, 1);
 
   q.submit([&](handler &h) {
     auto out_accessor = out_buf.get_access<access::mode::write>(h);
@@ -66,7 +67,7 @@ int main() {
   INTEL::fpga_selector device_selector;
 #endif
 
-  std::array<float, 1> out_data = {0};
+  float out_data = 0;
 
   try {
     queue q(device_selector, dpc_common::exception_handler);
@@ -77,12 +78,12 @@ int main() {
     // Enqueue the chain of kEngines compute units
     // Compute unit must take a single argument, its ID
     submit_compute_units<kEngines, ChainComputeUnit>(q, [=](auto ID) {
-      float f = Pipes::PipeAt<ID>::read();
+      auto f = Pipes::PipeAt<ID>::read();
       Pipes::PipeAt<ID + 1>::write(f);
     });
 
     // Enqueue the Sink kernel
-    SinkKernel(q, out_data);
+    SinkKernel(q, &out_data);
 
   } catch (sycl::exception const &e) {
     // Catches exceptions in the host code
@@ -99,9 +100,9 @@ int main() {
   }
 
   // Verify result
-  if (out_data[0] != kTestData) {
+  if (out_data != kTestData) {
     std::cout << "FAILED: The results are incorrect\n";
-    std::cout << "Expected: " << kTestData << " Got: " << out_data[0] << "\n";
+    std::cout << "Expected: " << kTestData << " Got: " << out_data << "\n";
     return 1;
   }
 

DirectProgramming/DPC++FPGA/Tutorials/DesignPatterns/compute_units/src/pipe_array.hpp

@@ -14,6 +14,7 @@
   namespace INTEL = sycl::intel;  // Namespace alias for backward compatibility
 #else
   #include <CL/sycl/INTEL/fpga_extensions.hpp>
+  using namespace sycl;
 #endif