[SYCL][FPGA] Add device query for pipe feature support

Pipe feature is enabled under SYCL_INTEL_data_flow_pipes
extension, to decide if a device support it we query for
device type (should be accelerator, later FPGA) and vendor
string should be "Intel(R) Corporation".

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

Author: MrSidims

sycl/include/CL/sycl/detail/device_info.hpp

@@ -289,6 +289,27 @@ struct get_device_info<vector_class<size_t>,
   }
 };
 
+// Specialization for kernel to kernel pipes.
+// Here we step away from OpenCL, since there is no appropriate cl_device_info
+// enum for global pipes feature.
+template <>
+struct get_device_info<bool, info::device::kernel_kernel_pipe_support> {
+  static bool _(RT::PiDevice dev) {
+    // We claim, that all Intel FPGA devices support kernel to kernel pipe
+    // feature (at least at the scope of SYCL_INTEL_data_flow_pipes extension).
+    platform plt = get_device_info<platform, info::device::platform>::_(dev);
+    string_class platform_name = plt.get_info<info::platform::name>();
+    if (platform_name == "Intel(R) FPGA Emulation Platform for OpenCL(TM)" ||
+        platform_name == "Intel(R) FPGA SDK for OpenCL(TM)")
+      return true;
+
+    // TODO: a better way is to query for supported SPIR-V capabilities when
+    // it's started to be possible. Also, if a device's backend supports
+    // SPIR-V 1.1 (where Pipe Storage feature was defined), than it supports
+    // the feature as well.
+    return false;
+  }
+};
 
 // SYCL host device information
 

sycl/include/CL/sycl/info/device_traits.def

@@ -79,3 +79,4 @@ PARAM_TRAITS_SPEC(device, reference_count, pi_uint32)
 PARAM_TRAITS_SPEC(device, max_num_sub_groups, pi_uint32)
 PARAM_TRAITS_SPEC(device, sub_group_independent_forward_progress, bool)
 PARAM_TRAITS_SPEC(device, sub_group_sizes, vector_class<size_t>)
+PARAM_TRAITS_SPEC(device, kernel_kernel_pipe_support, bool)

sycl/include/CL/sycl/info/info_desc.hpp

@@ -121,7 +121,8 @@ enum class device : cl_device_info {
   sub_group_independent_forward_progress =
       CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS,
   sub_group_sizes = CL_DEVICE_SUB_GROUP_SIZES_INTEL,
-  partition_type_property
+  partition_type_property,
+  kernel_kernel_pipe_support
 };
 
 enum class device_type : pi_uint64 {

sycl/source/detail/device_info.cpp

@@ -496,6 +496,11 @@ bool get_device_info_host<
   throw runtime_error("Sub-group feature is not supported on HOST device.");
 }
 
+template <>
+bool get_device_info_host<info::device::kernel_kernel_pipe_support>() {
+  return false;
+}
+
 } // namespace detail
 } // namespace sycl
 } // namespace cl

sycl/test/fpga_tests/fpga_pipes.cpp

@@ -1,4 +1,6 @@
 // RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 //==------------- fpga_pipes.cpp - SYCL FPGA pipes test --------------------==//
 //
@@ -292,6 +294,13 @@ int test_array_th_bl_pipe(cl::sycl::queue Queue) {
 int main() {
   cl::sycl::queue Queue;
 
+  if (!Queue.get_device()
+           .get_info<cl::sycl::info::device::kernel_kernel_pipe_support>()) {
+    std::cout << "SYCL_INTEL_data_flow_pipes not supported, skipping"
+              << std::endl;
+    return 0;
+  }
+
   // Non-blocking pipes
   int Result = test_simple_nb_pipe<some_nb_pipe, /*test number*/ 1>(Queue);
   Result &= test_simple_nb_pipe<some::nb_pipe, /*test number*/ 2>(Queue);

sycl/test/fpga_tests/pipes_info.cpp

@@ -0,0 +1,36 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+//==--------- pipes_info.cpp - SYCL device pipe info test --*- C++ -*-------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl.hpp>
+
+int main() {
+  cl::sycl::queue Queue;
+  cl::sycl::device Device = Queue.get_device();
+  cl::sycl::platform Platform = Device.get_platform();
+
+  // Query if the device supports kernel to kernel pipe feature
+  bool IsSupported =
+    Device.get_info<cl::sycl::info::device::kernel_kernel_pipe_support>();
+
+  // Query for platform string. We expect only Intel FPGA platforms to support
+  // SYCL_INTEL_data_flow_pipes extension.
+  std::string platform_name =
+    Platform.get_info<cl::sycl::info::platform::name>();
+  bool SupposedToBeSupported =
+    (platform_name == "Intel(R) FPGA Emulation Platform for OpenCL(TM)" ||
+     platform_name == "Intel(R) FPGA SDK for OpenCL(TM)")
+        ? true
+        : false;
+
+  return (SupposedToBeSupported != IsSupported);
+}