[SYCL] Move free function queries to experimental namespace

sycl/doc/extensions/README.md

@@ -43,6 +43,7 @@ DPC++ extensions status:
 | [Uniform](Uniform/Uniform.asciidoc)                                                                                         | Proposal                                  | |
 | [Assert](Assert/SYCL_ONEAPI_ASSERT.asciidoc) | Proposal | |
 | [Matrix](Matrix/dpcpp-joint-matrix.asciidoc)                                                                        | Partially supported(AMX AOT)               | Not supported: dynamic-extent, wg and wi scopes, layouts other than packed|
+| [SYCL_INTEL_free_function_queries](FreeFunctionQueries/SYCL_INTEL_free_function_queries.asciidoc)                           | Supported (experimental)                  | |
 
 Legend:
 

sycl/include/CL/sycl/group.hpp

@@ -428,12 +428,28 @@ template <int Dims> group<Dims> store_group(const group<Dims> *g) {
 }
 } // namespace detail
 
-template <int Dims> group<Dims> this_group() {
+template <int Dims>
+__SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_group() instead")
+group<Dims> this_group() {
 #ifdef __SYCL_DEVICE_ONLY__
   return detail::Builder::getElement(detail::declptr<group<Dims>>());
 #else
   return detail::store_group<Dims>(nullptr);
 #endif
 }
+
+namespace ext {
+namespace oneapi {
+namespace experimental {
+template <int Dims> group<Dims> this_group() {
+#ifdef __SYCL_DEVICE_ONLY__
+  return sycl::detail::Builder::getElement(detail::declptr<group<Dims>>());
+#else
+  return sycl::detail::store_group<Dims>(nullptr);
+#endif
+}
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/id.hpp

@@ -289,13 +289,28 @@ template <int Dims> id<Dims> store_id(const id<Dims> *i) {
 }
 } // namespace detail
 
-template <int Dims> id<Dims> this_id() {
+template <int Dims>
+__SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_id() instead")
+id<Dims> this_id() {
 #ifdef __SYCL_DEVICE_ONLY__
   return detail::Builder::getElement(detail::declptr<id<Dims>>());
 #else
   return detail::store_id<Dims>(nullptr);
 #endif
 }
 
+namespace ext {
+namespace oneapi {
+namespace experimental {
+template <int Dims> id<Dims> this_id() {
+#ifdef __SYCL_DEVICE_ONLY__
+  return sycl::detail::Builder::getElement(detail::declptr<id<Dims>>());
+#else
+  return sycl::detail::store_id<Dims>(nullptr);
+#endif
+}
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/item.hpp

@@ -132,13 +132,28 @@ template <int Dims> item<Dims> store_item(const item<Dims> *i) {
 }
 } // namespace detail
 
-template <int Dims> item<Dims> this_item() {
+template <int Dims>
+__SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_item() instead")
+item<Dims> this_item() {
 #ifdef __SYCL_DEVICE_ONLY__
   return detail::Builder::getElement(detail::declptr<item<Dims>>());
 #else
   return detail::store_item<Dims>(nullptr);
 #endif
 }
 
+namespace ext {
+namespace oneapi {
+namespace experimental {
+template <int Dims> item<Dims> this_item() {
+#ifdef __SYCL_DEVICE_ONLY__
+  return sycl::detail::Builder::getElement(detail::declptr<item<Dims>>());
+#else
+  return sycl::detail::store_item<Dims>(nullptr);
+#endif
+}
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/nd_item.hpp

@@ -202,13 +202,28 @@ template <int Dims> nd_item<Dims> store_nd_item(const nd_item<Dims> *nd_i) {
 }
 } // namespace detail
 
-template <int Dims> nd_item<Dims> this_nd_item() {
+template <int Dims>
+__SYCL_DEPRECATED("use sycl::ext::oneapi::experimental::this_nd_item() instead")
+nd_item<Dims> this_nd_item() {
 #ifdef __SYCL_DEVICE_ONLY__
   return detail::Builder::getElement(detail::declptr<nd_item<Dims>>());
 #else
   return detail::store_nd_item<Dims>(nullptr);
 #endif
 }
 
+namespace ext {
+namespace oneapi {
+namespace experimental {
+template <int Dims> nd_item<Dims> this_nd_item() {
+#ifdef __SYCL_DEVICE_ONLY__
+  return sycl::detail::Builder::getElement(detail::declptr<nd_item<Dims>>());
+#else
+  return sycl::detail::store_nd_item<Dims>(nullptr);
+#endif
+}
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/sub_group.hpp

@@ -16,5 +16,20 @@ namespace sycl {
 using ext::oneapi::sub_group;
 // TODO move the entire sub_group class implementation to this file once
 // breaking changes are allowed.
+
+namespace ext {
+namespace oneapi {
+namespace experimental {
+inline sub_group this_sub_group() {
+#ifdef __SYCL_DEVICE_ONLY__
+  return sub_group();
+#else
+  throw runtime_error("Sub-groups are not supported on host device.",
+                      PI_INVALID_DEVICE);
+#endif
+}
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/sycl/ext/oneapi/sub_group.hpp

@@ -100,6 +100,11 @@ void store(multi_ptr<T, Space> dst, const vec<T, N> &x) {
 namespace ext {
 namespace oneapi {
 
+struct sub_group;
+namespace experimental {
+inline sub_group this_sub_group();
+}
+
 struct sub_group {
 
   using id_type = id<1>;
@@ -733,9 +738,12 @@ struct sub_group {
 protected:
   template <int dimensions> friend class cl::sycl::nd_item;
   friend sub_group this_sub_group();
+  friend sub_group experimental::this_sub_group();
   sub_group() = default;
 };
 
+__SYCL_DEPRECATED(
+    "use sycl::ext::oneapi::experimental::this_sub_group() instead")
 inline sub_group this_sub_group() {
 #ifdef __SYCL_DEVICE_ONLY__
   return sub_group();

sycl/test/basic_tests/free_function_queries/free_function_queries_interface.cpp

@@ -13,26 +13,30 @@
 #include <type_traits>
 
 template <int Dims> struct this_id_caller {
-  auto operator()() const -> decltype(sycl::this_id<Dims>()) {
-    return sycl::this_id<Dims>();
+  auto operator()() const
+      -> decltype(sycl::ext::oneapi::experimental::this_id<Dims>()) {
+    return sycl::ext::oneapi::experimental::this_id<Dims>();
   }
 };
 
 template <int Dims> struct this_item_caller {
-  auto operator()() const -> decltype(sycl::this_item<Dims>()) {
-    return sycl::this_item<Dims>();
+  auto operator()() const
+      -> decltype(sycl::ext::oneapi::experimental::this_item<Dims>()) {
+    return sycl::ext::oneapi::experimental::this_item<Dims>();
   }
 };
 
 template <int Dims> struct this_nd_item_caller {
-  auto operator()() const -> decltype(sycl::this_nd_item<Dims>()) {
-    return sycl::this_nd_item<Dims>();
+  auto operator()() const
+      -> decltype(sycl::ext::oneapi::experimental::this_nd_item<Dims>()) {
+    return sycl::ext::oneapi::experimental::this_nd_item<Dims>();
   }
 };
 
 template <int Dims> struct this_group_caller {
-  auto operator()() const -> decltype(sycl::this_group<Dims>()) {
-    return sycl::this_group<Dims>();
+  auto operator()() const
+      -> decltype(sycl::ext::oneapi::experimental::this_group<Dims>()) {
+    return sycl::ext::oneapi::experimental::this_group<Dims>();
   }
 };
 
@@ -66,7 +70,8 @@ int main() {
   test<sycl::group>(this_group_caller<2>{});
   test<sycl::group>(this_group_caller<3>{});
 
-  static_assert(std::is_same<decltype(sycl::ext::oneapi::this_sub_group()),
-                             sycl::ext::oneapi::sub_group>::value,
-                "Wrong return type of free function query for Sub Group");
+  static_assert(
+      std::is_same<decltype(sycl::ext::oneapi::experimental::this_sub_group()),
+                   sycl::ext::oneapi::sub_group>::value,
+      "Wrong return type of free function query for Sub Group");
 }

sycl/test/warnings/free_functions_deprecation.cpp

@@ -0,0 +1,18 @@
+// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s -o %t.out
+
+#include <CL/sycl.hpp>
+
+int main() {
+  // expected-warning@+1 {{'this_group<1>' is deprecated: use sycl::ext::oneapi::experimental::this_group() instead}}
+  (void)sycl::this_group<1>();
+  // expected-warning@+1 {{'this_item<1>' is deprecated: use sycl::ext::oneapi::experimental::this_item() instead}}
+  (void)sycl::this_item<1>();
+  // expected-warning@+1 {{'this_nd_item<1>' is deprecated: use sycl::ext::oneapi::experimental::this_nd_item() instead}}
+  (void)sycl::this_nd_item<1>();
+  // expected-warning@+1 {{'this_id<1>' is deprecated: use sycl::ext::oneapi::experimental::this_id() instead}}
+  (void)sycl::this_id<1>();
+  // expected-warning@+1 {{'this_sub_group' is deprecated: use sycl::ext::oneapi::experimental::this_sub_group() instead}}
+  (void)sycl::ext::oneapi::this_sub_group();
+
+  return 0;
+}