[SYCL][E2E] Remove subgroup supported checks from e2e tests (#14313)

Subgroups are core sycl functionality which should be tested on all
backends.

Author: ayylol

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

@@ -9,7 +9,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "../SubGroup/helper.hpp"
 #include <algorithm>
 #include <cstdio>
 #include <cstdlib>
@@ -20,10 +19,6 @@ using namespace sycl;
 
 int main(int argc, char *argv[]) {
   queue q;
-  if (!core_sg_supported(q.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
 
   // Fill output array with sub-group IDs
   const uint32_t outer = 2;

sycl/test-e2e/Regression/get_subgroup_sizes.cpp

@@ -1,3 +1,9 @@
+// UNSUPPORTED: accelerator
+// TODO: FPGAs currently report `sub_group_sizes` as non-empty list,
+// despite not having extension `cl_intel_required_subgroup_size`
+// UNSUPPORTED: cuda || hip
+// TODO: Similar issue to FPGAs
+
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
@@ -18,13 +24,15 @@ int main() {
   queue Q;
   auto Dev = Q.get_device();
   auto Vec = Dev.get_info<info::device::extensions>();
+  std::vector<size_t> SubGroupSizes =
+      Dev.get_info<sycl::info::device::sub_group_sizes>();
   if (std::find(Vec.begin(), Vec.end(), "cl_intel_required_subgroup_size") !=
       std::end(Vec)) {
-    std::vector<size_t> SubGroupSizes =
-        Dev.get_info<sycl::info::device::sub_group_sizes>();
-    std::vector<size_t>::const_iterator MaxIter =
-        std::max_element(SubGroupSizes.begin(), SubGroupSizes.end());
-    int MaxSubGroup_size = *MaxIter;
+    assert(!SubGroupSizes.empty() &&
+           "Required sub-group size list should not be empty");
+  } else {
+    assert(SubGroupSizes.empty() &&
+           "Required sub-group size list should be empty");
   }
   return 0;
 }

sycl/test-e2e/SubGroup/attributes.cpp

@@ -1,3 +1,10 @@
+// UNSUPPORTED: accelerator
+// TODO: FPGAs currently report supported subgroups as {4,8,16,32,64}, causing
+// this test to fail
+// UNSUPPORTED: cuda || hip
+// TODO: Device subgroup sizes reports {32}, but when we try to use it with a
+// kernel attribute and check it, we get a subgroup size of 0.
+
 // RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out
 // RUN: %{run} %t.out
 //==------- attributes.cpp - SYCL sub_group attributes test ----*- C++ -*---==//
@@ -13,7 +20,7 @@
 #define KERNEL_FUNCTOR_WITH_SIZE(SIZE)                                         \
   class KernelFunctor##SIZE {                                                  \
   public:                                                                      \
-    [[intel::reqd_sub_group_size(SIZE)]] void                                  \
+    [[sycl::reqd_sub_group_size(SIZE)]] void                                   \
     operator()(sycl::nd_item<1> Item) const {                                  \
       const auto GID = Item.get_global_id();                                   \
     }                                                                          \
@@ -49,19 +56,6 @@ int main() {
   queue Queue;
   device Device = Queue.get_device();
 
-  // According to specification, this kernel query requires `cl_khr_subgroups`
-  // or `cl_intel_subgroups`, and also `cl_intel_required_subgroup_size`
-  auto Vec = Device.get_info<info::device::extensions>();
-  if (std::find(Vec.begin(), Vec.end(), "cl_intel_subgroups") ==
-              std::end(Vec) &&
-          std::find(Vec.begin(), Vec.end(), "cl_khr_subgroups") ==
-              std::end(Vec) ||
-      std::find(Vec.begin(), Vec.end(), "cl_intel_required_subgroup_size") ==
-          std::end(Vec)) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
-
   try {
     const auto SGSizes = Device.get_info<info::device::sub_group_sizes>();
 

sycl/test-e2e/SubGroup/helper.hpp

@@ -164,24 +164,3 @@ void exit_if_not_equal_vec(vec<T, N> val, vec<T, N> ref, const char *name) {
     exit(1);
   }
 }
-
-bool core_sg_supported(const device &Device) {
-  auto Vec = Device.get_info<info::device::extensions>();
-  if (std::find(Vec.begin(), Vec.end(), "cl_khr_subgroups") != std::end(Vec))
-    return true;
-
-  if (std::find(Vec.begin(), Vec.end(), "cl_intel_subgroups") != std::end(Vec))
-    return true;
-
-  if (Device.get_backend() == sycl::backend::opencl) {
-    // Extract the numerical version from the version string, OpenCL version
-    // string have the format "OpenCL <major>.<minor> <vendor specific data>".
-    std::string ver = Device.get_info<info::device::version>().substr(7, 3);
-
-    // cl_khr_subgroups was core in OpenCL 2.1 and 2.2, but went back to
-    // optional in 3.0
-    return ver >= "2.1" && ver < "3.0";
-  }
-
-  return false;
-}

sycl/test-e2e/SubGroup/info.cpp

@@ -1,3 +1,7 @@
+// UNSUPPORTED: accelerator
+// TODO: FPGAs currently report supported subgroups as {4,8,16,32,64}, causing
+// this test to fail. Additionally, the kernel max_sub_group_size checks
+// crash on FPGAs
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
@@ -17,14 +21,20 @@ int main() {
   queue Queue;
   device Device = Queue.get_device();
 
-  /* Basic sub-group functionality is supported as part of cl_khr_subgroups
-   * extension or as core OpenCL 2.1 feature. */
-  if (!core_sg_supported(Device)) {
-    std::cout << "Skipping test\n";
-    return 0;
+  bool old_opencl = false;
+  if (Device.get_backend() == sycl::backend::opencl) {
+    // Extract the numerical version from the version string, OpenCL version
+    // string have the format "OpenCL <major>.<minor> <vendor specific data>".
+    std::string ver = Device.get_info<info::device::version>().substr(7, 3);
+    old_opencl = (ver < "2.1");
   }
+
   /* Check info::device parameters. */
-  Device.get_info<info::device::sub_group_independent_forward_progress>();
+  if (!old_opencl) {
+    // Independent forward progress is missing on OpenCL backend prior to
+    // version 2.1
+    Device.get_info<info::device::sub_group_independent_forward_progress>();
+  }
   Device.get_info<info::device::max_num_sub_groups>();
 
   try {
@@ -49,30 +59,24 @@ int main() {
     });
     uint32_t Res = 0;
 
-    /* sub_group_sizes can be queried only if cl_intel_required_subgroup_size
-     * extension is supported by device*/
-    auto Vec = Device.get_info<info::device::extensions>();
-    if (std::find(Vec.begin(), Vec.end(), "cl_intel_required_subgroup_size") !=
-        std::end(Vec)) {
-      auto sg_sizes = Device.get_info<info::device::sub_group_sizes>();
+    auto sg_sizes = Device.get_info<info::device::sub_group_sizes>();
+
+    // Max sub-group size for a particular kernel might not be the max
+    // supported size on the device in general. Can only check that it is
+    // contained in list of valid sizes.
+    Res = Kernel.get_info<info::kernel_device_specific::max_sub_group_size>(
+        Device);
+    bool Expected =
+        std::find(sg_sizes.begin(), sg_sizes.end(), Res) != sg_sizes.end();
+    exit_if_not_equal<bool>(Expected, true, "max_sub_group_size");
 
-      // Max sub-group size for a particular kernel might not be the max
-      // supported size on the device in general. Can only check that it is
-      // contained in list of valid sizes.
+    for (auto r : {range<3>(3, 4, 5), range<3>(1, 1, 1), range<3>(4, 2, 1),
+                   range<3>(32, 3, 4), range<3>(7, 9, 11)}) {
       Res = Kernel.get_info<info::kernel_device_specific::max_sub_group_size>(
           Device);
-      bool Expected =
+      Expected =
           std::find(sg_sizes.begin(), sg_sizes.end(), Res) != sg_sizes.end();
       exit_if_not_equal<bool>(Expected, true, "max_sub_group_size");
-
-      for (auto r : {range<3>(3, 4, 5), range<3>(1, 1, 1), range<3>(4, 2, 1),
-                     range<3>(32, 3, 4), range<3>(7, 9, 11)}) {
-        Res = Kernel.get_info<info::kernel_device_specific::max_sub_group_size>(
-            Device);
-        Expected =
-            std::find(sg_sizes.begin(), sg_sizes.end(), Res) != sg_sizes.end();
-        exit_if_not_equal<bool>(Expected, true, "max_sub_group_size");
-      }
     }
 
     Res = Kernel.get_info<info::kernel_device_specific::compile_num_sub_groups>(
@@ -81,21 +85,11 @@ int main() {
     /* Sub-group size is not specified in kernel or IL*/
     exit_if_not_equal<uint32_t>(Res, 0, "compile_num_sub_groups");
 
-    // According to specification, this kernel query requires `cl_khr_subgroups`
-    // or `cl_intel_subgroups`
-    if ((std::find(Vec.begin(), Vec.end(), "cl_khr_subgroups") !=
-         std::end(Vec)) ||
-        std::find(Vec.begin(), Vec.end(), "cl_intel_subgroups") !=
-                std::end(Vec) &&
-            std::find(Vec.begin(), Vec.end(),
-                      "cl_intel_required_subgroup_size") != std::end(Vec)) {
-      Res =
-          Kernel.get_info<info::kernel_device_specific::compile_sub_group_size>(
-              Device);
-
-      /* Required sub-group size is not specified in kernel or IL*/
-      exit_if_not_equal<uint32_t>(Res, 0, "compile_sub_group_size");
-    }
+    Res = Kernel.get_info<info::kernel_device_specific::compile_sub_group_size>(
+        Device);
+
+    /* Required sub-group size is not specified in kernel or IL*/
+    exit_if_not_equal<uint32_t>(Res, 0, "compile_sub_group_size");
 
   } catch (exception e) {
     std::cout << "SYCL exception caught: " << e.what();

sycl/test-e2e/SubGroup/reduce.cpp

@@ -13,10 +13,6 @@
 #include <iostream>
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check<class KernelName_AJprOaCZgUmsYFRTTGNw, int>(Queue);
   check<class KernelName_ShKFIYTqaI, unsigned int>(Queue);
   check<class KernelName_TovsKTk, long>(Queue);

sycl/test-e2e/SubGroup/reduce_fp16.cpp

@@ -10,10 +10,6 @@
 
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check<class KernelName_oMg, sycl::half>(Queue);
   std::cout << "Test passed." << std::endl;
   return 0;

sycl/test-e2e/SubGroup/reduce_fp64.cpp

@@ -8,10 +8,6 @@
 
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check<class KernelName_alTnImqzYasRyHjYg, double>(Queue);
   std::cout << "Test passed." << std::endl;
   return 0;

sycl/test-e2e/SubGroup/reduce_spirv13.cpp

@@ -8,10 +8,6 @@
 #include <iostream>
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
 
   check_mul<class MulA, int>(Queue);
   check_mul<class MulB, unsigned int>(Queue);

sycl/test-e2e/SubGroup/reduce_spirv13_fp16.cpp

@@ -11,10 +11,6 @@
 
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check_mul<class MulHalf, sycl::half>(Queue);
   std::cout << "Test passed." << std::endl;
   return 0;

sycl/test-e2e/SubGroup/reduce_spirv13_fp64.cpp

@@ -10,10 +10,6 @@
 #include <iostream>
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check_mul<class MulDouble, double>(Queue);
   std::cout << "Test passed." << std::endl;
   return 0;

sycl/test-e2e/SubGroup/scan.cpp

@@ -14,10 +14,6 @@
 
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check<class KernelName_QTbNYAsEmawQ, int>(Queue);
   check<class KernelName_FQFNSdcVGrCLUbn, unsigned int>(Queue);
   check<class KernelName_kWYnyHJx, long>(Queue);

sycl/test-e2e/SubGroup/scan_fp16.cpp

@@ -11,10 +11,6 @@
 #include <iostream>
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check<class KernelName_dlpo, sycl::half>(Queue);
   std::cout << "Test passed." << std::endl;
   return 0;

sycl/test-e2e/SubGroup/scan_fp64.cpp

@@ -9,10 +9,6 @@
 #include <iostream>
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check<class KernelName_cYZflKkIXS, double>(Queue);
   std::cout << "Test passed." << std::endl;
   return 0;

sycl/test-e2e/SubGroup/scan_spirv13.cpp

@@ -9,10 +9,6 @@
 
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check_mul<class MulA, int>(Queue);
   check_mul<class MulB, unsigned int>(Queue);
   check_mul<class MulC, long>(Queue);

sycl/test-e2e/SubGroup/scan_spirv13_fp16.cpp

@@ -12,10 +12,6 @@
 
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check_mul<class MulHalf, sycl::half>(Queue);
   std::cout << "Test passed." << std::endl;
   return 0;

sycl/test-e2e/SubGroup/scan_spirv13_fp64.cpp

@@ -11,10 +11,6 @@
 
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check<class MulDouble, double>(Queue);
   std::cout << "Test passed." << std::endl;
   return 0;

sycl/test-e2e/SubGroup/vote.cpp

@@ -69,10 +69,6 @@ void check(queue Queue, const int G, const int L, const int D, const int R) {
 }
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
-    std::cout << "Skipping test\n";
-    return 0;
-  }
   check(Queue, 240, 80, 3, 1);
   check(Queue, 24, 12, 3, 4);
   check(Queue, 1024, 256, 3, 1);