diff --git a/sycl/test-e2e/WorkGroupMemory/Dynamic/copy_dynamic_size.cpp b/sycl/test-e2e/WorkGroupScratchMemory/copy_dynamic_size.cpp similarity index 68% rename from sycl/test-e2e/WorkGroupMemory/Dynamic/copy_dynamic_size.cpp rename to sycl/test-e2e/WorkGroupScratchMemory/copy_dynamic_size.cpp index 1f61653efc44e..411fd51011dea 100644 --- a/sycl/test-e2e/WorkGroupMemory/Dynamic/copy_dynamic_size.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/copy_dynamic_size.cpp @@ -10,7 +10,6 @@ #include #include -constexpr size_t Size = 1024; using DataType = int; namespace sycl_ext = sycl::ext::oneapi::experimental; @@ -31,23 +30,27 @@ void copy_via_smem(DataType *a, DataType *b, sycl::nd_item<1> it) { int main() { sycl::queue queue; - DataType *a = sycl::malloc_device(Size, queue); - DataType *b = sycl::malloc_device(Size, queue); - std::vector a_host(Size, 1.0); - std::vector b_host(Size, -5.0); + auto size = std::min( + queue.get_device().get_info(), + size_t{1024}); - queue.copy(a_host.data(), a, Size).wait_and_throw(); + DataType *a = sycl::malloc_device(size, queue); + DataType *b = sycl::malloc_device(size, queue); + std::vector a_host(size, 1.0); + std::vector 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}), + cgh.parallel_for(sycl::nd_range<1>({size}, {size}), sycl_ext::properties{sycl_ext::work_group_scratch_size( - Size * sizeof(DataType))}, + size * sizeof(DataType))}, [=](sycl::nd_item<1> it) { copy_via_smem(a, b, it); }); }) .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]); } diff --git a/sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_alloc_local_accessor.cpp b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp similarity index 98% rename from sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_alloc_local_accessor.cpp rename to sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp index 04d8a85a808ff..f133865ec2833 100644 --- a/sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_alloc_local_accessor.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp @@ -2,7 +2,7 @@ // RUN: %{run} %t.out // -// UNSUPPORTED: gpu-intel-gen12 +// UNSUPPORTED: gpu-intel-gen12, cpu // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/16072 // Test work_group_dynamic extension with allocation size specified at runtime diff --git a/sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_alloc_ptr_alias.cpp b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp similarity index 98% rename from sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_alloc_ptr_alias.cpp rename to sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp index 46346d5f2ee85..969db5061fb67 100644 --- a/sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_alloc_ptr_alias.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp @@ -2,7 +2,7 @@ // RUN: %{run} %t.out // -// UNSUPPORTED: gpu-intel-gen12 +// UNSUPPORTED: gpu-intel-gen12, cpu // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/16072 // Test work_group_dynamic extension with allocation size specified at runtime diff --git a/sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_allocation.cpp b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_allocation.cpp similarity index 98% rename from sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_allocation.cpp rename to sycl/test-e2e/WorkGroupScratchMemory/dynamic_allocation.cpp index 224bf2607f772..1787d406f8ce3 100644 --- a/sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_allocation.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_allocation.cpp @@ -2,7 +2,7 @@ // RUN: %{run} %t.out // -// UNSUPPORTED: gpu-intel-gen12 +// UNSUPPORTED: gpu-intel-gen12, cpu // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/16072 // Test work_group_dynamic extension with allocation size specified at runtime. diff --git a/sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_unused.cpp b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_unused.cpp similarity index 61% rename from sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_unused.cpp rename to sycl/test-e2e/WorkGroupScratchMemory/dynamic_unused.cpp index e427305c18ed3..3944966e62e77 100644 --- a/sycl/test-e2e/WorkGroupMemory/Dynamic/dynamic_unused.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_unused.cpp @@ -9,25 +9,28 @@ #include #include -constexpr size_t Size = 1024; using DataType = int; namespace sycl_ext = sycl::ext::oneapi::experimental; int main() { sycl::queue queue; - DataType *a = sycl::malloc_device(Size, queue); - DataType *b = sycl::malloc_device(Size, queue); - std::vector a_host(Size, 1.0); - std::vector b_host(Size, -5.0); + auto size = std::min( + queue.get_device().get_info(), + size_t{1024}); - queue.copy(a_host.data(), a, Size).wait_and_throw(); + DataType *a = sycl::malloc_device(size, queue); + DataType *b = sycl::malloc_device(size, queue); + std::vector a_host(size, 1.0); + std::vector 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}), + cgh.parallel_for(sycl::nd_range<1>({size}, {size}), sycl_ext::properties{sycl_ext::work_group_scratch_size( - Size * sizeof(DataType))}, + size * sizeof(DataType))}, [=](sycl::nd_item<1> it) { b[it.get_local_linear_id()] = a[it.get_local_linear_id()]; @@ -35,7 +38,7 @@ int main() { }) .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]); } diff --git a/sycl/test-e2e/WorkGroupMemory/lit.local.cfg b/sycl/test-e2e/WorkGroupScratchMemory/lit.local.cfg similarity index 100% rename from sycl/test-e2e/WorkGroupMemory/lit.local.cfg rename to sycl/test-e2e/WorkGroupScratchMemory/lit.local.cfg diff --git a/sycl/test-e2e/WorkGroupStatic/lit.local.cfg b/sycl/test-e2e/WorkGroupStatic/lit.local.cfg new file mode 100644 index 0000000000000..bafde197c494f --- /dev/null +++ b/sycl/test-e2e/WorkGroupStatic/lit.local.cfg @@ -0,0 +1,3 @@ + +# https://github.com/intel/llvm/issues/16072 +config.unsupported_features += ['hip'] diff --git a/sycl/test-e2e/WorkGroupMemory/Static/work_group_static_memory_block_scope.cpp b/sycl/test-e2e/WorkGroupStatic/work_group_static_memory_block_scope.cpp similarity index 100% rename from sycl/test-e2e/WorkGroupMemory/Static/work_group_static_memory_block_scope.cpp rename to sycl/test-e2e/WorkGroupStatic/work_group_static_memory_block_scope.cpp diff --git a/sycl/test-e2e/WorkGroupMemory/Static/work_group_static_memory_class_scope.cpp b/sycl/test-e2e/WorkGroupStatic/work_group_static_memory_class_scope.cpp similarity index 100% rename from sycl/test-e2e/WorkGroupMemory/Static/work_group_static_memory_class_scope.cpp rename to sycl/test-e2e/WorkGroupStatic/work_group_static_memory_class_scope.cpp diff --git a/sycl/test-e2e/WorkGroupMemory/Static/work_group_static_memory_namespace_scope.cpp b/sycl/test-e2e/WorkGroupStatic/work_group_static_memory_namespace_scope.cpp similarity index 100% rename from sycl/test-e2e/WorkGroupMemory/Static/work_group_static_memory_namespace_scope.cpp rename to sycl/test-e2e/WorkGroupStatic/work_group_static_memory_namespace_scope.cpp