From b326450d6b65083e87bd7c6cc4e66f50a69ca41e Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 18 Sep 2024 16:13:23 +0200 Subject: [PATCH 1/7] Add memory copy to device memory --- dpnp/backend/kernels/dpnp_krnl_indexing.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp index 5400da817581..ebb81eef0e61 100644 --- a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp @@ -60,19 +60,19 @@ DPCTLSyclEventRef dpnp_choose_c(DPCTLSyclQueueRef q_ref, sycl::queue q = *(reinterpret_cast(q_ref)); - DPNPC_ptr_adapter<_DataType1> input1_ptr(q_ref, array1_in, size); + DPNPC_ptr_adapter<_DataType1> input1_ptr(q_ref, array1_in, size, true); _DataType1 *array_in = input1_ptr.get_ptr(); - DPNPC_ptr_adapter<_DataType2 *> choices_ptr(q_ref, choices1, choices_size); + 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) { DPNPC_ptr_adapter<_DataType2> choice_ptr(q_ref, choices[i], - choice_size); + choice_size, true); choices[i] = choice_ptr.get_ptr(); } - DPNPC_ptr_adapter<_DataType2> result1_ptr(q_ref, result1, size, false, + DPNPC_ptr_adapter<_DataType2> result1_ptr(q_ref, result1, size, true, true); _DataType2 *result = result1_ptr.get_ptr(); @@ -88,6 +88,7 @@ DPCTLSyclEventRef dpnp_choose_c(DPCTLSyclQueueRef q_ref, }; sycl::event event = q.submit(kernel_func); + result1_ptr.depends_on(event); event_ref = reinterpret_cast(&event); From be09527fd0e7b4231578788d375a27e24eee1f1e Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Thu, 19 Sep 2024 10:31:22 +0200 Subject: [PATCH 2/7] Keep extra memcopy only for the result array --- dpnp/backend/kernels/dpnp_krnl_indexing.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp index ebb81eef0e61..2f16d706223e 100644 --- a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp @@ -60,15 +60,15 @@ DPCTLSyclEventRef dpnp_choose_c(DPCTLSyclQueueRef q_ref, sycl::queue q = *(reinterpret_cast(q_ref)); - DPNPC_ptr_adapter<_DataType1> input1_ptr(q_ref, array1_in, size, true); + 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, true); + DPNPC_ptr_adapter<_DataType2 *> choices_ptr(q_ref, choices1, choices_size); _DataType2 **choices = choices_ptr.get_ptr(); for (size_t i = 0; i < choices_size; ++i) { DPNPC_ptr_adapter<_DataType2> choice_ptr(q_ref, choices[i], - choice_size, true); + choice_size); choices[i] = choice_ptr.get_ptr(); } From 544fb8874100a070b52109c93b1235a75eec25f5 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Thu, 19 Sep 2024 12:03:50 +0200 Subject: [PATCH 3/7] Add missing memory copy from host to device memory --- dpnp/backend/kernels/dpnp_krnl_indexing.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp index 2f16d706223e..e98780e18145 100644 --- a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp @@ -63,7 +63,9 @@ 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) { @@ -72,7 +74,7 @@ DPCTLSyclEventRef dpnp_choose_c(DPCTLSyclQueueRef q_ref, choices[i] = choice_ptr.get_ptr(); } - DPNPC_ptr_adapter<_DataType2> result1_ptr(q_ref, result1, size, true, + DPNPC_ptr_adapter<_DataType2> result1_ptr(q_ref, result1, size, false, true); _DataType2 *result = result1_ptr.get_ptr(); @@ -88,7 +90,7 @@ DPCTLSyclEventRef dpnp_choose_c(DPCTLSyclQueueRef q_ref, }; sycl::event event = q.submit(kernel_func); - result1_ptr.depends_on(event); + choices_ptr.depends_on(event); event_ref = reinterpret_cast(&event); From 948213b6da9bf7d1dd39a8c0384d3e165d96c0e3 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Thu, 19 Sep 2024 12:04:43 +0200 Subject: [PATCH 4/7] Applied pre-commit formatting --- dpnp/backend/kernels/dpnp_krnl_indexing.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp index e98780e18145..eda36511e048 100644 --- a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp @@ -65,7 +65,8 @@ DPCTLSyclEventRef dpnp_choose_c(DPCTLSyclQueueRef q_ref, // 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); + 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) { From a3b43cdf2c72393ad3581e677662862b3c6f3409 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Thu, 19 Sep 2024 12:06:53 +0200 Subject: [PATCH 5/7] Updated changelog --- CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index b697e4a227ca..115d8685c8d4 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -120,7 +120,7 @@ In addition, this release completes implementation of `dpnp.fft` module and adds * Resolved a possible race condition in `dpnp.inv` [#1940](https://github.com/IntelPython/dpnp/pull/1940) * 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) - +* Fixed a crash in `dpnp.choose` caused by missing memory copying from host to device memory [#2063](https://github.com/IntelPython/dpnp/pull/2063) ## [0.15.0] - 05/25/2024 From 9b336642a5fc871c74fae3fe3c886723f7c84a12 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Thu, 19 Sep 2024 12:07:24 +0200 Subject: [PATCH 6/7] Add empty line in changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 115d8685c8d4..9ffa41b4b715 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -122,6 +122,7 @@ In addition, this release completes implementation of `dpnp.fft` module and adds * Resolved an issue with input array of `usm_ndarray` passed into `dpnp.ix_` [#2047](https://github.com/IntelPython/dpnp/pull/2047) * Fixed a crash in `dpnp.choose` caused by missing memory copying from host to device memory [#2063](https://github.com/IntelPython/dpnp/pull/2063) + ## [0.15.0] - 05/25/2024 This release completes implementation of `dpnp.linalg` module and array creation routine, adds cumulative reductions and histogram functions. From f579bcec568cda7f708f90a7e189f4bfe1c77f0b Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Thu, 19 Sep 2024 12:20:47 +0200 Subject: [PATCH 7/7] Correct the description of issue in changelog --- CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 65dffa3b390a..ca2da727486f 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -121,7 +121,7 @@ In addition, this release completes implementation of `dpnp.fft` module and adds * Resolved a possible race condition in `dpnp.inv` [#1940](https://github.com/IntelPython/dpnp/pull/1940) * 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) -* Fixed a crash in `dpnp.choose` caused by missing memory copying from host to device memory [#2063](https://github.com/IntelPython/dpnp/pull/2063) +* 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