From 4845abdbf57e38e7b0b9e9bc84e01917a283c066 Mon Sep 17 00:00:00 2001 From: Vladislav Markin Date: Fri, 1 Apr 2022 17:01:50 +0300 Subject: [PATCH 1/3] remove mixed host\dev impl from dpnp_all dpnp_any --- dpnp/backend/kernels/dpnp_krnl_logic.cpp | 30 ++++++++++-------------- 1 file changed, 12 insertions(+), 18 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp index cb323734aebf..692640cbcf71 100644 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_logic.cpp @@ -51,17 +51,17 @@ DPCTLSyclEventRef dpnp_all_c(DPCTLSyclQueueRef q_ref, } sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size); DPNPC_ptr_adapter<_ResultType> result1_ptr(q_ref, result1, 1, true, true); const _DataType* array_in = input1_ptr.get_ptr(); _ResultType* result = result1_ptr.get_ptr(); - result[0] = true; + auto init_mem_event = q.fill<_ResultType>(result, true, 1); if (!size) { + init_mem_event.wait(); return event_ref; } @@ -74,14 +74,11 @@ DPCTLSyclEventRef dpnp_all_c(DPCTLSyclQueueRef q_ref, result[0] = false; } }; + auto parallel_for_event = + q.parallel_for>( + gws, init_mem_event, kernel_parallel_for_func); - auto kernel_func = [&](sycl::handler& cgh) { - cgh.parallel_for>(gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - - event.wait(); + parallel_for_event.wait(); return event_ref; } @@ -228,17 +225,17 @@ DPCTLSyclEventRef dpnp_any_c(DPCTLSyclQueueRef q_ref, } sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size); DPNPC_ptr_adapter<_ResultType> result1_ptr(q_ref, result1, 1, true, true); const _DataType* array_in = input1_ptr.get_ptr(); _ResultType* result = result1_ptr.get_ptr(); - result[0] = false; + auto init_mem_event = q.fill<_ResultType>(result, false, 1); if (!size) { + init_mem_event.wait(); return event_ref; } @@ -251,14 +248,11 @@ DPCTLSyclEventRef dpnp_any_c(DPCTLSyclQueueRef q_ref, result[0] = true; } }; + auto parallel_for_event = + q.parallel_for>( + gws, init_mem_event, kernel_parallel_for_func); - auto kernel_func = [&](sycl::handler& cgh) { - cgh.parallel_for>(gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - - event.wait(); + parallel_for_event.wait(); return event_ref; } From ad5f6f618ffcd949bbe00c9106e8f7220f5b0254 Mon Sep 17 00:00:00 2001 From: Vladislav Markin Date: Mon, 4 Apr 2022 14:13:49 +0300 Subject: [PATCH 2/3] remove DPNPC_ptr_adapter usage from dpnp_all dpnp_any --- dpnp/backend/kernels/dpnp_krnl_logic.cpp | 14 ++++---------- 1 file changed, 4 insertions(+), 10 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp index 692640cbcf71..b98ddb3a0312 100644 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_logic.cpp @@ -51,11 +51,8 @@ DPCTLSyclEventRef dpnp_all_c(DPCTLSyclQueueRef q_ref, } sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size); - DPNPC_ptr_adapter<_ResultType> result1_ptr(q_ref, result1, 1, true, true); - const _DataType* array_in = input1_ptr.get_ptr(); - _ResultType* result = result1_ptr.get_ptr(); + const _DataType* array_in = reinterpret_cast(array1_in); + _ResultType* result = reinterpret_cast<_ResultType*>(result1); auto init_mem_event = q.fill<_ResultType>(result, true, 1); @@ -225,11 +222,8 @@ DPCTLSyclEventRef dpnp_any_c(DPCTLSyclQueueRef q_ref, } sycl::queue q = *(reinterpret_cast(q_ref)); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size); - DPNPC_ptr_adapter<_ResultType> result1_ptr(q_ref, result1, 1, true, true); - const _DataType* array_in = input1_ptr.get_ptr(); - _ResultType* result = result1_ptr.get_ptr(); + const _DataType* array_in = reinterpret_cast(array1_in); + _ResultType* result = reinterpret_cast<_ResultType*>(result1); auto init_mem_event = q.fill<_ResultType>(result, false, 1); From 8091acaf5c57dad9a4eb46114e7d293642be4002 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 17 Aug 2022 12:28:31 -0500 Subject: [PATCH 3/3] Rework CFD implementation for 3 functions: dpnp_all_c(), dpnp_allclose_c(), dpnp_any_c() --- dpnp/backend/kernels/dpnp_krnl_logic.cpp | 51 ++++++++++++------------ 1 file changed, 25 insertions(+), 26 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp index d9fbfe76042c..1f2464b10747 100644 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_logic.cpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2016-2020, Intel Corporation +// Copyright (c) 2016-2022, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without @@ -51,15 +51,16 @@ DPCTLSyclEventRef dpnp_all_c(DPCTLSyclQueueRef q_ref, } sycl::queue q = *(reinterpret_cast(q_ref)); + const _DataType* array_in = reinterpret_cast(array1_in); _ResultType* result = reinterpret_cast<_ResultType*>(result1); - auto init_mem_event = q.fill<_ResultType>(result, true, 1); + auto fill_event = q.fill<_ResultType>(result, true, 1); if (!size) { - init_mem_event.wait(); - return event_ref; + event_ref = reinterpret_cast(&fill_event); + return DPCTLEvent_Copy(event_ref); } sycl::range<1> gws(size); @@ -73,13 +74,12 @@ DPCTLSyclEventRef dpnp_all_c(DPCTLSyclQueueRef q_ref, }; auto kernel_func = [&](sycl::handler& cgh) { + cgh.depends_on(fill_event); cgh.parallel_for>(gws, kernel_parallel_for_func); }; - event = q.submit(kernel_func); - + auto event = q.submit(kernel_func); event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); } @@ -94,6 +94,7 @@ void dpnp_all_c(const void* array1_in, void* result1, const size_t size) size, dep_event_vec_ref); DPCTLEvent_WaitAndThrow(event_ref); + DPCTLEvent_Delete(event_ref); } template @@ -124,26 +125,23 @@ DPCTLSyclEventRef dpnp_allclose_c(DPCTLSyclQueueRef q_ref, DPCTLSyclEventRef event_ref = nullptr; - if (!array1_in || !result1) + if (!array1_in || !array2_in || !result1) { return event_ref; } sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; - DPNPC_ptr_adapter<_DataType1> input1_ptr(q_ref, array1_in, size); - DPNPC_ptr_adapter<_DataType2> input2_ptr(q_ref, array2_in, size); - DPNPC_ptr_adapter<_ResultType> result1_ptr(q_ref, result1, 1, true, true); - const _DataType1* array1 = input1_ptr.get_ptr(); - const _DataType2* array2 = input2_ptr.get_ptr(); - _ResultType* result = result1_ptr.get_ptr(); + const _DataType1* array1 = reinterpret_cast(array1_in); + const _DataType2* array2 = reinterpret_cast(array2_in); + _ResultType* result = reinterpret_cast<_ResultType*>(result1); - result[0] = true; + auto fill_event = q.fill<_ResultType>(result, true, 1); if (!size) { - return event_ref; + event_ref = reinterpret_cast(&fill_event); + return DPCTLEvent_Copy(event_ref); } sycl::range<1> gws(size); @@ -157,14 +155,13 @@ DPCTLSyclEventRef dpnp_allclose_c(DPCTLSyclQueueRef q_ref, }; auto kernel_func = [&](sycl::handler& cgh) { + cgh.depends_on(fill_event); cgh.parallel_for>(gws, kernel_parallel_for_func); }; - event = q.submit(kernel_func); - + auto event = q.submit(kernel_func); event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); } @@ -183,6 +180,7 @@ void dpnp_allclose_c( atol_val, dep_event_vec_ref); DPCTLEvent_WaitAndThrow(event_ref); + DPCTLEvent_Delete(event_ref); } template @@ -225,15 +223,16 @@ DPCTLSyclEventRef dpnp_any_c(DPCTLSyclQueueRef q_ref, } sycl::queue q = *(reinterpret_cast(q_ref)); + const _DataType* array_in = reinterpret_cast(array1_in); _ResultType* result = reinterpret_cast<_ResultType*>(result1); - auto init_mem_event = q.fill<_ResultType>(result, false, 1); + auto fill_event = q.fill<_ResultType>(result, false, 1); if (!size) { - init_mem_event.wait(); - return event_ref; + event_ref = reinterpret_cast(&fill_event); + return DPCTLEvent_Copy(event_ref); } sycl::range<1> gws(size); @@ -247,13 +246,12 @@ DPCTLSyclEventRef dpnp_any_c(DPCTLSyclQueueRef q_ref, }; auto kernel_func = [&](sycl::handler& cgh) { + cgh.depends_on(fill_event); cgh.parallel_for>(gws, kernel_parallel_for_func); }; - event = q.submit(kernel_func); - + auto event = q.submit(kernel_func); event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); } @@ -268,6 +266,7 @@ void dpnp_any_c(const void* array1_in, void* result1, const size_t size) size, dep_event_vec_ref); DPCTLEvent_WaitAndThrow(event_ref); + DPCTLEvent_Delete(event_ref); } template