[SYCL][E2E] Limit work group size in WorkGroupScratchMemory tests (#17046)

Some devices don't support work group sizes of 1024, so use the maximum
size if it is smaller in the copy_dynamic_size.cpp and
dynamic_unused.cpp tests.

Author: RossBrunton

sycl/test-e2e/WorkGroupScratchMemory/copy_dynamic_size.cpp

@@ -12,7 +12,6 @@
 #include <sycl/group_barrier.hpp>
 #include <sycl/usm.hpp>
 
-constexpr size_t Size = 1024;
 using DataType = int;
 
 namespace sycl_ext = sycl::ext::oneapi::experimental;
@@ -44,25 +43,29 @@ template <typename T> struct KernelFunctor {
 
 int main() {
   sycl::queue queue;
-  DataType *a = sycl::malloc_device<DataType>(Size, queue);
-  DataType *b = sycl::malloc_device<DataType>(Size, queue);
-  std::vector<DataType> a_host(Size, 1.0);
-  std::vector<DataType> b_host(Size, -5.0);
+  auto size = std::min(
+      queue.get_device().get_info<sycl::info::device::max_work_group_size>(),
+      1024ul);
 
-  queue.copy(a_host.data(), a, Size).wait_and_throw();
+  DataType *a = sycl::malloc_device<DataType>(size, queue);
+  DataType *b = sycl::malloc_device<DataType>(size, queue);
+  std::vector<DataType> a_host(size, 1.0);
+  std::vector<DataType> b_host(size, -5.0);
+
+  queue.copy(a_host.data(), a, size).wait_and_throw();
 
   queue
       .submit([&](sycl::handler &cgh) {
         cgh.parallel_for(
-            sycl::nd_range<1>({Size}, {Size}),
+            sycl::nd_range<1>({size}, {size}),
             KernelFunctor(
                 sycl_ext::properties{
-                    sycl_ext::work_group_scratch_size(Size * sizeof(DataType))},
+                    sycl_ext::work_group_scratch_size(size * sizeof(DataType))},
                 a, b));
       })
       .wait_and_throw();
 
-  queue.copy(b, b_host.data(), Size).wait_and_throw();
+  queue.copy(b, b_host.data(), size).wait_and_throw();
   for (size_t i = 0; i < b_host.size(); i++) {
     assert(b_host[i] == a_host[i]);
   }

sycl/test-e2e/WorkGroupScratchMemory/dynamic_unused.cpp

@@ -9,7 +9,6 @@
 #include <sycl/ext/oneapi/work_group_scratch_memory.hpp>
 #include <sycl/usm.hpp>
 
-constexpr size_t Size = 1024;
 using DataType = int;
 
 namespace sycl_ext = sycl::ext::oneapi::experimental;
@@ -29,25 +28,29 @@ template <typename T> struct KernelFunctor {
 
 int main() {
   sycl::queue queue;
-  DataType *a = sycl::malloc_device<DataType>(Size, queue);
-  DataType *b = sycl::malloc_device<DataType>(Size, queue);
-  std::vector<DataType> a_host(Size, 1.0);
-  std::vector<DataType> b_host(Size, -5.0);
+  auto size = std::min(
+      queue.get_device().get_info<sycl::info::device::max_work_group_size>(),
+      1024ul);
 
-  queue.copy(a_host.data(), a, Size).wait_and_throw();
+  DataType *a = sycl::malloc_device<DataType>(size, queue);
+  DataType *b = sycl::malloc_device<DataType>(size, queue);
+  std::vector<DataType> a_host(size, 1.0);
+  std::vector<DataType> b_host(size, -5.0);
+
+  queue.copy(a_host.data(), a, size).wait_and_throw();
 
   queue
       .submit([&](sycl::handler &cgh) {
         cgh.parallel_for(
-            sycl::nd_range<1>({Size}, {Size}),
+            sycl::nd_range<1>({size}, {size}),
             KernelFunctor(
                 sycl_ext::properties{
-                    sycl_ext::work_group_scratch_size(Size * sizeof(DataType))},
+                    sycl_ext::work_group_scratch_size(size * sizeof(DataType))},
                 a, b));
       })
       .wait_and_throw();
 
-  queue.copy(b, b_host.data(), Size).wait_and_throw();
+  queue.copy(b, b_host.data(), size).wait_and_throw();
   for (size_t i = 0; i < b_host.size(); i++) {
     assert(b_host[i] == a_host[i]);
   }