[SYCL] Update sycl_ext_oneapi_device_architecture device headers to support JIT

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;
+}