Skip to content

Commit f1f3d7e

Browse files
samir-nasibliAlexander-Makaryev
samir-nasibli
and
Alexander-Makaryev
authored
syle 9/14/2021 (#931)
Co-authored-by: Alexander-Makaryev <[email protected]>
1 parent 69557ee commit f1f3d7e

11 files changed

+91
-61
lines changed

‎dpnp/backend/include/dpnp_iface.hpp

Lines changed: 27 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -253,8 +253,12 @@ INP_DLLEXPORT void dpnp_nanvar_c(void* array, void* mask_arr, void* result, cons
253253
* @param [in] j Number input array.
254254
*/
255255
template <typename _DataType>
256-
INP_DLLEXPORT void
257-
dpnp_nonzero_c(const void* array1, void* result1, const size_t result_size, const size_t* shape, const size_t ndim, const size_t j);
256+
INP_DLLEXPORT void dpnp_nonzero_c(const void* array1,
257+
void* result1,
258+
const size_t result_size,
259+
const size_t* shape,
260+
const size_t ndim,
261+
const size_t j);
258262

259263
/**
260264
* @ingroup BACKEND_API
@@ -634,8 +638,13 @@ INP_DLLEXPORT void dpnp_diag_indices_c(void* result1, size_t size);
634638
* @param [in] ndim Number of elements in shape.
635639
*/
636640
template <typename _DataType>
637-
INP_DLLEXPORT void dpnp_diagonal_c(
638-
void* array1_in, const size_t input1_size, void* result1, const size_t offset, size_t* shape, size_t* res_shape, const size_t res_ndim);
641+
INP_DLLEXPORT void dpnp_diagonal_c(void* array1_in,
642+
const size_t input1_size,
643+
void* result1,
644+
const size_t offset,
645+
size_t* shape,
646+
size_t* res_shape,
647+
const size_t res_ndim);
639648

640649
/**
641650
* @ingroup BACKEND_API
@@ -695,8 +704,13 @@ INP_DLLEXPORT void dpnp_matrix_rank_c(void* array1_in, void* result1, size_t* sh
695704
* @param [in] naxis Number of elements in axis.
696705
*/
697706
template <typename _DataType>
698-
INP_DLLEXPORT void
699-
dpnp_max_c(void* array1_in, void* result1, const size_t result_size, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis);
707+
INP_DLLEXPORT void dpnp_max_c(void* array1_in,
708+
void* result1,
709+
const size_t result_size,
710+
const size_t* shape,
711+
size_t ndim,
712+
const size_t* axis,
713+
size_t naxis);
700714

701715
/**
702716
* @ingroup BACKEND_API
@@ -741,8 +755,13 @@ INP_DLLEXPORT void
741755
* @param [in] naxis Number of elements in axis.
742756
*/
743757
template <typename _DataType>
744-
INP_DLLEXPORT void
745-
dpnp_min_c(void* array, void* result, const size_t result_size, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis);
758+
INP_DLLEXPORT void dpnp_min_c(void* array,
759+
void* result,
760+
const size_t result_size,
761+
const size_t* shape,
762+
size_t ndim,
763+
const size_t* axis,
764+
size_t naxis);
746765

747766
/**
748767
* @ingroup BACKEND_API

‎dpnp/backend/kernels/dpnp_krnl_indexing.cpp

Lines changed: 13 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -45,8 +45,13 @@ template <typename _DataType>
4545
class dpnp_diagonal_c_kernel;
4646

4747
template <typename _DataType>
48-
void dpnp_diagonal_c(
49-
void* array1_in, const size_t input1_size, void* result1, const size_t offset, size_t* shape, size_t* res_shape, const size_t res_ndim)
48+
void dpnp_diagonal_c(void* array1_in,
49+
const size_t input1_size,
50+
void* result1,
51+
const size_t offset,
52+
size_t* shape,
53+
size_t* res_shape,
54+
const size_t res_ndim)
5055
{
5156
const size_t res_size = std::accumulate(res_shape, res_shape + res_ndim, 1, std::multiplies<size_t>());
5257
if (!(res_size && input1_size))
@@ -196,7 +201,12 @@ void dpnp_fill_diagonal_c(void* array1_in, void* val_in, size_t* shape, const si
196201
}
197202

198203
template <typename _DataType>
199-
void dpnp_nonzero_c(const void* in_array1, void* result1, const size_t result_size, const size_t* shape, const size_t ndim, const size_t j)
204+
void dpnp_nonzero_c(const void* in_array1,
205+
void* result1,
206+
const size_t result_size,
207+
const size_t* shape,
208+
const size_t ndim,
209+
const size_t j)
200210
{
201211
if ((in_array1 == nullptr) || (result1 == nullptr))
202212
{
@@ -215,7 +225,6 @@ void dpnp_nonzero_c(const void* in_array1, void* result1, const size_t result_si
215225
const _DataType* arr = input1_ptr.get_ptr();
216226
long* result = result_ptr.get_ptr();
217227

218-
219228
size_t idx = 0;
220229
for (size_t i = 0; i < input1_size; ++i)
221230
{

‎dpnp/backend/kernels/dpnp_krnl_mathematical.cpp

Lines changed: 6 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -374,8 +374,8 @@ void dpnp_remainder_c(void* result_out,
374374
_DataType_input2* input2_data = input2_ptr.get_ptr();
375375
_DataType_output* result = reinterpret_cast<_DataType_output*>(result_out);
376376

377-
std::vector<size_t> result_shape = get_result_shape(input1_shape, input1_shape_ndim,
378-
input2_shape, input2_shape_ndim);
377+
std::vector<size_t> result_shape =
378+
get_result_shape(input1_shape, input1_shape_ndim, input2_shape, input2_shape_ndim);
379379

380380
DPNPC_id<_DataType_input1>* input1_it;
381381
const size_t input1_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input1>);
@@ -390,18 +390,17 @@ void dpnp_remainder_c(void* result_out,
390390
new (input2_it) DPNPC_id<_DataType_input2>(input2_data, input2_shape, input2_shape_ndim);
391391

392392
input2_it->broadcast_to_shape(result_shape);
393-
393+
394394
const size_t result_size = input1_it->get_output_size();
395395

396-
cl::sycl::range<1> gws(result_size);
396+
cl::sycl::range<1> gws(result_size);
397397
auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
398398
const size_t i = global_id[0];
399399
const _DataType_output input1_elem = (*input1_it)[i];
400400
const _DataType_output input2_elem = (*input2_it)[i];
401401
double fmod_res = cl::sycl::fmod((double)input1_elem, (double)input2_elem);
402402
double add = fmod_res + input2_elem;
403403
result[i] = cl::sycl::fmod(add, (double)input2_elem);
404-
405404
};
406405
auto kernel_func = [&](cl::sycl::handler& cgh) {
407406
cgh.parallel_for<class dpnp_remainder_c_kernel<_DataType_output, _DataType_input1, _DataType_input2>>(
@@ -412,9 +411,8 @@ void dpnp_remainder_c(void* result_out,
412411

413412
if (input1_size == input2_size)
414413
{
415-
if constexpr ((std::is_same<_DataType_input1, double>::value ||
416-
std::is_same<_DataType_input1, float>::value) &&
417-
std::is_same<_DataType_input2, _DataType_input1>::value)
414+
if constexpr ((std::is_same<_DataType_input1, double>::value || std::is_same<_DataType_input1, float>::value) &&
415+
std::is_same<_DataType_input2, _DataType_input1>::value)
418416
{
419417
event = oneapi::mkl::vm::fmod(DPNP_QUEUE, input1_size, input1_data, input2_data, result);
420418
event.wait();
@@ -436,7 +434,6 @@ void dpnp_remainder_c(void* result_out,
436434

437435
input1_it->~DPNPC_id();
438436
input2_it->~DPNPC_id();
439-
440437
}
441438

442439
template <typename _KernelNameSpecialization1, typename _KernelNameSpecialization2, typename _KernelNameSpecialization3>

‎dpnp/backend/kernels/dpnp_krnl_random.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1253,7 +1253,7 @@ void dpnp_rng_vonmises_large_kappa_c(void* result, const _DataType mu, const _Da
12531253
{
12541254
return;
12551255
}
1256-
1256+
12571257
DPNPC_ptr_adapter<_DataType> result1_ptr(result, size, true, true);
12581258
_DataType* result1 = result1_ptr.get_ptr();
12591259

@@ -1547,7 +1547,7 @@ void dpnp_rng_zipf_c(void* result, const _DataType a, const size_t size)
15471547
long X;
15481548
const _DataType d_zero = 0.0;
15491549
const _DataType d_one = 1.0;
1550-
1550+
15511551
DPNPC_ptr_adapter<_DataType> result1_ptr(result, size, true, true);
15521552
_DataType* result1 = result1_ptr.get_ptr();
15531553

‎dpnp/backend/kernels/dpnp_krnl_reduction.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -124,7 +124,8 @@ void dpnp_sum_c(void* result_out,
124124
policy, input_it.begin(output_id), input_it.end(output_id), init, std::plus<_DataType_output>());
125125
policy.queue().wait(); // TODO move out of the loop
126126

127-
dpnp_memory_memcpy_c(result + output_id, &accumulator, sizeof(_DataType_output)); // result[output_id] = accumulator;
127+
dpnp_memory_memcpy_c(
128+
result + output_id, &accumulator, sizeof(_DataType_output)); // result[output_id] = accumulator;
128129
}
129130

130131
return;
@@ -184,7 +185,8 @@ void dpnp_prod_c(void* result_out,
184185
policy, input_it.begin(output_id), input_it.end(output_id), init, std::multiplies<_DataType_output>());
185186
policy.queue().wait(); // TODO move out of the loop
186187

187-
dpnp_memory_memcpy_c(result + output_id, &accumulator, sizeof(_DataType_output)); // result[output_id] = accumulator;
188+
dpnp_memory_memcpy_c(
189+
result + output_id, &accumulator, sizeof(_DataType_output)); // result[output_id] = accumulator;
188190
}
189191

190192
return;

‎dpnp/backend/kernels/dpnp_krnl_sorting.cpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -118,8 +118,6 @@ void dpnp_partition_c(
118118
_DataType* arr2 = input2_ptr.get_ptr();
119119
_DataType* result = result1_ptr.get_ptr();
120120

121-
122-
123121
auto arr_to_result_event = DPNP_QUEUE.memcpy(result, arr, size * sizeof(_DataType));
124122
arr_to_result_event.wait();
125123

‎dpnp/backend/kernels/dpnp_krnl_statistics.cpp

Lines changed: 16 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -52,8 +52,8 @@ void dpnp_correlate_c(void* result_out,
5252
(void)where;
5353

5454
dpnp_dot_c<_DataType_output, _DataType_input1, _DataType_input2>(result_out,
55-
42, // dummy result_size
56-
42, // dummy result_ndim
55+
42, // dummy result_size
56+
42, // dummy result_ndim
5757
NULL, // dummy result_shape
5858
NULL, // dummy result_strides
5959
input1_in,
@@ -153,7 +153,13 @@ template <typename _DataType>
153153
class dpnp_max_c_kernel;
154154

155155
template <typename _DataType>
156-
void dpnp_max_c(void* array1_in, void* result1, const size_t result_size, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis)
156+
void dpnp_max_c(void* array1_in,
157+
void* result1,
158+
const size_t result_size,
159+
const size_t* shape,
160+
size_t ndim,
161+
const size_t* axis,
162+
size_t naxis)
157163
{
158164
const size_t size_input = std::accumulate(shape, shape + ndim, 1, std::multiplies<size_t>());
159165
if (!size_input)
@@ -416,7 +422,13 @@ template <typename _DataType>
416422
class dpnp_min_c_kernel;
417423

418424
template <typename _DataType>
419-
void dpnp_min_c(void* array1_in, void* result1, const size_t result_size, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis)
425+
void dpnp_min_c(void* array1_in,
426+
void* result1,
427+
const size_t result_size,
428+
const size_t* shape,
429+
size_t ndim,
430+
const size_t* axis,
431+
size_t naxis)
420432
{
421433
__attribute__((unused)) void* tmp = (void*)(axis + naxis);
422434

‎dpnp/backend/src/constants.cpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -30,14 +30,14 @@
3030
void* python_constants::py_none = nullptr;
3131
void* python_constants::py_nan = nullptr;
3232

33-
void dpnp_python_constants_initialize_c(void *_py_none, void *_py_nan)
33+
void dpnp_python_constants_initialize_c(void* _py_none, void* _py_nan)
3434
{
3535
python_constants::py_none = _py_none;
3636
python_constants::py_nan = _py_nan;
37-
// std::cout << "========dpnp_python_constants_initialize_c=============" << std::endl;
38-
// std::cout << "\t None=" << _py_none
39-
// << "\n\t NaN=" << _py_nan
40-
// << "\n\t py_none=" << python_constants::py_none
41-
// << "\n\t py_nan=" << python_constants::py_nan
42-
// << std::endl;
37+
// std::cout << "========dpnp_python_constants_initialize_c=============" << std::endl;
38+
// std::cout << "\t None=" << _py_none
39+
// << "\n\t NaN=" << _py_nan
40+
// << "\n\t py_none=" << python_constants::py_none
41+
// << "\n\t py_nan=" << python_constants::py_nan
42+
// << std::endl;
4343
}

‎dpnp/backend/src/constants.hpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,6 @@
3333
#define INP_DLLEXPORT
3434
#endif
3535

36-
3736
/**
3837
* This is container for the constants from Python interpreter and other modules. These constants are subject to use
3938
* in algorithms.
@@ -53,6 +52,6 @@ struct python_constants
5352
* @param [in] py_none Python NONE representation
5453
* @param [in] py_nan Python NAN representation
5554
*/
56-
INP_DLLEXPORT void dpnp_python_constants_initialize_c(void *py_none, void *py_nan);
55+
INP_DLLEXPORT void dpnp_python_constants_initialize_c(void* py_none, void* py_nan);
5756

5857
#endif // CONSTANTS_H

‎dpnp/backend/src/dpnpc_memory_adapter.hpp

Lines changed: 11 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -43,12 +43,12 @@
4343
template <typename _DataType>
4444
class DPNPC_ptr_adapter final
4545
{
46-
void* aux_ptr = nullptr; /**< pointer to allocated memory by this adapter */
47-
void* orig_ptr = nullptr; /**< original pointer to memory given by parameters */
48-
size_t size_in_bytes = 0; /**< size of bytes of the memory */
49-
bool allocated = false; /**< True if the memory allocated by this procedure and needs to be free */
50-
bool target_no_queue = false; /**< Indicates that original memory will be accessed from non SYCL environment */
51-
bool copy_back = false; /**< If the memory is 'result' it needs to be copied back to original */
46+
void* aux_ptr = nullptr; /**< pointer to allocated memory by this adapter */
47+
void* orig_ptr = nullptr; /**< original pointer to memory given by parameters */
48+
size_t size_in_bytes = 0; /**< size of bytes of the memory */
49+
bool allocated = false; /**< True if the memory allocated by this procedure and needs to be free */
50+
bool target_no_queue = false; /**< Indicates that original memory will be accessed from non SYCL environment */
51+
bool copy_back = false; /**< If the memory is 'result' it needs to be copied back to original */
5252
const bool verbose = false;
5353

5454
public:
@@ -67,8 +67,8 @@ class DPNPC_ptr_adapter final
6767
// enum class alloc { host = 0, device = 1, shared = 2, unknown = 3 };
6868
cl::sycl::usm::alloc src_ptr_type = cl::sycl::usm::alloc::unknown;
6969
src_ptr_type = cl::sycl::get_pointer_type(src_ptr, DPNP_QUEUE.get_context());
70-
if (verbose)
71-
{
70+
if (verbose)
71+
{
7272
std::cerr << "DPNPC_ptr_converter:";
7373
std::cerr << "\n\t target_no_queue=" << target_no_queue;
7474
std::cerr << "\n\t copy_back=" << copy_back;
@@ -83,7 +83,7 @@ class DPNPC_ptr_adapter final
8383
std::cerr << "\n\t queue device is_gpu=" << DPNP_QUEUE.get_device().is_gpu();
8484
std::cerr << "\n\t queue device is_accelerator=" << DPNP_QUEUE.get_device().is_accelerator();
8585
std::cerr << std::endl;
86-
}
86+
}
8787

8888
if (is_memcpy_required(src_ptr_type))
8989
{
@@ -93,9 +93,7 @@ class DPNPC_ptr_adapter final
9393
if (verbose)
9494
{
9595
std::cerr << "DPNPC_ptr_converter::alloc and copy memory"
96-
<< " from=" << src_ptr
97-
<< " to=" << aux_ptr
98-
<< " size_in_bytes=" << size_in_bytes
96+
<< " from=" << src_ptr << " to=" << aux_ptr << " size_in_bytes=" << size_in_bytes
9997
<< std::endl;
10098
}
10199
}
@@ -150,10 +148,7 @@ class DPNPC_ptr_adapter final
150148
if (verbose)
151149
{
152150
std::cerr << "DPNPC_ptr_converter::copy_data_back:"
153-
<< " from=" << aux_ptr
154-
<< " to=" << orig_ptr
155-
<< " size_in_bytes=" << size_in_bytes
156-
<< std::endl;
151+
<< " from=" << aux_ptr << " to=" << orig_ptr << " size_in_bytes=" << size_in_bytes << std::endl;
157152
}
158153

159154
dpnp_memory_memcpy_c(orig_ptr, aux_ptr, size_in_bytes);

0 commit comments

Comments
 (0)