Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
35 changes: 27 additions & 8 deletions dpnp/backend/include/dpnp_iface.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -253,8 +253,12 @@ INP_DLLEXPORT void dpnp_nanvar_c(void* array, void* mask_arr, void* result, cons
* @param [in] j Number input array.
*/
template <typename _DataType>
INP_DLLEXPORT void
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);
INP_DLLEXPORT void 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);

/**
* @ingroup BACKEND_API
Expand Down Expand Up @@ -634,8 +638,13 @@ INP_DLLEXPORT void dpnp_diag_indices_c(void* result1, size_t size);
* @param [in] ndim Number of elements in shape.
*/
template <typename _DataType>
INP_DLLEXPORT void dpnp_diagonal_c(
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);
INP_DLLEXPORT void dpnp_diagonal_c(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);

/**
* @ingroup BACKEND_API
Expand Down Expand Up @@ -695,8 +704,13 @@ INP_DLLEXPORT void dpnp_matrix_rank_c(void* array1_in, void* result1, size_t* sh
* @param [in] naxis Number of elements in axis.
*/
template <typename _DataType>
INP_DLLEXPORT 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);
INP_DLLEXPORT 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);

/**
* @ingroup BACKEND_API
Expand Down Expand Up @@ -741,8 +755,13 @@ INP_DLLEXPORT void
* @param [in] naxis Number of elements in axis.
*/
template <typename _DataType>
INP_DLLEXPORT void
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);
INP_DLLEXPORT void 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);

/**
* @ingroup BACKEND_API
Expand Down
17 changes: 13 additions & 4 deletions dpnp/backend/kernels/dpnp_krnl_indexing.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -45,8 +45,13 @@ template <typename _DataType>
class dpnp_diagonal_c_kernel;

template <typename _DataType>
void dpnp_diagonal_c(
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)
void dpnp_diagonal_c(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)
{
const size_t res_size = std::accumulate(res_shape, res_shape + res_ndim, 1, std::multiplies<size_t>());
if (!(res_size && input1_size))
Expand Down Expand Up @@ -196,7 +201,12 @@ void dpnp_fill_diagonal_c(void* array1_in, void* val_in, size_t* shape, const si
}

template <typename _DataType>
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)
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)
{
if ((in_array1 == nullptr) || (result1 == nullptr))
{
Expand All @@ -215,7 +225,6 @@ void dpnp_nonzero_c(const void* in_array1, void* result1, const size_t result_si
const _DataType* arr = input1_ptr.get_ptr();
long* result = result_ptr.get_ptr();


size_t idx = 0;
for (size_t i = 0; i < input1_size; ++i)
{
Expand Down
15 changes: 6 additions & 9 deletions dpnp/backend/kernels/dpnp_krnl_mathematical.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -374,8 +374,8 @@ void dpnp_remainder_c(void* result_out,
_DataType_input2* input2_data = input2_ptr.get_ptr();
_DataType_output* result = reinterpret_cast<_DataType_output*>(result_out);

std::vector<size_t> result_shape = get_result_shape(input1_shape, input1_shape_ndim,
input2_shape, input2_shape_ndim);
std::vector<size_t> result_shape =
get_result_shape(input1_shape, input1_shape_ndim, input2_shape, input2_shape_ndim);

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

input2_it->broadcast_to_shape(result_shape);

const size_t result_size = input1_it->get_output_size();

cl::sycl::range<1> gws(result_size);
cl::sycl::range<1> gws(result_size);
auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
const size_t i = global_id[0];
const _DataType_output input1_elem = (*input1_it)[i];
const _DataType_output input2_elem = (*input2_it)[i];
double fmod_res = cl::sycl::fmod((double)input1_elem, (double)input2_elem);
double add = fmod_res + input2_elem;
result[i] = cl::sycl::fmod(add, (double)input2_elem);

};
auto kernel_func = [&](cl::sycl::handler& cgh) {
cgh.parallel_for<class dpnp_remainder_c_kernel<_DataType_output, _DataType_input1, _DataType_input2>>(
Expand All @@ -412,9 +411,8 @@ void dpnp_remainder_c(void* result_out,

if (input1_size == input2_size)
{
if constexpr ((std::is_same<_DataType_input1, double>::value ||
std::is_same<_DataType_input1, float>::value) &&
std::is_same<_DataType_input2, _DataType_input1>::value)
if constexpr ((std::is_same<_DataType_input1, double>::value || std::is_same<_DataType_input1, float>::value) &&
std::is_same<_DataType_input2, _DataType_input1>::value)
{
event = oneapi::mkl::vm::fmod(DPNP_QUEUE, input1_size, input1_data, input2_data, result);
event.wait();
Expand All @@ -436,7 +434,6 @@ void dpnp_remainder_c(void* result_out,

input1_it->~DPNPC_id();
input2_it->~DPNPC_id();

}

template <typename _KernelNameSpecialization1, typename _KernelNameSpecialization2, typename _KernelNameSpecialization3>
Expand Down
4 changes: 2 additions & 2 deletions dpnp/backend/kernels/dpnp_krnl_random.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1253,7 +1253,7 @@ void dpnp_rng_vonmises_large_kappa_c(void* result, const _DataType mu, const _Da
{
return;
}

DPNPC_ptr_adapter<_DataType> result1_ptr(result, size, true, true);
_DataType* result1 = result1_ptr.get_ptr();

Expand Down Expand Up @@ -1547,7 +1547,7 @@ void dpnp_rng_zipf_c(void* result, const _DataType a, const size_t size)
long X;
const _DataType d_zero = 0.0;
const _DataType d_one = 1.0;

DPNPC_ptr_adapter<_DataType> result1_ptr(result, size, true, true);
_DataType* result1 = result1_ptr.get_ptr();

Expand Down
6 changes: 4 additions & 2 deletions dpnp/backend/kernels/dpnp_krnl_reduction.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -124,7 +124,8 @@ void dpnp_sum_c(void* result_out,
policy, input_it.begin(output_id), input_it.end(output_id), init, std::plus<_DataType_output>());
policy.queue().wait(); // TODO move out of the loop

dpnp_memory_memcpy_c(result + output_id, &accumulator, sizeof(_DataType_output)); // result[output_id] = accumulator;
dpnp_memory_memcpy_c(
result + output_id, &accumulator, sizeof(_DataType_output)); // result[output_id] = accumulator;
}

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

dpnp_memory_memcpy_c(result + output_id, &accumulator, sizeof(_DataType_output)); // result[output_id] = accumulator;
dpnp_memory_memcpy_c(
result + output_id, &accumulator, sizeof(_DataType_output)); // result[output_id] = accumulator;
}

return;
Expand Down
2 changes: 0 additions & 2 deletions dpnp/backend/kernels/dpnp_krnl_sorting.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -118,8 +118,6 @@ void dpnp_partition_c(
_DataType* arr2 = input2_ptr.get_ptr();
_DataType* result = result1_ptr.get_ptr();



auto arr_to_result_event = DPNP_QUEUE.memcpy(result, arr, size * sizeof(_DataType));
arr_to_result_event.wait();

Expand Down
20 changes: 16 additions & 4 deletions dpnp/backend/kernels/dpnp_krnl_statistics.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -52,8 +52,8 @@ void dpnp_correlate_c(void* result_out,
(void)where;

dpnp_dot_c<_DataType_output, _DataType_input1, _DataType_input2>(result_out,
42, // dummy result_size
42, // dummy result_ndim
42, // dummy result_size
42, // dummy result_ndim
NULL, // dummy result_shape
NULL, // dummy result_strides
input1_in,
Expand Down Expand Up @@ -153,7 +153,13 @@ template <typename _DataType>
class dpnp_max_c_kernel;

template <typename _DataType>
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)
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)
{
const size_t size_input = std::accumulate(shape, shape + ndim, 1, std::multiplies<size_t>());
if (!size_input)
Expand Down Expand Up @@ -416,7 +422,13 @@ template <typename _DataType>
class dpnp_min_c_kernel;

template <typename _DataType>
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)
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)
{
__attribute__((unused)) void* tmp = (void*)(axis + naxis);

Expand Down
14 changes: 7 additions & 7 deletions dpnp/backend/src/constants.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -30,14 +30,14 @@
void* python_constants::py_none = nullptr;
void* python_constants::py_nan = nullptr;

void dpnp_python_constants_initialize_c(void *_py_none, void *_py_nan)
void dpnp_python_constants_initialize_c(void* _py_none, void* _py_nan)
{
python_constants::py_none = _py_none;
python_constants::py_nan = _py_nan;
// std::cout << "========dpnp_python_constants_initialize_c=============" << std::endl;
// std::cout << "\t None=" << _py_none
// << "\n\t NaN=" << _py_nan
// << "\n\t py_none=" << python_constants::py_none
// << "\n\t py_nan=" << python_constants::py_nan
// << std::endl;
// std::cout << "========dpnp_python_constants_initialize_c=============" << std::endl;
// std::cout << "\t None=" << _py_none
// << "\n\t NaN=" << _py_nan
// << "\n\t py_none=" << python_constants::py_none
// << "\n\t py_nan=" << python_constants::py_nan
// << std::endl;
}
3 changes: 1 addition & 2 deletions dpnp/backend/src/constants.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -33,7 +33,6 @@
#define INP_DLLEXPORT
#endif


/**
* This is container for the constants from Python interpreter and other modules. These constants are subject to use
* in algorithms.
Expand All @@ -53,6 +52,6 @@ struct python_constants
* @param [in] py_none Python NONE representation
* @param [in] py_nan Python NAN representation
*/
INP_DLLEXPORT void dpnp_python_constants_initialize_c(void *py_none, void *py_nan);
INP_DLLEXPORT void dpnp_python_constants_initialize_c(void* py_none, void* py_nan);

#endif // CONSTANTS_H
27 changes: 11 additions & 16 deletions dpnp/backend/src/dpnpc_memory_adapter.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -43,12 +43,12 @@
template <typename _DataType>
class DPNPC_ptr_adapter final
{
void* aux_ptr = nullptr; /**< pointer to allocated memory by this adapter */
void* orig_ptr = nullptr; /**< original pointer to memory given by parameters */
size_t size_in_bytes = 0; /**< size of bytes of the memory */
bool allocated = false; /**< True if the memory allocated by this procedure and needs to be free */
bool target_no_queue = false; /**< Indicates that original memory will be accessed from non SYCL environment */
bool copy_back = false; /**< If the memory is 'result' it needs to be copied back to original */
void* aux_ptr = nullptr; /**< pointer to allocated memory by this adapter */
void* orig_ptr = nullptr; /**< original pointer to memory given by parameters */
size_t size_in_bytes = 0; /**< size of bytes of the memory */
bool allocated = false; /**< True if the memory allocated by this procedure and needs to be free */
bool target_no_queue = false; /**< Indicates that original memory will be accessed from non SYCL environment */
bool copy_back = false; /**< If the memory is 'result' it needs to be copied back to original */
const bool verbose = false;

public:
Expand All @@ -67,8 +67,8 @@ class DPNPC_ptr_adapter final
// enum class alloc { host = 0, device = 1, shared = 2, unknown = 3 };
cl::sycl::usm::alloc src_ptr_type = cl::sycl::usm::alloc::unknown;
src_ptr_type = cl::sycl::get_pointer_type(src_ptr, DPNP_QUEUE.get_context());
if (verbose)
{
if (verbose)
{
std::cerr << "DPNPC_ptr_converter:";
std::cerr << "\n\t target_no_queue=" << target_no_queue;
std::cerr << "\n\t copy_back=" << copy_back;
Expand All @@ -83,7 +83,7 @@ class DPNPC_ptr_adapter final
std::cerr << "\n\t queue device is_gpu=" << DPNP_QUEUE.get_device().is_gpu();
std::cerr << "\n\t queue device is_accelerator=" << DPNP_QUEUE.get_device().is_accelerator();
std::cerr << std::endl;
}
}

if (is_memcpy_required(src_ptr_type))
{
Expand All @@ -93,9 +93,7 @@ class DPNPC_ptr_adapter final
if (verbose)
{
std::cerr << "DPNPC_ptr_converter::alloc and copy memory"
<< " from=" << src_ptr
<< " to=" << aux_ptr
<< " size_in_bytes=" << size_in_bytes
<< " from=" << src_ptr << " to=" << aux_ptr << " size_in_bytes=" << size_in_bytes
<< std::endl;
}
}
Expand Down Expand Up @@ -150,10 +148,7 @@ class DPNPC_ptr_adapter final
if (verbose)
{
std::cerr << "DPNPC_ptr_converter::copy_data_back:"
<< " from=" << aux_ptr
<< " to=" << orig_ptr
<< " size_in_bytes=" << size_in_bytes
<< std::endl;
<< " from=" << aux_ptr << " to=" << orig_ptr << " size_in_bytes=" << size_in_bytes << std::endl;
}

dpnp_memory_memcpy_c(orig_ptr, aux_ptr, size_in_bytes);
Expand Down
9 changes: 4 additions & 5 deletions dpnp/backend/src/memory_sycl.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,6 @@
#include <dpnp_iface.hpp>
#include "queue_sycl.hpp"


static bool use_sycl_device_memory()
{
// TODO need to move all getenv() into common dpnpc place
Expand Down Expand Up @@ -70,10 +69,10 @@ char* dpnp_memory_alloc_c(size_t size_in_bytes)
#if not defined(NDEBUG)
if (memory_type != cl::sycl::usm::alloc::device)
{
for (size_t i = 0; i < size_in_bytes / sizeof(char); ++i)
{
array[i] = 0; // type dependant is better. set double(42.42) instead zero
}
for (size_t i = 0; i < size_in_bytes / sizeof(char); ++i)
{
array[i] = 0; // type dependant is better. set double(42.42) instead zero
}
}
// std::cout << ") -> ptr=" << (void*)array << std::endl;
#endif
Expand Down