Resolve crash identified in backend implementation of dpnp.choose (#2063) (#2068)

* 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

Co-authored-by: Anton <[email protected]>

Author: vtavana

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

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<DPCTLSyclEventRef>(&event);