From 4006375830642e1e1ab1c4fa705d1e69149e3b02 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Fri, 30 Jun 2023 12:41:48 -0700
Subject: [PATCH 1/2] [SYCL] Implement sycl_ext_oneapi_if_device

Specification PR is at https://github.com/intel/llvm/pull/8917.
---
 .../ext/oneapi/experimental/if_device.hpp     | 42 +++++++++++++++++
 sycl/include/sycl/sycl.hpp                    |  1 +
 sycl/test-e2e/Basic/if_device.cpp             | 45 +++++++++++++++++++
 3 files changed, 88 insertions(+)
 create mode 100644 sycl/include/sycl/ext/oneapi/experimental/if_device.hpp
 create mode 100644 sycl/test-e2e/Basic/if_device.cpp

diff --git a/sycl/include/sycl/ext/oneapi/experimental/if_device.hpp b/sycl/include/sycl/ext/oneapi/experimental/if_device.hpp
new file mode 100644
index 000000000000..77001af5b5da
--- /dev/null
+++ b/sycl/include/sycl/ext/oneapi/experimental/if_device.hpp
@@ -0,0 +1,42 @@
+#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
diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp
index f972b95908f8..932bb92f8e97 100644
--- a/sycl/include/sycl/sycl.hpp
+++ b/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>
diff --git a/sycl/test-e2e/Basic/if_device.cpp b/sycl/test-e2e/Basic/if_device.cpp
new file mode 100644
index 000000000000..c6d98daea337
--- /dev/null
+++ b/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;
+}

From 6050fe4739c3a097942cf3a9087707cd304f89e1 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Wed, 5 Jul 2023 08:43:02 -0700
Subject: [PATCH 2/2] Add licence header

---
 sycl/include/sycl/ext/oneapi/experimental/if_device.hpp | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/sycl/include/sycl/ext/oneapi/experimental/if_device.hpp b/sycl/include/sycl/ext/oneapi/experimental/if_device.hpp
index 77001af5b5da..0143ae4c9414 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/if_device.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/if_device.hpp
@@ -1,3 +1,11 @@
+//==------ 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>