From 9e1bfc65fd9163e9a4e8271b1eb1aca41c0fbed4 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Wed, 20 Nov 2024 22:10:31 -0800 Subject: [PATCH 01/30] [SYCL] Deprecate parallel_for and single_task overloads in the sycl_ext_oneapi_kernel_properties extension Signed-off-by: Hu, Peisen --- sycl/include/sycl/handler.hpp | 71 +++++++++++++++++++++++------------ sycl/include/sycl/queue.hpp | 43 +++++++++++++-------- 2 files changed, 76 insertions(+), 38 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index d0a9867ec4c40..092eea6bd1411 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -2314,9 +2314,12 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - single_task(PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { + __SYCL2020_DEPRECATED( + "Use sycl::ext::oneapi::experimental::single_task (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> single_task(PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { single_task_lambda_impl(Props, KernelFunc); } @@ -2364,11 +2367,15 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - (sizeof...(RestT) > 1) && - detail::AreAllButLastReductions::value && - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<1> Range, PropertiesT Properties, RestT &&...Rest) { + __SYCL2020_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t<(sizeof...(RestT) > 1) && + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list< + PropertiesT>::value> parallel_for(range<1> Range, + PropertiesT Properties, + RestT &&...Rest) { #ifndef __SYCL_DEVICE_ONLY__ throwIfGraphAssociated(); @@ -2379,11 +2386,15 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - (sizeof...(RestT) > 1) && - detail::AreAllButLastReductions::value && - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<2> Range, PropertiesT Properties, RestT &&...Rest) { + __SYCL2020_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t<(sizeof...(RestT) > 1) && + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list< + PropertiesT>::value> parallel_for(range<2> Range, + PropertiesT Properties, + RestT &&...Rest) { #ifndef __SYCL_DEVICE_ONLY__ throwIfGraphAssociated(); @@ -2394,11 +2405,15 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - (sizeof...(RestT) > 1) && - detail::AreAllButLastReductions::value && - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<3> Range, PropertiesT Properties, RestT &&...Rest) { + __SYCL2020_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t<(sizeof...(RestT) > 1) && + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list< + PropertiesT>::value> parallel_for(range<3> Range, + PropertiesT Properties, + RestT &&...Rest) { #ifndef __SYCL_DEVICE_ONLY__ throwIfGraphAssociated(); @@ -2433,11 +2448,15 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - (sizeof...(RestT) > 1) && - detail::AreAllButLastReductions::value && - ext::oneapi::experimental::is_property_list::value> - parallel_for(nd_range Range, PropertiesT Properties, RestT &&...Rest) { + __SYCL2020_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t<(sizeof...(RestT) > 1) && + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list< + PropertiesT>::value> parallel_for(nd_range Range, + PropertiesT Properties, + RestT &&...Rest) { #ifndef __SYCL_DEVICE_ONLY__ throwIfGraphAssociated(); @@ -2459,6 +2478,9 @@ class __SYCL_EXPORT handler { template + __SYCL2020_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") void parallel_for_work_group(range NumWorkGroups, PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { parallel_for_work_group_lambda_impl + __SYCL2020_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") void parallel_for_work_group(range NumWorkGroups, range WorkGroupSize, PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 9e530604ce84e..0d10160b07fb4 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -2089,11 +2089,14 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param CodeLoc contains the code location of user code template + __SYCL2020_DEPRECATED( + "Use sycl::ext::oneapi::experimental::single_task (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< - ext::oneapi::experimental::is_property_list::value, event> - single_task( - PropertiesT Properties, _KERNELFUNCPARAM(KernelFunc), - const detail::code_location &CodeLoc = detail::code_location::current()) { + ext::oneapi::experimental::is_property_list::value, + event> single_task(PropertiesT Properties, _KERNELFUNCPARAM(KernelFunc), + const detail::code_location &CodeLoc = + detail::code_location::current()) { static_assert( (detail::check_fn_signature, void()>::value || @@ -2131,11 +2134,15 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param CodeLoc contains the code location of user code template + __SYCL2020_DEPRECATED( + "Use sycl::ext::oneapi::experimental::single_task (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< - ext::oneapi::experimental::is_property_list::value, event> - single_task( - event DepEvent, PropertiesT Properties, _KERNELFUNCPARAM(KernelFunc), - const detail::code_location &CodeLoc = detail::code_location::current()) { + ext::oneapi::experimental::is_property_list::value, + event> single_task(event DepEvent, PropertiesT Properties, + _KERNELFUNCPARAM(KernelFunc), + const detail::code_location &CodeLoc = + detail::code_location::current()) { static_assert( (detail::check_fn_signature, void()>::value || @@ -2177,12 +2184,15 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param CodeLoc contains the code location of user code template + __SYCL2020_DEPRECATED( + "Use sycl::ext::oneapi::experimental::single_task (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< - ext::oneapi::experimental::is_property_list::value, event> - single_task( - const std::vector &DepEvents, PropertiesT Properties, - _KERNELFUNCPARAM(KernelFunc), - const detail::code_location &CodeLoc = detail::code_location::current()) { + ext::oneapi::experimental::is_property_list::value, + event> single_task(const std::vector &DepEvents, + PropertiesT Properties, _KERNELFUNCPARAM(KernelFunc), + const detail::code_location &CodeLoc = + detail::code_location::current()) { static_assert( (detail::check_fn_signature, void()>::value || @@ -2428,11 +2438,14 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// const KernelType &KernelFunc". template + __SYCL2020_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< detail::AreAllButLastReductions::value && ext::oneapi::experimental::is_property_list::value, - event> - parallel_for(nd_range Range, PropertiesT Properties, RestT &&...Rest) { + event> parallel_for(nd_range Range, PropertiesT Properties, + RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( From 21fd2d9fa4b51634a5305666f2c41fff0dd82c91 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Thu, 21 Nov 2024 16:10:53 -0800 Subject: [PATCH 02/30] [SYCL] Update affected llvm-lit test cases to ignore deprecation warnings Signed-off-by: Hu, Peisen --- .../extensions/properties/properties_kernel_device_has.cpp | 4 ++-- .../extensions/properties/properties_kernel_launch_bounds.cpp | 4 ++-- .../properties/properties_kernel_max_work_group_size.cpp | 4 ++-- .../properties/properties_kernel_sub_group_size.cpp | 4 ++-- .../properties/properties_kernel_work_group_size.cpp | 4 ++-- .../properties/properties_kernel_work_group_size_hint.cpp | 4 ++-- .../properties/properties_kernel_device_has_warning.cpp | 2 +- .../properties/properties_kernel_negative_device.cpp | 2 +- sycl/test/virtual-functions/diagnostics-positive.cpp | 2 +- sycl/test/virtual-functions/properties-positive.cpp | 2 +- 10 files changed, 16 insertions(+), 16 deletions(-) diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp index 055b25b920b8b..5e263326113b4 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp @@ -1,5 +1,5 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm -Xclang -disable-llvm-passes %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm -Xclang -disable-llvm-passes %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp index 96ac3da42a504..1b85f1b6fed55 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp @@ -1,5 +1,5 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp index 924270bb6cafe..87eab697d2405 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp @@ -1,5 +1,5 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp index a179c134749e9..5878f5c5d80fd 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp @@ -1,5 +1,5 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp index 932b92fab9009..f3912f28c87ed 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp @@ -1,5 +1,5 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp index a844b484b8b51..6ed9f6b9b1feb 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp @@ -1,5 +1,5 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp index f24e089bed4b1..267f19a2aad9e 100644 --- a/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp +++ b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s +// RUN: %clangxx -fsycl-device-only -Wno-deprecated-declarations -Xclang -verify -Xclang -verify-ignore-unexpected=note %s // Tests for warnings when propagated aspects do not match the aspects available // in a function, as specified through the 'sycl::device_has' property. diff --git a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp index d451e319a2670..5990fe96c2ced 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -ferror-limit=0 -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s +// RUN: %clangxx -fsycl-device-only -Wno-deprecated-declarations -Xclang -fsycl-is-device -ferror-limit=0 -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s #include diff --git a/sycl/test/virtual-functions/diagnostics-positive.cpp b/sycl/test/virtual-functions/diagnostics-positive.cpp index 1b1e346678e8f..a0a9de1f819aa 100644 --- a/sycl/test/virtual-functions/diagnostics-positive.cpp +++ b/sycl/test/virtual-functions/diagnostics-positive.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -c -Xclang -verify %s -o %t.ignored +// RUN: %clangxx -fsycl -c -Wno-deprecated-declarations -Xclang -verify %s -o %t.ignored // // This test is intended to check that no diagnostics are emitted when a kernel // performing virtual function calls is submitted with the right properties. diff --git a/sycl/test/virtual-functions/properties-positive.cpp b/sycl/test/virtual-functions/properties-positive.cpp index 970e348457388..f4b66f322968d 100644 --- a/sycl/test/virtual-functions/properties-positive.cpp +++ b/sycl/test/virtual-functions/properties-positive.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // // This test is intended to check that we can successfully compile code that // uses new properties from the virtual functions extension. From ed7dec29b5409011560e0d8a7306a71915f3eab4 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Thu, 21 Nov 2024 17:30:12 -0800 Subject: [PATCH 03/30] [SYCL] Fix affected e2e test cases Signed-off-by: Hu, Peisen --- sycl/test-e2e/Basic/work_group_size_prop.cpp | 2 +- sycl/test-e2e/Graph/RecordReplay/work_group_size_prop.cpp | 2 +- sycl/test-e2e/GroupAlgorithm/root_group.cpp | 2 +- sycl/test-e2e/Properties/cache_config.cpp | 2 +- sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp | 2 +- sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp | 2 +- sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp | 2 +- .../VirtualFunctions/2/2/single-construct-single-use.cpp | 2 +- sycl/test-e2e/VirtualFunctions/misc/math.cpp | 2 +- .../forward_progress/forward_progress_kernel_param_L0_gpu.cpp | 2 +- .../forward_progress/forward_progress_kernel_param_ocl_cpu.cpp | 2 +- 11 files changed, 11 insertions(+), 11 deletions(-) diff --git a/sycl/test-e2e/Basic/work_group_size_prop.cpp b/sycl/test-e2e/Basic/work_group_size_prop.cpp index ac8400dcc31b8..a373a6c5477ef 100644 --- a/sycl/test-e2e/Basic/work_group_size_prop.cpp +++ b/sycl/test-e2e/Basic/work_group_size_prop.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out +// RUN: %{build} -o %t.out -Wno-deprecated-declarations // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/Graph/RecordReplay/work_group_size_prop.cpp b/sycl/test-e2e/Graph/RecordReplay/work_group_size_prop.cpp index 3a0f21f4f94a7..8f0e53a24305e 100644 --- a/sycl/test-e2e/Graph/RecordReplay/work_group_size_prop.cpp +++ b/sycl/test-e2e/Graph/RecordReplay/work_group_size_prop.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out +// RUN: %{build} -o %t.out -Wno-deprecated-declarations // RUN: %{run} %t.out // Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} diff --git a/sycl/test-e2e/GroupAlgorithm/root_group.cpp b/sycl/test-e2e/GroupAlgorithm/root_group.cpp index 92e5d69ffcab4..63bb9f523661d 100644 --- a/sycl/test-e2e/GroupAlgorithm/root_group.cpp +++ b/sycl/test-e2e/GroupAlgorithm/root_group.cpp @@ -1,6 +1,6 @@ // Fails with opencl non-cpu, enable when fixed. // XFAIL: (opencl && !cpu && !accelerator) -// RUN: %{build} -I . -o %t.out %if any-device-is-cuda %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %} +// RUN: %{build} -I . -o %t.out %if any-device-is-cuda %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %} -Wno-deprecated-declarations // RUN: %{run} %t.out // Disabled temporarily while investigation into the failure is ongoing. diff --git a/sycl/test-e2e/Properties/cache_config.cpp b/sycl/test-e2e/Properties/cache_config.cpp index 0cda3e97a5d1f..b62c50094584b 100644 --- a/sycl/test-e2e/Properties/cache_config.cpp +++ b/sycl/test-e2e/Properties/cache_config.cpp @@ -1,6 +1,6 @@ // REQUIRES: gpu, level_zero -// RUN: %{build} -o %t.out +// RUN: %{build} -o %t.out -Wno-deprecated-declarations // RUN: env UR_L0_DEBUG=1 %{run} %t.out 2>&1 | FileCheck %s #include diff --git a/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp b/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp index f198bc94f855f..fb8b1f305fba3 100644 --- a/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp +++ b/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out %helper-includes +// RUN: %{build} -o %t.out %helper-includes -Wno-deprecated-declarations // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp b/sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp index bb334972c3f77..46d407ff81074 100644 --- a/sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp +++ b/sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out %helper-includes +// RUN: %{build} -o %t.out %helper-includes -Wno-deprecated-declarations // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp b/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp index 2bfb3dd0f010d..7919a18f09785 100644 --- a/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp +++ b/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out %helper-includes +// RUN: %{build} -o %t.out %helper-includes -Wno-deprecated-declarations // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp b/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp index ccf0c77036085..23561a96c9abe 100644 --- a/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp +++ b/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out %helper-includes +// RUN: %{build} -o %t.out %helper-includes -Wno-deprecated-declarations // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/VirtualFunctions/misc/math.cpp b/sycl/test-e2e/VirtualFunctions/misc/math.cpp index 71b34c23cef1f..1d3482f671243 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/math.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/math.cpp @@ -3,7 +3,7 @@ // This test checks that SYCL math built-in functions work correctly // inside virtual functions. // -// RUN: %{build} -o %t.out %helper-includes +// RUN: %{build} -o %t.out %helper-includes -Wno-deprecated-declarations // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/forward_progress/forward_progress_kernel_param_L0_gpu.cpp b/sycl/test-e2e/forward_progress/forward_progress_kernel_param_L0_gpu.cpp index 003840a8c1299..4f9a7eaf79fbc 100644 --- a/sycl/test-e2e/forward_progress/forward_progress_kernel_param_L0_gpu.cpp +++ b/sycl/test-e2e/forward_progress/forward_progress_kernel_param_L0_gpu.cpp @@ -1,5 +1,5 @@ // REQUIRES: level_zero, gpu -// RUN: %{build} -o %t.out +// RUN: %{build} -o %t.out -Wno-deprecated-declarations // RUN: %{run} %t.out // The purpose of this test is to check that the forward_progress_guarantee diff --git a/sycl/test-e2e/forward_progress/forward_progress_kernel_param_ocl_cpu.cpp b/sycl/test-e2e/forward_progress/forward_progress_kernel_param_ocl_cpu.cpp index ffdd99184d233..364517f3ac411 100644 --- a/sycl/test-e2e/forward_progress/forward_progress_kernel_param_ocl_cpu.cpp +++ b/sycl/test-e2e/forward_progress/forward_progress_kernel_param_ocl_cpu.cpp @@ -1,5 +1,5 @@ // REQUIRES: opencl, cpu -// RUN: %{build} -o %t.out +// RUN: %{build} -o %t.out -Wno-deprecated-declarations // RUN: %{run} %t.out // The purpose of this test is to check that the forward_progress_guarantee From e9e77e1f2b79b491f0e87e18ddec054e3369900a Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Thu, 21 Nov 2024 18:52:43 -0800 Subject: [PATCH 04/30] [SYCL] Fix affected CUDA e2e test cases Signed-off-by: Hu, Peisen --- sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp | 2 +- sycl/test-e2e/Basic/max_work_group_size_props.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp b/sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp index 7009ca367d8e9..b6f584da3896c 100644 --- a/sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp +++ b/sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out +// RUN: %{build} -o %t.out -Wno-deprecated-declarations // RUN: %{run} %t.out // This property is not yet supported by all UR adapters diff --git a/sycl/test-e2e/Basic/max_work_group_size_props.cpp b/sycl/test-e2e/Basic/max_work_group_size_props.cpp index 96439971d904a..7946aef87c659 100644 --- a/sycl/test-e2e/Basic/max_work_group_size_props.cpp +++ b/sycl/test-e2e/Basic/max_work_group_size_props.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out +// RUN: %{build} -o %t.out -Wno-deprecated-declarations // RUN: %{run} %t.out // This property is not yet supported by all UR adapters From 22ae374169ca539a258b579147534e4c3a6daa42 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Fri, 22 Nov 2024 07:42:52 -0800 Subject: [PATCH 05/30] [SYCL] Update Graph/Explicit/work_group_size_prop.cpp Signed-off-by: Hu, Peisen --- sycl/test-e2e/Graph/Explicit/work_group_size_prop.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/Graph/Explicit/work_group_size_prop.cpp b/sycl/test-e2e/Graph/Explicit/work_group_size_prop.cpp index 3c97c3c351516..dc42a1360541e 100644 --- a/sycl/test-e2e/Graph/Explicit/work_group_size_prop.cpp +++ b/sycl/test-e2e/Graph/Explicit/work_group_size_prop.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out +// RUN: %{build} -o %t.out -Wno-deprecated-declarations // RUN: %{run} %t.out // Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} From 241d2712fa5989cfe115d7045169bd8016713819 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Mon, 25 Nov 2024 22:20:02 -0800 Subject: [PATCH 06/30] [SYCL] Use proper deprecation warning type Signed-off-by: Hu, Peisen --- sycl/include/sycl/handler.hpp | 14 +++++++------- sycl/include/sycl/queue.hpp | 8 ++++---- 2 files changed, 11 insertions(+), 11 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 092eea6bd1411..6b59cb5c497c7 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -2314,7 +2314,7 @@ class __SYCL_EXPORT handler { template - __SYCL2020_DEPRECATED( + __SYCL_DEPRECATED( "Use sycl::ext::oneapi::experimental::single_task (provided in the " "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t - __SYCL2020_DEPRECATED( + __SYCL_DEPRECATED( "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t<(sizeof...(RestT) > 1) && @@ -2386,7 +2386,7 @@ class __SYCL_EXPORT handler { template - __SYCL2020_DEPRECATED( + __SYCL_DEPRECATED( "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t<(sizeof...(RestT) > 1) && @@ -2405,7 +2405,7 @@ class __SYCL_EXPORT handler { template - __SYCL2020_DEPRECATED( + __SYCL_DEPRECATED( "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t<(sizeof...(RestT) > 1) && @@ -2448,7 +2448,7 @@ class __SYCL_EXPORT handler { template - __SYCL2020_DEPRECATED( + __SYCL_DEPRECATED( "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t<(sizeof...(RestT) > 1) && @@ -2478,7 +2478,7 @@ class __SYCL_EXPORT handler { template - __SYCL2020_DEPRECATED( + __SYCL_DEPRECATED( "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " "sycl_ext_oneapi_enqueue_functions extension) instead.") void parallel_for_work_group(range NumWorkGroups, PropertiesT Props, @@ -2490,7 +2490,7 @@ class __SYCL_EXPORT handler { template - __SYCL2020_DEPRECATED( + __SYCL_DEPRECATED( "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " "sycl_ext_oneapi_enqueue_functions extension) instead.") void parallel_for_work_group(range NumWorkGroups, diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 0d10160b07fb4..8193e53a60e44 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -2089,7 +2089,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param CodeLoc contains the code location of user code template - __SYCL2020_DEPRECATED( + __SYCL_DEPRECATED( "Use sycl::ext::oneapi::experimental::single_task (provided in the " "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< @@ -2134,7 +2134,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param CodeLoc contains the code location of user code template - __SYCL2020_DEPRECATED( + __SYCL_DEPRECATED( "Use sycl::ext::oneapi::experimental::single_task (provided in the " "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< @@ -2184,7 +2184,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param CodeLoc contains the code location of user code template - __SYCL2020_DEPRECATED( + __SYCL_DEPRECATED( "Use sycl::ext::oneapi::experimental::single_task (provided in the " "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< @@ -2438,7 +2438,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// const KernelType &KernelFunc". template - __SYCL2020_DEPRECATED( + __SYCL_DEPRECATED( "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< From f68f4568644d2efe58c58de709c2a4f83f0999cb Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Mon, 16 Dec 2024 13:34:31 -0800 Subject: [PATCH 07/30] [SYCL] Revert changes to E2E tests Signed-off-by: Hu, Peisen --- .../max_linear_work_group_size_props.cpp | 2 +- .../Basic/max_work_group_size_props.cpp | 2 +- sycl/test-e2e/Basic/work_group_size_prop.cpp | 2 +- .../Graph/Explicit/work_group_size_prop.cpp | 2 +- .../RecordReplay/work_group_size_prop.cpp | 2 +- sycl/test-e2e/GroupAlgorithm/root_group.cpp | 40 ++++++++++++------- sycl/test-e2e/Properties/cache_config.cpp | 2 +- .../2/1/1/missing-overrides.cpp | 2 +- .../2/1/1/more-complex-hierarchy.cpp | 2 +- .../2/1/1/simple-hierarchy.cpp | 2 +- .../2/2/single-construct-single-use.cpp | 2 +- sycl/test-e2e/VirtualFunctions/misc/math.cpp | 2 +- .../forward_progress_kernel_param_L0_gpu.cpp | 2 +- .../forward_progress_kernel_param_ocl_cpu.cpp | 2 +- 14 files changed, 39 insertions(+), 27 deletions(-) diff --git a/sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp b/sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp index b6f584da3896c..7009ca367d8e9 100644 --- a/sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp +++ b/sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out -Wno-deprecated-declarations +// RUN: %{build} -o %t.out // RUN: %{run} %t.out // This property is not yet supported by all UR adapters diff --git a/sycl/test-e2e/Basic/max_work_group_size_props.cpp b/sycl/test-e2e/Basic/max_work_group_size_props.cpp index 7946aef87c659..96439971d904a 100644 --- a/sycl/test-e2e/Basic/max_work_group_size_props.cpp +++ b/sycl/test-e2e/Basic/max_work_group_size_props.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out -Wno-deprecated-declarations +// RUN: %{build} -o %t.out // RUN: %{run} %t.out // This property is not yet supported by all UR adapters diff --git a/sycl/test-e2e/Basic/work_group_size_prop.cpp b/sycl/test-e2e/Basic/work_group_size_prop.cpp index a373a6c5477ef..ac8400dcc31b8 100644 --- a/sycl/test-e2e/Basic/work_group_size_prop.cpp +++ b/sycl/test-e2e/Basic/work_group_size_prop.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out -Wno-deprecated-declarations +// RUN: %{build} -o %t.out // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/Graph/Explicit/work_group_size_prop.cpp b/sycl/test-e2e/Graph/Explicit/work_group_size_prop.cpp index dc42a1360541e..3c97c3c351516 100644 --- a/sycl/test-e2e/Graph/Explicit/work_group_size_prop.cpp +++ b/sycl/test-e2e/Graph/Explicit/work_group_size_prop.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out -Wno-deprecated-declarations +// RUN: %{build} -o %t.out // RUN: %{run} %t.out // Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} diff --git a/sycl/test-e2e/Graph/RecordReplay/work_group_size_prop.cpp b/sycl/test-e2e/Graph/RecordReplay/work_group_size_prop.cpp index 8f0e53a24305e..3a0f21f4f94a7 100644 --- a/sycl/test-e2e/Graph/RecordReplay/work_group_size_prop.cpp +++ b/sycl/test-e2e/Graph/RecordReplay/work_group_size_prop.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out -Wno-deprecated-declarations +// RUN: %{build} -o %t.out // RUN: %{run} %t.out // Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} diff --git a/sycl/test-e2e/GroupAlgorithm/root_group.cpp b/sycl/test-e2e/GroupAlgorithm/root_group.cpp index 63bb9f523661d..2e50634fd21c8 100644 --- a/sycl/test-e2e/GroupAlgorithm/root_group.cpp +++ b/sycl/test-e2e/GroupAlgorithm/root_group.cpp @@ -1,6 +1,8 @@ // Fails with opencl non-cpu, enable when fixed. // XFAIL: (opencl && !cpu && !accelerator) -// RUN: %{build} -I . -o %t.out %if any-device-is-cuda %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %} -Wno-deprecated-declarations +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/14641 + +// RUN: %{build} -I . -o %t.out %if any-device-is-cuda %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %} // RUN: %{run} %t.out // Disabled temporarily while investigation into the failure is ongoing. @@ -14,6 +16,7 @@ #include #include #include +#include static constexpr int WorkGroupSize = 32; @@ -27,13 +30,18 @@ void testQueriesAndProperties() { const auto bundle = sycl::get_kernel_bundle(q.get_context()); const auto kernel = bundle.get_kernel(); - const auto maxWGs = kernel.ext_oneapi_get_info< - sycl::ext::oneapi::experimental::info::kernel_queue_specific:: - max_num_work_group_sync>(q); - const auto wgRange = sycl::range{WorkGroupSize, 1, 1}; - const auto maxWGsWithLimits = kernel.ext_oneapi_get_info< - sycl::ext::oneapi::experimental::info::kernel_queue_specific:: - max_num_work_group_sync>(q, wgRange, wgRange.size() * sizeof(int)); + const auto local_range = sycl::range<1>(1); + const auto maxWGs = + kernel + .ext_oneapi_get_info( + q, local_range, 0); + const auto wgRange = sycl::range<3>{WorkGroupSize, 1, 1}; + const auto maxWGsWithLimits = + kernel + .ext_oneapi_get_info( + q, wgRange, wgRange.size() * sizeof(int)); const auto props = sycl::ext::oneapi::experimental::properties{ sycl::ext::oneapi::experimental::use_root_sync}; q.single_task(props, []() {}); @@ -52,9 +60,11 @@ void testRootGroup() { const auto bundle = sycl::get_kernel_bundle(q.get_context()); const auto kernel = bundle.get_kernel(); - const auto maxWGs = kernel.ext_oneapi_get_info< - sycl::ext::oneapi::experimental::info::kernel_queue_specific:: - max_num_work_group_sync>(q); + const auto maxWGs = + kernel + .ext_oneapi_get_info( + q, WorkGroupSize, 0); const auto props = sycl::ext::oneapi::experimental::properties{ sycl::ext::oneapi::experimental::use_root_sync}; sycl::buffer dataBuf{sycl::range{maxWGs * WorkGroupSize}}; @@ -94,9 +104,11 @@ void testRootGroupFunctions() { const auto bundle = sycl::get_kernel_bundle(q.get_context()); const auto kernel = bundle.get_kernel(); - const auto maxWGs = kernel.ext_oneapi_get_info< - sycl::ext::oneapi::experimental::info::kernel_queue_specific:: - max_num_work_group_sync>(q); + const auto maxWGs = + kernel + .ext_oneapi_get_info( + q, WorkGroupSize, 0); const auto props = sycl::ext::oneapi::experimental::properties{ sycl::ext::oneapi::experimental::use_root_sync}; diff --git a/sycl/test-e2e/Properties/cache_config.cpp b/sycl/test-e2e/Properties/cache_config.cpp index b62c50094584b..0cda3e97a5d1f 100644 --- a/sycl/test-e2e/Properties/cache_config.cpp +++ b/sycl/test-e2e/Properties/cache_config.cpp @@ -1,6 +1,6 @@ // REQUIRES: gpu, level_zero -// RUN: %{build} -o %t.out -Wno-deprecated-declarations +// RUN: %{build} -o %t.out // RUN: env UR_L0_DEBUG=1 %{run} %t.out 2>&1 | FileCheck %s #include diff --git a/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp b/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp index fb8b1f305fba3..f198bc94f855f 100644 --- a/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp +++ b/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out %helper-includes -Wno-deprecated-declarations +// RUN: %{build} -o %t.out %helper-includes // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp b/sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp index 46d407ff81074..bb334972c3f77 100644 --- a/sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp +++ b/sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out %helper-includes -Wno-deprecated-declarations +// RUN: %{build} -o %t.out %helper-includes // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp b/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp index 7919a18f09785..2bfb3dd0f010d 100644 --- a/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp +++ b/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out %helper-includes -Wno-deprecated-declarations +// RUN: %{build} -o %t.out %helper-includes // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp b/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp index 23561a96c9abe..ccf0c77036085 100644 --- a/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp +++ b/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp @@ -1,4 +1,4 @@ -// RUN: %{build} -o %t.out %helper-includes -Wno-deprecated-declarations +// RUN: %{build} -o %t.out %helper-includes // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/VirtualFunctions/misc/math.cpp b/sycl/test-e2e/VirtualFunctions/misc/math.cpp index 1d3482f671243..71b34c23cef1f 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/math.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/math.cpp @@ -3,7 +3,7 @@ // This test checks that SYCL math built-in functions work correctly // inside virtual functions. // -// RUN: %{build} -o %t.out %helper-includes -Wno-deprecated-declarations +// RUN: %{build} -o %t.out %helper-includes // RUN: %{run} %t.out #include diff --git a/sycl/test-e2e/forward_progress/forward_progress_kernel_param_L0_gpu.cpp b/sycl/test-e2e/forward_progress/forward_progress_kernel_param_L0_gpu.cpp index 4f9a7eaf79fbc..003840a8c1299 100644 --- a/sycl/test-e2e/forward_progress/forward_progress_kernel_param_L0_gpu.cpp +++ b/sycl/test-e2e/forward_progress/forward_progress_kernel_param_L0_gpu.cpp @@ -1,5 +1,5 @@ // REQUIRES: level_zero, gpu -// RUN: %{build} -o %t.out -Wno-deprecated-declarations +// RUN: %{build} -o %t.out // RUN: %{run} %t.out // The purpose of this test is to check that the forward_progress_guarantee diff --git a/sycl/test-e2e/forward_progress/forward_progress_kernel_param_ocl_cpu.cpp b/sycl/test-e2e/forward_progress/forward_progress_kernel_param_ocl_cpu.cpp index 364517f3ac411..ffdd99184d233 100644 --- a/sycl/test-e2e/forward_progress/forward_progress_kernel_param_ocl_cpu.cpp +++ b/sycl/test-e2e/forward_progress/forward_progress_kernel_param_ocl_cpu.cpp @@ -1,5 +1,5 @@ // REQUIRES: opencl, cpu -// RUN: %{build} -o %t.out -Wno-deprecated-declarations +// RUN: %{build} -o %t.out // RUN: %{run} %t.out // The purpose of this test is to check that the forward_progress_guarantee From e7666a0cc5d10d8fec45fbb57f0dcb86ef52a65e Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Mon, 16 Dec 2024 13:48:43 -0800 Subject: [PATCH 08/30] [SYCL] Revert changes to llvm-lit tests (without reductions) Signed-off-by: Hu, Peisen --- .../properties/properties_kernel_launch_bounds.cpp | 4 ++-- .../properties/properties_kernel_max_work_group_size.cpp | 4 ++-- .../properties/properties_kernel_device_has_warning.cpp | 2 +- .../properties/properties_kernel_negative_device.cpp | 2 +- sycl/test/virtual-functions/diagnostics-positive.cpp | 2 +- sycl/test/virtual-functions/properties-positive.cpp | 6 +----- 6 files changed, 8 insertions(+), 12 deletions(-) diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp index 1b85f1b6fed55..96ac3da42a504 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp @@ -1,5 +1,5 @@ -// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s +// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp index 87eab697d2405..924270bb6cafe 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp @@ -1,5 +1,5 @@ -// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s +// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp index 267f19a2aad9e..f24e089bed4b1 100644 --- a/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp +++ b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl-device-only -Wno-deprecated-declarations -Xclang -verify -Xclang -verify-ignore-unexpected=note %s +// RUN: %clangxx -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s // Tests for warnings when propagated aspects do not match the aspects available // in a function, as specified through the 'sycl::device_has' property. diff --git a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp index 5990fe96c2ced..d451e319a2670 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl-device-only -Wno-deprecated-declarations -Xclang -fsycl-is-device -ferror-limit=0 -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s +// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -ferror-limit=0 -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s #include diff --git a/sycl/test/virtual-functions/diagnostics-positive.cpp b/sycl/test/virtual-functions/diagnostics-positive.cpp index a0a9de1f819aa..1b1e346678e8f 100644 --- a/sycl/test/virtual-functions/diagnostics-positive.cpp +++ b/sycl/test/virtual-functions/diagnostics-positive.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -c -Wno-deprecated-declarations -Xclang -verify %s -o %t.ignored +// RUN: %clangxx -fsycl -c -Xclang -verify %s -o %t.ignored // // This test is intended to check that no diagnostics are emitted when a kernel // performing virtual function calls is submitted with the right properties. diff --git a/sycl/test/virtual-functions/properties-positive.cpp b/sycl/test/virtual-functions/properties-positive.cpp index f4b66f322968d..3441e91a4d45c 100644 --- a/sycl/test/virtual-functions/properties-positive.cpp +++ b/sycl/test/virtual-functions/properties-positive.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s +// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s // // This test is intended to check that we can successfully compile code that // uses new properties from the virtual functions extension. @@ -44,10 +44,6 @@ class SubSubDerived : public SubDerived { int main() { sycl::queue q; - static_assert( - oneapi::is_property_key::value); - static_assert(oneapi::is_property_key::value); - oneapi::properties props_empty{oneapi::assume_indirect_calls}; oneapi::properties props_void{oneapi::assume_indirect_calls_to}; oneapi::properties props_int{oneapi::assume_indirect_calls_to}; From 677cf8e0bd85c5fec842899c4a88dd4d6221cc24 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Mon, 16 Dec 2024 14:12:58 -0800 Subject: [PATCH 09/30] [SYCL] Add TODO comments for tests with -Wno-deprecated-declarations flag Signed-off-by: Hu, Peisen --- .../extensions/properties/properties_kernel_device_has.cpp | 3 +++ .../extensions/properties/properties_kernel_sub_group_size.cpp | 3 +++ .../properties/properties_kernel_work_group_size.cpp | 3 +++ .../properties/properties_kernel_work_group_size_hint.cpp | 3 +++ 4 files changed, 12 insertions(+) diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp index 5e263326113b4..3d1c528744afd 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp @@ -1,3 +1,6 @@ +// TODO: Currently using the -Wno-deprecated-declarations flag due to issue +// https://github.com/intel/llvm/issues/16320. Remove the flag once the issue is +// resolved. // RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm -Xclang -disable-llvm-passes %s -o - | FileCheck %s --check-prefix CHECK-IR // RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp index 5878f5c5d80fd..ad81d1db1fe0b 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp @@ -1,3 +1,6 @@ +// TODO: Currently using the -Wno-deprecated-declarations flag due to issue +// https://github.com/intel/llvm/issues/16320. Remove the flag once the issue is +// resolved. // RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR // RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp index f3912f28c87ed..63280fcc638f3 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp @@ -1,3 +1,6 @@ +// TODO: Currently using the -Wno-deprecated-declarations flag due to issue +// https://github.com/intel/llvm/issues/16320. Remove the flag once the issue is +// resolved. // RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR // RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp index 6ed9f6b9b1feb..a0bae31ad8004 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp @@ -1,3 +1,6 @@ +// TODO: Currently using the -Wno-deprecated-declarations flag due to issue +// https://github.com/intel/llvm/issues/16320. Remove the flag once the issue is +// resolved. // RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR // RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics From 914304b38ebe6de747f9828df91cc1e0edd171df Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Mon, 16 Dec 2024 14:28:23 -0800 Subject: [PATCH 10/30] [SYCL] Revert unrelated changes Signed-off-by: Hu, Peisen --- sycl/test-e2e/GroupAlgorithm/root_group.cpp | 38 +++++++------------ .../virtual-functions/properties-positive.cpp | 4 ++ 2 files changed, 17 insertions(+), 25 deletions(-) diff --git a/sycl/test-e2e/GroupAlgorithm/root_group.cpp b/sycl/test-e2e/GroupAlgorithm/root_group.cpp index 2e50634fd21c8..92e5d69ffcab4 100644 --- a/sycl/test-e2e/GroupAlgorithm/root_group.cpp +++ b/sycl/test-e2e/GroupAlgorithm/root_group.cpp @@ -1,7 +1,5 @@ // Fails with opencl non-cpu, enable when fixed. // XFAIL: (opencl && !cpu && !accelerator) -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/14641 - // RUN: %{build} -I . -o %t.out %if any-device-is-cuda %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %} // RUN: %{run} %t.out @@ -16,7 +14,6 @@ #include #include #include -#include static constexpr int WorkGroupSize = 32; @@ -30,18 +27,13 @@ void testQueriesAndProperties() { const auto bundle = sycl::get_kernel_bundle(q.get_context()); const auto kernel = bundle.get_kernel(); - const auto local_range = sycl::range<1>(1); - const auto maxWGs = - kernel - .ext_oneapi_get_info( - q, local_range, 0); - const auto wgRange = sycl::range<3>{WorkGroupSize, 1, 1}; - const auto maxWGsWithLimits = - kernel - .ext_oneapi_get_info( - q, wgRange, wgRange.size() * sizeof(int)); + const auto maxWGs = kernel.ext_oneapi_get_info< + sycl::ext::oneapi::experimental::info::kernel_queue_specific:: + max_num_work_group_sync>(q); + const auto wgRange = sycl::range{WorkGroupSize, 1, 1}; + const auto maxWGsWithLimits = kernel.ext_oneapi_get_info< + sycl::ext::oneapi::experimental::info::kernel_queue_specific:: + max_num_work_group_sync>(q, wgRange, wgRange.size() * sizeof(int)); const auto props = sycl::ext::oneapi::experimental::properties{ sycl::ext::oneapi::experimental::use_root_sync}; q.single_task(props, []() {}); @@ -60,11 +52,9 @@ void testRootGroup() { const auto bundle = sycl::get_kernel_bundle(q.get_context()); const auto kernel = bundle.get_kernel(); - const auto maxWGs = - kernel - .ext_oneapi_get_info( - q, WorkGroupSize, 0); + const auto maxWGs = kernel.ext_oneapi_get_info< + sycl::ext::oneapi::experimental::info::kernel_queue_specific:: + max_num_work_group_sync>(q); const auto props = sycl::ext::oneapi::experimental::properties{ sycl::ext::oneapi::experimental::use_root_sync}; sycl::buffer dataBuf{sycl::range{maxWGs * WorkGroupSize}}; @@ -104,11 +94,9 @@ void testRootGroupFunctions() { const auto bundle = sycl::get_kernel_bundle(q.get_context()); const auto kernel = bundle.get_kernel(); - const auto maxWGs = - kernel - .ext_oneapi_get_info( - q, WorkGroupSize, 0); + const auto maxWGs = kernel.ext_oneapi_get_info< + sycl::ext::oneapi::experimental::info::kernel_queue_specific:: + max_num_work_group_sync>(q); const auto props = sycl::ext::oneapi::experimental::properties{ sycl::ext::oneapi::experimental::use_root_sync}; diff --git a/sycl/test/virtual-functions/properties-positive.cpp b/sycl/test/virtual-functions/properties-positive.cpp index 3441e91a4d45c..970e348457388 100644 --- a/sycl/test/virtual-functions/properties-positive.cpp +++ b/sycl/test/virtual-functions/properties-positive.cpp @@ -44,6 +44,10 @@ class SubSubDerived : public SubDerived { int main() { sycl::queue q; + static_assert( + oneapi::is_property_key::value); + static_assert(oneapi::is_property_key::value); + oneapi::properties props_empty{oneapi::assume_indirect_calls}; oneapi::properties props_void{oneapi::assume_indirect_calls_to}; oneapi::properties props_int{oneapi::assume_indirect_calls_to}; From 370a34e8d3336d61ae3a00aeeb576e4c2706e708 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 7 Jan 2025 05:55:30 -0800 Subject: [PATCH 11/30] [SYCL] Add test for newly added depre. warnings Signed-off-by: Hu, Peisen --- sycl/include/sycl/handler.hpp | 65 ++++++++----- sycl/include/sycl/queue.hpp | 27 ++++-- ...ed_single_task_parallel_for_with_props.cpp | 93 +++++++++++++++++++ 3 files changed, 154 insertions(+), 31 deletions(-) create mode 100644 sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index b80df4e304b3c..298b64fc344d2 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -2303,40 +2303,52 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<1> NumWorkItems, PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> parallel_for(range<1> NumWorkItems, + PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<2> NumWorkItems, PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> parallel_for(range<2> NumWorkItems, + PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<3> NumWorkItems, PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> parallel_for(range<3> NumWorkItems, + PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(nd_range Range, PropertiesT Properties, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> parallel_for(nd_range Range, + PropertiesT Properties, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_impl(Range, Properties, std::move(KernelFunc)); } @@ -2400,24 +2412,33 @@ class __SYCL_EXPORT handler { } template - std::enable_if_t::value> - parallel_for(range<1> Range, RestT &&...Rest) { + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> parallel_for(range<1> Range, RestT &&...Rest) { parallel_for(Range, ext::oneapi::experimental::empty_properties_t{}, std::forward(Rest)...); } template - std::enable_if_t::value> - parallel_for(range<2> Range, RestT &&...Rest) { + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> parallel_for(range<2> Range, RestT &&...Rest) { parallel_for(Range, ext::oneapi::experimental::empty_properties_t{}, std::forward(Rest)...); } template - std::enable_if_t::value> - parallel_for(range<3> Range, RestT &&...Rest) { + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> parallel_for(range<3> Range, RestT &&...Rest) { parallel_for(Range, ext::oneapi::experimental::empty_properties_t{}, std::forward(Rest)...); diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 68c176cc4d583..513c3410cf032 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -2907,12 +2907,14 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param KernelFunc is the Kernel functor or lambda template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< detail::AreAllButLastReductions::value && ext::oneapi::experimental::is_property_list::value, - event> - parallel_for_impl(range Range, PropertiesT Properties, - RestT &&...Rest) { + event> parallel_for_impl(range Range, PropertiesT Properties, + RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( @@ -2943,10 +2945,13 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param KernelFunc is the Kernel functor or lambda template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< - ext::oneapi::experimental::is_property_list::value, event> - parallel_for_impl(range Range, event DepEvent, PropertiesT Properties, - RestT &&...Rest) { + ext::oneapi::experimental::is_property_list::value, + event> parallel_for_impl(range Range, event DepEvent, + PropertiesT Properties, RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( @@ -2980,10 +2985,14 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param KernelFunc is the Kernel functor or lambda template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< - ext::oneapi::experimental::is_property_list::value, event> - parallel_for_impl(range Range, const std::vector &DepEvents, - PropertiesT Properties, RestT &&...Rest) { + ext::oneapi::experimental::is_property_list::value, + event> parallel_for_impl(range Range, + const std::vector &DepEvents, + PropertiesT Properties, RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( diff --git a/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp b/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp new file mode 100644 index 0000000000000..e6aaa78097f10 --- /dev/null +++ b/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp @@ -0,0 +1,93 @@ +// Ignore unexpected warnings because for some reason the warnings are emitted +// twice, e.g. once for `single_task`, then for `single_task>>`. +// RUN: %clangxx -fsycl -sycl-std=2020 -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=warning -Xclang -verify-ignore-unexpected=note %s -fsyntax-only -Wall -Wextra +#include + +using namespace sycl; +int main() { + queue Q; + event Ev; + range<1> R1{1}; + range<2> R2(1, 1); + range<3> R3(1, 1, 1); + nd_range<1> NDR1{R1, R1}; + nd_range<2> NDR2{R2, R2}; + nd_range<3> NDR3{R3, R3}; + constexpr auto Props = sycl::ext::oneapi::experimental::properties{}; + + // expected-warning@+1{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.single_task(Props, []() {}); + // expected-warning@+1{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.single_task(Ev, Props, []() {}); + // expected-warning@+1{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.single_task({Ev}, Props, []() {}); + + // expected-warning@+1{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.parallel_for(NDR1, Props, [](nd_item<1>) {}); + + // expected-warning@+2{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.single_task(Props, []() {}); + }); + + // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for(R1, Props, [](id<1>) {}); + }); + // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for(R2, Props, [](id<2>) {}); + }); + // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for(R3, Props, [](id<3>) {}); + }); + + // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for(NDR1, Props, [](nd_item<1>) {}); + }); + // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for(NDR2, Props, [](nd_item<2>) {}); + }); + // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for(NDR3, Props, [](nd_item<3>) {}); + }); + + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R1, Props, + [](sycl::group<1>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R2, Props, + [](sycl::group<2>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R3, Props, + [](sycl::group<3>) {}); + }); + + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R1, R1, Props, + [](sycl::group<1>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R2, R2, Props, + [](sycl::group<2>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R3, R3, Props, + [](sycl::group<3>) {}); + }); + return 0; +} From 870f92f9dce00593d017d60329b660605b9ad9d6 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 7 Jan 2025 07:36:57 -0800 Subject: [PATCH 12/30] Revert "[SYCL] Add test for newly added depre. warnings" This reverts commit 370a34e8d3336d61ae3a00aeeb576e4c2706e708. --- sycl/include/sycl/handler.hpp | 65 +++++-------- sycl/include/sycl/queue.hpp | 27 ++---- ...ed_single_task_parallel_for_with_props.cpp | 93 ------------------- 3 files changed, 31 insertions(+), 154 deletions(-) delete mode 100644 sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 298b64fc344d2..b80df4e304b3c 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -2303,52 +2303,40 @@ class __SYCL_EXPORT handler { template - __SYCL_DEPRECATED( - "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " - "sycl_ext_oneapi_enqueue_functions extension) instead.") - std::enable_if_t::value> parallel_for(range<1> NumWorkItems, - PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value> + parallel_for(range<1> NumWorkItems, PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - __SYCL_DEPRECATED( - "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " - "sycl_ext_oneapi_enqueue_functions extension) instead.") - std::enable_if_t::value> parallel_for(range<2> NumWorkItems, - PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value> + parallel_for(range<2> NumWorkItems, PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - __SYCL_DEPRECATED( - "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " - "sycl_ext_oneapi_enqueue_functions extension) instead.") - std::enable_if_t::value> parallel_for(range<3> NumWorkItems, - PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value> + parallel_for(range<3> NumWorkItems, PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - __SYCL_DEPRECATED( - "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " - "sycl_ext_oneapi_enqueue_functions extension) instead.") - std::enable_if_t::value> parallel_for(nd_range Range, - PropertiesT Properties, - _KERNELFUNCPARAM(KernelFunc)) { + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value> + parallel_for(nd_range Range, PropertiesT Properties, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_impl(Range, Properties, std::move(KernelFunc)); } @@ -2412,33 +2400,24 @@ class __SYCL_EXPORT handler { } template - __SYCL_DEPRECATED( - "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " - "sycl_ext_oneapi_enqueue_functions extension) instead.") - std::enable_if_t::value> parallel_for(range<1> Range, RestT &&...Rest) { + std::enable_if_t::value> + parallel_for(range<1> Range, RestT &&...Rest) { parallel_for(Range, ext::oneapi::experimental::empty_properties_t{}, std::forward(Rest)...); } template - __SYCL_DEPRECATED( - "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " - "sycl_ext_oneapi_enqueue_functions extension) instead.") - std::enable_if_t::value> parallel_for(range<2> Range, RestT &&...Rest) { + std::enable_if_t::value> + parallel_for(range<2> Range, RestT &&...Rest) { parallel_for(Range, ext::oneapi::experimental::empty_properties_t{}, std::forward(Rest)...); } template - __SYCL_DEPRECATED( - "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " - "sycl_ext_oneapi_enqueue_functions extension) instead.") - std::enable_if_t::value> parallel_for(range<3> Range, RestT &&...Rest) { + std::enable_if_t::value> + parallel_for(range<3> Range, RestT &&...Rest) { parallel_for(Range, ext::oneapi::experimental::empty_properties_t{}, std::forward(Rest)...); diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 513c3410cf032..68c176cc4d583 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -2907,14 +2907,12 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param KernelFunc is the Kernel functor or lambda template - __SYCL_DEPRECATED( - "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " - "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< detail::AreAllButLastReductions::value && ext::oneapi::experimental::is_property_list::value, - event> parallel_for_impl(range Range, PropertiesT Properties, - RestT &&...Rest) { + event> + parallel_for_impl(range Range, PropertiesT Properties, + RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( @@ -2945,13 +2943,10 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param KernelFunc is the Kernel functor or lambda template - __SYCL_DEPRECATED( - "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " - "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< - ext::oneapi::experimental::is_property_list::value, - event> parallel_for_impl(range Range, event DepEvent, - PropertiesT Properties, RestT &&...Rest) { + ext::oneapi::experimental::is_property_list::value, event> + parallel_for_impl(range Range, event DepEvent, PropertiesT Properties, + RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( @@ -2985,14 +2980,10 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param KernelFunc is the Kernel functor or lambda template - __SYCL_DEPRECATED( - "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " - "sycl_ext_oneapi_enqueue_functions extension) instead.") std::enable_if_t< - ext::oneapi::experimental::is_property_list::value, - event> parallel_for_impl(range Range, - const std::vector &DepEvents, - PropertiesT Properties, RestT &&...Rest) { + ext::oneapi::experimental::is_property_list::value, event> + parallel_for_impl(range Range, const std::vector &DepEvents, + PropertiesT Properties, RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( diff --git a/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp b/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp deleted file mode 100644 index e6aaa78097f10..0000000000000 --- a/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp +++ /dev/null @@ -1,93 +0,0 @@ -// Ignore unexpected warnings because for some reason the warnings are emitted -// twice, e.g. once for `single_task`, then for `single_task>>`. -// RUN: %clangxx -fsycl -sycl-std=2020 -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=warning -Xclang -verify-ignore-unexpected=note %s -fsyntax-only -Wall -Wextra -#include - -using namespace sycl; -int main() { - queue Q; - event Ev; - range<1> R1{1}; - range<2> R2(1, 1); - range<3> R3(1, 1, 1); - nd_range<1> NDR1{R1, R1}; - nd_range<2> NDR2{R2, R2}; - nd_range<3> NDR3{R3, R3}; - constexpr auto Props = sycl::ext::oneapi::experimental::properties{}; - - // expected-warning@+1{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.single_task(Props, []() {}); - // expected-warning@+1{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.single_task(Ev, Props, []() {}); - // expected-warning@+1{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.single_task({Ev}, Props, []() {}); - - // expected-warning@+1{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.parallel_for(NDR1, Props, [](nd_item<1>) {}); - - // expected-warning@+2{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.single_task(Props, []() {}); - }); - - // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for(R1, Props, [](id<1>) {}); - }); - // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for(R2, Props, [](id<2>) {}); - }); - // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for(R3, Props, [](id<3>) {}); - }); - - // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for(NDR1, Props, [](nd_item<1>) {}); - }); - // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for(NDR2, Props, [](nd_item<2>) {}); - }); - // expected-warning@+2{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for(NDR3, Props, [](nd_item<3>) {}); - }); - - // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for_work_group(R1, Props, - [](sycl::group<1>) {}); - }); - // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for_work_group(R2, Props, - [](sycl::group<2>) {}); - }); - // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for_work_group(R3, Props, - [](sycl::group<3>) {}); - }); - - // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for_work_group(R1, R1, Props, - [](sycl::group<1>) {}); - }); - // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for_work_group(R2, R2, Props, - [](sycl::group<2>) {}); - }); - // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} - Q.submit([&](handler &CGH) { - CGH.parallel_for_work_group(R3, R3, Props, - [](sycl::group<3>) {}); - }); - return 0; -} From 14943ba836ea3b14f685b3412a7ed9bacf52c06c Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 7 Jan 2025 07:42:51 -0800 Subject: [PATCH 13/30] [SYCL] Add new warning test Signed-off-by: Hu, Peisen --- ...ed_single_task_parallel_for_with_props.cpp | 67 +++++++++++++++++++ 1 file changed, 67 insertions(+) create mode 100644 sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp diff --git a/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp b/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp new file mode 100644 index 0000000000000..0550abb3f25b8 --- /dev/null +++ b/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp @@ -0,0 +1,67 @@ +// Ignore unexpected warnings because for some reason the warnings are emitted +// twice, e.g. once for `single_task`, then for `single_task>>`. +// RUN: %clangxx -fsycl -sycl-std=2020 -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=warning -Xclang -verify-ignore-unexpected=note %s -fsyntax-only -Wall -Wextra +#include + +using namespace sycl; +int main() { + queue Q; + event Ev; + range<1> R1{1}; + range<2> R2(1, 1); + range<3> R3(1, 1, 1); + nd_range<1> NDR1{R1, R1}; + nd_range<2> NDR2{R2, R2}; + nd_range<3> NDR3{R3, R3}; + constexpr auto Props = sycl::ext::oneapi::experimental::properties{}; + + // expected-warning@+1{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.single_task(Props, []() {}); + // expected-warning@+1{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.single_task(Ev, Props, []() {}); + // expected-warning@+1{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.single_task({Ev}, Props, []() {}); + + // expected-warning@+1{{'parallel_for' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.parallel_for(NDR1, Props, [](nd_item<1>) {}); + + // expected-warning@+2{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.single_task(Props, []() {}); + }); + + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R1, Props, + [](sycl::group<1>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R2, Props, + [](sycl::group<2>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R3, Props, + [](sycl::group<3>) {}); + }); + + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R1, R1, Props, + [](sycl::group<1>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R2, R2, Props, + [](sycl::group<2>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: Use sycl::ext::oneapi::experimental::parallel_for (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R3, R3, Props, + [](sycl::group<3>) {}); + }); + return 0; +} From d56a99b97e4b2a9eca5496394439f216c590ed21 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 7 Jan 2025 07:44:59 -0800 Subject: [PATCH 14/30] [SYCL] Remove unused variables Signed-off-by: Hu, Peisen --- .../warnings/deprecated_single_task_parallel_for_with_props.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp b/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp index 0550abb3f25b8..e4bbf16091808 100644 --- a/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp +++ b/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp @@ -13,8 +13,6 @@ int main() { range<2> R2(1, 1); range<3> R3(1, 1, 1); nd_range<1> NDR1{R1, R1}; - nd_range<2> NDR2{R2, R2}; - nd_range<3> NDR3{R3, R3}; constexpr auto Props = sycl::ext::oneapi::experimental::properties{}; // expected-warning@+1{{'single_task' is deprecated: Use sycl::ext::oneapi::experimental::single_task (provided in the sycl_ext_oneapi_enqueue_functions extension) instead.}} From 9616ef65fa791102f0ce280f21c5b774dcd0765a Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Wed, 15 Jan 2025 18:25:18 -0800 Subject: [PATCH 15/30] [SYCL] Rewrite multiple-translation-units tests Signed-off-by: Hu, Peisen --- .../Inputs/call.cpp | 22 ++++++++++++++----- 1 file changed, 16 insertions(+), 6 deletions(-) diff --git a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp index 9ce59931405d6..b3310b31793ee 100644 --- a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp +++ b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp @@ -1,17 +1,27 @@ #include "declarations.hpp" +template struct KernelFunctor { + T1 mDeviceStorage; + T2 mDataAcc; + KernelFunctor(T1 &DeviceStorage, T2 &DataAcc) + : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc) {} + + void operator()() const { + auto *Ptr = mDeviceStorage->getAs(); + Ptr->increment(mDataAcc.get_multi_ptr().get()); + } + auto get(oneapi::properties_tag) const { + return oneapi::properties{oneapi::assume_indirect_calls}; + } +}; + int call(sycl::queue Q, storage_t *DeviceStorage, int Init) { int Data = Init; { sycl::buffer DataStorage(&Data, sycl::range{1}); - constexpr oneapi::properties props{oneapi::assume_indirect_calls}; Q.submit([&](sycl::handler &CGH) { sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only); - CGH.single_task(props, [=]() { - auto *Ptr = DeviceStorage->getAs(); - Ptr->increment( - DataAcc.get_multi_ptr().get()); - }); + CGH.single_task(KernelFunctor(DeviceStorage, DataAcc)); }); } From 6335da7ab20b94b68c9f711685a749c840c07c83 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Wed, 15 Jan 2025 18:44:46 -0800 Subject: [PATCH 16/30] [SYCL] Add template keywords Signed-off-by: Hu, Peisen --- .../multiple-translation-units/Inputs/call.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp index b3310b31793ee..b8c934109ab44 100644 --- a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp +++ b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp @@ -7,8 +7,8 @@ template struct KernelFunctor { : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc) {} void operator()() const { - auto *Ptr = mDeviceStorage->getAs(); - Ptr->increment(mDataAcc.get_multi_ptr().get()); + auto *Ptr = mDeviceStorage->template getAs(); + Ptr->increment(mDataAcc.template get_multi_ptr().get()); } auto get(oneapi::properties_tag) const { return oneapi::properties{oneapi::assume_indirect_calls}; From 1fdcd879877544955ac4fd25cf9a00ee6f3445ac Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Wed, 15 Jan 2025 19:12:47 -0800 Subject: [PATCH 17/30] [SYCL] Fix formatting issue Signed-off-by: Hu, Peisen --- .../multiple-translation-units/Inputs/call.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp index b8c934109ab44..4ea4e7cf125b7 100644 --- a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp +++ b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp @@ -8,7 +8,8 @@ template struct KernelFunctor { void operator()() const { auto *Ptr = mDeviceStorage->template getAs(); - Ptr->increment(mDataAcc.template get_multi_ptr().get()); + Ptr->increment( + mDataAcc.template get_multi_ptr().get()); } auto get(oneapi::properties_tag) const { return oneapi::properties{oneapi::assume_indirect_calls}; From 7280d801c9f592eb8557c05b5da38930112bd92d Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 21 Jan 2025 09:40:25 -0800 Subject: [PATCH 18/30] [SYCL] Add missing deprecation warnings in handler.hpp Signed-off-by: Hu, Peisen --- sycl/include/sycl/handler.hpp | 44 ++++++++++++++++++++++------------- 1 file changed, 28 insertions(+), 16 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 1a63b4c8fdc63..4fef8be3da705 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -2304,40 +2304,52 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<1> NumWorkItems, PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::single_task (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> parallel_for(range<1> NumWorkItems, + PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<2> NumWorkItems, PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::single_task (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> parallel_for(range<2> NumWorkItems, + PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<3> NumWorkItems, PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::single_task (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> parallel_for(range<3> NumWorkItems, + PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(nd_range Range, PropertiesT Properties, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::single_task (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t::value> parallel_for(nd_range Range, + PropertiesT Properties, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_impl(Range, Properties, std::move(KernelFunc)); } From 1c6354b269bcf2f9df029eae488b3fca3c60042a Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 21 Jan 2025 12:00:11 -0800 Subject: [PATCH 19/30] [SYCL] Deprecate parallel_for(nd_range with props) in queue.hpp Signed-off-by: Hu, Peisen --- sycl/include/sycl/queue.hpp | 75 ++++++++++++++++++++++++++++++++++--- 1 file changed, 70 insertions(+), 5 deletions(-) diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index cbd2f531f6c78..a01f7213f8183 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -2479,8 +2479,41 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { typename... RestT> std::enable_if_t::value, event> parallel_for(nd_range Range, RestT &&...Rest) { - return parallel_for( - Range, ext::oneapi::experimental::empty_properties_t{}, Rest...); + constexpr detail::code_location CodeLoc = getCodeLocation(); + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.template parallel_for(Range, Rest...); + }, + TlsCodeLocCapture.query()); + } + + /// parallel_for version with a kernel represented as a lambda + nd_range that + /// specifies global, local sizes and offset. + /// + /// \param Range specifies the global and local work spaces of the kernel + /// \param DepEvent is an event that specifies the kernel dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t< + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(nd_range Range, event DepEvent, + PropertiesT Properties, RestT &&...Rest) { + constexpr detail::code_location CodeLoc = getCodeLocation(); + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvent); + CGH.template parallel_for(Range, Properties, Rest...); + }, + TlsCodeLocCapture.query()); } /// parallel_for version with a kernel represented as a lambda + nd_range that @@ -2492,7 +2525,8 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// const KernelType &KernelFunc". template - event parallel_for(nd_range Range, event DepEvent, RestT &&...Rest) { + std::enable_if_t::value, event> + parallel_for(nd_range Range, event DepEvent, RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( @@ -2503,6 +2537,36 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { TlsCodeLocCapture.query()); } + /// parallel_for version with a kernel represented as a lambda + nd_range that + /// specifies global, local sizes and offset. + /// + /// \param Range specifies the global and local work spaces of the kernel + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t< + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(nd_range Range, + const std::vector &DepEvents, + PropertiesT Properties, RestT &&...Rest) { + constexpr detail::code_location CodeLoc = getCodeLocation(); + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvents); + CGH.template parallel_for(Range, Properties, Rest...); + }, + TlsCodeLocCapture.query()); + } + /// parallel_for version with a kernel represented as a lambda + nd_range that /// specifies global, local sizes and offset. /// @@ -2513,8 +2577,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// const KernelType &KernelFunc". template - event parallel_for(nd_range Range, const std::vector &DepEvents, - RestT &&...Rest) { + std::enable_if_t::value, event> + parallel_for(nd_range Range, const std::vector &DepEvents, + RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( From bab42403c47c2004e33f22b6f13bb86e6fefd1f0 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 21 Jan 2025 13:26:46 -0800 Subject: [PATCH 20/30] [SYCL] Deprecate parallel_for(range with props) in queue.hpp Signed-off-by: Hu, Peisen --- sycl/include/sycl/queue.hpp | 180 ++++++++++++++++++++++++++++++++++++ 1 file changed, 180 insertions(+) diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index a01f7213f8183..f04b8a1e88960 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -2240,6 +2240,25 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { CodeLoc); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<1> Range, PropertiesT Properties, + RestT &&...Rest) { + return parallel_for_impl(Range, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2251,6 +2270,25 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<2> Range, PropertiesT Properties, + RestT &&...Rest) { + return parallel_for_impl(Range, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2262,6 +2300,25 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<3> Range, PropertiesT Properties, + RestT &&...Rest) { + return parallel_for_impl(Range, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2273,6 +2330,26 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvent is an event that specifies the kernel dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<1> Range, event DepEvent, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvent, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2285,6 +2362,26 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, DepEvent, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvent is an event that specifies the kernel dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<2> Range, event DepEvent, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvent, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2297,6 +2394,26 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, DepEvent, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvent is an event that specifies the kernel dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<3> Range, event DepEvent, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvent, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2309,6 +2426,27 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, DepEvent, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<1> Range, const std::vector &DepEvents, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvents, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2323,6 +2461,27 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, DepEvents, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<2> Range, const std::vector &DepEvents, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvents, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2337,6 +2496,27 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, DepEvents, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED( + "Use sycl::ext::oneapi::experimental::parallel_for (provided in the " + "sycl_ext_oneapi_enqueue_functions extension) instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<3> Range, const std::vector &DepEvents, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvents, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// From 4c126d28a3b9d33e61c554e5e98ee724c7f854c4 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Wed, 22 Jan 2025 22:22:04 -0800 Subject: [PATCH 21/30] [SYCL] E2E Test fix for last update #1 Signed-off-by: Hu, Peisen --- sycl/test-e2e/Basic/kernel_max_wg_size.cpp | 14 +++- sycl/test-e2e/Basic/sub_group_size_prop.cpp | 37 --------- .../cluster_launch_parallel_for.cpp | 76 +++++++++++-------- 3 files changed, 54 insertions(+), 73 deletions(-) diff --git a/sycl/test-e2e/Basic/kernel_max_wg_size.cpp b/sycl/test-e2e/Basic/kernel_max_wg_size.cpp index 55c993734df6a..eb0ff5483bae4 100644 --- a/sycl/test-e2e/Basic/kernel_max_wg_size.cpp +++ b/sycl/test-e2e/Basic/kernel_max_wg_size.cpp @@ -29,6 +29,16 @@ __attribute__((noinline)) void f(int *result, nd_item<1> &index) { result[index.get_global_id()] = index.get_global_id(); } +struct KernelFunctor { + int *mResult; + KernelFunctor(int *result) : mResult(result) {} + + void operator()(nd_item<1> index) const { f(mResult, index); } + auto get(syclex::properties_tag) const { + return syclex::properties{intelex::grf_size<256>}; + } +}; + int main() { queue myQueue; auto myContext = myQueue.get_context(); @@ -46,11 +56,9 @@ int main() { nd_range myRange{range{maxWgSize}, range{maxWgSize}}; int *result = sycl::malloc_shared(maxWgSize, myQueue); - syclex::properties kernelProperties{intelex::grf_size<256>}; myQueue.submit([&](handler &cgh) { cgh.use_kernel_bundle(myBundle); - cgh.parallel_for(myRange, kernelProperties, - ([=](nd_item<1> index) { f(result, index); })); + cgh.parallel_for(myRange, KernelFunctor(result)); }); myQueue.wait(); diff --git a/sycl/test-e2e/Basic/sub_group_size_prop.cpp b/sycl/test-e2e/Basic/sub_group_size_prop.cpp index 6da86acd09c45..ae8281903a92b 100644 --- a/sycl/test-e2e/Basic/sub_group_size_prop.cpp +++ b/sycl/test-e2e/Basic/sub_group_size_prop.cpp @@ -44,33 +44,12 @@ void test(queue &Queue, const std::vector SupportedSGSizes) { return; } - auto Props = ext::oneapi::experimental::properties{ - ext::oneapi::experimental::sub_group_size}; - nd_range<1> NdRange(SGSize * 4, SGSize * 2); size_t ReadSubGroupSize = 0; { buffer ReadSubGroupSizeBuf(&ReadSubGroupSize, range(1)); - Queue.submit([&](handler &CGH) { - accessor ReadSubGroupSizeBufAcc{ReadSubGroupSizeBuf, CGH, - sycl::write_only, sycl::no_init}; - - CGH.parallel_for>( - NdRange, Props, [=](nd_item<1> NdItem) { - auto SG = NdItem.get_sub_group(); - if (NdItem.get_global_linear_id() == 0) - ReadSubGroupSizeBufAcc[0] = SG.get_local_linear_range(); - }); - }); - } - assert(ReadSubGroupSize == SGSize && "Failed check for function."); - - ReadSubGroupSize = 0; - { - buffer ReadSubGroupSizeBuf(&ReadSubGroupSize, range(1)); - Queue.submit([&](handler &CGH) { accessor ReadSubGroupSizeBufAcc{ReadSubGroupSizeBuf, CGH, sycl::write_only, sycl::no_init}; @@ -81,22 +60,6 @@ void test(queue &Queue, const std::vector SupportedSGSizes) { }); } assert(ReadSubGroupSize == SGSize && "Failed check for functor."); - - ReadSubGroupSize = 0; - { - buffer ReadSubGroupSizeBuf(&ReadSubGroupSize, range(1)); - - Queue.submit([&](handler &CGH) { - accessor ReadSubGroupSizeBufAcc{ReadSubGroupSizeBuf, CGH, - sycl::write_only, sycl::no_init}; - KernelFunctorWithSGSizeProp KernelFunctor{ReadSubGroupSizeBufAcc}; - - CGH.parallel_for>(NdRange, Props, - KernelFunctor); - }); - } - assert(ReadSubGroupSize == SGSize && - "Failed check for functor and properties."); } int main() { diff --git a/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp b/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp index e37d4ea1f1fb3..48273bffe277a 100644 --- a/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp +++ b/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp @@ -10,6 +10,46 @@ #include +template struct KernelFunctor { + int *mCorrectResultFlag; + cuda::cluster_size mClusterDims; + sycl::range mClusterRange; + KernelFunctor(int *CorrectResultFlag, cuda::cluster_size ClusterDims, + sycl::range ClusterRange) + : mCorrectResultFlag(CorrectResultFlag), mClusterDims(ClusterDims), + mClusterRange(ClusterRange) {} + + void operator()(sycl::nd_item It) const { + uint32_t ClusterDimX, ClusterDimY, ClusterDimZ; +// Temporary solution till cluster group class is implemented +#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_CUDA_ARCH__) && \ + (__SYCL_CUDA_ARCH__ >= 900) + asm volatile("\n\t" + "mov.u32 %0, %%cluster_nctaid.x; \n\t" + "mov.u32 %1, %%cluster_nctaid.y; \n\t" + "mov.u32 %2, %%cluster_nctaid.z; \n\t" + : "=r"(ClusterDimZ), "=r"(ClusterDimY), "=r"(ClusterDimX)); +#endif + if constexpr (Dim == 1) { + if (ClusterDimZ == mClusterRange[0] && ClusterDimY == 1 && + ClusterDimX == 1) { + *mCorrectResultFlag = 1; + } + } else if constexpr (Dim == 2) { + if (ClusterDimZ == mClusterRange[1] && ClusterDimY == mClusterRange[0] && + ClusterDimX == 1) { + *mCorrectResultFlag = 1; + } + } else { + if (ClusterDimZ == mClusterRange[2] && ClusterDimY == mClusterRange[1] && + ClusterDimX == mClusterRange[0]) { + *mCorrectResultFlag = 1; + } + } + } + auto get(properties_tag) const { return properties{mClusterDims}; } +}; + template int test_cluster_launch_parallel_for(sycl::queue &Queue, sycl::range GlobalRange, @@ -18,45 +58,15 @@ int test_cluster_launch_parallel_for(sycl::queue &Queue, using namespace sycl::ext::oneapi::experimental; cuda::cluster_size ClusterDims(ClusterRange); - properties ClusterLaunchProperty{ClusterDims}; int *CorrectResultFlag = sycl::malloc_device(1, Queue); Queue.memset(CorrectResultFlag, 0, sizeof(int)).wait(); Queue .submit([&](sycl::handler &CGH) { - CGH.parallel_for(sycl::nd_range(GlobalRange, LocalRange), - ClusterLaunchProperty, [=](sycl::nd_item It) { - uint32_t ClusterDimX, ClusterDimY, ClusterDimZ; -// Temporary solution till cluster group class is implemented -#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_CUDA_ARCH__) && \ - (__SYCL_CUDA_ARCH__ >= 900) - asm volatile("\n\t" - "mov.u32 %0, %%cluster_nctaid.x; \n\t" - "mov.u32 %1, %%cluster_nctaid.y; \n\t" - "mov.u32 %2, %%cluster_nctaid.z; \n\t" - : "=r"(ClusterDimZ), "=r"(ClusterDimY), - "=r"(ClusterDimX)); -#endif - if constexpr (Dim == 1) { - if (ClusterDimZ == ClusterRange[0] && - ClusterDimY == 1 && ClusterDimX == 1) { - *CorrectResultFlag = 1; - } - } else if constexpr (Dim == 2) { - if (ClusterDimZ == ClusterRange[1] && - ClusterDimY == ClusterRange[0] && - ClusterDimX == 1) { - *CorrectResultFlag = 1; - } - } else { - if (ClusterDimZ == ClusterRange[2] && - ClusterDimY == ClusterRange[1] && - ClusterDimX == ClusterRange[0]) { - *CorrectResultFlag = 1; - } - } - }); + CGH.parallel_for( + sycl::nd_range(GlobalRange, LocalRange), + KernelFunctor(CorrectResultFlag, ClusterDims, ClusterRange)); }) .wait_and_throw(); From ce067bbedd143146d5e538b2d86c3b46dd788803 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Wed, 22 Jan 2025 22:28:32 -0800 Subject: [PATCH 22/30] [SYCL] Minor fix Signed-off-by: Hu, Peisen --- sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp b/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp index 48273bffe277a..2a2343d585bed 100644 --- a/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp +++ b/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp @@ -11,6 +11,8 @@ #include template struct KernelFunctor { + using namespace sycl::ext::oneapi::experimental; + int *mCorrectResultFlag; cuda::cluster_size mClusterDims; sycl::range mClusterRange; From 7cb9d8e1bfe9c395572a26e0384b59fc38a676d0 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Wed, 22 Jan 2025 23:12:37 -0800 Subject: [PATCH 23/30] [SYCL] E2E Test fix for last update #2 Signed-off-by: Hu, Peisen --- .../enqueueLaunchCustom_check_event_deps.cpp | 25 +++++-- sycl/test-e2e/DeviceCodeSplit/grf.cpp | 13 +++- sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp | 68 ------------------- .../VirtualFunctions/misc/group-barrier.cpp | 27 ++++++-- .../misc/range-non-uniform-vf-2.cpp | 29 +++++--- .../misc/range-non-uniform-vf.cpp | 25 +++++-- .../misc/range-uniform-vf.cpp | 21 ++++-- 7 files changed, 103 insertions(+), 105 deletions(-) diff --git a/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp b/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp index 8900d10328871..622bf8f856ed9 100644 --- a/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp +++ b/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp @@ -24,6 +24,22 @@ template void dummy_kernel(T *Input, int N, sycl::nd_item<1> It) { #endif } +template struct KernelFunctor { + using namespace sycl::ext::oneapi::experimental; + + T mAcc; + properties mClusterLaunchProperty; + KernelFunctor(properties ClusterLaunchProperty, T Acc) + : mClusterLaunchProperty(ClusterLaunchProperty), mAcc(Acc) {} + + void operator()(sycl::nd_item<1> It) const { + dummy_kernel( + mAcc.template get_multi_ptr().get(), 4096, + It); + } + auto get(properties_tag) const { return mClusterLaunchProperty; } +}; + int main() { std::vector HostArray(4096, -20); @@ -46,13 +62,8 @@ int main() { cuda::cluster_size ClusterDims(sycl::range{2}); properties ClusterLaunchProperty{ClusterDims}; auto Acc = Buff.template get_access(CGH); - CGH.parallel_for( - sycl::nd_range({4096}, {32}), ClusterLaunchProperty, - [=](sycl::nd_item<1> It) { - dummy_kernel( - Acc.get_multi_ptr().get(), 4096, - It); - }); + CGH.parallel_for(sycl::nd_range({4096}, {32}), + KernelFunctor(ClusterLaunchProperty, Acc)); }); Queue.submit([&](sycl::handler &CGH) { auto Acc = Buff.template get_access(CGH); diff --git a/sycl/test-e2e/DeviceCodeSplit/grf.cpp b/sycl/test-e2e/DeviceCodeSplit/grf.cpp index 62f1a76a5f017..4717d6fa71ea0 100644 --- a/sycl/test-e2e/DeviceCodeSplit/grf.cpp +++ b/sycl/test-e2e/DeviceCodeSplit/grf.cpp @@ -67,6 +67,15 @@ bool checkResult(const std::vector &A, int Inc) { return true; } +template struct KernelFunctor { + T mPA; + properties mProp; + KernelFunctor(properties Prop, T PA) : mProp(Prop), mPA(PA) {} + + void operator()(id<1> i) const { PA[i] += 2; } + auto get(properties_tag) const { return mProp; } +}; + int main(void) { constexpr unsigned Size = 32; constexpr unsigned VL = 16; @@ -122,8 +131,8 @@ int main(void) { auto e = q.submit([&](handler &cgh) { auto PA = bufa.get_access(cgh); - cgh.parallel_for( - Size, prop, [=](id<1> i) { PA[i] += 2; }); + cgh.parallel_for(Size, + KernelFunctor(prop, PA)); }); e.wait(); } catch (sycl::exception const &e) { diff --git a/sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp b/sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp index 7c0bfe5161530..adaf6e1977ea4 100644 --- a/sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp +++ b/sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp @@ -39,9 +39,6 @@ void test(queue &Queue, const std::vector SupportedSGSizes) { return; } - auto Props = ext::oneapi::experimental::properties{ - ext::oneapi::experimental::sub_group_size}; - nd_range<1> NdRange(SGSize * 4, SGSize * 2); size_t ReadSubGroupSize = 0; @@ -49,39 +46,6 @@ void test(queue &Queue, const std::vector SupportedSGSizes) { buffer ReadSubGroupSizeBuf(&ReadSubGroupSize, range(1)); ReadSubGroupSizeBuf.set_write_back(false); - { - exp_ext::command_graph Graph{ - Queue.get_context(), - Queue.get_device(), - {exp_ext::property::graph::assume_buffer_outlives_graph{}}}; - - add_node(Graph, Queue, [&](handler &CGH) { - accessor ReadSubGroupSizeBufAcc{ReadSubGroupSizeBuf, CGH, - sycl::write_only, sycl::no_init}; - - CGH.parallel_for>( - NdRange, Props, [=](nd_item<1> NdItem) { - auto SG = NdItem.get_sub_group(); - if (NdItem.get_global_linear_id() == 0) - ReadSubGroupSizeBufAcc[0] = SG.get_local_linear_range(); - }); - }); - - auto ExecGraph = Graph.finalize(); - Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); }); - Queue.wait_and_throw(); - } - - host_accessor HostAcc(ReadSubGroupSizeBuf); - ReadSubGroupSize = HostAcc[0]; - } - assert(ReadSubGroupSize == SGSize && "Failed check for function."); - - ReadSubGroupSize = 0; - { - buffer ReadSubGroupSizeBuf(&ReadSubGroupSize, range(1)); - ReadSubGroupSizeBuf.set_write_back(false); - { exp_ext::command_graph Graph{ Queue.get_context(), @@ -107,38 +71,6 @@ void test(queue &Queue, const std::vector SupportedSGSizes) { ReadSubGroupSize = HostAcc[0]; } assert(ReadSubGroupSize == SGSize && "Failed check for functor."); - - ReadSubGroupSize = 0; - { - buffer ReadSubGroupSizeBuf(&ReadSubGroupSize, range(1)); - ReadSubGroupSizeBuf.set_write_back(false); - - { - exp_ext::command_graph Graph{ - Queue.get_context(), - Queue.get_device(), - {exp_ext::property::graph::assume_buffer_outlives_graph{}}}; - - add_node(Graph, Queue, [&](handler &CGH) { - accessor ReadSubGroupSizeBufAcc{ReadSubGroupSizeBuf, CGH, - sycl::write_only, sycl::no_init}; - KernelFunctorWithSGSizeProp KernelFunctor{ - ReadSubGroupSizeBufAcc}; - - CGH.parallel_for>( - NdRange, Props, KernelFunctor); - }); - - auto ExecGraph = Graph.finalize(); - Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); }); - Queue.wait_and_throw(); - } - - host_accessor HostAcc(ReadSubGroupSizeBuf); - ReadSubGroupSize = HostAcc[0]; - } - assert(ReadSubGroupSize == SGSize && - "Failed check for functor and properties."); } int main() { diff --git a/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp b/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp index f01a25d4179f4..48db619d94081 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp @@ -101,6 +101,24 @@ class MultiplyOp : public BaseOp { } }; +template struct KernelFunctor { + T1 mDeviceStorage; + T2 mDataAcc; + T3 mLocalAcc; + KernelFunctor(T1 DeviceStorage, T2 DataAcc, T3 LocalAcc) + : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc), mLocalAcc(LocalAcc) {} + + void operator()(sycl::nd_item<1> It) const { + auto *Ptr = mDeviceStorage->template getAs(); + mDataAcc[It.get_global_id()] = Ptr->apply( + mLocalAcc.template get_multi_ptr().get(), + It.get_group()); + } + auto get(oneapi::properties_tag) const { + return oneapi::properties{oneapi::assume_indirect_calls}; + } +}; + int main() try { using storage_t = obj_storage_t; @@ -113,7 +131,6 @@ int main() try { sycl::range G{16}; sycl::range L{4}; - constexpr oneapi::properties props{oneapi::assume_indirect_calls}; for (unsigned TestCase = 0; TestCase < 2; ++TestCase) { sycl::buffer DataStorage(G); @@ -126,12 +143,8 @@ int main() try { q.submit([&](sycl::handler &CGH) { sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write); sycl::local_accessor LocalAcc(L, CGH); - CGH.parallel_for(sycl::nd_range{G, L}, props, [=](auto It) { - auto *Ptr = DeviceStorage->getAs(); - DataAcc[It.get_global_id()] = Ptr->apply( - LocalAcc.get_multi_ptr().get(), - It.get_group()); - }); + CGH.parallel_for(sycl::nd_range{G, L}, + KernelFunctor(DeviceStorage, DataAcc, LocalAcc)); }).wait_and_throw(); auto *Ptr = HostStorage.construct(TestCase); diff --git a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp index 45b56916a5c1d..6a6e5df031a2d 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp @@ -44,6 +44,25 @@ class OpB : public BaseOp { virtual int bar(int V) { return V / 2; } }; +template struct KernelFunctor { + T1 mDeviceStorage; + T2 mDataAcc; + KernelFunctor(T1 DeviceStorage, T2 DataAcc) + : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc) {} + + void operator()(sycl::id<1> It) const { + // Select method that corresponds to this work-item + auto *Ptr = mDeviceStorage->template getAs(); + if (It % 2) + mDataAcc[It] = Ptr->foo(mDataAcc[It]); + else + mDataAcc[It] = Ptr->bar(mDataAcc[It]); + } + auto get(oneapi::properties_tag) const { + return oneapi::properties{oneapi::assume_indirect_calls}; + } +}; + int main() try { using storage_t = obj_storage_t; @@ -54,7 +73,6 @@ int main() try { auto *DeviceStorage = sycl::malloc_shared(1, q); sycl::range R{1024}; - constexpr oneapi::properties props{oneapi::assume_indirect_calls}; for (size_t TestCase = 0; TestCase < 2; ++TestCase) { std::vector HostData(R.size()); std::iota(HostData.begin(), HostData.end(), 0); @@ -69,14 +87,7 @@ int main() try { q.submit([&](sycl::handler &CGH) { sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write); - CGH.parallel_for(R, props, [=](auto It) { - // Select method that corresponds to this work-item - auto *Ptr = DeviceStorage->template getAs(); - if (It % 2) - DataAcc[It] = Ptr->foo(DataAcc[It]); - else - DataAcc[It] = Ptr->bar(DataAcc[It]); - }); + CGH.parallel_for(R, KernelFunctor(DeviceStorage, DataAcc)); }); BaseOp *Ptr = HostStorage.construct(TestCase); diff --git a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp index 453a3aee81fa6..e689cde95a9bf 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp @@ -41,6 +41,23 @@ class RoundOp : public BaseOp { virtual float apply(float V) { return sycl::round(V); } }; +template struct KernelFunctor { + T1 mDeviceStorage; + T2 mDataAcc; + KernelFunctor(T1 DeviceStorage, T2 DataAcc) + : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc) {} + + void operator()(sycl::id<1> It) const { + // Select an object that corresponds to this work-item + auto Ind = It % 3; + auto *Ptr = mDeviceStorage[Ind].template getAs(); + mDataAcc[It] = Ptr->apply(mDataAcc[It]); + } + auto get(oneapi::properties_tag) const { + return oneapi::properties{oneapi::assume_indirect_calls}; + } +}; + int main() try { using storage_t = obj_storage_t; @@ -51,7 +68,6 @@ int main() try { auto *DeviceStorage = sycl::malloc_shared(3, q); sycl::range R{1024}; - constexpr oneapi::properties props{oneapi::assume_indirect_calls}; { std::vector HostData(R.size()); for (size_t I = 1; I < HostData.size(); ++I) @@ -69,12 +85,7 @@ int main() try { q.submit([&](sycl::handler &CGH) { sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write); - CGH.parallel_for(R, props, [=](auto it) { - // Select an object that corresponds to this work-item - auto Ind = it % 3; - auto *Ptr = DeviceStorage[Ind].template getAs(); - DataAcc[it] = Ptr->apply(DataAcc[it]); - }); + CGH.parallel_for(R, KernelFunctor(DeviceStorage, DataAcc)); }); BaseOp *Ptr[] = {HostStorage[0].construct(0), diff --git a/sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp b/sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp index 66db6a0c5af7a..2d5facfd35a81 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp @@ -41,6 +41,21 @@ class RoundOp : public BaseOp { virtual float apply(float V) { return sycl::round(V); } }; +template struct KernelFunctor { + T1 mDeviceStorage; + T2 mDataAcc; + KernelFunctor(T1 DeviceStorage, T2 DataAcc) + : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc) {} + + void operator()(sycl::id<1> It) const { + auto *Ptr = mDeviceStorage->getAs(); + mDataAcc[It] = Ptr->apply(mDataAcc[It]); + } + auto get(oneapi::properties_tag) const { + return oneapi::properties{oneapi::assume_indirect_calls}; + } +}; + int main() try { using storage_t = obj_storage_t; @@ -51,7 +66,6 @@ int main() try { auto *DeviceStorage = sycl::malloc_shared(1, q); sycl::range R{1024}; - constexpr oneapi::properties props{oneapi::assume_indirect_calls}; for (unsigned TestCase = 0; TestCase < 3; ++TestCase) { std::vector HostData(R.size()); for (size_t I = 1; I < HostData.size(); ++I) @@ -67,10 +81,7 @@ int main() try { q.submit([&](sycl::handler &CGH) { sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write); - CGH.parallel_for(R, props, [=](auto it) { - auto *Ptr = DeviceStorage->getAs(); - DataAcc[it] = Ptr->apply(DataAcc[it]); - }); + CGH.parallel_for(R, KernelFunctor(DeviceStorage, DataAcc)); }); auto *Ptr = HostStorage.construct(TestCase); From 4314e4b43f15c5e6f75496a6caf6e35eb37674af Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Wed, 22 Jan 2025 23:20:11 -0800 Subject: [PATCH 24/30] [SYCL] Minor fix 2 Signed-off-by: Hu, Peisen --- .../ClusterLaunch/cluster_launch_parallel_for.cpp | 11 ++++++----- .../enqueueLaunchCustom_check_event_deps.cpp | 11 ++++++----- 2 files changed, 12 insertions(+), 10 deletions(-) diff --git a/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp b/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp index 2a2343d585bed..05ddc30ec3e24 100644 --- a/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp +++ b/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp @@ -11,12 +11,11 @@ #include template struct KernelFunctor { - using namespace sycl::ext::oneapi::experimental; - int *mCorrectResultFlag; - cuda::cluster_size mClusterDims; + sycl::ext::oneapi::experimental::cuda::cluster_size mClusterDims; sycl::range mClusterRange; - KernelFunctor(int *CorrectResultFlag, cuda::cluster_size ClusterDims, + KernelFunctor(int *CorrectResultFlag, + sycl::ext::oneapi::experimental::cuda::cluster_size ClusterDims, sycl::range ClusterRange) : mCorrectResultFlag(CorrectResultFlag), mClusterDims(ClusterDims), mClusterRange(ClusterRange) {} @@ -49,7 +48,9 @@ template struct KernelFunctor { } } } - auto get(properties_tag) const { return properties{mClusterDims}; } + auto get(sycl::ext::oneapi::experimental::properties_tag) const { + return sycl::ext::oneapi::experimental::properties{mClusterDims}; + } }; template diff --git a/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp b/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp index 622bf8f856ed9..7b777b1c0fbe6 100644 --- a/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp +++ b/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp @@ -25,11 +25,10 @@ template void dummy_kernel(T *Input, int N, sycl::nd_item<1> It) { } template struct KernelFunctor { - using namespace sycl::ext::oneapi::experimental; - T mAcc; - properties mClusterLaunchProperty; - KernelFunctor(properties ClusterLaunchProperty, T Acc) + sycl::ext::oneapi::experimental::properties mClusterLaunchProperty; + KernelFunctor( + sycl::ext::oneapi::experimental::properties ClusterLaunchProperty, T Acc) : mClusterLaunchProperty(ClusterLaunchProperty), mAcc(Acc) {} void operator()(sycl::nd_item<1> It) const { @@ -37,7 +36,9 @@ template struct KernelFunctor { mAcc.template get_multi_ptr().get(), 4096, It); } - auto get(properties_tag) const { return mClusterLaunchProperty; } + auto get(sycl::ext::oneapi::experimental::properties_tag) const { + return mClusterLaunchProperty; + } }; int main() { From 82ec02668120fa930d06fb45172e839d01424c27 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Wed, 22 Jan 2025 23:58:02 -0800 Subject: [PATCH 25/30] [SYCL] Minor fix 3 Signed-off-by: Hu, Peisen --- .../ClusterLaunch/cluster_launch_parallel_for.cpp | 7 +++---- .../enqueueLaunchCustom_check_event_deps.cpp | 9 ++++----- sycl/test-e2e/DeviceCodeSplit/grf.cpp | 12 ++++++------ .../VirtualFunctions/misc/range-non-uniform-vf-2.cpp | 2 +- .../VirtualFunctions/misc/range-non-uniform-vf.cpp | 2 +- .../VirtualFunctions/misc/range-uniform-vf.cpp | 2 +- 6 files changed, 16 insertions(+), 18 deletions(-) diff --git a/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp b/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp index 05ddc30ec3e24..a98f1fad23300 100644 --- a/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp +++ b/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp @@ -10,12 +10,11 @@ #include -template struct KernelFunctor { +template struct KernelFunctor { int *mCorrectResultFlag; - sycl::ext::oneapi::experimental::cuda::cluster_size mClusterDims; + T mClusterDims; sycl::range mClusterRange; - KernelFunctor(int *CorrectResultFlag, - sycl::ext::oneapi::experimental::cuda::cluster_size ClusterDims, + KernelFunctor(int *CorrectResultFlag, T ClusterDims, sycl::range ClusterRange) : mCorrectResultFlag(CorrectResultFlag), mClusterDims(ClusterDims), mClusterRange(ClusterRange) {} diff --git a/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp b/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp index 7b777b1c0fbe6..0460defa72104 100644 --- a/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp +++ b/sycl/test-e2e/ClusterLaunch/enqueueLaunchCustom_check_event_deps.cpp @@ -24,11 +24,10 @@ template void dummy_kernel(T *Input, int N, sycl::nd_item<1> It) { #endif } -template struct KernelFunctor { - T mAcc; - sycl::ext::oneapi::experimental::properties mClusterLaunchProperty; - KernelFunctor( - sycl::ext::oneapi::experimental::properties ClusterLaunchProperty, T Acc) +template struct KernelFunctor { + T1 mAcc; + T2 mClusterLaunchProperty; + KernelFunctor(T2 ClusterLaunchProperty, T1 Acc) : mClusterLaunchProperty(ClusterLaunchProperty), mAcc(Acc) {} void operator()(sycl::nd_item<1> It) const { diff --git a/sycl/test-e2e/DeviceCodeSplit/grf.cpp b/sycl/test-e2e/DeviceCodeSplit/grf.cpp index 4717d6fa71ea0..28ff4c0c4c80c 100644 --- a/sycl/test-e2e/DeviceCodeSplit/grf.cpp +++ b/sycl/test-e2e/DeviceCodeSplit/grf.cpp @@ -67,12 +67,12 @@ bool checkResult(const std::vector &A, int Inc) { return true; } -template struct KernelFunctor { - T mPA; - properties mProp; - KernelFunctor(properties Prop, T PA) : mProp(Prop), mPA(PA) {} +template struct KernelFunctor { + T1 mPA; + T2 mProp; + KernelFunctor(T1 PA, T2 Prop) : mPA(PA), mProp(Prop) {} - void operator()(id<1> i) const { PA[i] += 2; } + void operator()(id<1> i) const { mPA[i] += 2; } auto get(properties_tag) const { return mProp; } }; @@ -132,7 +132,7 @@ int main(void) { auto e = q.submit([&](handler &cgh) { auto PA = bufa.get_access(cgh); cgh.parallel_for(Size, - KernelFunctor(prop, PA)); + KernelFunctor(PA, prop)); }); e.wait(); } catch (sycl::exception const &e) { diff --git a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp index 6a6e5df031a2d..93af9fa9692d3 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp @@ -50,7 +50,7 @@ template struct KernelFunctor { KernelFunctor(T1 DeviceStorage, T2 DataAcc) : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc) {} - void operator()(sycl::id<1> It) const { + void operator()(sycl::item<1> It) const { // Select method that corresponds to this work-item auto *Ptr = mDeviceStorage->template getAs(); if (It % 2) diff --git a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp index e689cde95a9bf..56b233dbff8cb 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp @@ -47,7 +47,7 @@ template struct KernelFunctor { KernelFunctor(T1 DeviceStorage, T2 DataAcc) : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc) {} - void operator()(sycl::id<1> It) const { + void operator()(sycl::item<1> It) const { // Select an object that corresponds to this work-item auto Ind = It % 3; auto *Ptr = mDeviceStorage[Ind].template getAs(); diff --git a/sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp b/sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp index 2d5facfd35a81..f624dcb26d66a 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp @@ -48,7 +48,7 @@ template struct KernelFunctor { : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc) {} void operator()(sycl::id<1> It) const { - auto *Ptr = mDeviceStorage->getAs(); + auto *Ptr = mDeviceStorage->template getAs(); mDataAcc[It] = Ptr->apply(mDataAcc[It]); } auto get(oneapi::properties_tag) const { From 91af441f9d95dd9b1f4a07f4c087f088e579eb3c Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Thu, 23 Jan 2025 07:44:55 -0800 Subject: [PATCH 26/30] [SYCL] Minor fix 3 Signed-off-by: Hu, Peisen --- .../ClusterLaunch/cluster_launch_parallel_for.cpp | 13 ++++++++----- sycl/test-e2e/DeviceCodeSplit/grf.cpp | 2 +- 2 files changed, 9 insertions(+), 6 deletions(-) diff --git a/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp b/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp index a98f1fad23300..01db70b11464a 100644 --- a/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp +++ b/sycl/test-e2e/ClusterLaunch/cluster_launch_parallel_for.cpp @@ -12,11 +12,12 @@ template struct KernelFunctor { int *mCorrectResultFlag; - T mClusterDims; + T mClusterLaunchProperty; sycl::range mClusterRange; - KernelFunctor(int *CorrectResultFlag, T ClusterDims, + KernelFunctor(int *CorrectResultFlag, T ClusterLaunchProperty, sycl::range ClusterRange) - : mCorrectResultFlag(CorrectResultFlag), mClusterDims(ClusterDims), + : mCorrectResultFlag(CorrectResultFlag), + mClusterLaunchProperty(ClusterLaunchProperty), mClusterRange(ClusterRange) {} void operator()(sycl::nd_item It) const { @@ -48,7 +49,7 @@ template struct KernelFunctor { } } auto get(sycl::ext::oneapi::experimental::properties_tag) const { - return sycl::ext::oneapi::experimental::properties{mClusterDims}; + return mClusterLaunchProperty; } }; @@ -60,6 +61,7 @@ int test_cluster_launch_parallel_for(sycl::queue &Queue, using namespace sycl::ext::oneapi::experimental; cuda::cluster_size ClusterDims(ClusterRange); + properties ClusterLaunchProperty{ClusterDims}; int *CorrectResultFlag = sycl::malloc_device(1, Queue); Queue.memset(CorrectResultFlag, 0, sizeof(int)).wait(); @@ -68,7 +70,8 @@ int test_cluster_launch_parallel_for(sycl::queue &Queue, .submit([&](sycl::handler &CGH) { CGH.parallel_for( sycl::nd_range(GlobalRange, LocalRange), - KernelFunctor(CorrectResultFlag, ClusterDims, ClusterRange)); + KernelFunctor( + CorrectResultFlag, ClusterLaunchProperty, ClusterRange)); }) .wait_and_throw(); diff --git a/sycl/test-e2e/DeviceCodeSplit/grf.cpp b/sycl/test-e2e/DeviceCodeSplit/grf.cpp index 28ff4c0c4c80c..4080049f665af 100644 --- a/sycl/test-e2e/DeviceCodeSplit/grf.cpp +++ b/sycl/test-e2e/DeviceCodeSplit/grf.cpp @@ -67,7 +67,7 @@ bool checkResult(const std::vector &A, int Inc) { return true; } -template struct KernelFunctor { +template struct KernelFunctor { T1 mPA; T2 mProp; KernelFunctor(T1 PA, T2 Prop) : mPA(PA), mProp(Prop) {} From 6e9f228c8dc8ec10f45f34caa6b95ecf129ee51f Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Thu, 23 Jan 2025 09:44:06 -0800 Subject: [PATCH 27/30] [SYCL] Minor fix 4 Signed-off-by: Hu, Peisen --- sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp index 93af9fa9692d3..c62e65c7b9d69 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp @@ -50,7 +50,7 @@ template struct KernelFunctor { KernelFunctor(T1 DeviceStorage, T2 DataAcc) : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc) {} - void operator()(sycl::item<1> It) const { + template void operator()(T It) const { // Select method that corresponds to this work-item auto *Ptr = mDeviceStorage->template getAs(); if (It % 2) From 12514e7c1fb7473987b6c089ce56a3b7026a3363 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Thu, 23 Jan 2025 13:12:01 -0800 Subject: [PATCH 28/30] [SYCL] Test Signed-off-by: Hu, Peisen --- sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp b/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp index 48db619d94081..e332174567cca 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp @@ -108,7 +108,7 @@ template struct KernelFunctor { KernelFunctor(T1 DeviceStorage, T2 DataAcc, T3 LocalAcc) : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc), mLocalAcc(LocalAcc) {} - void operator()(sycl::nd_item<1> It) const { + void operator()(sycl::item<1> It) const { auto *Ptr = mDeviceStorage->template getAs(); mDataAcc[It.get_global_id()] = Ptr->apply( mLocalAcc.template get_multi_ptr().get(), From 1e3e47ad034b775abc95ea55efa20609845ead08 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Thu, 23 Jan 2025 15:01:51 -0800 Subject: [PATCH 29/30] [SYCL] Revert commit Test Signed-off-by: Hu, Peisen --- sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp b/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp index e332174567cca..48db619d94081 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp @@ -108,7 +108,7 @@ template struct KernelFunctor { KernelFunctor(T1 DeviceStorage, T2 DataAcc, T3 LocalAcc) : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc), mLocalAcc(LocalAcc) {} - void operator()(sycl::item<1> It) const { + void operator()(sycl::nd_item<1> It) const { auto *Ptr = mDeviceStorage->template getAs(); mDataAcc[It.get_global_id()] = Ptr->apply( mLocalAcc.template get_multi_ptr().get(), From 0b54ead030a55ec942e698385f6b9fa0cbffbcf6 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Thu, 23 Jan 2025 20:04:11 -0800 Subject: [PATCH 30/30] [SYCL] E2E Test fix for last update #3 Signed-off-by: Hu, Peisen --- .../copy_dynamic_size.cpp | 21 ++++-- .../dynamic_alloc_local_accessor.cpp | 64 +++++++++++-------- .../dynamic_alloc_ptr_alias.cpp | 56 ++++++++-------- .../dynamic_allocation.cpp | 47 ++++++++------ .../WorkGroupScratchMemory/dynamic_unused.cpp | 26 ++++++-- 5 files changed, 130 insertions(+), 84 deletions(-) diff --git a/sycl/test-e2e/WorkGroupScratchMemory/copy_dynamic_size.cpp b/sycl/test-e2e/WorkGroupScratchMemory/copy_dynamic_size.cpp index 1f61653efc44e..e1716cff85c67 100644 --- a/sycl/test-e2e/WorkGroupScratchMemory/copy_dynamic_size.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/copy_dynamic_size.cpp @@ -29,6 +29,17 @@ void copy_via_smem(DataType *a, DataType *b, sycl::nd_item<1> it) { b[threadIdx_x] = smem_ptr[threadIdx_x]; } +template struct KernelFunctor { + T m_props; + DataType *m_a; + DataType *m_b; + KernelFunctor(T props, DataType *a, DataType *b) + : m_props(props), m_a(a), m_b(b) {} + + void operator()(sycl::nd_item<1> it) const { copy_via_smem(m_a, m_b, it); } + auto get(sycl_ext::properties_tag) const { return m_props; } +}; + int main() { sycl::queue queue; DataType *a = sycl::malloc_device(Size, queue); @@ -40,10 +51,12 @@ int main() { queue .submit([&](sycl::handler &cgh) { - cgh.parallel_for(sycl::nd_range<1>({Size}, {Size}), - sycl_ext::properties{sycl_ext::work_group_scratch_size( - Size * sizeof(DataType))}, - [=](sycl::nd_item<1> it) { copy_via_smem(a, b, it); }); + cgh.parallel_for( + sycl::nd_range<1>({Size}, {Size}), + KernelFunctor( + sycl_ext::properties{ + sycl_ext::work_group_scratch_size(Size * sizeof(DataType))}, + a, b)); }) .wait_and_throw(); diff --git a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp index 04d8a85a808ff..ebcc17855cbf1 100644 --- a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_local_accessor.cpp @@ -23,6 +23,41 @@ using namespace sycl; namespace sycl_ext = sycl::ext::oneapi::experimental; +template struct KernelFunctor { + T1 m_props; + T2 mLocalAccessor; + T3 mAcc; + KernelFunctor(T1 props, T2 LocalAccessor, T3 Acc) + : m_props(props), mLocalAccessor(LocalAccessor), mAcc(Acc) {} + + void operator()(nd_item<1> Item) const { + int *Ptr = + reinterpret_cast(sycl_ext::get_work_group_scratch_memory()); + size_t GroupOffset = Item.get_group_linear_id() * ElemPerWG; + for (size_t I = 0; I < RepeatWG; ++I) { + Ptr[WgSize * I + Item.get_local_linear_id()] = Item.get_local_linear_id(); + } + Item.barrier(); + + for (size_t I = 0; I < RepeatWG; ++I) { + // Check that the local accessor works. + size_t LocalIdx = Item.get_local_linear_id() ^ 1; + mLocalAccessor[WgSize * I + LocalIdx] = Ptr[WgSize * I + LocalIdx] + 1; + } + Item.barrier(); + + for (size_t I = 0; I < RepeatWG; ++I) { + // Check that the memory is accessible from other + // work-items + size_t BaseIdx = GroupOffset + (I * WgSize); + size_t LocalIdx = Item.get_local_linear_id(); + size_t GlobalIdx = BaseIdx + LocalIdx; + mAcc[GlobalIdx] = mLocalAccessor[WgSize * I + LocalIdx]; + } + } + auto get(sycl_ext::properties_tag) const { return m_props; } +}; + int main() { queue Q; std::vector Vec(Size, 0); @@ -36,34 +71,7 @@ int main() { auto LocalAccessor = sycl::local_accessor(WgSize * RepeatWG * sizeof(int), Cgh); Cgh.parallel_for(nd_range<1>(range<1>(WgSize * WgCount), range<1>(WgSize)), - properties, [=](nd_item<1> Item) { - int *Ptr = reinterpret_cast( - sycl_ext::get_work_group_scratch_memory()); - size_t GroupOffset = - Item.get_group_linear_id() * ElemPerWG; - for (size_t I = 0; I < RepeatWG; ++I) { - Ptr[WgSize * I + Item.get_local_linear_id()] = - Item.get_local_linear_id(); - } - Item.barrier(); - - for (size_t I = 0; I < RepeatWG; ++I) { - // Check that the local accessor works. - size_t LocalIdx = Item.get_local_linear_id() ^ 1; - LocalAccessor[WgSize * I + LocalIdx] = - Ptr[WgSize * I + LocalIdx] + 1; - } - Item.barrier(); - - for (size_t I = 0; I < RepeatWG; ++I) { - // Check that the memory is accessible from other - // work-items - size_t BaseIdx = GroupOffset + (I * WgSize); - size_t LocalIdx = Item.get_local_linear_id(); - size_t GlobalIdx = BaseIdx + LocalIdx; - Acc[GlobalIdx] = LocalAccessor[WgSize * I + LocalIdx]; - } - }); + KernelFunctor(properties, LocalAccessor, Acc)); }); host_accessor Acc(Buf, read_only); diff --git a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp index 46346d5f2ee85..2aba3369ada2f 100644 --- a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_alloc_ptr_alias.cpp @@ -23,6 +23,36 @@ using namespace sycl; namespace sycl_ext = sycl::ext::oneapi::experimental; +template struct KernelFunctor { + T1 m_props; + T2 mAcc; + KernelFunctor(T1 props, T2 Acc) : m_props(props), mAcc(Acc) {} + + void operator()(nd_item<1> Item) const { + int *Ptr = + reinterpret_cast(sycl_ext::get_work_group_scratch_memory()); + size_t GroupOffset = Item.get_group_linear_id() * ElemPerWG; + for (size_t I = 0; I < RepeatWG; ++I) { + Ptr[WgSize * I + Item.get_local_linear_id()] = Item.get_local_linear_id(); + } + + Item.barrier(); + // Check that multiple calls return the same pointer. + unsigned int *PtrAlias = reinterpret_cast( + sycl_ext::get_work_group_scratch_memory()); + + for (size_t I = 0; I < RepeatWG; ++I) { + // Check that the memory is accessible from other + // work-items + size_t BaseIdx = GroupOffset + (I * WgSize); + size_t LocalIdx = Item.get_local_linear_id() ^ 1; + size_t GlobalIdx = BaseIdx + LocalIdx; + mAcc[GlobalIdx] = PtrAlias[WgSize * I + LocalIdx]; + } + } + auto get(sycl_ext::properties_tag) const { return m_props; } +}; + int main() { queue Q; std::vector Vec(Size, 0); @@ -34,31 +64,7 @@ int main() { sizeof(int)); sycl_ext::properties properties{static_size}; Cgh.parallel_for(nd_range<1>(range<1>(WgSize * WgCount), range<1>(WgSize)), - properties, [=](nd_item<1> Item) { - int *Ptr = reinterpret_cast( - sycl_ext::get_work_group_scratch_memory()); - size_t GroupOffset = - Item.get_group_linear_id() * ElemPerWG; - for (size_t I = 0; I < RepeatWG; ++I) { - Ptr[WgSize * I + Item.get_local_linear_id()] = - Item.get_local_linear_id(); - } - - Item.barrier(); - // Check that multiple calls return the same pointer. - unsigned int *PtrAlias = - reinterpret_cast( - sycl_ext::get_work_group_scratch_memory()); - - for (size_t I = 0; I < RepeatWG; ++I) { - // Check that the memory is accessible from other - // work-items - size_t BaseIdx = GroupOffset + (I * WgSize); - size_t LocalIdx = Item.get_local_linear_id() ^ 1; - size_t GlobalIdx = BaseIdx + LocalIdx; - Acc[GlobalIdx] = PtrAlias[WgSize * I + LocalIdx]; - } - }); + KernelFunctor(properties, Acc)); }); host_accessor Acc(Buf, read_only); diff --git a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_allocation.cpp b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_allocation.cpp index 224bf2607f772..bf61ddd51a4b3 100644 --- a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_allocation.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_allocation.cpp @@ -22,6 +22,32 @@ using namespace sycl; namespace sycl_ext = sycl::ext::oneapi::experimental; +template struct KernelFunctor { + T1 m_props; + T2 mAcc; + KernelFunctor(T1 props, T2 Acc) : m_props(props), mAcc(Acc) {} + + void operator()(nd_item<1> Item) const { + int *Ptr = + reinterpret_cast(sycl_ext::get_work_group_scratch_memory()); + size_t GroupOffset = Item.get_group_linear_id() * ElemPerWG; + for (size_t I = 0; I < RepeatWG; ++I) { + Ptr[WgSize * I + Item.get_local_linear_id()] = Item.get_local_linear_id(); + } + + Item.barrier(); + for (size_t I = 0; I < RepeatWG; ++I) { + // Check that the memory is accessible from other + // work-items + size_t BaseIdx = GroupOffset + (I * WgSize); + size_t LocalIdx = Item.get_local_linear_id() ^ 1; + size_t GlobalIdx = BaseIdx + LocalIdx; + mAcc[GlobalIdx] = Ptr[WgSize * I + LocalIdx]; + } + } + auto get(sycl_ext::properties_tag) const { return m_props; } +}; + int main() { queue Q; std::vector Vec(Size, 0); @@ -33,26 +59,7 @@ int main() { sizeof(int)); sycl_ext::properties properties{static_size}; Cgh.parallel_for(nd_range<1>(range<1>(WgSize * WgCount), range<1>(WgSize)), - properties, [=](nd_item<1> Item) { - int *Ptr = reinterpret_cast( - sycl_ext::get_work_group_scratch_memory()); - size_t GroupOffset = - Item.get_group_linear_id() * ElemPerWG; - for (size_t I = 0; I < RepeatWG; ++I) { - Ptr[WgSize * I + Item.get_local_linear_id()] = - Item.get_local_linear_id(); - } - - Item.barrier(); - for (size_t I = 0; I < RepeatWG; ++I) { - // Check that the memory is accessible from other - // work-items - size_t BaseIdx = GroupOffset + (I * WgSize); - size_t LocalIdx = Item.get_local_linear_id() ^ 1; - size_t GlobalIdx = BaseIdx + LocalIdx; - Acc[GlobalIdx] = Ptr[WgSize * I + LocalIdx]; - } - }); + KernelFunctor(properties, Acc)); }); host_accessor Acc(Buf, read_only); diff --git a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_unused.cpp b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_unused.cpp index e427305c18ed3..6608eed567633 100644 --- a/sycl/test-e2e/WorkGroupScratchMemory/dynamic_unused.cpp +++ b/sycl/test-e2e/WorkGroupScratchMemory/dynamic_unused.cpp @@ -14,6 +14,19 @@ using DataType = int; namespace sycl_ext = sycl::ext::oneapi::experimental; +template struct KernelFunctor { + T m_props; + DataType *m_a; + DataType *m_b; + KernelFunctor(T props, DataType *a, DataType *b) + : m_props(props), m_a(a), m_b(b) {} + + void operator()(sycl::nd_item<1> it) const { + m_b[it.get_local_linear_id()] = m_a[it.get_local_linear_id()]; + } + auto get(sycl_ext::properties_tag) const { return m_props; } +}; + int main() { sycl::queue queue; DataType *a = sycl::malloc_device(Size, queue); @@ -25,13 +38,12 @@ int main() { queue .submit([&](sycl::handler &cgh) { - cgh.parallel_for(sycl::nd_range<1>({Size}, {Size}), - sycl_ext::properties{sycl_ext::work_group_scratch_size( - Size * sizeof(DataType))}, - [=](sycl::nd_item<1> it) { - b[it.get_local_linear_id()] = - a[it.get_local_linear_id()]; - }); + cgh.parallel_for( + sycl::nd_range<1>({Size}, {Size}), + KernelFunctor( + sycl_ext::properties{ + sycl_ext::work_group_scratch_size(Size * sizeof(DataType))}, + a, b)); }) .wait_and_throw();