[SYCL][FPGA] Fix success code for non-blocking pipes

SYCL pipe built-ins are mapped on SPIR-V pipe instructions.
From SPIR-V spec (3.32.23 Pipe Instructions):
OpReadPipe
  Result is 0 if the operation is successful and a negative value
  if the pipe is empty.

Same for OpWritePipe.

The bug was that SYCL API was expecting the opposite behaviour.

Signed-off-by: Dmitry Sidorov <[email protected]>

Author: MrSidims

sycl/include/CL/sycl/pipes.hpp

@@ -25,7 +25,7 @@ template <class name, class dataT, int32_t min_capacity = 0> class pipe {
     RPipeTy<dataT> RPipe =
       __spirv_CreatePipeFromPipeStorage_read<dataT>(&m_Storage);
     dataT TempData;
-    Success = static_cast<bool>(
+    Success = !static_cast<bool>(
         __spirv_ReadPipe(RPipe, &TempData, m_Size, m_Alignment));
     return TempData;
 #else
@@ -39,7 +39,7 @@ template <class name, class dataT, int32_t min_capacity = 0> class pipe {
 #ifdef __SYCL_DEVICE_ONLY__
     WPipeTy<dataT> WPipe =
       __spirv_CreatePipeFromPipeStorage_write<dataT>(&m_Storage);
-    Success = static_cast<bool>(
+    Success = !static_cast<bool>(
         __spirv_WritePipe(WPipe, &Data, m_Size, m_Alignment));
 #else
     assert(!"Pipes are not supported on a host device!");

sycl/test/fpga_tests/fpga_pipes.cpp

@@ -10,6 +10,9 @@
 #include <CL/sycl.hpp>
 #include <iostream>
 
+// Size of an array passing through a pipe
+constexpr size_t N = 10;
+
 // For simple non-blocking pipes with explicit type
 class some_nb_pipe;
 
@@ -140,6 +143,47 @@ int test_multiple_nb_pipe(cl::sycl::queue Queue) {
   return 0;
 }
 
+// Test for array passing through a non-blocking pipe
+template<int TestNumber>
+int test_array_th_nb_pipe(cl::sycl::queue Queue) {
+  int data[N] = {0};
+  using AnotherNbPipe = cl::sycl::pipe<class another_nb_pipe, int>;
+
+  Queue.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task<class writer<TestNumber>>([=]() {
+      bool SuccessCode = false;
+      for (size_t i = 0; i != N; ++i) {
+        do {
+          AnotherNbPipe::write(i, SuccessCode);
+        } while (!SuccessCode);
+      }
+    });
+  });
+
+  cl::sycl::buffer<int, 1> writeBuf(data, N);
+  Queue.submit([&](cl::sycl::handler &cgh) {
+    auto write_acc = writeBuf.get_access<cl::sycl::access::mode::write>(cgh);
+    cgh.single_task<class reader<TestNumber>>([=]() {
+      for (size_t i = 0; i != N; ++i) {
+        bool SuccessCode = false;
+        do {
+          write_acc[i] = AnotherNbPipe::read(SuccessCode);
+        } while (!SuccessCode);
+      }
+    });
+  });
+
+  auto readHostBuffer = writeBuf.get_access<cl::sycl::access::mode::read>();
+  for (size_t i = 0; i != N; ++i) {
+    if (readHostBuffer[i] != i)
+      std::cout << "Test: " << TestNumber << "\nResult mismatches "
+                << readHostBuffer[i] << " Vs expected " << i << std::endl;
+    return -1;
+  }
+
+  return 0;
+}
+
 // Test for simple blocking pipes
 template<typename PipeName, int TestNumber>
 int test_simple_bl_pipe(cl::sycl::queue Queue) {
@@ -211,6 +255,39 @@ int test_multiple_bl_pipe(cl::sycl::queue Queue) {
   return 0;
 }
 
+// Test for array passing through a blocking pipe
+template<int TestNumber>
+int test_array_th_bl_pipe(cl::sycl::queue Queue) {
+  int data[N] = {0};
+  using AnotherBlPipe = cl::sycl::pipe<class another_bl_pipe, int>;
+
+  Queue.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task<class writer<TestNumber>>([=]() {
+      for (size_t i = 0; i != N; ++i)
+        AnotherBlPipe::write(i);
+    });
+  });
+
+  cl::sycl::buffer<int, 1> writeBuf(data, N);
+  Queue.submit([&](cl::sycl::handler &cgh) {
+    auto write_acc = writeBuf.get_access<cl::sycl::access::mode::write>(cgh);
+    cgh.single_task<class reader<TestNumber>>([=]() {
+      for (size_t i = 0; i != N; ++i)
+        write_acc[i] = AnotherBlPipe::read();
+    });
+  });
+
+  auto readHostBuffer = writeBuf.get_access<cl::sycl::access::mode::read>();
+  for (size_t i = 0; i != N; ++i) {
+    if (readHostBuffer[i] != i)
+      std::cout << "Test: " << TestNumber << "\nResult mismatches "
+                << readHostBuffer[i] << " Vs expected " << i << std::endl;
+    return -1;
+  }
+
+  return 0;
+}
+
 int main() {
   cl::sycl::queue Queue;
 
@@ -230,5 +307,9 @@ int main() {
   Result &= test_simple_bl_pipe<templ_bl_pipe<0>, /*test number*/ 9>(Queue);
   Result &= test_multiple_bl_pipe</*test number*/ 10>(Queue);
 
+  // Test for an array data passing through a pipe
+  Result &= test_array_th_nb_pipe</*test number*/ 11>(Queue);
+  Result &= test_array_th_bl_pipe</*test number*/ 12>(Queue);
+
   return Result;
 }