[SYCL] Update sycl_ext_oneapi_device_architecture device headers to support JIT (#13434)

This patch updates device headers of sycl_ext_oneapi_device_architecture
extension to prepare to support API in JIT. New functionality is always
disabled now as it breaks the existing functionality, and will be
enabled when the feature is fully implemented by updating other
components.

Author: dm-vodopyanov

sycl/include/sycl/ext/oneapi/experimental/device_architecture.hpp

@@ -10,6 +10,7 @@
 
 #include <cstdint> // for uint64_t
 #include <optional>
+#include <utility> // for std::integer_sequence
 
 namespace sycl {
 inline namespace _V1 {
@@ -1009,6 +1010,163 @@ template <bool MakeCall> class if_architecture_helper {
 
 namespace ext::oneapi::experimental {
 
+namespace detail {
+// Call the callable object "fn" only when this code runs on a device which
+// has a certain set of aspects or a particular architecture.
+//
+// Condition is a parameter pack of int's that define a simple expression
+// language which tells the set of aspects or architectures that the device
+// must have in order to enable the call.  See the "Condition*" values below.
+template <typename T, typename... Condition>
+#ifdef __SYCL_DEVICE_ONLY__
+[[__sycl_detail__::add_ir_attributes_function(
+    "sycl-call-if-on-device-conditionally", true)]]
+#endif
+void call_if_on_device_conditionally(T fn, Condition...) {
+  fn();
+}
+
+// The "Condition" parameter pack above is a sequence of int's that define an
+// expression tree.  Each node represents a boolean subexpression:
+//
+// ConditionAspect -       Next int is a value from "enum aspect".  The
+//                           subexpression is true if the device has this
+//                           aspect.
+// ConditionArchitecture - Next int is a value from "enum architecture".  The
+//                           subexpression is true if the device has this
+//                           architecture.
+// ConditionNot -          Next int is the root of another subexpression S1.
+//                           This subexpression is true if S1 is false.
+// ConditionAnd -          Next int is the root of another subexpression S1.
+//                           The int following that subexpression is the root
+//                           of another subexpression S2.  This subexpression
+//                           is true if both S1 and S2 are true.
+// ConditionOr -           Next int is the root of another subexpression S1.
+//                           The int following that subexpression is the root
+//                           of another subexpression S2.  This subexpression
+//                           is true if either S1 or S2 are true.
+//
+// These values are stored in the application's executable, so they are
+// effectively part of the ABI.  Therefore, any change to an existing value
+// is an ABI break.
+//
+// There is no programmatic reason for the values to be negative.  They are
+// negative only by convention to make it easier for humans to distinguish them
+// from aspect or architecture values (which are positive).
+static constexpr int ConditionAspect = -1;
+static constexpr int ConditionArchitecture = -2;
+static constexpr int ConditionNot = -3;
+static constexpr int ConditionAnd = -4;
+static constexpr int ConditionOr = -5;
+
+// Metaprogramming helper to construct a ConditionOr expression for a sequence
+// of architectures.  "ConditionAnyArchitectureBuilder<Archs...>::seq" is an
+// "std::integer_sequence" representing the expression.
+template <architecture... Archs> struct ConditionAnyArchitectureBuilder;
+
+template <architecture Arch, architecture... Archs>
+struct ConditionAnyArchitectureBuilder<Arch, Archs...> {
+  template <int I1, int I2, int I3, int... Is>
+  static auto append(std::integer_sequence<int, Is...>) {
+    return std::integer_sequence<int, I1, I2, I3, Is...>{};
+  }
+  using rest = typename ConditionAnyArchitectureBuilder<Archs...>::seq;
+  static constexpr int arch = static_cast<int>(Arch);
+  using seq =
+      decltype(append<ConditionOr, ConditionArchitecture, arch>(rest{}));
+};
+
+template <architecture Arch> struct ConditionAnyArchitectureBuilder<Arch> {
+  static constexpr int arch = static_cast<int>(Arch);
+  using seq = std::integer_sequence<int, ConditionArchitecture, arch>;
+};
+
+// Metaprogramming helper to construct a ConditionNot expression.
+// ConditionNotBuilder<Exp>::seq" is an "std::integer_sequence" representing
+// the expression.
+template <typename Exp> struct ConditionNotBuilder {
+  template <int I, int... Is>
+  static auto append(std::integer_sequence<int, Is...>) {
+    return std::integer_sequence<int, I, Is...>{};
+  }
+  using rest = typename Exp::seq;
+  using seq = decltype(append<ConditionNot>(rest{}));
+};
+
+// Metaprogramming helper to construct a ConditionAnd expression.
+// "ConditionAndBuilder<Exp1, Exp2>::seq" is an "std::integer_sequence"
+// representing the expression.
+template <typename Exp1, typename Exp2> struct ConditionAndBuilder {
+  template <int I, int... I1s, int... I2s>
+  static auto append(std::integer_sequence<int, I1s...>,
+                     std::integer_sequence<int, I2s...>) {
+    return std::integer_sequence<int, I, I1s..., I2s...>{};
+  }
+  using rest1 = typename Exp1::seq;
+  using rest2 = typename Exp2::seq;
+  using seq = decltype(append<ConditionAnd>(rest1{}, rest2{}));
+};
+
+// Metaprogramming helper to construct a ConditionOr expression.
+// "ConditionOrBuilder<Exp1, Exp2>::seq" is an "std::integer_sequence"
+// representing the expression.
+template <typename Exp1, typename Exp2> struct ConditionOrBuilder {
+  template <int I, int... I1s, int... I2s>
+  static auto append(std::integer_sequence<int, I1s...>,
+                     std::integer_sequence<int, I2s...>) {
+    return std::integer_sequence<int, I, I1s..., I2s...>{};
+  }
+  using rest1 = typename Exp1::seq;
+  using rest2 = typename Exp2::seq;
+  using seq = decltype(append<ConditionOr>(rest1{}, rest2{}));
+};
+
+// Helper function to call call_if_on_device_conditionally() while converting
+// the "std::integer_sequence" for a condition expression into individual
+// arguments of type int.
+template <typename T, int... Is>
+void call_if_on_device_conditionally_helper(T fn,
+                                            std::integer_sequence<int, Is...>) {
+  call_if_on_device_conditionally(fn, Is...);
+}
+
+// Same sort of helper object for "else_if_architecture_is".
+template <typename MakeCallIf> class if_architecture_is_helper {
+public:
+  template <architecture... Archs, typename T,
+            typename = std::enable_if<std::is_invocable_v<T>>>
+  auto else_if_architecture_is(T fn) {
+    using make_call_if =
+        ConditionAndBuilder<MakeCallIf,
+                            ConditionAnyArchitectureBuilder<Archs...>>;
+    using make_else_call_if = ConditionAndBuilder<
+        MakeCallIf,
+        ConditionNotBuilder<ConditionAnyArchitectureBuilder<Archs...>>>;
+
+    using cond = typename make_call_if::seq;
+    call_if_on_device_conditionally_helper(fn, cond{});
+    return if_architecture_is_helper<make_else_call_if>{};
+  }
+
+  template <typename T> void otherwise(T fn) {
+    using cond = typename MakeCallIf::seq;
+    call_if_on_device_conditionally_helper(fn, cond{});
+  }
+};
+
+} // namespace detail
+
+#ifdef SYCL_EXT_ONEAPI_DEVICE_ARCHITECTURE_NEW_DESIGN_IMPL
+template <architecture... Archs, typename T>
+static auto if_architecture_is(T fn) {
+  using make_call_if = detail::ConditionAnyArchitectureBuilder<Archs...>;
+  using make_else_call_if = detail::ConditionNotBuilder<make_call_if>;
+
+  using cond = typename make_call_if::seq;
+  detail::call_if_on_device_conditionally_helper(fn, cond{});
+  return detail::if_architecture_is_helper<make_else_call_if>{};
+}
+#else
 /// The condition is `true` only if the device which executes the
 /// `if_architecture_is` function has any one of the architectures listed in the
 /// @tparam Archs pack.
@@ -1026,6 +1184,7 @@ constexpr static auto if_architecture_is(T fn) {
     return sycl::detail::if_architecture_helper<true>{};
   }
 }
+#endif // SYCL_EXT_ONEAPI_DEVICE_ARCHITECTURE_NEW_DESIGN_IMPL
 
 /// The condition is `true` only if the device which executes the
 /// `if_architecture_is` function has an architecture that is in any one of the

sycl/test/extensions/device_architecture/device_architecture.cpp

@@ -0,0 +1,35 @@
+// The goal of this test is to check that new design of
+// sycl_ext_oneapi_device_architecture extension can be compiled successfullly.
+// During binary run there are some errors, this is expected, so there is no run
+// line yet for this test.
+
+// RUN: %clangxx -fsycl -DSYCL_EXT_ONEAPI_DEVICE_ARCHITECTURE_NEW_DESIGN_IMPL %s -o %t.out
+
+#include <sycl/ext/oneapi/experimental/device_architecture.hpp>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi::experimental;
+
+int main() {
+  std::vector<int> vec(4);
+  {
+    buffer<int> buf(vec.data(), vec.size());
+
+    queue q(gpu_selector_v);
+
+    // test if_architecture_is
+    q.submit([&](handler &cgh) {
+      auto acc = buf.get_access<access::mode::read_write>(cgh);
+      cgh.single_task([=]() {
+        if_architecture_is<architecture::intel_gpu_pvc>([&]() {
+          acc[0] = 2;
+        }).otherwise([&]() { acc[0] = 1; });
+      });
+    });
+  }
+
+  assert(vec[0] == 1);
+
+  return 0;
+}