[SYCL] [E2E] Fix free funcs tests (intel#17742)

jchlanda · web-flow · commit cb5ef368e53e · 2025-04-09T10:56:11.000+01:00
This patch:
* enables free function tests
* re-writes them in terms of `sycl::ext::oneapi::this_work_item`
(instead of, now deprecated, `sycl::ext::oneapi::experimental::this_id`
and friends)
* fixes build failure caused by the removal of `core_sg_supported`
* removes a test from `free_function_queries.cpp` as it was only
checking correct translation between old and new spelling of the
extension
diff --git a/sycl/test-e2e/Basic/free_function_queries/free_function_queries.cpp b/sycl/test-e2e/Basic/free_function_queries/free_function_queries.cpp
@@ -1,10 +1,7 @@
-// FIXME: Investigate OS-agnostic failures
-// UNSUPPORTED: true
 // RUN: %{build} -o %t.out
-// RUN: %{run} %t.out
-
-// UNSUPPORTED: windows
-// The failure is caused by intel/llvm#5213
+// With awkward sizes (such as 1030) make sure that there is no range rounding,
+// so `get_global_id` returns correct values.
+// RUN: env SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING=1 %{run} %t.out
 
 //==- free_function_queries.cpp - SYCL free function queries test -=//
 //
@@ -29,7 +26,7 @@ int main() {
 
   {
     constexpr int checks_number{3};
-    int results[checks_number]{};
+    int results[checks_number]{41, 42, 43};
     {
       sycl::buffer<int> buf(data, sycl::range<1>(n));
       sycl::buffer<int> results_buf(results, sycl::range<1>(checks_number));
@@ -42,12 +39,12 @@ int main() {
                        sycl::access::target::device>
             results_acc(results_buf.get_access<sycl::access::mode::write>(cgh));
         cgh.parallel_for<class IdTest>(n, [=](sycl::id<1> i) {
-          auto that_id = sycl::ext::oneapi::experimental::this_id<1>();
-          results_acc[0] = that_id == i;
+          auto that_id = sycl::ext::oneapi::this_work_item::get_nd_item<1>();
+          results_acc[0] = that_id.get_global_id() == i;
 
-          auto that_item = sycl::ext::oneapi::experimental::this_item<1>();
-          results_acc[1] = that_item.get_id() == i;
-          results_acc[2] = that_item.get_range() == sycl::range<1>(n);
+          auto that_item = sycl::ext::oneapi::this_work_item::get_nd_item<1>();
+          results_acc[1] = that_item.get_global_id() == i;
+          results_acc[2] = that_item.get_global_range() == sycl::range<1>(n);
           acc[i]++;
         });
       });
@@ -63,7 +60,7 @@ int main() {
 
   {
     constexpr int checks_number{2};
-    int results[checks_number]{};
+    int results[checks_number]{41, 42};
     {
       sycl::buffer<int> buf(data, sycl::range<1>(n));
       sycl::buffer<int> results_buf(results, sycl::range<1>(checks_number));
@@ -78,10 +75,10 @@ int main() {
         cgh.parallel_for<class ItemTest>(n, [=](auto i) {
           static_assert(std::is_same<decltype(i), sycl::item<1>>::value,
                         "lambda arg type is unexpected");
-          auto that_id = sycl::ext::oneapi::experimental::this_id<1>();
-          results_acc[0] = i.get_id() == that_id;
-          auto that_item = sycl::ext::oneapi::experimental::this_item<1>();
-          results_acc[1] = i == that_item;
+          auto that_id = sycl::ext::oneapi::this_work_item::get_nd_item<1>();
+          results_acc[0] = that_id.get_global_id() == i;
+          auto that_item = sycl::ext::oneapi::this_work_item::get_nd_item<1>();
+          results_acc[1] = that_item.get_global_id() == i;
           acc[i]++;
         });
       });
@@ -96,8 +93,11 @@ int main() {
   }
 
   {
+// Make sure that we ignore global offset warning.
+#pragma GCC diagnostic push
+#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
     constexpr int checks_number{2};
-    int results[checks_number]{};
+    int results[checks_number]{41, 42};
     {
       sycl::buffer<int> buf(data, sycl::range<1>(n));
       sycl::buffer<int> results_buf(results, sycl::range<1>(checks_number));
@@ -112,11 +112,13 @@ int main() {
             results_acc(results_buf.get_access<sycl::access::mode::write>(cgh));
         cgh.parallel_for<class ItemOffsetTest>(
             sycl::range<1>{n}, offset, [=](sycl::item<1, true> i) {
-              auto that_id = sycl::ext::oneapi::experimental::this_id<1>();
-              results_acc[0] = i.get_id() == that_id;
-              auto that_item = sycl::ext::oneapi::experimental::this_item<1>();
-              results_acc[1] = i == that_item;
-              acc[that_item.get_linear_id()]++;
+              auto that_id =
+                  sycl::ext::oneapi::this_work_item::get_nd_item<1>();
+              results_acc[0] = i.get_id() == that_id.get_global_id();
+              auto that_item =
+                  sycl::ext::oneapi::this_work_item::get_nd_item<1>();
+              results_acc[1] = i == that_item.get_global_id();
+              acc[that_item.get_global_linear_id()]++;
             });
       });
     }
@@ -127,51 +129,6 @@ int main() {
     for (auto val : results) {
       assert(val == 1);
     }
-  }
-
-  {
-    constexpr int checks_number{5};
-    int results[checks_number]{};
-    {
-      sycl::buffer<int> buf(data, sycl::range<1>(n));
-      sycl::buffer<int> results_buf(results, sycl::range<1>(checks_number));
-      sycl::queue q;
-      sycl::nd_range<1> NDR(sycl::range<1>{n}, sycl::range<1>{2});
-      q.submit([&](sycl::handler &cgh) {
-        sycl::accessor<int, 1, sycl::access::mode::write,
-                       sycl::access::target::device>
-            acc(buf.get_access<sycl::access::mode::write>(cgh));
-        sycl::accessor<int, 1, sycl::access::mode::write,
-                       sycl::access::target::device>
-            results_acc(results_buf.get_access<sycl::access::mode::write>(cgh));
-        cgh.parallel_for<class NdItemTest>(NDR, [=](auto nd_i) {
-          static_assert(std::is_same<decltype(nd_i), sycl::nd_item<1>>::value,
-                        "lambda arg type is unexpected");
-          auto that_nd_item =
-              sycl::ext::oneapi::this_work_item::get_nd_item<1>();
-          results_acc[0] = that_nd_item == nd_i;
-          auto nd_item_group = that_nd_item.get_group();
-          results_acc[1] =
-              nd_item_group ==
-              sycl::ext::oneapi::this_work_item::get_work_group<1>();
-          auto nd_item_id = that_nd_item.get_global_id();
-          results_acc[2] =
-              nd_item_id == sycl::ext::oneapi::experimental::this_id<1>();
-          auto that_item = sycl::ext::oneapi::experimental::this_item<1>();
-          results_acc[3] = nd_item_id == that_item.get_id();
-          results_acc[4] =
-              that_nd_item.get_global_range() == that_item.get_range();
-
-          acc[that_nd_item.get_global_id(0)]++;
-        });
-      });
-    }
-    ++counter;
-    for (int i = 0; i < n; i++) {
-      assert(data[i] == counter);
-    }
-    for (auto val : results) {
-      assert(val == 1);
-    }
+#pragma GCC diagnostic pop
   }
 }
diff --git a/sycl/test-e2e/Basic/free_function_queries/free_function_queries_sub_group.cpp b/sycl/test-e2e/Basic/free_function_queries/free_function_queries_sub_group.cpp
@@ -1,14 +1,6 @@
-// FIXME: Investigate OS-agnostic failures
-// UNSUPPORTED: true
-// UNSUPPORTED: cuda || hip
-// CUDA and HIP compilation and runtime do not yet support sub-groups.
-//
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// UNSUPPORTED: windows
-// The failure is caused by intel/llvm#5213
-
 //==- free_function_queries_sub_group.cpp - SYCL free queries for sub group -=//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
@@ -31,15 +23,11 @@ int main() {
   int counter{0};
   {
     constexpr int checks_number{4};
-    int results[checks_number]{};
+    int results[checks_number]{41, 42, 43, 44};
     {
       sycl::buffer<int> buf(data, sycl::range<1>(n));
       sycl::buffer<int> results_buf(results, sycl::range<1>(checks_number));
       sycl::queue q;
-      if (!core_sg_supported(q.get_device())) {
-        std::cout << "Skipping test" << std::endl;
-        return 0;
-      }
       sycl::nd_range<1> NDR(sycl::range<1>{n}, sycl::range<1>{2});
       q.submit([&](sycl::handler &cgh) {
         sycl::accessor<int, 1, sycl::access::mode::write,
diff --git a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp
@@ -54,7 +54,7 @@
 // tests to match the required format and in that case you should just update
 // (i.e. reduce) the number and the list below.
 //
-// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 280
+// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 276
 //
 // List of improperly UNSUPPORTED tests.
 // Remove the CHECK once the test has been properly UNSUPPORTED.
@@ -86,10 +86,6 @@
 // CHECK-NEXT: Assert/check_resource_leak.cpp
 // CHECK-NEXT: Basic/buffer/buffer_create.cpp
 // CHECK-NEXT: Basic/build_log.cpp
-// CHECK-NEXT: Basic/free_function_queries/free_function_queries.cpp
-// CHECK-NEXT: Basic/free_function_queries/free_function_queries.cpp
-// CHECK-NEXT: Basic/free_function_queries/free_function_queries_sub_group.cpp
-// CHECK-NEXT: Basic/free_function_queries/free_function_queries_sub_group.cpp
 // CHECK-NEXT: Basic/gpu_max_wgs_error.cpp
 // CHECK-NEXT: Basic/group_async_copy.cpp
 // CHECK-NEXT: Basic/host-task-dependency.cpp
