diff --git a/dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp b/dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp index b1af79e019d0..90c3d9a1ee53 100644 --- a/dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2016-2023, Intel Corporation +// Copyright (c) 2016-2024, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without @@ -855,11 +855,12 @@ DPCTLSyclEventRef dpnp_tril_c(DPCTLSyclQueueRef q_ref, _DataType *array_m = input1_ptr.get_ptr(); _DataType *result = result_ptr.get_ptr(); + int *ids = new int[res_ndim]; + if (ndim == 1) { for (size_t i = 0; i < res_size; ++i) { size_t n = res_size; size_t val = i; - int ids[res_ndim]; for (size_t j = 0; j < res_ndim; ++j) { n /= res_shape[j]; size_t p = val / n; @@ -886,7 +887,6 @@ DPCTLSyclEventRef dpnp_tril_c(DPCTLSyclQueueRef q_ref, for (size_t i = 0; i < res_size; ++i) { size_t n = res_size; size_t val = i; - int ids[res_ndim]; for (size_t j = 0; j < res_ndim; ++j) { n /= res_shape[j]; size_t p = val / n; @@ -909,6 +909,8 @@ DPCTLSyclEventRef dpnp_tril_c(DPCTLSyclQueueRef q_ref, } } } + + delete[] ids; return DPCTLEvent_Copy(event_ref); } @@ -989,11 +991,12 @@ DPCTLSyclEventRef dpnp_triu_c(DPCTLSyclQueueRef q_ref, _DataType *array_m = input1_ptr.get_ptr(); _DataType *result = result_ptr.get_ptr(); + int *ids = new int[res_ndim]; + if (ndim == 1) { for (size_t i = 0; i < res_size; ++i) { size_t n = res_size; size_t val = i; - int ids[res_ndim]; for (size_t j = 0; j < res_ndim; ++j) { n /= res_shape[j]; size_t p = val / n; @@ -1020,7 +1023,6 @@ DPCTLSyclEventRef dpnp_triu_c(DPCTLSyclQueueRef q_ref, for (size_t i = 0; i < res_size; ++i) { size_t n = res_size; size_t val = i; - int ids[res_ndim]; for (size_t j = 0; j < res_ndim; ++j) { n /= res_shape[j]; size_t p = val / n; @@ -1043,6 +1045,8 @@ DPCTLSyclEventRef dpnp_triu_c(DPCTLSyclQueueRef q_ref, } } } + + delete[] ids; return DPCTLEvent_Copy(event_ref); } diff --git a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp index e9addf36b707..7dc35fb5a803 100644 --- a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2016-2023, Intel Corporation +// Copyright (c) 2016-2024, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without @@ -233,18 +233,14 @@ DPCTLSyclEventRef dpnp_diagonal_c(DPCTLSyclQueueRef q_ref, continue; } else { - size_t ind_input_size = ind_list.size() + 2; - size_t ind_input_[ind_input_size]; - ind_input_[0] = i; - ind_input_[1] = i + offset; - size_t ind_output_size = ind_list.size() + 1; - size_t ind_output_[ind_output_size]; - for (size_t k = 0; k < ind_list.size(); k++) { - ind_input_[k + 2] = ind_list.at(k); - ind_output_[k] = ind_list.at(k); - } - ind_output_[ind_list.size()] = i; + std::vector ind_input_{i, i + offset}; + ind_input_.insert(ind_input_.end(), ind_list.begin(), + ind_list.end()); + + std::vector ind_output_ = ind_list; + ind_output_.push_back(i); + const size_t ind_output_size = ind_output_.size(); size_t ind_output = 0; size_t n = 1; for (size_t k = 0; k < ind_output_size; k++) { @@ -253,6 +249,7 @@ DPCTLSyclEventRef dpnp_diagonal_c(DPCTLSyclQueueRef q_ref, n *= res_shape[ind]; } + const size_t ind_input_size = ind_input_.size(); size_t ind_input = 0; size_t m = 1; for (size_t k = 0; k < ind_input_size; k++) { @@ -423,11 +420,13 @@ DPCTLSyclEventRef dpnp_nonzero_c(DPCTLSyclQueueRef q_ref, long *result = result_ptr.get_ptr(); size_t idx = 0; + size_t *ids = new size_t[ndim]; + for (size_t i = 0; i < input1_size; ++i) { if (arr[i] != 0) { - size_t ids[ndim]; size_t ind1 = input1_size; size_t ind2 = i; + for (size_t k = 0; k < ndim; ++k) { ind1 = ind1 / shape[k]; ids[k] = ind2 / ind1; @@ -438,6 +437,7 @@ DPCTLSyclEventRef dpnp_nonzero_c(DPCTLSyclQueueRef q_ref, idx += 1; } } + delete[] ids; return event_ref; } @@ -621,8 +621,6 @@ DPCTLSyclEventRef DPCTLSyclEventRef event_ref = nullptr; sycl::queue q = *(reinterpret_cast(q_ref)); - size_t res_ndim = ndim - 1; - size_t res_shape[res_ndim]; const size_t size_arr = std::accumulate(shape, shape + ndim, 1, std::multiplies()); @@ -635,14 +633,14 @@ DPCTLSyclEventRef _DataType *values = input2_ptr.get_ptr(); _DataType *arr = result_ptr.get_ptr(); - if (axis != res_ndim) { - int ind = 0; + if (axis != (ndim - 1)) { + std::vector res_shape; for (size_t i = 0; i < ndim; i++) { if (axis != i) { - res_shape[ind] = shape[i]; - ind++; + res_shape.push_back(shape[i]); } } + size_t res_ndim = res_shape.size(); size_t prod = 1; for (size_t i = 0; i < res_ndim; ++i) { @@ -651,12 +649,13 @@ DPCTLSyclEventRef } } - size_t ind_array[prod]; - bool bool_ind_array[prod]; + size_t *ind_array = new size_t[prod]; + bool *bool_ind_array = new bool[prod]; for (size_t i = 0; i < prod; ++i) { bool_ind_array[i] = true; } - size_t arr_shape_offsets[ndim]; + + size_t *arr_shape_offsets = new size_t[ndim]; size_t acc = 1; for (size_t i = ndim - 1; i > 0; --i) { arr_shape_offsets[i] = acc; @@ -664,7 +663,7 @@ DPCTLSyclEventRef } arr_shape_offsets[0] = acc; - size_t output_shape_offsets[res_ndim]; + size_t *output_shape_offsets = new size_t[res_ndim]; acc = 1; if (res_ndim > 0) { for (size_t i = res_ndim - 1; i > 0; --i) { @@ -680,31 +679,31 @@ DPCTLSyclEventRef } // init result array + size_t *xyz = new size_t[res_ndim]; for (size_t result_idx = 0; result_idx < size_result; ++result_idx) { - size_t xyz[res_ndim]; size_t remainder = result_idx; for (size_t i = 0; i < res_ndim; ++i) { xyz[i] = remainder / output_shape_offsets[i]; remainder = remainder - xyz[i] * output_shape_offsets[i]; } - size_t source_axis[ndim]; - size_t result_axis_idx = 0; - for (size_t idx = 0; idx < ndim; ++idx) { - bool found = false; - if (axis == idx) { - found = true; - } - if (found) { - source_axis[idx] = 0; - } - else { - source_axis[idx] = xyz[result_axis_idx]; - result_axis_idx++; - } - } + // FIXME: computed and unused. Commented out per compiler warning + // size_t source_axis[ndim]; + // size_t result_axis_idx = 0; + // for (size_t idx = 0; idx < ndim; ++idx) { + // bool found = false; + // if (axis == idx) { + // found = true; + // } + // if (found) { + // source_axis[idx] = 0; + // } + // else { + // source_axis[idx] = xyz[result_axis_idx]; + // result_axis_idx++; + // } + // } - // FIXME: computed, but unused. Commented out per compiler warning // size_t source_idx = 0; // for (size_t i = 0; i < static_cast(ndim); ++i) // { @@ -714,7 +713,6 @@ DPCTLSyclEventRef for (size_t source_idx = 0; source_idx < size_arr; ++source_idx) { // reconstruct x,y,z from linear source_idx - size_t xyz[ndim]; size_t remainder = source_idx; for (size_t i = 0; i < ndim; ++i) { xyz[i] = remainder / arr_shape_offsets[i]; @@ -722,17 +720,11 @@ DPCTLSyclEventRef } // extract result axis - size_t result_axis[res_ndim]; - size_t result_idx = 0; + std::vector result_axis; for (size_t idx = 0; idx < ndim; ++idx) { // try to find current idx in axis array - bool found = false; - if (axis == idx) { - found = true; - } - if (!found) { - result_axis[result_idx] = xyz[idx]; - result_idx++; + if (axis != idx) { + result_axis.push_back(xyz[idx]); } } @@ -756,6 +748,12 @@ DPCTLSyclEventRef arr[source_idx] = values[source_idx % values_size]; } } + + delete[] ind_array; + delete[] bool_ind_array; + delete[] arr_shape_offsets; + delete[] output_shape_offsets; + delete[] xyz; } else { for (size_t i = 0; i < size_arr; ++i) { diff --git a/dpnp/backend/kernels/dpnp_krnl_linalg.cpp b/dpnp/backend/kernels/dpnp_krnl_linalg.cpp index 5bf54c2b84d3..5e78a7cda176 100644 --- a/dpnp/backend/kernels/dpnp_krnl_linalg.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_linalg.cpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2016-2023, Intel Corporation +// Copyright (c) 2016-2024, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without @@ -172,50 +172,52 @@ DPCTLSyclEventRef dpnp_det_c(DPCTLSyclQueueRef q_ref, _DataType *array_1 = input1_ptr.get_ptr(); _DataType *result = result_ptr.get_ptr(); + _DataType *matrix = new _DataType[n * n]; + _DataType *elems = new _DataType[n * n]; + for (size_t i = 0; i < size_out; i++) { - _DataType matrix[n][n]; if (size_out > 1) { - _DataType elems[n * n]; for (size_t j = i * n * n; j < (i + 1) * n * n; j++) { elems[j - i * n * n] = array_1[j]; } for (size_t j = 0; j < n; j++) { for (size_t k = 0; k < n; k++) { - matrix[j][k] = elems[j * n + k]; + matrix[j * n + k] = elems[j * n + k]; } } } else { for (size_t j = 0; j < n; j++) { for (size_t k = 0; k < n; k++) { - matrix[j][k] = array_1[j * n + k]; + matrix[j * n + k] = array_1[j * n + k]; } } } _DataType det_val = 1; for (size_t l = 0; l < n; l++) { - if (matrix[l][l] == 0) { + if (matrix[l * n + l] == 0) { for (size_t j = l; j < n; j++) { - if (matrix[j][l] != 0) { + if (matrix[j * n + l] != 0) { for (size_t k = l; k < n; k++) { - _DataType c = matrix[l][k]; - matrix[l][k] = -1 * matrix[j][k]; - matrix[j][k] = c; + _DataType c = matrix[l * n + k]; + matrix[l * n + k] = -1 * matrix[j * n + k]; + matrix[j * n + k] = c; } break; } - if (j == n - 1 and matrix[j][l] == 0) { + if (j == n - 1 and matrix[j * n + l] == 0) { det_val = 0; } } } if (det_val != 0) { for (size_t j = l + 1; j < n; j++) { - _DataType quotient = -(matrix[j][l] / matrix[l][l]); + _DataType quotient = + -(matrix[j * n + l] / matrix[l * n + l]); for (size_t k = l + 1; k < n; k++) { - matrix[j][k] += quotient * matrix[l][k]; + matrix[j * n + k] += quotient * matrix[l * n + k]; } } } @@ -223,13 +225,15 @@ DPCTLSyclEventRef dpnp_det_c(DPCTLSyclQueueRef q_ref, if (det_val != 0) { for (size_t l = 0; l < n; l++) { - det_val *= matrix[l][l]; + det_val *= matrix[l * n + l]; } } result[i] = det_val; } + delete[] elems; + delete[] matrix; return event_ref; } @@ -291,50 +295,50 @@ DPCTLSyclEventRef dpnp_inv_c(DPCTLSyclQueueRef q_ref, size_t n = shape[0]; - _ResultType a_arr[n][n]; - _ResultType e_arr[n][n]; + _ResultType *a_arr = new _ResultType[n * n]; + _ResultType *e_arr = new _ResultType[n * n]; for (size_t i = 0; i < n; ++i) { for (size_t j = 0; j < n; ++j) { - a_arr[i][j] = array_1[i * n + j]; + a_arr[i * n + j] = array_1[i * n + j]; if (i == j) { - e_arr[i][j] = 1; + e_arr[i * n + j] = 1; } else { - e_arr[i][j] = 0; + e_arr[i * n + j] = 0; } } } for (size_t k = 0; k < n; ++k) { - if (a_arr[k][k] == 0) { + if (a_arr[k * n + k] == 0) { for (size_t i = k; i < n; ++i) { - if (a_arr[i][k] != 0) { + if (a_arr[i * n + k] != 0) { for (size_t j = 0; j < n; ++j) { - float c = a_arr[k][j]; - a_arr[k][j] = a_arr[i][j]; - a_arr[i][j] = c; - float c_e = e_arr[k][j]; - e_arr[k][j] = e_arr[i][j]; - e_arr[i][j] = c_e; + float c = a_arr[k * n + j]; + a_arr[k * n + j] = a_arr[i * n + j]; + a_arr[i * n + j] = c; + float c_e = e_arr[k * n + j]; + e_arr[k * n + j] = e_arr[i * n + j]; + e_arr[i * n + j] = c_e; } break; } } } - float temp = a_arr[k][k]; + float temp = a_arr[k * n + k]; for (size_t j = 0; j < n; ++j) { - a_arr[k][j] = a_arr[k][j] / temp; - e_arr[k][j] = e_arr[k][j] / temp; + a_arr[k * n + j] = a_arr[k * n + j] / temp; + e_arr[k * n + j] = e_arr[k * n + j] / temp; } for (size_t i = k + 1; i < n; ++i) { - temp = a_arr[i][k]; + temp = a_arr[i * n + k]; for (size_t j = 0; j < n; j++) { - a_arr[i][j] = a_arr[i][j] - a_arr[k][j] * temp; - e_arr[i][j] = e_arr[i][j] - e_arr[k][j] * temp; + a_arr[i * n + j] = a_arr[i * n + j] - a_arr[k * n + j] * temp; + e_arr[i * n + j] = e_arr[i * n + j] - e_arr[k * n + j] * temp; } } } @@ -344,20 +348,24 @@ DPCTLSyclEventRef dpnp_inv_c(DPCTLSyclQueueRef q_ref, for (size_t i = 0; i < ind_k; ++i) { size_t ind_i = ind_k - 1 - i; - float temp = a_arr[ind_i][ind_k]; + float temp = a_arr[ind_i * n + ind_k]; for (size_t j = 0; j < n; ++j) { - a_arr[ind_i][j] = a_arr[ind_i][j] - a_arr[ind_k][j] * temp; - e_arr[ind_i][j] = e_arr[ind_i][j] - e_arr[ind_k][j] * temp; + a_arr[ind_i * n + j] = + a_arr[ind_i * n + j] - a_arr[ind_k * n + j] * temp; + e_arr[ind_i * n + j] = + e_arr[ind_i * n + j] - e_arr[ind_k * n + j] * temp; } } } for (size_t i = 0; i < n; ++i) { for (size_t j = 0; j < n; ++j) { - result[i * n + j] = e_arr[i][j]; + result[i * n + j] = e_arr[i * n + j]; } } + delete[] a_arr; + delete[] e_arr; return event_ref; } diff --git a/dpnp/backend/kernels/dpnp_krnl_sorting.cpp b/dpnp/backend/kernels/dpnp_krnl_sorting.cpp index 8a5a8ea69aaf..ac4992466e2f 100644 --- a/dpnp/backend/kernels/dpnp_krnl_sorting.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_sorting.cpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2016-2023, Intel Corporation +// Copyright (c) 2016-2024, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without @@ -192,11 +192,12 @@ DPCTLSyclEventRef dpnp_partition_c(DPCTLSyclQueueRef q_ref, auto arr_to_result_event = q.memcpy(result, arr, size * sizeof(_DataType)); arr_to_result_event.wait(); + _DataType *matrix = new _DataType[shape_[ndim - 1]]; + for (size_t i = 0; i < size_; ++i) { size_t ind_begin = i * shape_[ndim - 1]; size_t ind_end = (i + 1) * shape_[ndim - 1] - 1; - _DataType matrix[shape_[ndim - 1]]; for (size_t j = ind_begin; j < ind_end + 1; ++j) { size_t ind = j - ind_begin; matrix[ind] = arr2[j]; @@ -242,6 +243,7 @@ DPCTLSyclEventRef dpnp_partition_c(DPCTLSyclQueueRef q_ref, event.wait(); + delete[] matrix; sycl::free(shape, q); return event_ref; diff --git a/dpnp/backend/kernels/dpnp_krnl_statistics.cpp b/dpnp/backend/kernels/dpnp_krnl_statistics.cpp index 8f685c97cb38..97df3c2d7f11 100644 --- a/dpnp/backend/kernels/dpnp_krnl_statistics.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_statistics.cpp @@ -358,9 +358,7 @@ DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, } } else { - size_t res_ndim = ndim - naxis; - size_t res_shape[res_ndim]; - int ind = 0; + std::vector res_shape; for (size_t i = 0; i < ndim; ++i) { bool found = false; for (size_t j = 0; j < naxis; ++j) { @@ -370,28 +368,24 @@ DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, } } if (!found) { - res_shape[ind] = shape[i]; - ind++; + res_shape.push_back(shape[i]); } } + const size_t res_ndim = res_shape.size(); - size_t input_shape_offsets[ndim]; size_t acc = 1; - for (size_t i = ndim - 1; i > 0; --i) { - input_shape_offsets[i] = acc; + std::vector input_shape_offsets{acc}; + for (size_t i = ndim - 2; i > 0; --i) { acc *= shape[i]; + input_shape_offsets.insert(input_shape_offsets.begin(), acc); } - input_shape_offsets[0] = acc; - size_t output_shape_offsets[res_ndim]; acc = 1; - if (res_ndim > 0) { - for (size_t i = res_ndim - 1; i > 0; --i) { - output_shape_offsets[i] = acc; - acc *= res_shape[i]; - } + std::vector output_shape_offsets{acc}; + for (size_t i = res_ndim - 2; i > 0; --i) { + acc *= res_shape[i]; + output_shape_offsets.insert(output_shape_offsets.begin(), acc); } - output_shape_offsets[0] = acc; size_t size_result = 1; for (size_t i = 0; i < res_ndim; ++i) { @@ -399,15 +393,16 @@ DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, } // init result array + size_t *xyz = new size_t[res_ndim]; + size_t *source_axis = new size_t[ndim]; + size_t *result_axis = new size_t[res_ndim]; for (size_t result_idx = 0; result_idx < size_result; ++result_idx) { - size_t xyz[res_ndim]; size_t remainder = result_idx; for (size_t i = 0; i < res_ndim; ++i) { xyz[i] = remainder / output_shape_offsets[i]; remainder = remainder - xyz[i] * output_shape_offsets[i]; } - size_t source_axis[ndim]; size_t result_axis_idx = 0; for (size_t idx = 0; idx < ndim; ++idx) { bool found = false; @@ -436,7 +431,6 @@ DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, for (size_t source_idx = 0; source_idx < size_input; ++source_idx) { // reconstruct x,y,z from linear source_idx - size_t xyz[ndim]; size_t remainder = source_idx; for (size_t i = 0; i < ndim; ++i) { xyz[i] = remainder / input_shape_offsets[i]; @@ -444,7 +438,6 @@ DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, } // extract result axis - size_t result_axis[res_ndim]; size_t result_idx = 0; for (size_t idx = 0; idx < ndim; ++idx) { // try to find current idx in axis array @@ -471,11 +464,57 @@ DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, result[result_offset] = array_1[source_idx]; } } + + delete[] xyz; + delete[] source_axis; + delete[] result_axis; } return event_ref; } +// Explicit instantiation of the function, since dpnp_max_c() is used by +// other template functions, but implicit instantiation is not applied anymore. +template DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, + void *, + void *, + const size_t, + const shape_elem_type *, + size_t, + const shape_elem_type *, + size_t, + const DPCTLEventVectorRef); + +template DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, + void *, + void *, + const size_t, + const shape_elem_type *, + size_t, + const shape_elem_type *, + size_t, + const DPCTLEventVectorRef); + +template DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, + void *, + void *, + const size_t, + const shape_elem_type *, + size_t, + const shape_elem_type *, + size_t, + const DPCTLEventVectorRef); + +template DPCTLSyclEventRef dpnp_max_c(DPCTLSyclQueueRef q_ref, + void *, + void *, + const size_t, + const shape_elem_type *, + size_t, + const shape_elem_type *, + size_t, + const DPCTLEventVectorRef); + template void dpnp_max_c(void *array1_in, void *result1, @@ -730,9 +769,7 @@ DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, } } else { - size_t res_ndim = ndim - naxis; - size_t res_shape[res_ndim]; - int ind = 0; + std::vector res_shape; for (size_t i = 0; i < ndim; i++) { bool found = false; for (size_t j = 0; j < naxis; j++) { @@ -742,28 +779,24 @@ DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, } } if (!found) { - res_shape[ind] = shape[i]; - ind++; + res_shape.push_back(shape[i]); } } + const size_t res_ndim = res_shape.size(); - size_t input_shape_offsets[ndim]; size_t acc = 1; - for (size_t i = ndim - 1; i > 0; --i) { - input_shape_offsets[i] = acc; + std::vector input_shape_offsets{acc}; + for (size_t i = ndim - 2; i > 0; --i) { acc *= shape[i]; + input_shape_offsets.insert(input_shape_offsets.begin(), acc); } - input_shape_offsets[0] = acc; - size_t output_shape_offsets[res_ndim]; acc = 1; - if (res_ndim > 0) { - for (size_t i = res_ndim - 1; i > 0; --i) { - output_shape_offsets[i] = acc; - acc *= res_shape[i]; - } + std::vector output_shape_offsets{acc}; + for (size_t i = res_ndim - 2; i > 0; --i) { + acc *= res_shape[i]; + output_shape_offsets.insert(output_shape_offsets.begin(), acc); } - output_shape_offsets[0] = acc; size_t size_result = 1; for (size_t i = 0; i < res_ndim; ++i) { @@ -771,15 +804,16 @@ DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, } // init result array + size_t *xyz = new size_t[res_ndim]; + size_t *source_axis = new size_t[ndim]; + size_t *result_axis = new size_t[res_ndim]; for (size_t result_idx = 0; result_idx < size_result; ++result_idx) { - size_t xyz[res_ndim]; size_t remainder = result_idx; for (size_t i = 0; i < res_ndim; ++i) { xyz[i] = remainder / output_shape_offsets[i]; remainder = remainder - xyz[i] * output_shape_offsets[i]; } - size_t source_axis[ndim]; size_t result_axis_idx = 0; for (size_t idx = 0; idx < ndim; ++idx) { bool found = false; @@ -808,7 +842,6 @@ DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, for (size_t source_idx = 0; source_idx < size_input; ++source_idx) { // reconstruct x,y,z from linear source_idx - size_t xyz[ndim]; size_t remainder = source_idx; for (size_t i = 0; i < ndim; ++i) { xyz[i] = remainder / input_shape_offsets[i]; @@ -816,7 +849,6 @@ DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, } // extract result axis - size_t result_axis[res_ndim]; size_t result_idx = 0; for (size_t idx = 0; idx < ndim; ++idx) { // try to find current idx in axis array @@ -843,11 +875,57 @@ DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, result[result_offset] = array_1[source_idx]; } } + + delete[] xyz; + delete[] source_axis; + delete[] result_axis; } return event_ref; } +// Explicit instantiation of the function, since dpnp_min_c() is used by +// other template functions, but implicit instantiation is not applied anymore. +template DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, + void *, + void *, + const size_t, + const shape_elem_type *, + size_t, + const shape_elem_type *, + size_t, + const DPCTLEventVectorRef); + +template DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, + void *, + void *, + const size_t, + const shape_elem_type *, + size_t, + const shape_elem_type *, + size_t, + const DPCTLEventVectorRef); + +template DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, + void *, + void *, + const size_t, + const shape_elem_type *, + size_t, + const shape_elem_type *, + size_t, + const DPCTLEventVectorRef); + +template DPCTLSyclEventRef dpnp_min_c(DPCTLSyclQueueRef q_ref, + void *, + void *, + const size_t, + const shape_elem_type *, + size_t, + const shape_elem_type *, + size_t, + const DPCTLEventVectorRef); + template void dpnp_min_c(void *array1_in, void *result1,