Skip to content

Commit 3f3d10f

Browse files
vtavanavtavana
authored
Merge branch 'master' into reciprocal_angle
2 parents 1927e8b + 5b25163 commit 3f3d10f

File tree

5 files changed

+224
-134
lines changed

5 files changed

+224
-134
lines changed

‎dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp

Lines changed: 9 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
//*****************************************************************************
2-
// Copyright (c) 2016-2023, Intel Corporation
2+
// Copyright (c) 2016-2024, Intel Corporation
33
// All rights reserved.
44
//
55
// Redistribution and use in source and binary forms, with or without
@@ -855,11 +855,12 @@ DPCTLSyclEventRef dpnp_tril_c(DPCTLSyclQueueRef q_ref,
855855
_DataType *array_m = input1_ptr.get_ptr();
856856
_DataType *result = result_ptr.get_ptr();
857857

858+
int *ids = new int[res_ndim];
859+
858860
if (ndim == 1) {
859861
for (size_t i = 0; i < res_size; ++i) {
860862
size_t n = res_size;
861863
size_t val = i;
862-
int ids[res_ndim];
863864
for (size_t j = 0; j < res_ndim; ++j) {
864865
n /= res_shape[j];
865866
size_t p = val / n;
@@ -886,7 +887,6 @@ DPCTLSyclEventRef dpnp_tril_c(DPCTLSyclQueueRef q_ref,
886887
for (size_t i = 0; i < res_size; ++i) {
887888
size_t n = res_size;
888889
size_t val = i;
889-
int ids[res_ndim];
890890
for (size_t j = 0; j < res_ndim; ++j) {
891891
n /= res_shape[j];
892892
size_t p = val / n;
@@ -909,6 +909,8 @@ DPCTLSyclEventRef dpnp_tril_c(DPCTLSyclQueueRef q_ref,
909909
}
910910
}
911911
}
912+
913+
delete[] ids;
912914
return DPCTLEvent_Copy(event_ref);
913915
}
914916

@@ -989,11 +991,12 @@ DPCTLSyclEventRef dpnp_triu_c(DPCTLSyclQueueRef q_ref,
989991
_DataType *array_m = input1_ptr.get_ptr();
990992
_DataType *result = result_ptr.get_ptr();
991993

994+
int *ids = new int[res_ndim];
995+
992996
if (ndim == 1) {
993997
for (size_t i = 0; i < res_size; ++i) {
994998
size_t n = res_size;
995999
size_t val = i;
996-
int ids[res_ndim];
9971000
for (size_t j = 0; j < res_ndim; ++j) {
9981001
n /= res_shape[j];
9991002
size_t p = val / n;
@@ -1020,7 +1023,6 @@ DPCTLSyclEventRef dpnp_triu_c(DPCTLSyclQueueRef q_ref,
10201023
for (size_t i = 0; i < res_size; ++i) {
10211024
size_t n = res_size;
10221025
size_t val = i;
1023-
int ids[res_ndim];
10241026
for (size_t j = 0; j < res_ndim; ++j) {
10251027
n /= res_shape[j];
10261028
size_t p = val / n;
@@ -1043,6 +1045,8 @@ DPCTLSyclEventRef dpnp_triu_c(DPCTLSyclQueueRef q_ref,
10431045
}
10441046
}
10451047
}
1048+
1049+
delete[] ids;
10461050
return DPCTLEvent_Copy(event_ref);
10471051
}
10481052

‎dpnp/backend/kernels/dpnp_krnl_indexing.cpp

Lines changed: 48 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
//*****************************************************************************
2-
// Copyright (c) 2016-2023, Intel Corporation
2+
// Copyright (c) 2016-2024, Intel Corporation
33
// All rights reserved.
44
//
55
// Redistribution and use in source and binary forms, with or without
@@ -233,18 +233,14 @@ DPCTLSyclEventRef dpnp_diagonal_c(DPCTLSyclQueueRef q_ref,
233233
continue;
234234
}
235235
else {
236-
size_t ind_input_size = ind_list.size() + 2;
237-
size_t ind_input_[ind_input_size];
238-
ind_input_[0] = i;
239-
ind_input_[1] = i + offset;
240-
size_t ind_output_size = ind_list.size() + 1;
241-
size_t ind_output_[ind_output_size];
242-
for (size_t k = 0; k < ind_list.size(); k++) {
243-
ind_input_[k + 2] = ind_list.at(k);
244-
ind_output_[k] = ind_list.at(k);
245-
}
246-
ind_output_[ind_list.size()] = i;
236+
std::vector<size_t> ind_input_{i, i + offset};
237+
ind_input_.insert(ind_input_.end(), ind_list.begin(),
238+
ind_list.end());
239+
240+
std::vector<size_t> ind_output_ = ind_list;
241+
ind_output_.push_back(i);
247242

243+
const size_t ind_output_size = ind_output_.size();
248244
size_t ind_output = 0;
249245
size_t n = 1;
250246
for (size_t k = 0; k < ind_output_size; k++) {
@@ -253,6 +249,7 @@ DPCTLSyclEventRef dpnp_diagonal_c(DPCTLSyclQueueRef q_ref,
253249
n *= res_shape[ind];
254250
}
255251

252+
const size_t ind_input_size = ind_input_.size();
256253
size_t ind_input = 0;
257254
size_t m = 1;
258255
for (size_t k = 0; k < ind_input_size; k++) {
@@ -423,11 +420,13 @@ DPCTLSyclEventRef dpnp_nonzero_c(DPCTLSyclQueueRef q_ref,
423420
long *result = result_ptr.get_ptr();
424421

425422
size_t idx = 0;
423+
size_t *ids = new size_t[ndim];
424+
426425
for (size_t i = 0; i < input1_size; ++i) {
427426
if (arr[i] != 0) {
428-
size_t ids[ndim];
429427
size_t ind1 = input1_size;
430428
size_t ind2 = i;
429+
431430
for (size_t k = 0; k < ndim; ++k) {
432431
ind1 = ind1 / shape[k];
433432
ids[k] = ind2 / ind1;
@@ -438,6 +437,7 @@ DPCTLSyclEventRef dpnp_nonzero_c(DPCTLSyclQueueRef q_ref,
438437
idx += 1;
439438
}
440439
}
440+
delete[] ids;
441441

442442
return event_ref;
443443
}
@@ -621,8 +621,6 @@ DPCTLSyclEventRef
621621
DPCTLSyclEventRef event_ref = nullptr;
622622
sycl::queue q = *(reinterpret_cast<sycl::queue *>(q_ref));
623623

624-
size_t res_ndim = ndim - 1;
625-
size_t res_shape[res_ndim];
626624
const size_t size_arr = std::accumulate(shape, shape + ndim, 1,
627625
std::multiplies<shape_elem_type>());
628626

@@ -635,14 +633,14 @@ DPCTLSyclEventRef
635633
_DataType *values = input2_ptr.get_ptr();
636634
_DataType *arr = result_ptr.get_ptr();
637635

638-
if (axis != res_ndim) {
639-
int ind = 0;
636+
if (axis != (ndim - 1)) {
637+
std::vector<size_t> res_shape;
640638
for (size_t i = 0; i < ndim; i++) {
641639
if (axis != i) {
642-
res_shape[ind] = shape[i];
643-
ind++;
640+
res_shape.push_back(shape[i]);
644641
}
645642
}
643+
size_t res_ndim = res_shape.size();
646644

647645
size_t prod = 1;
648646
for (size_t i = 0; i < res_ndim; ++i) {
@@ -651,20 +649,21 @@ DPCTLSyclEventRef
651649
}
652650
}
653651

654-
size_t ind_array[prod];
655-
bool bool_ind_array[prod];
652+
size_t *ind_array = new size_t[prod];
653+
bool *bool_ind_array = new bool[prod];
656654
for (size_t i = 0; i < prod; ++i) {
657655
bool_ind_array[i] = true;
658656
}
659-
size_t arr_shape_offsets[ndim];
657+
658+
size_t *arr_shape_offsets = new size_t[ndim];
660659
size_t acc = 1;
661660
for (size_t i = ndim - 1; i > 0; --i) {
662661
arr_shape_offsets[i] = acc;
663662
acc *= shape[i];
664663
}
665664
arr_shape_offsets[0] = acc;
666665

667-
size_t output_shape_offsets[res_ndim];
666+
size_t *output_shape_offsets = new size_t[res_ndim];
668667
acc = 1;
669668
if (res_ndim > 0) {
670669
for (size_t i = res_ndim - 1; i > 0; --i) {
@@ -680,31 +679,31 @@ DPCTLSyclEventRef
680679
}
681680

682681
// init result array
682+
size_t *xyz = new size_t[res_ndim];
683683
for (size_t result_idx = 0; result_idx < size_result; ++result_idx) {
684-
size_t xyz[res_ndim];
685684
size_t remainder = result_idx;
686685
for (size_t i = 0; i < res_ndim; ++i) {
687686
xyz[i] = remainder / output_shape_offsets[i];
688687
remainder = remainder - xyz[i] * output_shape_offsets[i];
689688
}
690689

691-
size_t source_axis[ndim];
692-
size_t result_axis_idx = 0;
693-
for (size_t idx = 0; idx < ndim; ++idx) {
694-
bool found = false;
695-
if (axis == idx) {
696-
found = true;
697-
}
698-
if (found) {
699-
source_axis[idx] = 0;
700-
}
701-
else {
702-
source_axis[idx] = xyz[result_axis_idx];
703-
result_axis_idx++;
704-
}
705-
}
690+
// FIXME: computed and unused. Commented out per compiler warning
691+
// size_t source_axis[ndim];
692+
// size_t result_axis_idx = 0;
693+
// for (size_t idx = 0; idx < ndim; ++idx) {
694+
// bool found = false;
695+
// if (axis == idx) {
696+
// found = true;
697+
// }
698+
// if (found) {
699+
// source_axis[idx] = 0;
700+
// }
701+
// else {
702+
// source_axis[idx] = xyz[result_axis_idx];
703+
// result_axis_idx++;
704+
// }
705+
// }
706706

707-
// FIXME: computed, but unused. Commented out per compiler warning
708707
// size_t source_idx = 0;
709708
// for (size_t i = 0; i < static_cast<size_t>(ndim); ++i)
710709
// {
@@ -714,25 +713,18 @@ DPCTLSyclEventRef
714713

715714
for (size_t source_idx = 0; source_idx < size_arr; ++source_idx) {
716715
// reconstruct x,y,z from linear source_idx
717-
size_t xyz[ndim];
718716
size_t remainder = source_idx;
719717
for (size_t i = 0; i < ndim; ++i) {
720718
xyz[i] = remainder / arr_shape_offsets[i];
721719
remainder = remainder - xyz[i] * arr_shape_offsets[i];
722720
}
723721

724722
// extract result axis
725-
size_t result_axis[res_ndim];
726-
size_t result_idx = 0;
723+
std::vector<size_t> result_axis;
727724
for (size_t idx = 0; idx < ndim; ++idx) {
728725
// try to find current idx in axis array
729-
bool found = false;
730-
if (axis == idx) {
731-
found = true;
732-
}
733-
if (!found) {
734-
result_axis[result_idx] = xyz[idx];
735-
result_idx++;
726+
if (axis != idx) {
727+
result_axis.push_back(xyz[idx]);
736728
}
737729
}
738730

@@ -756,6 +748,12 @@ DPCTLSyclEventRef
756748
arr[source_idx] = values[source_idx % values_size];
757749
}
758750
}
751+
752+
delete[] ind_array;
753+
delete[] bool_ind_array;
754+
delete[] arr_shape_offsets;
755+
delete[] output_shape_offsets;
756+
delete[] xyz;
759757
}
760758
else {
761759
for (size_t i = 0; i < size_arr; ++i) {

0 commit comments

Comments
 (0)