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 7e30406189e29..afe3ebd0d2557 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 @@ -58,17 +58,15 @@ template struct KernelFunctorWithMaxWGSizeProp { } }; -template -int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { +template +int test(queue &Q, KernelType KernelFunc) { constexpr size_t Dims = 1; // Positive test case: Specify local size that matches required size. try { Q.submit([&](handler &CGH) { CGH.parallel_for>( - nd_range(repeatRange(8), range(I)), Props, - KernelFunc); + nd_range(repeatRange(8), range(I)), KernelFunc); }); Q.wait_and_throw(); } catch (exception &E) { @@ -81,8 +79,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { // Same as above but using the queue shortcuts. try { Q.parallel_for>( - nd_range(repeatRange(8), range(I)), Props, - KernelFunc); + nd_range(repeatRange(8), range(I)), KernelFunc); Q.wait_and_throw(); } catch (exception &E) { std::cerr @@ -97,7 +94,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { try { Q.submit([&](handler &CGH) { CGH.parallel_for>( - repeatRange(16), Props, KernelFunc); + repeatRange(16), KernelFunc); }); Q.wait_and_throw(); } catch (exception &E) { @@ -109,7 +106,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { try { Q.parallel_for>( - repeatRange(16), Props, KernelFunc); + repeatRange(16), KernelFunc); Q.wait_and_throw(); } catch (exception &E) { std::cerr << "Test case MaxLinearWGSizeNoLocalPositive shortcut failed: " @@ -122,7 +119,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { try { Q.submit([&](handler &CGH) { CGH.parallel_for>( - nd_range(repeatRange(16), repeatRange(8)), Props, + nd_range(repeatRange(16), repeatRange(8)), KernelFunc); }); Q.wait_and_throw(); @@ -147,7 +144,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { // Same as above but using the queue shortcuts. try { Q.parallel_for>( - nd_range(repeatRange(16), repeatRange(8)), Props, + nd_range(repeatRange(16), repeatRange(8)), KernelFunc); Q.wait_and_throw(); std::cerr @@ -174,17 +171,10 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { } template int test_max(queue &Q) { - auto Props = ext::oneapi::experimental::properties{ - ext::oneapi::experimental::max_linear_work_group_size}; - auto KernelFunction = [](auto) {}; - - auto EmptyProps = ext::oneapi::experimental::properties{}; KernelFunctorWithMaxWGSizeProp KernelFunctor; int Res = 0; - Res += test(Q, Props, KernelFunction); - Res += test(Q, EmptyProps, KernelFunctor); - Res += test(Q, Props, KernelFunctor); + Res += test(Q, KernelFunctor); return Res; } 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 6376aa0a10392..6694cb1d35d3f 100644 --- a/sycl/test-e2e/Basic/max_work_group_size_props.cpp +++ b/sycl/test-e2e/Basic/max_work_group_size_props.cpp @@ -49,17 +49,15 @@ template struct KernelFunctorWithMaxWGSizeProp { } }; -template -int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { +template +int test(queue &Q, KernelType KernelFunc) { constexpr size_t Dims = sizeof...(Is); // Positive test case: Specify local size that matches required size. try { Q.submit([&](handler &CGH) { CGH.parallel_for>( - nd_range(repeatRange(8), range(Is...)), Props, - KernelFunc); + nd_range(repeatRange(8), range(Is...)), KernelFunc); }); Q.wait_and_throw(); } catch (exception &E) { @@ -71,8 +69,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { // Same as above but using the queue shortcuts. try { Q.parallel_for>( - nd_range(repeatRange(8), range(Is...)), Props, - KernelFunc); + nd_range(repeatRange(8), range(Is...)), KernelFunc); Q.wait_and_throw(); } catch (exception &E) { std::cerr << "Test case MaxWGSizePositive shortcut failed: unexpected " @@ -86,7 +83,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { try { Q.submit([&](handler &CGH) { CGH.parallel_for>( - repeatRange(16), Props, KernelFunc); + repeatRange(16), KernelFunc); }); Q.wait_and_throw(); } catch (exception &E) { @@ -98,7 +95,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { try { Q.parallel_for>( - repeatRange(16), Props, KernelFunc); + repeatRange(16), KernelFunc); Q.wait_and_throw(); } catch (exception &E) { std::cerr << "Test case MaxWGSizeNoLocalPositive shortcut failed: " @@ -111,7 +108,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { try { Q.submit([&](handler &CGH) { CGH.parallel_for>( - nd_range(repeatRange(16), repeatRange(8)), Props, + nd_range(repeatRange(16), repeatRange(8)), KernelFunc); }); Q.wait_and_throw(); @@ -134,7 +131,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { // Same as above but using the queue shortcuts. try { Q.parallel_for>( - nd_range(repeatRange(16), repeatRange(8)), Props, + nd_range(repeatRange(16), repeatRange(8)), KernelFunc); Q.wait_and_throw(); std::cerr << "Test case MaxWGSizeNegative shortcut failed: no exception " @@ -159,17 +156,10 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { } template int test_max(queue &Q) { - auto Props = ext::oneapi::experimental::properties{ - ext::oneapi::experimental::max_work_group_size}; - auto KernelFunction = [](auto) {}; - - auto EmptyProps = ext::oneapi::experimental::properties{}; KernelFunctorWithMaxWGSizeProp KernelFunctor; int Res = 0; - Res += test(Q, Props, KernelFunction); - Res += test(Q, EmptyProps, KernelFunctor); - Res += test(Q, Props, KernelFunctor); + Res += test(Q, KernelFunctor); return Res; } diff --git a/sycl/test-e2e/Basic/work_group_size_prop.cpp b/sycl/test-e2e/Basic/work_group_size_prop.cpp index ac8400dcc31b8..9cf04c4d2ea66 100644 --- a/sycl/test-e2e/Basic/work_group_size_prop.cpp +++ b/sycl/test-e2e/Basic/work_group_size_prop.cpp @@ -45,9 +45,8 @@ template struct KernelFunctorWithWGSizeProp { } }; -template -int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { +template +int test(queue &Q, KernelType KernelFunc) { constexpr size_t Dims = sizeof...(Is); bool IsOpenCL = (Q.get_backend() == backend::opencl); @@ -56,8 +55,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { try { Q.submit([&](handler &CGH) { CGH.parallel_for>( - nd_range(repeatRange(8), range(Is...)), Props, - KernelFunc); + nd_range(repeatRange(8), range(Is...)), KernelFunc); }); Q.wait_and_throw(); } catch (exception &E) { @@ -69,8 +67,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { // Same as above but using the queue shortcuts. try { Q.parallel_for>( - nd_range(repeatRange(8), range(Is...)), Props, - KernelFunc); + nd_range(repeatRange(8), range(Is...)), KernelFunc); Q.wait_and_throw(); } catch (exception &E) { std::cerr << "Test case ReqdWGSizePositiveA shortcut failed: unexpected " @@ -87,7 +84,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { Q.submit([&](handler &CGH) { CGH.parallel_for< ReqdWGSizeNoLocalPositive>( - repeatRange(16), Props, KernelFunc); + repeatRange(16), KernelFunc); }); Q.wait_and_throw(); } catch (exception &E) { @@ -99,7 +96,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { try { Q.parallel_for>( - repeatRange(16), Props, KernelFunc); + repeatRange(16), KernelFunc); Q.wait_and_throw(); } catch (exception &E) { std::cerr << "Test case ReqdWGSizeNoLocalPositive shortcut failed: " @@ -113,7 +110,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { try { Q.submit([&](handler &CGH) { CGH.parallel_for>( - nd_range(repeatRange(16), repeatRange(8)), Props, + nd_range(repeatRange(16), repeatRange(8)), KernelFunc); }); Q.wait_and_throw(); @@ -137,7 +134,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { // Same as above but using the queue shortcuts. try { Q.parallel_for>( - nd_range(repeatRange(16), repeatRange(8)), Props, + nd_range(repeatRange(16), repeatRange(8)), KernelFunc); Q.wait_and_throw(); std::cerr << "Test case ReqdWGSizeNegativeA shortcut failed: no exception " @@ -162,17 +159,10 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) { } template int test(queue &Q) { - auto Props = ext::oneapi::experimental::properties{ - ext::oneapi::experimental::work_group_size}; - auto KernelFunction = [](auto) {}; - - auto EmptyProps = ext::oneapi::experimental::properties{}; KernelFunctorWithWGSizeProp KernelFunctor; int Res = 0; - Res += test(Q, Props, KernelFunction); - Res += test(Q, EmptyProps, KernelFunctor); - Res += test(Q, Props, KernelFunctor); + Res += test(Q, KernelFunctor); return Res; } diff --git a/sycl/test-e2e/Graph/Inputs/work_group_size_prop.cpp b/sycl/test-e2e/Graph/Inputs/work_group_size_prop.cpp index e052ab5acb3bf..7fd3d8eef1856 100644 --- a/sycl/test-e2e/Graph/Inputs/work_group_size_prop.cpp +++ b/sycl/test-e2e/Graph/Inputs/work_group_size_prop.cpp @@ -40,9 +40,8 @@ template struct KernelFunctorWithWGSizeProp { } }; -template -int test(queue &Queue, PropertiesT Props, KernelType KernelFunc) { +template +int test(queue &Queue, KernelType KernelFunc) { constexpr size_t Dims = sizeof...(Is); // Positive test case: Specify local size that matches required size. @@ -52,15 +51,13 @@ int test(queue &Queue, PropertiesT Props, KernelType KernelFunc) { add_node(Graph, Queue, [&](handler &CGH) { CGH.parallel_for>( - nd_range(repeatRange(8), range(Is...)), Props, - KernelFunc); + nd_range(repeatRange(8), range(Is...)), KernelFunc); }); #ifdef GRAPH_E2E_RECORD_REPLAY Graph.begin_recording(Queue); Queue.parallel_for>( - nd_range(repeatRange(8), range(Is...)), Props, - KernelFunc); + nd_range(repeatRange(8), range(Is...)), KernelFunc); Graph.end_recording(Queue); #endif @@ -83,7 +80,7 @@ int test(queue &Queue, PropertiesT Props, KernelType KernelFunc) { try { add_node(GraphN, Queue, [&](handler &CGH) { CGH.parallel_for>( - nd_range(repeatRange(16), repeatRange(8)), Props, + nd_range(repeatRange(16), repeatRange(8)), KernelFunc); }); auto ExecGraph = GraphN.finalize(); @@ -119,7 +116,7 @@ int test(queue &Queue, PropertiesT Props, KernelType KernelFunc) { GraphN.begin_recording(Queue); Queue.parallel_for>( - nd_range(repeatRange(16), repeatRange(8)), Props, + nd_range(repeatRange(16), repeatRange(8)), KernelFunc); GraphN.end_recording(Queue); @@ -156,17 +153,10 @@ int test(queue &Queue, PropertiesT Props, KernelType KernelFunc) { } template int test(queue &Queue) { - auto Props = sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::work_group_size}; - auto KernelFunction = [](auto) {}; - - auto EmptyProps = sycl::ext::oneapi::experimental::properties{}; KernelFunctorWithWGSizeProp KernelFunctor; int Res = 0; - Res += test(Queue, Props, KernelFunction); - Res += test(Queue, EmptyProps, KernelFunctor); - Res += test(Queue, Props, KernelFunctor); + Res += test(Queue, KernelFunctor); return Res; } diff --git a/sycl/test-e2e/GroupAlgorithm/root_group.cpp b/sycl/test-e2e/GroupAlgorithm/root_group.cpp index 2e50634fd21c8..257b5a4e4457f 100644 --- a/sycl/test-e2e/GroupAlgorithm/root_group.cpp +++ b/sycl/test-e2e/GroupAlgorithm/root_group.cpp @@ -2,7 +2,10 @@ // 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 %} +// TODO: Currently using the -Wno-deprecated-declarations flag due to issue +// https://github.com/intel/llvm/issues/16451. Rewrite testRootGroup() amd +// remove the flag once the issue is resolved. +// RUN: %{build} -I . -o %t.out -Wno-deprecated-declarations %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. @@ -42,9 +45,14 @@ void testQueriesAndProperties() { .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, []() {}); + struct TestKernel0 { + void operator()() const {} + auto get(sycl::ext::oneapi::experimental::properties_tag) { + return sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::use_root_sync}; + } + }; + q.single_task(TestKernel0{}); static auto check_max_num_work_group_sync = [](auto Result) { static_assert(std::is_same_v, size_t>, @@ -99,6 +107,32 @@ void testRootGroup() { } } +template struct TestKernel2 { + T m_testResults; + TestKernel2(T &testResults_) : m_testResults(testResults_) {} + void operator()(sycl::nd_item<1> it) const { + const auto root = it.ext_oneapi_get_root_group(); + if (root.leader() || root.get_local_id() == 3) { + m_testResults[0] = root.get_group_id() == sycl::id<1>(0); + m_testResults[1] = root.leader() ? root.get_local_id() == sycl::id<1>(0) + : root.get_local_id() == sycl::id<1>(3); + m_testResults[2] = root.get_group_range() == sycl::range<1>(1); + m_testResults[3] = root.get_local_range() == it.get_global_range(); + m_testResults[4] = root.get_max_local_range() == root.get_local_range(); + m_testResults[5] = root.get_group_linear_id() == 0; + m_testResults[6] = + root.get_local_linear_id() == root.get_local_id().get(0); + m_testResults[7] = root.get_group_linear_range() == 1; + m_testResults[8] = + root.get_local_linear_range() == root.get_local_range().size(); + } + } + auto get(sycl::ext::oneapi::experimental::properties_tag) { + return sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::use_root_sync}; + } +}; + void testRootGroupFunctions() { sycl::queue q; const auto bundle = @@ -109,34 +143,13 @@ void testRootGroupFunctions() { .ext_oneapi_get_info( q, WorkGroupSize, 0); - const auto props = sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::use_root_sync}; - constexpr int testCount = 9; sycl::buffer testResultsBuf{sycl::range{testCount}}; const auto range = sycl::nd_range<1>{maxWGs * WorkGroupSize, WorkGroupSize}; q.submit([&](sycl::handler &h) { sycl::accessor testResults{testResultsBuf, h}; - h.parallel_for( - range, props, [=](sycl::nd_item<1> it) { - const auto root = it.ext_oneapi_get_root_group(); - if (root.leader() || root.get_local_id() == 3) { - testResults[0] = root.get_group_id() == sycl::id<1>(0); - testResults[1] = root.leader() - ? root.get_local_id() == sycl::id<1>(0) - : root.get_local_id() == sycl::id<1>(3); - testResults[2] = root.get_group_range() == sycl::range<1>(1); - testResults[3] = root.get_local_range() == it.get_global_range(); - testResults[4] = - root.get_max_local_range() == root.get_local_range(); - testResults[5] = root.get_group_linear_id() == 0; - testResults[6] = - root.get_local_linear_id() == root.get_local_id().get(0); - testResults[7] = root.get_group_linear_range() == 1; - testResults[8] = - root.get_local_linear_range() == root.get_local_range().size(); - } - }); + h.parallel_for(range, + TestKernel2(testResults)); }); sycl::host_accessor testResults{testResultsBuf}; for (int i = 0; i < testCount; i++) { diff --git a/sycl/test-e2e/Properties/cache_config.cpp b/sycl/test-e2e/Properties/cache_config.cpp index 0cda3e97a5d1f..666f85631c9b1 100644 --- a/sycl/test-e2e/Properties/cache_config.cpp +++ b/sycl/test-e2e/Properties/cache_config.cpp @@ -1,6 +1,9 @@ // REQUIRES: gpu, level_zero -// RUN: %{build} -o %t.out +// 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: %{build} -o %t.out -Wno-deprecated-declarations // RUN: env UR_L0_DEBUG=1 %{run} %t.out 2>&1 | FileCheck %s #include @@ -36,6 +39,14 @@ struct NegativeKernelFunctor { auto get(properties_tag) const { return properties{}; } }; +struct RangeKernelFunctor { + + RangeKernelFunctor() {} + + void operator()(id<2> i) const {} + auto get(properties_tag) const { return properties{cache_config(large_slm)}; } +}; + int main() { sycl::property_list q_prop{sycl::property::queue::in_order()}; queue q{q_prop}; @@ -43,22 +54,10 @@ int main() { sycl::ext::oneapi::experimental::properties properties{ cache_config(large_slm)}; - // CHECK: single_task - // CHECK: ZE ---> zeKernelSetCacheConfig - std::cout << "single_task" << std::endl; - q.single_task(properties, [=]() {}).wait(); - // CHECK: parallel_for with sycl::range // CHECK: ZE ---> zeKernelSetCacheConfig std::cout << "parallel_for with sycl::range" << std::endl; - q.parallel_for(range<2>{16, 16}, properties, [=](id<2> i) {}).wait(); - - // CHECK: parallel_for with sycl::nd_range - // CHECK: ZE ---> zeKernelSetCacheConfig - std::cout << "parallel_for with sycl::nd_range" << std::endl; - q.parallel_for(nd_range<2>{range<2>(4, 4), range<2>(2, 2)}, properties, - [=](nd_item<2> i) {}) - .wait(); + q.parallel_for(range<2>{16, 16}, RangeKernelFunctor{}).wait(); // CHECK: parallel_for_work_group(range, func) // CHECK: ZE ---> zeKernelSetCacheConfig 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..b008effc626c0 100644 --- a/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp +++ b/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp @@ -66,6 +66,25 @@ void applyOp(int *DataPtr, Base *ObjPtr) { ObjPtr->multiply(DataPtr); } +template struct KernelFunctor { + T1 mStorageAcc; + T2 mDataAcc; + unsigned mTestCase; + KernelFunctor(T1 &StorageAcc, T2 &DataAcc, unsigned TestCase) + : mStorageAcc(StorageAcc), mDataAcc(DataAcc), mTestCase(TestCase) {} + + void operator()() const { + auto *Ptr = + mStorageAcc[0].template construct(mTestCase); + applyOp( + mDataAcc.template get_multi_ptr().get(), + Ptr); + } + auto get(oneapi::properties_tag) const { + return oneapi::properties{oneapi::assume_indirect_calls}; + } +}; + int main() try { using storage_t = obj_storage_t(TestCase); - applyOp(DataAcc.get_multi_ptr().get(), - Ptr); - }); + CGH.single_task(KernelFunctor(StorageAcc, DataAcc, TestCase)); }); Base *Ptr = HostStorage.construct(TestCase); 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..d1c2c5fd092a9 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 @@ -45,6 +45,25 @@ class IncrementBy8 : public IncrementOp { void applyOp(int *Data, AbstractOp *Obj) { Obj->applyOp(Data); } +template struct KernelFunctor { + T1 mStorageAcc; + T2 mDataAcc; + unsigned mTestCase; + KernelFunctor(T1 &StorageAcc, T2 &DataAcc, unsigned TestCase) + : mStorageAcc(StorageAcc), mDataAcc(DataAcc), mTestCase(TestCase) {} + + void operator()() const { + auto *Ptr = mStorageAcc[0].template construct( + mTestCase); + applyOp( + mDataAcc.template get_multi_ptr().get(), + Ptr); + } + auto get(oneapi::properties_tag) const { + return oneapi::properties{oneapi::assume_indirect_calls}; + } +}; + int main() try { using storage_t = obj_storage_t; @@ -59,7 +78,6 @@ int main() try { sycl::queue q(asyncHandler); - constexpr oneapi::properties props{oneapi::assume_indirect_calls}; for (unsigned TestCase = 0; TestCase < 4; ++TestCase) { int HostData = 42; int Data = HostData; @@ -68,12 +86,7 @@ int main() try { q.submit([&](sycl::handler &CGH) { sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::write_only); sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only); - CGH.single_task(props, [=]() { - auto *Ptr = - StorageAcc[0].construct(TestCase); - applyOp(DataAcc.get_multi_ptr().get(), - Ptr); - }); + CGH.single_task(KernelFunctor(StorageAcc, DataAcc, TestCase)); }); auto *Ptr = HostStorage.construct(TestCase); 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..aad1e1ccecffa 100644 --- a/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp +++ b/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp @@ -30,6 +30,24 @@ class IncrementBy8 : public BaseIncrement { void increment(int *Data) override { *Data += 8; } }; +template struct KernelFunctor { + T1 mStorageAcc; + T2 mDataAcc; + unsigned mTestCase; + KernelFunctor(T1 &StorageAcc, T2 &DataAcc, unsigned TestCase) + : mStorageAcc(StorageAcc), mDataAcc(DataAcc), mTestCase(TestCase) {} + void operator()() const { + auto *Ptr = + mStorageAcc[0].template construct( + mTestCase); + Ptr->increment( + mDataAcc.template get_multi_ptr().get()); + } + auto get(oneapi::properties_tag) const { + return oneapi::properties{oneapi::assume_indirect_calls}; + } +}; + int main() try { using storage_t = obj_storage_t; @@ -44,7 +62,6 @@ int main() try { sycl::queue q(asyncHandler); - constexpr oneapi::properties props{oneapi::assume_indirect_calls}; for (unsigned TestCase = 0; TestCase < 4; ++TestCase) { int HostData = 42; int Data = HostData; @@ -53,12 +70,7 @@ int main() try { q.submit([&](sycl::handler &CGH) { sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::write_only); sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only); - CGH.single_task(props, [=]() { - auto *Ptr = - StorageAcc[0].construct(TestCase); - Ptr->increment( - DataAcc.get_multi_ptr().get()); - }); + CGH.single_task(KernelFunctor(StorageAcc, DataAcc, TestCase)); }); auto *Ptr = HostStorage.construct(TestCase); 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..467d4e5b006c1 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 @@ -57,6 +57,22 @@ class IncrementBy16 : public BaseIncrement { void increment(int *Data) override { *Data += 16 + Mod; } }; +template struct KernelFunctor { + T1 mStorageAcc; + T2 mDataAcc; + KernelFunctor(T1 &StorageAcc, T2 &DataAcc) + : mStorageAcc(StorageAcc), mDataAcc(DataAcc) {} + void operator()() const { + auto *Ptr = mStorageAcc[0].template getAs(); + Ptr->increment( + mDataAcc.template get_multi_ptr().get()); + } + auto get(oneapi::properties_tag) const { + return oneapi::properties{ + oneapi::assume_indirect_calls_to}; + } +}; + int main() try { using storage_t = obj_storage_t; @@ -72,8 +88,6 @@ int main() try { sycl::queue q(asyncHandler); // TODO: cover uses case when objects are passed through USM - constexpr oneapi::properties props{ - oneapi::assume_indirect_calls_to}; for (unsigned TestCase = 0; TestCase < 5; ++TestCase) { int HostData = 42; int Data = HostData; @@ -90,11 +104,7 @@ int main() try { q.submit([&](sycl::handler &CGH) { sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::read_write); sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only); - CGH.single_task(props, [=]() { - auto *Ptr = StorageAcc[0].getAs(); - Ptr->increment( - DataAcc.get_multi_ptr().get()); - }); + CGH.single_task(KernelFunctor(StorageAcc, DataAcc)); }); auto *Ptr = diff --git a/sycl/test-e2e/VirtualFunctions/misc/math.cpp b/sycl/test-e2e/VirtualFunctions/misc/math.cpp index 71b34c23cef1f..da0570ce0291e 100644 --- a/sycl/test-e2e/VirtualFunctions/misc/math.cpp +++ b/sycl/test-e2e/VirtualFunctions/misc/math.cpp @@ -40,6 +40,21 @@ class RoundOp : public BaseOp { virtual float apply(float V) { return sycl::round(V); } }; +template struct KernelFunctor { + T1 mDataAcc; + T2 mDeviceStorage; + KernelFunctor(T1 &DataAcc, T2 &DeviceStorage) + : mDataAcc(DataAcc), mDeviceStorage(DeviceStorage) {} + + void operator()() const { + auto *Ptr = mDeviceStorage->template getAs(); + mDataAcc[0] = Ptr->apply(mDataAcc[0]); + } + auto get(oneapi::properties_tag) const { + return oneapi::properties{oneapi::assume_indirect_calls}; + } +}; + int main() try { using storage_t = obj_storage_t; @@ -49,7 +64,6 @@ int main() try { auto *DeviceStorage = sycl::malloc_shared(1, q); - constexpr oneapi::properties props{oneapi::assume_indirect_calls}; for (unsigned TestCase = 0; TestCase < 3; ++TestCase) { float HostData = 3.56; float Data = HostData; @@ -63,10 +77,7 @@ int main() try { q.submit([&](sycl::handler &CGH) { sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write); - CGH.single_task(props, [=]() { - auto *Ptr = DeviceStorage->getAs(); - DataAcc[0] = Ptr->apply(DataAcc[0]); - }); + CGH.single_task(KernelFunctor(DataAcc, DeviceStorage)); }); auto *Ptr = HostStorage.construct(TestCase); 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..b8b80b9541569 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 @@ -23,41 +23,42 @@ void check_props(sycl::queue &q) {} // Full specializations for each progress guarantee +template struct KernelFunctor { + T props; + KernelFunctor(const T &props_) : props(props_) {} + void operator()() const {} + auto get(properties_tag) const { return props; } +}; + template <> void check_props(sycl::queue &q) { constexpr auto guarantee = forward_progress_guarantee::parallel; // Check properties at execution_scope::root_group coordination level - q.single_task( - properties{work_group_progress}, - [=]() {}); - q.single_task( - properties{sub_group_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_group_progress})); + q.single_task(KernelFunctor( + properties{sub_group_progress})); try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + work_item_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } // Check properties at execution_scope::work_group coordination level - q.single_task( - properties{sub_group_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{sub_group_progress})); try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + work_item_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } // Check properties at execution_scope::sub_group coordination level try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_item_progress})); } catch (sycl::exception &ex) { } } @@ -66,66 +67,54 @@ template <> void check_props(sycl::queue &q) { constexpr auto guarantee = forward_progress_guarantee::weakly_parallel; // Check properties at execution_scope::root_group coordination level - q.single_task( - properties{work_group_progress}, - [=]() {}); - q.single_task( - properties{sub_group_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_group_progress})); + q.single_task(KernelFunctor( + properties{sub_group_progress})); - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_item_progress})); // Check properties at execution_scope::work_group coordination level - q.single_task( - properties{sub_group_progress}, - [=]() {}); - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{sub_group_progress})); + q.single_task(KernelFunctor( + properties{work_item_progress})); // Check properties at execution_scope::sub_group coordination level - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_item_progress})); } template <> void check_props(sycl::queue &q) { constexpr auto guarantee = forward_progress_guarantee::concurrent; // Check properties at execution_scope::root_group coordination level - q.single_task( - properties{work_group_progress}, - [=]() {}); - q.single_task( - properties{sub_group_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_group_progress})); + q.single_task(KernelFunctor( + properties{sub_group_progress})); try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + work_item_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } // Check properties at execution_scope::work_group coordination level - q.single_task( - properties{sub_group_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{sub_group_progress})); try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + work_item_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } // Check properties at execution_scope::sub_group coordination level try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_item_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } 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..8647d42ee24c3 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 @@ -22,49 +22,50 @@ void check_props(sycl::queue &q) {} // Full specializations for each progress guarantee +template struct KernelFunctor { + T props; + KernelFunctor(const T &props_) : props(props_) {} + void operator()() const {} + auto get(properties_tag) const { return props; } +}; + template <> void check_props(sycl::queue &q) { constexpr auto guarantee = forward_progress_guarantee::parallel; // Check properties at execution_scope::root_group coordination level - q.single_task( - properties{work_group_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_group_progress})); try { - q.single_task( - properties{sub_group_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + sub_group_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + work_item_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } // Check properties at execution_scope::work_group coordination level try { - q.single_task( - properties{sub_group_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + sub_group_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + work_item_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } // Check properties at execution_scope::sub_group coordination level try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_item_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } @@ -74,29 +75,23 @@ template <> void check_props(sycl::queue &q) { constexpr auto guarantee = forward_progress_guarantee::weakly_parallel; // Check properties at execution_scope::root_group coordination level - q.single_task( - properties{work_group_progress}, - [=]() {}); - q.single_task( - properties{sub_group_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_group_progress})); + q.single_task(KernelFunctor( + properties{sub_group_progress})); - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_item_progress})); // Check properties at execution_scope::work_group coordination level - q.single_task( - properties{sub_group_progress}, - [=]() {}); - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{sub_group_progress})); + q.single_task(KernelFunctor( + properties{work_item_progress})); // Check properties at execution_scope::sub_group coordination level - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_item_progress})); } template <> @@ -104,48 +99,42 @@ void check_props(sycl::queue &q) { constexpr auto guarantee = forward_progress_guarantee::concurrent; // Check properties at execution_scope::root_group coordination level try { - q.single_task( - properties{work_group_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + work_group_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } try { - q.single_task( - properties{sub_group_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + sub_group_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + work_item_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } // Check properties at execution_scope::work_group coordination level try { - q.single_task( - properties{sub_group_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + sub_group_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor(properties{ + work_item_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { } // Check properties at execution_scope::sub_group coordination level try { - q.single_task( - properties{work_item_progress}, - [=]() {}); + q.single_task(KernelFunctor( + properties{work_item_progress})); assert(false && "Expected exception not seen!"); } catch (sycl::exception &ex) { }