[SYCL] Fix assertion at host pipe initialization in case of multiple TUs (#16756)

If there are multiple translation units using the host pipe then
initialize function will be called multiple times with the same pointer
because special function for pipe registration and initialization
(__sycl_host_pipe_registration) has to be emitted by frontend for each
translation unit. Just assert that pointer is the same.

Author: againull

sycl/source/detail/host_pipe_map_entry.hpp

@@ -43,7 +43,18 @@ struct HostPipeMapEntry {
   }
 
   void initialize(const void *HostPipePtr) {
-    assert(!MHostPipePtr && "Host pipe pointer has already been initialized.");
+    // If there are multiple translation units using the host pipe then
+    // initialize function will be called multiple times with the same pointer
+    // because special function for pipe registration and initialization has to
+    // be emitted by frontend for each translation unit. Just make sure that
+    // pointer is the same.
+    if (MHostPipePtr) {
+      assert(MHostPipePtr == HostPipePtr &&
+             "Host pipe intializations disagree on address of the host pipe on "
+             "host.");
+      return;
+    }
+
     MHostPipePtr = HostPipePtr;
   }
 

sycl/test-e2e/Basic/fpga_tests/Inputs/kernel.cpp

@@ -0,0 +1,7 @@
+#include "mypipe.hpp"
+#include <cstdint>
+
+void KernelFunctor::operator()() const {
+  uint32_t data = 2;
+  my_pipe::write(data);
+}

sycl/test-e2e/Basic/fpga_tests/Inputs/mypipe.hpp

@@ -0,0 +1,11 @@
+#include <cstdint>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/intel/fpga_extensions.hpp>
+
+#pragma once
+
+struct KernelFunctor {
+  using my_pipe =
+      sycl::ext::intel::experimental::pipe<struct sample_host_pipe, uint32_t>;
+  SYCL_EXTERNAL void operator()() const;
+};

sycl/test-e2e/Basic/fpga_tests/fpga_host_pipe_multiple_tus.cpp

@@ -0,0 +1,20 @@
+// REQUIRES: accelerator
+// RUN: %clangxx -fsycl -fintelfpga %s %S/Inputs/kernel.cpp -I%S/Inputs -o %t.out
+// RUN: %{run} %t.out | FileCheck %s
+
+// Test checks that host pipe initialization doesn't fail if pipe is used in
+// multiple translation units.
+
+#include "mypipe.hpp"
+#include <iostream>
+
+int main() {
+  sycl::queue q{sycl::ext::intel::fpga_emulator_selector_v};
+  q.submit([&](sycl::handler &cgh) { cgh.single_task(KernelFunctor{}); });
+  uint32_t result = KernelFunctor::my_pipe::read(q);
+  q.wait();
+  // CHECK: 2
+  std::cout << result << std::endl;
+
+  return 0;
+}