[SYCL] Implement sycl_ext_oneapi_if_device

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

@@ -0,0 +1,50 @@
+//==------ if_device.hpp --- SYCL ext header file -----------=--*- 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
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/detail/common.hpp>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace ext::oneapi::experimental {
+namespace detail {
+// Helper object used to implement "otherwise".  The "MakeCall" template
+// parameter tells whether the previous call to "if_device" or "if_host" called
+// its "fn".  When "MakeCall" is true, the previous call to "fn" did not
+// happen, so the "otherwise" should call "fn".
+template <bool MakeCall> class if_device_or_host_helper {
+public:
+  template <typename T> void otherwise(T fn) {
+    if constexpr (MakeCall) {
+      fn();
+    }
+  }
+};
+} // namespace detail
+
+template <typename T> static auto if_device(T fn) {
+#ifdef __SYCL_DEVICE_ONLY__
+  fn();
+  return detail::if_device_or_host_helper<false>{};
+#else
+  return detail::if_device_or_host_helper<true>{};
+#endif
+}
+
+template <typename T> static auto if_host(T fn) {
+#ifdef __SYCL_DEVICE_ONLY__
+  return detail::if_device_or_host_helper<true>{};
+#else
+  fn();
+  return detail::if_device_or_host_helper<false>{};
+#endif
+}
+} // namespace ext::oneapi::experimental
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl

sycl/include/sycl/sycl.hpp

@@ -75,6 +75,7 @@
 #include <sycl/ext/oneapi/experimental/builtins.hpp>
 #include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
 #include <sycl/ext/oneapi/experimental/fixed_size_group.hpp>
+#include <sycl/ext/oneapi/experimental/if_device.hpp>
 #include <sycl/ext/oneapi/experimental/opportunistic_group.hpp>
 #include <sycl/ext/oneapi/experimental/tangle_group.hpp>
 #include <sycl/ext/oneapi/filter_selector.hpp>

sycl/test-e2e/Basic/if_device.cpp

@@ -0,0 +1,45 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/sycl.hpp>
+
+namespace syclex = sycl::ext::oneapi::experimental;
+
+int main() {
+  sycl::queue q;
+  sycl::buffer<int, 1> b(1);
+
+  {
+    sycl::host_accessor a{b};
+    a[0] = 1;
+    syclex::if_device([&]() { a[0] = 2; });
+    assert(a[0] == 1);
+    syclex::if_device([&]() { a[0] = 2; }).otherwise([&]() { a[0] = 3; });
+    assert(a[0] == 3);
+    syclex::if_host([&]() { a[0] = 2; });
+    assert(a[0] == 2);
+    syclex::if_host([&]() { a[0] = 1; }).otherwise([&]() { a[0] = 3; });
+    assert(a[0] == 1);
+  }
+  auto Do = [&](auto Fn) {
+    q.submit([&](sycl::handler &cgh) {
+      sycl::accessor a{b, cgh};
+      cgh.single_task([=]() { Fn(a); });
+    });
+  };
+
+  Do([&](auto a) { syclex::if_device([&]() { a[0] = 2; }); });
+  assert(sycl::host_accessor{b}[0] == 2);
+  Do([&](auto a) {
+    syclex::if_device([&]() { a[0] = 3; }).otherwise([&]() { a[0] = 1; });
+  });
+  assert(sycl::host_accessor{b}[0] == 3);
+  Do([&](auto a) { syclex::if_host([&]() { a[0] = 2; }); });
+  assert(sycl::host_accessor{b}[0] == 3);
+  Do([&](auto a) {
+    syclex::if_host([&]() { a[0] = 2; }).otherwise([&]() { a[0] = 1; });
+  });
+  assert(sycl::host_accessor{b}[0] == 1);
+
+  return 0;
+}