From 22b281a9269cf3e7cbd978246ce63fbe2d90cfb6 Mon Sep 17 00:00:00 2001 From: Anton <100830759+antonwolfy@users.noreply.github.com> Date: Fri, 20 Sep 2024 00:25:36 +0200 Subject: [PATCH] Resolve crash identified in backend implementation of `dpnp.choose` (#2063) * Add memory copy to device memory * Keep extra memcopy only for the result array * Add missing memory copy from host to device memory * Applied pre-commit formatting * Updated changelog * Add empty line in changelog * Correct the description of issue in changelog --- CHANGELOG.md | 1 + dpnp/backend/kernels/dpnp_krnl_indexing.cpp | 6 +++++- 2 files changed, 6 insertions(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 414d75a3f8f3..f77c86beb04a 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -113,6 +113,7 @@ In addition, this release completes implementation of `dpnp.fft` module and adds * Resolved an issue with failing tests for `dpnp.append` when running on a device without fp64 support [#2034](https://github.com/IntelPython/dpnp/pull/2034) * Resolved an issue with input array of `usm_ndarray` passed into `dpnp.ix_` [#2047](https://github.com/IntelPython/dpnp/pull/2047) * Added a workaround to prevent crash in tests on Windows in internal CI/CD (when running on either Lunar Lake or Arrow Lake) [#2062](https://github.com/IntelPython/dpnp/pull/2062) +* Fixed a crash in `dpnp.choose` caused by missing control of releasing temporary allocated device memory [#2063](https://github.com/IntelPython/dpnp/pull/2063) ## [0.15.0] - 05/25/2024 diff --git a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp index 5400da817581..eda36511e048 100644 --- a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp @@ -63,7 +63,10 @@ DPCTLSyclEventRef dpnp_choose_c(DPCTLSyclQueueRef q_ref, DPNPC_ptr_adapter<_DataType1> input1_ptr(q_ref, array1_in, size); _DataType1 *array_in = input1_ptr.get_ptr(); - DPNPC_ptr_adapter<_DataType2 *> choices_ptr(q_ref, choices1, choices_size); + // choices1 is a list of pointers to device memory, + // which is allocating on the host, so memcpy to device memory is required + DPNPC_ptr_adapter<_DataType2 *> choices_ptr(q_ref, choices1, choices_size, + true); _DataType2 **choices = choices_ptr.get_ptr(); for (size_t i = 0; i < choices_size; ++i) { @@ -88,6 +91,7 @@ DPCTLSyclEventRef dpnp_choose_c(DPCTLSyclQueueRef q_ref, }; sycl::event event = q.submit(kernel_func); + choices_ptr.depends_on(event); event_ref = reinterpret_cast(&event);