[SYCL] Only kernel 0 and 2

Signed-off-by: Hu, Peisen <[email protected]>

Author: HPS-1

sycl/test-e2e/GroupAlgorithm/root_group.cpp

@@ -104,6 +104,32 @@ void testRootGroup() {
   }
 }
 
+template <typename T> struct TestKernel2 {
+  T m_testResults;
+  TestKernel2(T &testResults_) : m_testResults(testResults_) {}
+  void operator()(sycl::nd_item<1> it) const {
+    const auto root = it.ext_oneapi_get_root_group();
+    if (root.leader() || root.get_local_id() == 3) {
+      m_testResults[0] = root.get_group_id() == sycl::id<1>(0);
+      m_testResults[1] = root.leader() ? root.get_local_id() == sycl::id<1>(0)
+                                       : root.get_local_id() == sycl::id<1>(3);
+      m_testResults[2] = root.get_group_range() == sycl::range<1>(1);
+      m_testResults[3] = root.get_local_range() == it.get_global_range();
+      m_testResults[4] = root.get_max_local_range() == root.get_local_range();
+      m_testResults[5] = root.get_group_linear_id() == 0;
+      m_testResults[6] =
+          root.get_local_linear_id() == root.get_local_id().get(0);
+      m_testResults[7] = root.get_group_linear_range() == 1;
+      m_testResults[8] =
+          root.get_local_linear_range() == root.get_local_range().size();
+    }
+  }
+  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+    return sycl::ext::oneapi::experimental::properties{
+        sycl::ext::oneapi::experimental::use_root_sync};
+  }
+};
+
 void testRootGroupFunctions() {
   sycl::queue q;
   const auto bundle =
@@ -114,34 +140,14 @@ void testRootGroupFunctions() {
           .ext_oneapi_get_info<sycl::ext::oneapi::experimental::info::
                                    kernel_queue_specific::max_num_work_groups>(
               q, WorkGroupSize, 0);
-  const auto props = sycl::ext::oneapi::experimental::properties{
-      sycl::ext::oneapi::experimental::use_root_sync};
 
   constexpr int testCount = 9;
   sycl::buffer<bool> testResultsBuf{sycl::range{testCount}};
   const auto range = sycl::nd_range<1>{maxWGs * WorkGroupSize, WorkGroupSize};
   q.submit([&](sycl::handler &h) {
     sycl::accessor testResults{testResultsBuf, h};
-    h.parallel_for<class RootGroupFunctionsKernel>(
-        range, props, [=](sycl::nd_item<1> it) {
-          const auto root = it.ext_oneapi_get_root_group();
-          if (root.leader() || root.get_local_id() == 3) {
-            testResults[0] = root.get_group_id() == sycl::id<1>(0);
-            testResults[1] = root.leader()
-                                 ? root.get_local_id() == sycl::id<1>(0)
-                                 : root.get_local_id() == sycl::id<1>(3);
-            testResults[2] = root.get_group_range() == sycl::range<1>(1);
-            testResults[3] = root.get_local_range() == it.get_global_range();
-            testResults[4] =
-                root.get_max_local_range() == root.get_local_range();
-            testResults[5] = root.get_group_linear_id() == 0;
-            testResults[6] =
-                root.get_local_linear_id() == root.get_local_id().get(0);
-            testResults[7] = root.get_group_linear_range() == 1;
-            testResults[8] =
-                root.get_local_linear_range() == root.get_local_range().size();
-          }
-        });
+    h.parallel_for<class RootGroupFunctionsKernel>(range,
+                                                   TestKernel2(testResults));
   });
   sycl::host_accessor testResults{testResultsBuf};
   for (int i = 0; i < testCount; i++) {