[SYCL] [FPGA] Create DSP control header (#5035)

Create header file that provides user API for FPGA DSP control feature.

Related links:

    SPIR-V spec: https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/INTEL/SPV_INTEL_fpga_dsp_control.asciidoc
    SPIR-V implementation: 

    [SPIRV] New FPGA function attribute for DSP control KhronosGroup/SPIRV-LLVM-Translator#1046

Test: intel/llvm-test-suite#591

Author: shuoniu-intel

sycl/include/sycl/ext/intel/fpga_dsp_control.hpp

@@ -0,0 +1,79 @@
+//==------------ fpga_dsp_control.hpp --- SYCL FPGA DSP Control ------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+#pragma once
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace ext {
+namespace intel {
+
+enum class Preference { DSP, Softlogic, Compiler_default };
+enum class Propagate { On, Off };
+
+template <typename Function>
+#ifdef __SYCL_DEVICE_ONLY__
+[[intel::prefer_dsp]]
+[[intel::propagate_dsp_preference]]
+#endif // __SYCL_DEVICE_ONLY__
+void math_prefer_dsp_propagate(Function f)
+{
+  f();
+}
+
+template <typename Function>
+#ifdef __SYCL_DEVICE_ONLY__
+[[intel::prefer_dsp]]
+#endif // __SYCL_DEVICE_ONLY__
+void math_prefer_dsp_no_propagate(Function f)
+{
+  f();
+}
+
+template <typename Function>
+#ifdef __SYCL_DEVICE_ONLY__
+[[intel::prefer_softlogic]]
+[[intel::propagate_dsp_preference]]
+#endif // __SYCL_DEVICE_ONLY__
+void math_prefer_softlogic_propagate(Function f)
+{
+  f();
+}
+
+template <typename Function>
+#ifdef __SYCL_DEVICE_ONLY__
+[[intel::prefer_softlogic]]
+#endif // __SYCL_DEVICE_ONLY__
+void math_prefer_softlogic_no_propagate(Function f)
+{
+  f();
+}
+
+template <Preference my_preference = Preference::DSP,
+          Propagate my_propagate = Propagate::On, typename Function>
+void math_dsp_control(Function f) {
+  if (my_preference == Preference::DSP) {
+    if (my_propagate == Propagate::On) {
+      math_prefer_dsp_propagate(f);
+    } else {
+      math_prefer_dsp_no_propagate(f);
+    }
+  } else if (my_preference == Preference::Softlogic) {
+    if (my_propagate == Propagate::On) {
+      math_prefer_softlogic_propagate(f);
+    } else {
+      math_prefer_softlogic_no_propagate(f);
+    }
+  } else { // my_preference == Preference::Compiler_default
+    math_prefer_dsp_no_propagate([&]() { f(); });
+  }
+}
+
+} // namespace intel
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/sycl/ext/intel/fpga_extensions.hpp

@@ -8,6 +8,7 @@
 
 #pragma once
 #include <sycl/ext/intel/fpga_device_selector.hpp>
+#include <sycl/ext/intel/fpga_dsp_control.hpp>
 #include <sycl/ext/intel/fpga_loop_fuse.hpp>
 #include <sycl/ext/intel/fpga_lsu.hpp>
 #include <sycl/ext/intel/fpga_reg.hpp>

sycl/test/extensions/fpga.cpp

@@ -128,5 +128,29 @@ int main() {
     }
   }
 
+  /*Check DSP control interface*/
+  cl::sycl::buffer<int, 1> output_buffer(1);
+  cl::sycl::buffer<int, 1> input_buffer(1);
+  Queue.submit([&](sycl::handler &cgh) {
+    auto output_accessor =
+        output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
+    auto input_accessor =
+        input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
+    cgh.single_task<class DSPControlKernel>([=]() {
+      float sum = input_accessor[0];
+      sycl::ext::intel::math_dsp_control<
+          sycl::ext::intel::Preference::Softlogic>([&] { sum += 1.23f; });
+      sycl::ext::intel::math_dsp_control<sycl::ext::intel::Preference::DSP>(
+          [&] { sum += 1.23f; });
+      sycl::ext::intel::math_dsp_control<
+          sycl::ext::intel::Preference::Softlogic,
+          sycl::ext::intel::Propagate::Off>([&] { sum += 4.56f; });
+      sycl::ext::intel::math_dsp_control<sycl::ext::intel::Preference::DSP,
+                                         sycl::ext::intel::Propagate::Off>(
+          [&] { sum += 4.56f; });
+      output_accessor[0] = sum;
+    });
+  });
+
   return 0;
 }