[SYCL] Deprecate ext::oneapi::sub_group (#10263)

sub_group extension is deprecated in favor of standard SYCL2020
sub_group.

Author: KornevNikita

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -50,7 +50,7 @@ static const char *LegalSYCLFunctions[] = {
     "^sycl::_V1::multi_ptr<.+>::.+",
     "^sycl::_V1::nd_item<.+>::.+",
     "^sycl::_V1::group<.+>::.+",
-    "^sycl::_V1::sub_group<.+>::.+",
+    "^sycl::_V1::sub_group::.+",
     "^sycl::_V1::range<.+>::.+",
     "^sycl::_V1::kernel_handler::.+",
     "^sycl::_V1::cos",

sycl/include/sycl/detail/spirv.hpp

@@ -22,6 +22,7 @@
 #ifdef __SYCL_DEVICE_ONLY__
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
+struct sub_group;
 namespace ext {
 namespace oneapi {
 struct sub_group;
@@ -63,6 +64,9 @@ template <int Dimensions> struct group_scope<group<Dimensions>> {
 template <> struct group_scope<::sycl::ext::oneapi::sub_group> {
   static constexpr __spv::Scope::Flag value = __spv::Scope::Flag::Subgroup;
 };
+template <> struct group_scope<::sycl::sub_group> {
+  static constexpr __spv::Scope::Flag value = __spv::Scope::Flag::Subgroup;
+};
 
 template <typename ParentGroup>
 struct group_scope<sycl::ext::oneapi::experimental::ballot_group<ParentGroup>> {
@@ -253,6 +257,9 @@ template <typename Group> struct GroupId {
 template <> struct GroupId<::sycl::ext::oneapi::sub_group> {
   using type = uint32_t;
 };
+template <> struct GroupId<::sycl::sub_group> {
+  using type = uint32_t;
+};
 template <typename Group, typename T, typename IdT>
 EnableIfNativeBroadcast<T, IdT> GroupBroadcast(Group, T x, IdT local_id) {
   using GroupIdT = typename GroupId<Group>::type;
@@ -341,7 +348,7 @@ GroupBroadcast(const ext::oneapi::experimental::opportunistic_group &g, T x,
   auto LocalId = detail::IdToMaskPosition(g, local_id);
 
   // TODO: Refactor to avoid duplication after design settles.
-  using GroupIdT = typename GroupId<sycl::ext::oneapi::sub_group>::type;
+  using GroupIdT = typename GroupId<::sycl::sub_group>::type;
   GroupIdT GroupLocalId = static_cast<GroupIdT>(LocalId);
   using OCLT = detail::ConvertToOpenCLType_t<T>;
   using WidenedT = WidenOpenCLTypeTo32_t<OCLT>;

sycl/include/sycl/detail/type_traits.hpp

@@ -28,6 +28,7 @@ inline constexpr bool is_fixed_size_group_v = is_fixed_size_group<T>::value;
 } // namespace detail
 
 template <int Dimensions> class group;
+struct sub_group;
 namespace ext::oneapi {
 struct sub_group;
 
@@ -50,6 +51,7 @@ struct is_fixed_topology_group<sycl::group<Dimensions>> : std::true_type {};
 template <>
 struct is_fixed_topology_group<sycl::ext::oneapi::sub_group> : std::true_type {
 };
+template <> struct is_fixed_topology_group<sycl::sub_group> : std::true_type {};
 
 template <class T> struct is_user_constructed_group : std::false_type {};
 
@@ -77,6 +79,7 @@ struct is_group<group<Dimensions>> : std::true_type {};
 template <typename T> struct is_sub_group : std::false_type {};
 
 template <> struct is_sub_group<ext::oneapi::sub_group> : std::true_type {};
+template <> struct is_sub_group<sycl::sub_group> : std::true_type {};
 
 template <typename T>
 struct is_generic_group

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

@@ -128,7 +128,7 @@ inline opportunistic_group get_opportunistic_group() {
 #ifdef __SYCL_DEVICE_ONLY__
 #if defined(__SPIR__)
   // TODO: It may be wiser to call the intrinsic than rely on this_group()
-  sycl::sub_group sg = sycl::ext::oneapi::this_sub_group();
+  sycl::sub_group sg = sycl::ext::oneapi::experimental::this_sub_group();
   sub_group_mask mask = sycl::ext::oneapi::group_ballot(sg, true);
   return opportunistic_group(mask);
 #elif defined(__NVPTX__)

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

@@ -93,7 +93,7 @@ group<Dimensions> get_child_group(root_group<Dimensions> g) {
   return this_group<Dimensions>();
 }
 
-template <int Dimensions> sub_group get_child_group(group<Dimensions> g) {
+template <int Dimensions> sycl::sub_group get_child_group(group<Dimensions> g) {
   (void)g;
   return this_sub_group();
 }

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

@@ -22,19 +22,15 @@
 // Forward declarations of types not allowed to be wrapped in uniform:
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
-namespace ext::oneapi {
 
 struct sub_group;
-
-} // namespace ext::oneapi
-
 template <int, bool> class item;
 template <int> class id;
 template <int> class nd_item;
 template <int> class h_item;
 template <int> class group;
 template <int> class nd_range;
-using ext::oneapi::sub_group;
+using sycl::sub_group;
 
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl