[SYCL] Make core.hpp independent from sub_group.hpp (#16044)

This PR will also allow us to remove `sub_group.hpp` from `sycl.hpp` to
make the former an opt-in header (for better compile-times).

Author: AlexeySachkov

sycl/include/sycl/group_algorithm.hpp

@@ -21,7 +21,6 @@
 #include <sycl/known_identity.hpp>     // for known_identity_v
 #include <sycl/nd_item.hpp>            // for nd_item
 #include <sycl/range.hpp>              // for range
-#include <sycl/sub_group.hpp>          // for sub_group
 #include <sycl/types.hpp>              // for vec
 
 #ifdef __SYCL_DEVICE_ONLY__
@@ -37,6 +36,7 @@
 
 namespace sycl {
 inline namespace _V1 {
+struct sub_group;
 namespace detail {
 
 // ---- linear_id_to_id

sycl/include/sycl/nd_item.hpp

@@ -25,14 +25,14 @@
 #include <sycl/nd_range.hpp>  // for nd_range
 #include <sycl/pointers.hpp>  // for decorated_global_ptr, decor...
 #include <sycl/range.hpp>     // for range
-#include <sycl/sub_group.hpp> // for sub_group
 
 #include <cstddef>     // for size_t
 #include <stdint.h>    // for uint32_t
 #include <type_traits> // for enable_if_t, remove_const_t
 
 namespace sycl {
 inline namespace _V1 {
+struct sub_group;
 namespace detail {
 class Builder;
 }
@@ -117,7 +117,8 @@ template <int Dimensions = 1> class nd_item {
                                         get_group_range(), get_group_id());
   }
 
-  sub_group get_sub_group() const { return sub_group(); }
+  // Out-of-class definition in sub_group.hpp
+  sub_group get_sub_group() const;
 
   size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const {
     size_t Id = get_group_id()[Dimension];

sycl/include/sycl/stream.hpp

@@ -24,12 +24,12 @@
 #include <sycl/half_type.hpp>                 // for half, operator-, operator<
 #include <sycl/handler.hpp>                   // for handler
 #include <sycl/item.hpp>                      // for item
-#include <sycl/nd_item.hpp>                   // for nd_item
-#include <sycl/nd_range.hpp>                  // for nd_range
-#include <sycl/property_list.hpp>             // for property_list
-#include <sycl/range.hpp>                     // for range
-#include <sycl/sub_group.hpp>                 // for multi_ptr
-#include <sycl/types.hpp>                     // for vec, SwizzleOp
+#include <sycl/multi_ptr.hpp>
+#include <sycl/nd_item.hpp>       // for nd_item
+#include <sycl/nd_range.hpp>      // for nd_range
+#include <sycl/property_list.hpp> // for property_list
+#include <sycl/range.hpp>         // for range
+#include <sycl/types.hpp>         // for vec, SwizzleOp
 
 #include <cstddef>     // for size_t, byte
 #include <memory>      // for hash, shared_ptr

sycl/include/sycl/sub_group.hpp

@@ -16,7 +16,8 @@
 #include <sycl/id.hpp>                         // for id
 #include <sycl/memory_enums.hpp>               // for memory_scope
 #include <sycl/multi_ptr.hpp>                  // for multi_ptr
-#include <sycl/range.hpp>                      // for range
+#include <sycl/nd_item.hpp>
+#include <sycl/range.hpp> // for range
 
 #include <stdint.h>    // for uint32_t
 #include <tuple>       // for _Swallow_assign, ignore
@@ -663,5 +664,10 @@ struct sub_group {
   friend sub_group ext::oneapi::this_work_item::get_sub_group();
   sub_group() = default;
 };
+
+template <int Dimensions> sub_group nd_item<Dimensions>::get_sub_group() const {
+  return sub_group();
+}
+
 } // namespace _V1
 } // namespace sycl

sycl/test-e2e/AOT/reqd-sg-size.cpp

@@ -14,6 +14,7 @@
 #include <iostream>
 
 #include <sycl/detail/core.hpp>
+#include <sycl/sub_group.hpp>
 
 using namespace sycl;
 

sycl/test-e2e/Basic/linear-sub_group.cpp

@@ -14,6 +14,7 @@
 #include <cstdlib>
 #include <iostream>
 #include <sycl/detail/core.hpp>
+#include <sycl/sub_group.hpp>
 
 using namespace sycl;
 

sycl/test-e2e/Basic/sub_group_size_prop.cpp

@@ -2,6 +2,7 @@
 // RUN: %{run} %t.out
 
 #include <sycl/detail/core.hpp>
+#include <sycl/sub_group.hpp>
 
 #include <iostream>
 

sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp

@@ -1,6 +1,7 @@
 // This test is adapted from "test-e2e/Basic/sub_group_size_prop.cpp"
 
 #include "../graph_common.hpp"
+#include <sycl/sub_group.hpp>
 
 enum class Variant { Function, Functor, FunctorAndProperty };
 

sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/common.hpp

@@ -1,6 +1,7 @@
 
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/experimental/group_sort.hpp>
+#include <sycl/sub_group.hpp>
 
 #pragma once
 

sycl/test-e2e/GroupAlgorithm/SYCL2020/permute_select.hpp

@@ -10,6 +10,7 @@
 #include <iostream>
 #include <sycl/detail/core.hpp>
 #include <sycl/group_algorithm.hpp>
+#include <sycl/sub_group.hpp>
 template <typename T, int N> class sycl_subgr;
 
 using namespace sycl;

sycl/test-e2e/GroupAlgorithm/SYCL2020/shift_left_right.hpp

@@ -9,6 +9,7 @@
 #include "helpers.hpp"
 #include <sycl/detail/core.hpp>
 #include <sycl/group_algorithm.hpp>
+#include <sycl/sub_group.hpp>
 template <typename T, int N> class sycl_subgr;
 
 using namespace sycl;

sycl/test-e2e/GroupAlgorithm/different_types.cpp

@@ -8,6 +8,7 @@
 #include <numeric>
 #include <sycl/detail/core.hpp>
 #include <sycl/group_algorithm.hpp>
+#include <sycl/sub_group.hpp>
 #include <sycl/sycl_span.hpp>
 
 using namespace sycl;

sycl/test-e2e/GroupAlgorithm/load_store/partial_sg.cpp

@@ -3,6 +3,7 @@
 
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/experimental/group_load_store.hpp>
+#include <sycl/sub_group.hpp>
 
 #include <numeric>
 

sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp

@@ -6,6 +6,9 @@
 // UNSUPPORTED: arch-intel_gpu_pvc && !igc-dev
 
 #include "support.h"
+
+#include <sycl/sub_group.hpp>
+
 #include <algorithm>
 #include <cassert>
 #include <complex>

sycl/test-e2e/GroupAlgorithm/shuffle_marray.cpp

@@ -6,6 +6,7 @@
 #include <sycl/detail/core.hpp>
 #include <sycl/group_algorithm.hpp>
 #include <sycl/marray.hpp>
+#include <sycl/sub_group.hpp>
 
 static constexpr size_t NumElems = 5;
 static constexpr size_t NumWorkItems = 64;

sycl/test-e2e/PerformanceTests/Reduction/reduce_over_sub_group.cpp

@@ -3,6 +3,7 @@
 // RUN: %{run} %t.out
 
 #include <sycl/detail/core.hpp>
+#include <sycl/sub_group.hpp>
 
 #include <sycl/ext/oneapi/experimental/user_defined_reductions.hpp>
 

sycl/test-e2e/SubGroup/helper.hpp

@@ -10,6 +10,7 @@
 #include <complex>
 #include <iostream>
 #include <sycl/detail/core.hpp>
+#include <sycl/sub_group.hpp>
 
 using namespace sycl;
 

sycl/test-e2e/SubGroup/sub_group_as.cpp

@@ -9,6 +9,7 @@
 #include <cstdio>
 #include <cstdlib>
 #include <sycl/detail/core.hpp>
+#include <sycl/sub_group.hpp>
 
 int main(int argc, char *argv[]) {
   sycl::queue queue;

sycl/test-e2e/SubGroup/sub_group_by_value_semantics.cpp

@@ -2,6 +2,7 @@
 // RUN: %{run} %t.out
 
 #include <sycl/detail/core.hpp>
+#include <sycl/sub_group.hpp>
 
 int main() {
   bool result = true;

sycl/test-e2e/SubGroup/sub_groups_sycl2020.cpp

@@ -2,6 +2,7 @@
 // RUN: %{run} %t.out
 
 #include <sycl/detail/core.hpp>
+#include <sycl/sub_group.hpp>
 
 int main() {
   sycl::queue Q;

sycl/test-e2e/SubGroupMask/Basic.cpp

@@ -16,6 +16,7 @@
 
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/sub_group_mask.hpp>
+#include <sycl/sub_group.hpp>
 
 #include <iostream>
 using namespace sycl;

sycl/test-e2e/SubGroupMask/GroupSize.cpp

@@ -15,6 +15,7 @@
 
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/sub_group_mask.hpp>
+#include <sycl/sub_group.hpp>
 
 #include <iostream>
 

sycl/test-e2e/SubGroupMask/sub_group_mask_ver2.cpp

@@ -4,6 +4,7 @@
 #include <iostream>
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/sub_group_mask.hpp>
+#include <sycl/sub_group.hpp>
 
 #define TEST_ON_DEVICE(TEST_BODY)                                              \
   {                                                                            \

sycl/test-e2e/XPTI/kernel/content.cpp

@@ -14,6 +14,7 @@
 #include <numeric>
 #include <sycl/detail/core.hpp>
 #include <sycl/reduction.hpp>
+#include <sycl/sub_group.hpp>
 
 using namespace sycl;
 int main() {

sycl/test/include_deps/sycl_detail_core.hpp.cpp

@@ -128,7 +128,6 @@
 // CHECK-NEXT: kernel_handler.hpp
 // CHECK-NEXT: nd_item.hpp
 // CHECK-NEXT: nd_range.hpp
-// CHECK-NEXT: sub_group.hpp
 // CHECK-NEXT: device.hpp
 // CHECK-NEXT: kernel_bundle_enums.hpp
 // CHECK-NEXT: event.hpp