diff --git a/dpnp/backend/CMakeLists.txt b/dpnp/backend/CMakeLists.txt index d96320bf0acd..7ed57fd929d8 100644 --- a/dpnp/backend/CMakeLists.txt +++ b/dpnp/backend/CMakeLists.txt @@ -30,7 +30,6 @@ set(DPNP_SRC kernels/dpnp_krnl_fft.cpp kernels/dpnp_krnl_indexing.cpp kernels/dpnp_krnl_logic.cpp - kernels/dpnp_krnl_manipulation.cpp kernels/dpnp_krnl_mathematical.cpp kernels/dpnp_krnl_random.cpp kernels/dpnp_krnl_reduction.cpp diff --git a/dpnp/backend/examples/example11.cpp b/dpnp/backend/examples/example11.cpp deleted file mode 100644 index 3a16991bae66..000000000000 --- a/dpnp/backend/examples/example11.cpp +++ /dev/null @@ -1,85 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -/** - * Example 11. - * - * This example shows simple usage of the DPNP C++ Backend library RNG shuffle - * function for one and ndim arrays. - * - * Possible compile line: - * g++ -g dpnp/backend/examples/example11.cpp -Idpnp -Idpnp/backend/include - * -Ldpnp -Wl,-rpath='$ORIGIN'/dpnp -ldpnp_backend_c -o example11 - * - */ - -#include - -#include - -template -void print_dpnp_array(T *arr, size_t size) -{ - std::cout << std::endl; - for (size_t i = 0; i < size; ++i) { - std::cout << arr[i] << ", "; - } - std::cout << std::endl; -} - -int main(int, char **) -{ - // Two cases: - // 1) array size = 100, ndim = 1, high_dim_size = 10 (aka ndarray with shape - // (100,) ) 2) array size = 100, ndim = 2, high_dim_size = 20 (e.g. ndarray - // with shape (20, 5) and len(array) = 20 ) - const size_t ndim_cases = 2; - const size_t itemsize = sizeof(double); - const size_t ndim[ndim_cases] = {1, 2}; - const size_t high_dim_size[ndim_cases] = {100, 20}; - const size_t size = 100; - const size_t seed = 1234; - - // DPNPC dpnp_rng_shuffle_c - // DPNPC interface - double *array_1 = - reinterpret_cast(dpnp_memory_alloc_c(size * sizeof(double))); - for (size_t i = 0; i < ndim_cases; i++) { - std::cout << "\nREPRODUCE: DPNPC dpnp_rng_shuffle_c:"; - std::cout << "\nDIMS: " << ndim[i] << std::endl; - // init array 0, 1, 2, 3, 4, 5, 6, .... - dpnp_arange_c(0, 1, array_1, size); - // print before shuffle - std::cout << "\nINPUT array:"; - print_dpnp_array(array_1, size); - dpnp_rng_srand_c(seed); - dpnp_rng_shuffle_c(array_1, itemsize, ndim[i], high_dim_size[i], - size); - // print shuffle result - std::cout << "\nSHUFFLE INPUT array:"; - print_dpnp_array(array_1, size); - } - dpnp_memory_free_c(array_1); -} diff --git a/dpnp/backend/examples/example3.cpp b/dpnp/backend/examples/example3.cpp deleted file mode 100644 index 2d516dc0b8de..000000000000 --- a/dpnp/backend/examples/example3.cpp +++ /dev/null @@ -1,79 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -/** - * Example 3. - * - * This example shows simple usage of the DPNP C++ Backend library - * to calculate cos of input vector elements - * - * Possible compile line: - * . /opt/intel/oneapi/setvars.sh - * g++ -g dpnp/backend/examples/example3.cpp -Idpnp -Idpnp/backend/include - * -Ldpnp -Wl,-rpath='$ORIGIN'/dpnp -ldpnp_backend_c -o example3 - * - */ - -#include - -#include "dpnp_iface.hpp" - -int main(int, char **) -{ - const size_t size = 256; - - std::cout << "SYCL queue is CPU: " << dpnp_queue_is_cpu_c() << std::endl; - - int *array1 = (int *)dpnp_memory_alloc_c(size * sizeof(int)); - double *result = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - - for (size_t i = 0; i < 10; ++i) { - array1[i] = i; - result[i] = 0; - std::cout << ", " << array1[i]; - } - std::cout << std::endl; - - const long ndim = 1; - shape_elem_type *shape = reinterpret_cast( - dpnp_memory_alloc_c(ndim * sizeof(shape_elem_type))); - shape[0] = size; - shape_elem_type *strides = reinterpret_cast( - dpnp_memory_alloc_c(ndim * sizeof(shape_elem_type))); - strides[0] = 1; - - dpnp_cos_c(result, size, ndim, shape, strides, array1, size, - ndim, shape, strides, NULL); - - for (size_t i = 0; i < 10; ++i) { - std::cout << ", " << result[i]; - } - std::cout << std::endl; - - dpnp_memory_free_c(result); - dpnp_memory_free_c(array1); - - return 0; -} diff --git a/dpnp/backend/examples/example7.cpp b/dpnp/backend/examples/example7.cpp deleted file mode 100644 index 49c12c5dd51c..000000000000 --- a/dpnp/backend/examples/example7.cpp +++ /dev/null @@ -1,77 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -/** - * Example 7. - * - * This example shows simple usage of the DPNP C++ Backend library - * to calculate eigenvalues and eigenvectors of a symmetric matrix - * - * Possible compile line: - * . /opt/intel/oneapi/setvars.sh - * g++ -g dpnp/backend/examples/example7.cpp -Idpnp -Idpnp/backend/include - * -Ldpnp -Wl,-rpath='$ORIGIN'/dpnp -ldpnp_backend_c -o example7 - * - */ - -#include - -#include "dpnp_iface.hpp" - -int main(int, char **) -{ - const size_t size = 2; - size_t len = size * size; - - float *array = (float *)dpnp_memory_alloc_c(len * sizeof(float)); - float *result1 = (float *)dpnp_memory_alloc_c(size * sizeof(float)); - float *result2 = (float *)dpnp_memory_alloc_c(len * sizeof(float)); - - /* init input diagonal array like: - 1, 0, 0, - 0, 2, 0, - 0, 0, 3 - */ - for (size_t i = 0; i < len; ++i) { - array[i] = 0; - } - for (size_t i = 0; i < size; ++i) { - array[size * i + i] = i + 1; - } - - dpnp_eig_c(array, result1, result2, size); - - std::cout << "eigen values" << std::endl; - for (size_t i = 0; i < size; ++i) { - std::cout << result1[i] << ", "; - } - std::cout << std::endl; - - dpnp_memory_free_c(result2); - dpnp_memory_free_c(result1); - dpnp_memory_free_c(array); - - return 0; -} diff --git a/dpnp/backend/examples/example_bs.cpp b/dpnp/backend/examples/example_bs.cpp deleted file mode 100644 index c20c6d27e297..000000000000 --- a/dpnp/backend/examples/example_bs.cpp +++ /dev/null @@ -1,282 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -/** - * Example BS. - * - * This example shows simple usage of the DPNP C++ Backend library - * to calculate black scholes algorithm like in Python version - * - * Possible compile line: - * . /opt/intel/oneapi/setvars.sh - * g++ -g dpnp/backend/examples/example_bs.cpp -Idpnp -Idpnp/backend/include - * -Ldpnp -Wl,-rpath='$ORIGIN'/dpnp -ldpnp_backend_c -o example_bs - */ - -#include -#include - -#include "dpnp_iface.hpp" - -void black_scholes(double *price, - double *strike, - double *t, - const double rate, - const double vol, - double *call, - double *put, - const size_t size) -{ - const size_t ndim = 1; - const size_t scalar_size = 1; - - double *mr = (double *)dpnp_memory_alloc_c(1 * sizeof(double)); - mr[0] = -rate; - - double *vol_vol_two = (double *)dpnp_memory_alloc_c(1 * sizeof(double)); - vol_vol_two[0] = vol * vol * 2; - - double *quarter = (double *)dpnp_memory_alloc_c(1 * sizeof(double)); - quarter[0] = 0.25; - - double *one = (double *)dpnp_memory_alloc_c(1 * sizeof(double)); - one[0] = 1.; - - double *half = (double *)dpnp_memory_alloc_c(1 * sizeof(double)); - half[0] = 0.5; - - double *P = price; - double *S = strike; - double *T = t; - - double *p_div_s = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // p_div_s = P / S - dpnp_divide_c(p_div_s, P, size, &size, ndim, S, - size, &size, ndim, NULL); - double *a = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - dpnp_log_c(p_div_s, a, size); // a = np.log(p_div_s) - dpnp_memory_free_c(p_div_s); - - double *b = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // b = T * mr - dpnp_multiply_c( - b, T, size, &size, ndim, mr, scalar_size, &scalar_size, ndim, NULL); - dpnp_memory_free_c(mr); - double *z = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // z = T * vol_vol_twos - dpnp_multiply_c(z, T, size, &size, ndim, - vol_vol_two, scalar_size, - &scalar_size, ndim, NULL); - dpnp_memory_free_c(vol_vol_two); - - double *c = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // c = quarters * z - dpnp_multiply_c(c, quarter, scalar_size, - &scalar_size, ndim, z, size, &size, - ndim, NULL); - dpnp_memory_free_c(quarter); - - double *sqrt_z = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - dpnp_sqrt_c(z, sqrt_z, size); // sqrt_z = np.sqrt(z) - dpnp_memory_free_c(z); - double *y = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // y = ones / np.sqrt(z) - dpnp_divide_c(y, one, scalar_size, &scalar_size, - ndim, sqrt_z, size, &size, ndim, - NULL); - dpnp_memory_free_c(sqrt_z); - dpnp_memory_free_c(one); - - double *a_sub_b = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // a_sub_b = a - b - dpnp_subtract_c(a_sub_b, a, size, &size, ndim, b, - size, &size, ndim, NULL); - dpnp_memory_free_c(a); - double *a_sub_b_add_c = - (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // a_sub_b_add_c = a_sub_b + c - dpnp_add_c(a_sub_b_add_c, a_sub_b, size, &size, - ndim, c, size, &size, ndim, NULL); - double *w1 = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // w1 = a_sub_b_add_c * y - dpnp_multiply_c(w1, a_sub_b_add_c, size, &size, - ndim, y, size, &size, ndim, NULL); - dpnp_memory_free_c(a_sub_b_add_c); - - double *a_sub_b_sub_c = - (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // a_sub_b_sub_c = a_sub_b - c - dpnp_subtract_c(a_sub_b_sub_c, a_sub_b, size, &size, - ndim, c, size, &size, ndim, NULL); - dpnp_memory_free_c(a_sub_b); - dpnp_memory_free_c(c); - double *w2 = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // w2 = a_sub_b_sub_c * y - dpnp_multiply_c(w2, a_sub_b_sub_c, size, &size, - ndim, y, size, &size, ndim, NULL); - dpnp_memory_free_c(a_sub_b_sub_c); - dpnp_memory_free_c(y); - - double *erf_w1 = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - dpnp_erf_c(w1, erf_w1, size); // erf_w1 = np.erf(w1) - dpnp_memory_free_c(w1); - double *halfs_mul_erf_w1 = - (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // halfs_mul_erf_w1 = half * erf_w1 - dpnp_multiply_c(halfs_mul_erf_w1, half, scalar_size, - &scalar_size, ndim, erf_w1, size, - &size, ndim, NULL); - dpnp_memory_free_c(erf_w1); - double *d1 = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // d1 = half + halfs_mul_erf_w1 - dpnp_add_c(d1, half, scalar_size, &scalar_size, - ndim, halfs_mul_erf_w1, size, &size, - ndim, NULL); - dpnp_memory_free_c(halfs_mul_erf_w1); - - double *erf_w2 = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - dpnp_erf_c(w2, erf_w2, size); // erf_w2 = np.erf(w2) - dpnp_memory_free_c(w2); - double *halfs_mul_erf_w2 = - (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // halfs_mul_erf_w2 = half * erf_w2 - dpnp_multiply_c(halfs_mul_erf_w2, half, scalar_size, - &scalar_size, ndim, erf_w2, size, - &size, ndim, NULL); - dpnp_memory_free_c(erf_w2); - double *d2 = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // d2 = half + halfs_mul_erf_w2 - dpnp_add_c(d2, half, scalar_size, &scalar_size, - ndim, halfs_mul_erf_w2, size, &size, - ndim, NULL); - dpnp_memory_free_c(halfs_mul_erf_w2); - dpnp_memory_free_c(half); - - double *exp_b = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - dpnp_exp_c(b, exp_b, size); // exp_b = np.exp(b) - double *Se = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // Se = exp_b * S - dpnp_multiply_c(Se, exp_b, size, &size, ndim, S, - size, &size, ndim, NULL); - dpnp_memory_free_c(exp_b); - dpnp_memory_free_c(b); - - double *P_mul_d1 = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // P_mul_d1 = P * d1 - dpnp_multiply_c(P_mul_d1, P, size, &size, ndim, d1, - size, &size, ndim, NULL); - dpnp_memory_free_c(d1); - double *Se_mul_d2 = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // Se_mul_d2 = Se * d2 - dpnp_multiply_c(Se_mul_d2, Se, size, &size, ndim, - d2, size, &size, ndim, NULL); - dpnp_memory_free_c(d2); - double *r = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // r = P_mul_d1 - Se_mul_d2 - dpnp_subtract_c(r, P_mul_d1, size, &size, ndim, - Se_mul_d2, size, &size, ndim, NULL); - dpnp_memory_free_c(Se_mul_d2); - dpnp_memory_free_c(P_mul_d1); - - dpnp_copyto_c(call, r, size); // call[:] = r - double *r_sub_P = (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // r_sub_P = r - P - dpnp_subtract_c(r_sub_P, r, size, &size, ndim, P, - size, &size, ndim, NULL); - dpnp_memory_free_c(r); - double *r_sub_P_add_Se = - (double *)dpnp_memory_alloc_c(size * sizeof(double)); - // r_sub_P_add_Se = r_sub_P + Se - dpnp_add_c(r_sub_P_add_Se, r_sub_P, size, &size, - ndim, Se, size, &size, ndim, NULL); - dpnp_memory_free_c(r_sub_P); - dpnp_memory_free_c(Se); - dpnp_copyto_c(put, r_sub_P_add_Se, - size); // put[:] = r_sub_P_add_Se - dpnp_memory_free_c(r_sub_P_add_Se); -} - -int main(int, char **) -{ - const size_t SIZE = 256; - - const size_t SEED = 7777777; - const long PL = 10, PH = 50; - const long SL = 10, SH = 50; - const long TL = 1, TH = 2; - const double RISK_FREE = 0.1; - const double VOLATILITY = 0.2; - - std::cout << "SYCL queue is CPU: " << dpnp_queue_is_cpu_c() << std::endl; - - double *price = (double *)dpnp_memory_alloc_c(SIZE * sizeof(double)); - double *strike = (double *)dpnp_memory_alloc_c(SIZE * sizeof(double)); - double *t = (double *)dpnp_memory_alloc_c(SIZE * sizeof(double)); - - dpnp_rng_srand_c(SEED); // np.random.seed(SEED) - dpnp_rng_uniform_c(price, PL, PH, - SIZE); // np.random.uniform(PL, PH, SIZE) - dpnp_rng_uniform_c(strike, SL, SH, - SIZE); // np.random.uniform(SL, SH, SIZE) - dpnp_rng_uniform_c(t, TL, TH, - SIZE); // np.random.uniform(TL, TH, SIZE) - - double *zero = (double *)dpnp_memory_alloc_c(1 * sizeof(double)); - zero[0] = 0.; - - double *mone = (double *)dpnp_memory_alloc_c(1 * sizeof(double)); - mone[0] = -1.; - - double *call = (double *)dpnp_memory_alloc_c(SIZE * sizeof(double)); - double *put = (double *)dpnp_memory_alloc_c(SIZE * sizeof(double)); - - dpnp_full_c(zero, call, SIZE); // np.full(SIZE, 0., dtype=DTYPE) - dpnp_full_c(mone, put, SIZE); // np.full(SIZE, -1., dtype=DTYPE) - - dpnp_memory_free_c(mone); - dpnp_memory_free_c(zero); - - black_scholes(price, strike, t, RISK_FREE, VOLATILITY, call, put, SIZE); - - std::cout << "call: "; - for (size_t i = 0; i < 10; ++i) { - std::cout << call[i] << ", "; - } - std::cout << "..." << std::endl; - std::cout << "put: "; - for (size_t i = 0; i < 10; ++i) { - std::cout << put[i] << ", "; - } - std::cout << "..." << std::endl; - - dpnp_memory_free_c(put); - dpnp_memory_free_c(call); - - dpnp_memory_free_c(t); - dpnp_memory_free_c(strike); - dpnp_memory_free_c(price); - - return 0; -} diff --git a/dpnp/backend/examples/example_experimental_iface.cpp b/dpnp/backend/examples/example_experimental_iface.cpp deleted file mode 100644 index 4454a34b9a45..000000000000 --- a/dpnp/backend/examples/example_experimental_iface.cpp +++ /dev/null @@ -1,63 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -/** - * Example of experimental interface. - * - * This example shows how to get a runtime pointer from DPNP C++ Backend library - * - * Possible compile line: - * . /opt/intel/oneapi/setvars.sh - * g++ -g dpnp/backend/examples/example_experimental_iface.cpp -Idpnp - * -Idpnp/backend/include -Ldpnp -Wl,-rpath='$ORIGIN'/dpnp -ldpnp_backend_c -o - * example_experimental_iface - */ - -#include - -#include -// TODO #include - -int main(int, char **) -{ - void *result = get_backend_function_name("dpnp_dot", "float"); - std::cout << "Result Dot() function pointer (by old interface): " << result - << std::endl; - - DPNPFuncData_t dpnp_dot_f = get_dpnp_function_ptr( - DPNPFuncName::DPNP_FN_DOT, DPNPFuncType::DPNP_FT_LONG); - std::cout << "Result Dot() function pointer: " << dpnp_dot_f.ptr - << " with return datatype " << (size_t)dpnp_dot_f.return_type - << std::endl; - - DPNPFuncData_t dpnp_add_f = get_dpnp_function_ptr( - DPNPFuncName::DPNP_FN_ADD, DPNPFuncType::DPNP_FT_FLOAT, - DPNPFuncType::DPNP_FT_INT); - std::cout << "Result Add() function pointer: " << dpnp_add_f.ptr - << " with return datatype " << (size_t)dpnp_add_f.return_type - << std::endl; - - return 0; -} diff --git a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp index dcec3f8192bb..e5a2c924653a 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -140,14 +140,4 @@ MACRO_2ARG_3TYPES_OP(dpnp_multiply_c, std::complex, std::complex)) -MACRO_2ARG_3TYPES_OP(dpnp_subtract_c, - input1_elem - input2_elem, - x1 - x2, - MACRO_UNPACK_TYPES(bool, std::int32_t, std::int64_t), - oneapi::mkl::vm::sub, - MACRO_UNPACK_TYPES(float, - double, - std::complex, - std::complex)) - #undef MACRO_2ARG_3TYPES_OP diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index 324e7a612b1a..0fc5595041c6 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -176,74 +176,6 @@ template INP_DLLEXPORT void dpnp_any_c(const void *array, void *result, const size_t size); -/** - * @ingroup BACKEND_API - * @brief Array initialization - * - * Input array, step based, initialization procedure. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] start Start of initialization sequence - * @param [in] step Step for initialization sequence - * @param [out] result1 Output array. - * @param [in] size Number of elements in input arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_arange_c(DPCTLSyclQueueRef q_ref, - size_t start, - size_t step, - void *result1, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_arange_c(size_t start, size_t step, void *result1, size_t size); - -/** - * @ingroup BACKEND_API - * @brief Implementation of full function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array_in Input one-element array. - * @param [out] result Output array. - * @param [in] size Number of elements in the output array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_full_c(DPCTLSyclQueueRef q_ref, - void *array_in, - void *result, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_full_c(void *array_in, void *result, const size_t size); - -/** - * @ingroup BACKEND_API - * @brief Implementation of full_like function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array_in Input one-element array. - * @param [out] result Output array. - * @param [in] size Number of elements in the output array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_full_like_c(DPCTLSyclQueueRef q_ref, - void *array_in, - void *result, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_full_like_c(void *array_in, void *result, size_t size); - /** * @ingroup BACKEND_API * @brief Compute the variance along the specified axis, while ignoring NaNs. @@ -591,56 +523,6 @@ INP_DLLEXPORT void dpnp_prod_c(void *result_out, const void *initial, const long *where); -/** - * @ingroup BACKEND_API - * @brief Range of values (maximum - minimum) along an axis. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result_out Output array. - * @param [in] result_size Size of output array. - * @param [in] result_ndim Number of output array dimensions. - * @param [in] result_shape Shape of output array. - * @param [in] result_strides Strides of output array. - * @param [in] input_in First input array. - * @param [in] input_size Size of first input array. - * @param [in] input_ndim Number of first input array dimensions. - * @param [in] input_shape Shape of first input array. - * @param [in] input_strides Strides of first input array. - * @param [in] axis Axis. - * @param [in] naxis Number of elements in axis. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_ptp_c(DPCTLSyclQueueRef q_ref, - void *result_out, - const size_t result_size, - const size_t result_ndim, - const shape_elem_type *result_shape, - const shape_elem_type *result_strides, - const void *input_in, - const size_t input_size, - const size_t input_ndim, - const shape_elem_type *input_shape, - const shape_elem_type *input_strides, - const shape_elem_type *axis, - const size_t naxis, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_ptp_c(void *result_out, - const size_t result_size, - const size_t result_ndim, - const shape_elem_type *result_shape, - const shape_elem_type *result_strides, - const void *input_in, - const size_t input_size, - const size_t input_ndim, - const shape_elem_type *input_shape, - const shape_elem_type *input_strides, - const shape_elem_type *axis, - const size_t naxis); - /** * @ingroup BACKEND_API * @brief Replaces specified elements of an array with given values. @@ -715,29 +597,7 @@ INP_DLLEXPORT void dpnp_put_along_axis_c(void *arr_in, /** * @ingroup BACKEND_API - * @brief Return a 2-D array with ones on the diagonal and zeros elsewhere. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result The eigenvalues, each repeated according to - * its multiplicity - * @param [in] k Index of the diagonal - * @param [in] shape Shape of result - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_eye_c(DPCTLSyclQueueRef q_ref, - void *result, - int k, - const shape_elem_type *res_shape, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_eye_c(void *result, int k, const shape_elem_type *res_shape); -/** - * @ingroup BACKEND_API * @brief math library implementation of argsort function * * @param [in] q_ref Reference to SYCL queue. @@ -916,60 +776,6 @@ INP_DLLEXPORT void dpnp_choose_c(void *result1, size_t choices_size, size_t choice_size); -/** - * @ingroup BACKEND_API - * @brief Extract a diagonal or construct a diagonal array. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result Output array. - * @param [in] k Diagonal in question. - * @param [in] shape Shape of input array. - * @param [in] res_shape Shape of result array. - * @param [in] ndim Number of elements in shape of input array. - * @param [in] res_ndim Number of elements in shape of result array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_diag_c(DPCTLSyclQueueRef q_ref, - void *array, - void *result, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_diag_c(void *array, - void *result, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim); - -/** - * @ingroup BACKEND_API - * @brief Return the indices to access the main diagonal of an array. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result1 Output array. - * @param [in] size Size of array. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_diag_indices_c(DPCTLSyclQueueRef q_ref, - void *result1, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_diag_indices_c(void *result1, size_t size); - /** * @ingroup BACKEND_API * @brief math library implementation of diagonal function @@ -1006,26 +812,6 @@ INP_DLLEXPORT void dpnp_diagonal_c(void *array1_in, shape_elem_type *res_shape, const size_t res_ndim); -/** - * @ingroup BACKEND_API - * @brief Implementation of identity function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result1 Output array. - * @param [in] n Number of rows (and columns) in n x n - * output. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_identity_c(DPCTLSyclQueueRef q_ref, - void *result1, - const size_t n, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_identity_c(void *result1, const size_t n); - /** * @ingroup BACKEND_API * @brief implementation of creating filled with value array function @@ -1287,128 +1073,6 @@ INP_DLLEXPORT void dpnp_take_c(void *array, void *result, size_t size); -/** - * @ingroup BACKEND_API - * @brief math library implementation of trace function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result Output array. - * @param [in] shape Shape of input array. - * @param [in] ndim Number of elements in array.shape. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_trace_c(DPCTLSyclQueueRef q_ref, - const void *array, - void *result, - const shape_elem_type *shape, - const size_t ndim, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_trace_c(const void *array, - void *result, - const shape_elem_type *shape, - const size_t ndim); - -/** - * @ingroup BACKEND_API - * @brief An array with ones at and below the given diagonal and zeros - * elsewhere. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [out] result Output array. - * @param [in] N Number of rows in the array. - * @param [in] M Number of columns in the array. - * @param [in] k The sub-diagonal at and below which the - * array is filled. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_tri_c(DPCTLSyclQueueRef q_ref, - void *result, - const size_t N, - const size_t M, - const int k, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_tri_c(void *result, const size_t N, const size_t M, const int k); - -/** - * @ingroup BACKEND_API - * @brief Lower triangle of an array. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result Output array. - * @param [in] k Diagonal above which to zero elements. - * @param [in] shape Shape of input array. - * @param [in] res_shape Shape of result array. - * @param [in] ndim Number of elements in array.shape. - * @param [in] res_ndim Number of elements in res_shape. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_tril_c(DPCTLSyclQueueRef q_ref, - void *array, - void *result, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_tril_c(void *array, - void *result, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim); - -/** - * @ingroup BACKEND_API - * @brief Upper triangle of an array. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array with data. - * @param [out] result Output array. - * @param [in] k Diagonal above which to zero elements. - * @param [in] shape Shape of input array. - * @param [in] res_shape Shape of result array. - * @param [in] ndim Number of elements in array.shape. - * @param [in] res_ndim Number of elements in res_shape. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_triu_c(DPCTLSyclQueueRef q_ref, - void *array, - void *result, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_triu_c(void *array, - void *result, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim); - /** * @ingroup BACKEND_API * @brief math library implementation of var function @@ -1609,62 +1273,6 @@ INP_DLLEXPORT DPCTLSyclEventRef template INP_DLLEXPORT void dpnp_ones_like_c(void *result, size_t size); -/** - * @ingroup BACKEND_API - * @brief repeat elements of an array. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array_in Input array. - * @param [out] result Output array. - * @param [in] repeats The number of repetitions for each element. - * @param [in] size Number of elements in input arrays. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_repeat_c(DPCTLSyclQueueRef q_ref, - const void *array_in, - void *result, - const size_t repeats, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_repeat_c(const void *array_in, - void *result, - const size_t repeats, - const size_t size); - -/** - * @ingroup BACKEND_API - * @brief Implementation of vander function - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array_in Input array. - * @param [out] result Output array. - * @param [in] size_in Number of elements in the input array. - * @param [in] N Number of columns in the output. - * @param [in] increasing Order of the powers of the columns. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - * - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_vander_c(DPCTLSyclQueueRef q_ref, - const void *array1_in, - void *result1, - const size_t size_in, - const size_t N, - const int increasing, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_vander_c(const void *array1_in, - void *result1, - const size_t size_in, - const size_t N, - const int increasing); - /** * @ingroup BACKEND_API * @brief Implementation of zeros function diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index a39174931fec..d62e5998583e 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -64,7 +64,6 @@ enum class DPNPFuncName : size_t DPNP_FN_ALLCLOSE_EXT, /**< Used in numpy.allclose() impl, requires extra parameters */ DPNP_FN_ANY, /**< Used in numpy.any() impl */ - DPNP_FN_ARANGE, /**< Used in numpy.arange() impl */ DPNP_FN_ARGMAX, /**< Used in numpy.argmax() impl */ DPNP_FN_ARGMIN, /**< Used in numpy.argmin() impl */ DPNP_FN_ARGSORT, /**< Used in numpy.argsort() impl */ @@ -82,8 +81,6 @@ enum class DPNPFuncName : size_t DPNP_FN_DEGREES, /**< Used in numpy.degrees() impl */ DPNP_FN_DEGREES_EXT, /**< Used in numpy.degrees() impl, requires extra parameters */ - DPNP_FN_DIAG, /**< Used in numpy.diag() impl */ - DPNP_FN_DIAG_INDICES, /**< Used in numpy.diag_indices() impl */ DPNP_FN_DIAGONAL, /**< Used in numpy.diagonal() impl */ DPNP_FN_DOT, /**< Used in numpy.dot() impl */ DPNP_FN_DOT_EXT, /**< Used in numpy.dot() impl, requires extra parameters */ @@ -93,7 +90,6 @@ enum class DPNPFuncName : size_t DPNP_FN_ERF, /**< Used in scipy.special.erf impl */ DPNP_FN_ERF_EXT, /**< Used in scipy.special.erf impl, requires extra parameters */ - DPNP_FN_EYE, /**< Used in numpy.eye() impl */ DPNP_FN_FFT_FFT, /**< Used in numpy.fft.fft() impl */ DPNP_FN_FFT_FFT_EXT, /**< Used in numpy.fft.fft() impl, requires extra parameters */ @@ -101,14 +97,10 @@ enum class DPNPFuncName : size_t DPNP_FN_FFT_RFFT_EXT, /**< Used in numpy.fft.rfft() impl, requires extra parameters */ DPNP_FN_FILL_DIAGONAL, /**< Used in numpy.fill_diagonal() impl */ - DPNP_FN_FULL, /**< Used in numpy.full() impl */ - DPNP_FN_FULL_LIKE, /**< Used in numpy.full_like() impl */ - DPNP_FN_IDENTITY, /**< Used in numpy.identity() impl */ DPNP_FN_INITVAL, /**< Used in numpy ones, ones_like, zeros, zeros_like impls */ DPNP_FN_INITVAL_EXT, /**< Used in numpy ones, ones_like, zeros, zeros_like impls */ - DPNP_FN_INVERT, /**< Used in numpy.invert() impl */ DPNP_FN_MAX, /**< Used in numpy.max() impl */ DPNP_FN_MAXIMUM_EXT, /**< Used in numpy.fmax() impl , requires extra parameters */ @@ -132,13 +124,11 @@ enum class DPNPFuncName : size_t parameters */ DPNP_FN_PLACE, /**< Used in numpy.place() impl */ DPNP_FN_PROD, /**< Used in numpy.prod() impl */ - DPNP_FN_PTP, /**< Used in numpy.ptp() impl */ DPNP_FN_PUT, /**< Used in numpy.put() impl */ DPNP_FN_PUT_ALONG_AXIS, /**< Used in numpy.put_along_axis() impl */ DPNP_FN_RADIANS, /**< Used in numpy.radians() impl */ DPNP_FN_RADIANS_EXT, /**< Used in numpy.radians() impl, requires extra parameters */ - DPNP_FN_REPEAT, /**< Used in numpy.repeat() impl */ DPNP_FN_RNG_BETA, /**< Used in numpy.random.beta() impl */ DPNP_FN_RNG_BETA_EXT, /**< Used in numpy.random.beta() impl, requires extra parameters */ @@ -262,22 +252,12 @@ enum class DPNPFuncName : size_t DPNP_FN_SQRT_EXT, /**< Used in numpy.sqrt() impl, requires extra parameters */ DPNP_FN_STD, /**< Used in numpy.std() impl */ - DPNP_FN_SUBTRACT_EXT, /**< Used in numpy.subtract() impl, requires extra - parameters */ - DPNP_FN_SUM, /**< Used in numpy.sum() impl */ - DPNP_FN_TAKE, /**< Used in numpy.take() impl */ - DPNP_FN_TRANSPOSE, /**< Used in numpy.transpose() impl */ - DPNP_FN_TRACE, /**< Used in numpy.trace() impl */ - DPNP_FN_TRAPZ_EXT, /**< Used in numpy.trapz() impl, requires extra - parameters */ - DPNP_FN_TRI, /**< Used in numpy.tri() impl */ - DPNP_FN_TRIL, /**< Used in numpy.tril() impl */ - DPNP_FN_TRIU, /**< Used in numpy.triu() impl */ - DPNP_FN_VANDER, /**< Used in numpy.vander() impl */ - DPNP_FN_VAR, /**< Used in numpy.var() impl */ - DPNP_FN_ZEROS, /**< Used in numpy.zeros() impl */ - DPNP_FN_ZEROS_LIKE, /**< Used in numpy.zeros_like() impl */ - DPNP_FN_LAST, /**< The latest element of the enumeration */ + DPNP_FN_SUM, /**< Used in numpy.sum() impl */ + DPNP_FN_TAKE, /**< Used in numpy.take() impl */ + DPNP_FN_VAR, /**< Used in numpy.var() impl */ + DPNP_FN_ZEROS, /**< Used in numpy.zeros() impl */ + DPNP_FN_ZEROS_LIKE, /**< Used in numpy.zeros_like() impl */ + DPNP_FN_LAST, /**< The latest element of the enumeration */ }; /** @@ -381,14 +361,4 @@ void *get_dpnp_function_ptr1( DPNPFuncType first_type, DPNPFuncType second_type = DPNPFuncType::DPNP_FT_NONE); -/** - * DEPRECATED. - * Experimental interface. DO NOT USE IT! - * - * parameter @ref type_name will be converted into var_args or char *[] with - * extra length parameter - */ -INP_DLLEXPORT -void *get_backend_function_name(const char *func_name, const char *type_name); - #endif // BACKEND_IFACE_FPTR_H diff --git a/dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp b/dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp index 175eb3d76987..ebcffa944c04 100644 --- a/dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp @@ -31,355 +31,6 @@ #include "dpnpc_memory_adapter.hpp" #include "queue_sycl.hpp" -template -class dpnp_arange_c_kernel; - -template -DPCTLSyclEventRef dpnp_arange_c(DPCTLSyclQueueRef q_ref, - size_t start, - size_t step, - void *result1, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // parameter `size` used instead `stop` to avoid dependency on array length - // calculation algorithm - // TODO: floating point (and negatives) types from `start` and `step` - - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; - - validate_type_for_device<_DataType>(q); - - _DataType *result = reinterpret_cast<_DataType *>(result1); - - sycl::range<1> gws(size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - - result[i] = start + i * step; - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - event_ref = reinterpret_cast(&event); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_arange_c(size_t start, size_t step, void *result1, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_arange_c<_DataType>( - q_ref, start, step, result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_arange_default_c)(size_t, size_t, void *, size_t) = - dpnp_arange_c<_DataType>; - -// Explicit instantiation of the function, since dpnp_arange_c() is used by -// other template functions, but implicit instantiation is not applied anymore. -template DPCTLSyclEventRef dpnp_arange_c(DPCTLSyclQueueRef, - size_t, - size_t, - void *, - size_t, - const DPCTLEventVectorRef); - -template DPCTLSyclEventRef dpnp_arange_c(DPCTLSyclQueueRef, - size_t, - size_t, - void *, - size_t, - const DPCTLEventVectorRef); - -template DPCTLSyclEventRef dpnp_arange_c(DPCTLSyclQueueRef, - size_t, - size_t, - void *, - size_t, - const DPCTLEventVectorRef); - -template DPCTLSyclEventRef dpnp_arange_c(DPCTLSyclQueueRef, - size_t, - size_t, - void *, - size_t, - const DPCTLEventVectorRef); - -template -DPCTLSyclEventRef dpnp_diag_c(DPCTLSyclQueueRef q_ref, - void *v_in, - void *result1, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)res_ndim; - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - sycl::queue q = *(reinterpret_cast(q_ref)); - - validate_type_for_device<_DataType>(q); - - const size_t input1_size = std::accumulate( - shape, shape + ndim, 1, std::multiplies()); - const size_t result_size = std::accumulate( - res_shape, res_shape + res_ndim, 1, std::multiplies()); - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, v_in, input1_size, true); - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result1, result_size, true, - true); - _DataType *v = input1_ptr.get_ptr(); - _DataType *result = result_ptr.get_ptr(); - - size_t init0 = std::max(0, -k); - size_t init1 = std::max(0, k); - - if (ndim == 1) { - for (size_t i = 0; i < static_cast(shape[0]); ++i) { - size_t ind = (init0 + i) * res_shape[1] + init1 + i; - result[ind] = v[i]; - } - } - else { - for (size_t i = 0; i < static_cast(res_shape[0]); ++i) { - size_t ind = (init0 + i) * shape[1] + init1 + i; - result[i] = v[ind]; - } - } - return event_ref; -} - -template -void dpnp_diag_c(void *v_in, - void *result1, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_diag_c<_DataType>(q_ref, v_in, result1, k, shape, res_shape, ndim, - res_ndim, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_diag_default_c)(void *, - void *, - const int, - shape_elem_type *, - shape_elem_type *, - const size_t, - const size_t) = dpnp_diag_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_eye_c(DPCTLSyclQueueRef q_ref, - void *result1, - int k, - const shape_elem_type *res_shape, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (result1 == nullptr) { - return event_ref; - } - - if (res_shape == nullptr) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - validate_type_for_device<_DataType>(q); - - size_t result_size = res_shape[0] * res_shape[1]; - - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result1, result_size, true, - true); - _DataType *result = result_ptr.get_ptr(); - - int diag_val_; - diag_val_ = std::min((int)res_shape[0], (int)res_shape[1]); - diag_val_ = std::min(diag_val_, ((int)res_shape[0] + k)); - diag_val_ = std::min(diag_val_, ((int)res_shape[1] - k)); - - size_t diag_val = (diag_val_ < 0) ? 0 : (size_t)diag_val_; - - for (size_t i = 0; i < result_size; ++i) { - result[i] = 0; - for (size_t j = 0; j < diag_val; ++j) { - size_t ind = (k >= 0) ? (j * res_shape[1] + j + k) - : (j - k) * res_shape[1] + j; - if (i == ind) { - result[i] = 1; - break; - } - } - } - - return event_ref; -} - -template -void dpnp_eye_c(void *result1, int k, const shape_elem_type *res_shape) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_eye_c<_DataType>(q_ref, result1, k, res_shape, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_eye_default_c)(void *, - int, - const shape_elem_type *) = dpnp_eye_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_full_c(DPCTLSyclQueueRef q_ref, - void *array_in, - void *result, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - return dpnp_initval_c<_DataType>(q_ref, result, array_in, size, - dep_event_vec_ref); -} - -template -void dpnp_full_c(void *array_in, void *result, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_full_c<_DataType>( - q_ref, array_in, result, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_full_default_c)(void *, - void *, - const size_t) = dpnp_full_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_full_like_c(DPCTLSyclQueueRef q_ref, - void *array_in, - void *result, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - return dpnp_full_c<_DataType>(q_ref, array_in, result, size, - dep_event_vec_ref); -} - -template -void dpnp_full_like_c(void *array_in, void *result, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_full_like_c<_DataType>( - q_ref, array_in, result, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_full_like_default_c)(void *, - void *, - const size_t) = dpnp_full_like_c<_DataType>; - -template -class dpnp_identity_c_kernel; - -template -DPCTLSyclEventRef dpnp_identity_c(DPCTLSyclQueueRef q_ref, - void *result1, - const size_t n, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (n == 0) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; - - validate_type_for_device<_DataType>(q); - - _DataType *result = static_cast<_DataType *>(result1); - - sycl::range<2> gws(n, n); - auto kernel_parallel_for_func = [=](sycl::id<2> global_id) { - size_t i = global_id[0]; - size_t j = global_id[1]; - result[i * n + j] = i == j; - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - event_ref = reinterpret_cast(&event); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_identity_c(void *result1, const size_t n) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_identity_c<_DataType>(q_ref, result1, n, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_identity_default_c)(void *, - const size_t) = dpnp_identity_c<_DataType>; - template class dpnp_ones_c_kernel; @@ -442,632 +93,6 @@ void dpnp_ones_like_c(void *result, size_t size) template void (*dpnp_ones_like_default_c)(void *, size_t) = dpnp_ones_like_c<_DataType>; -template -DPCTLSyclEventRef dpnp_ptp_c(DPCTLSyclQueueRef q_ref, - void *result1_out, - const size_t result_size, - const size_t result_ndim, - const shape_elem_type *result_shape, - const shape_elem_type *result_strides, - const void *input1_in, - const size_t input_size, - const size_t input_ndim, - const shape_elem_type *input_shape, - const shape_elem_type *input_strides, - const shape_elem_type *axis, - const size_t naxis, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)result_strides; - (void)input_strides; - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - DPCTLSyclEventRef e1_ref = nullptr; - DPCTLSyclEventRef e2_ref = nullptr; - DPCTLSyclEventRef e3_ref = nullptr; - - if ((input1_in == nullptr) || (result1_out == nullptr)) { - return event_ref; - } - - if (input_ndim < 1) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - validate_type_for_device<_DataType>(q); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, input1_in, input_size, true); - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result1_out, result_size, - false, true); - _DataType *arr = input1_ptr.get_ptr(); - _DataType *result = result_ptr.get_ptr(); - - _DataType *min_arr = reinterpret_cast<_DataType *>( - sycl::malloc_shared(result_size * sizeof(_DataType), q)); - _DataType *max_arr = reinterpret_cast<_DataType *>( - sycl::malloc_shared(result_size * sizeof(_DataType), q)); - - e1_ref = dpnp_min_c<_DataType>(q_ref, arr, min_arr, result_size, - input_shape, input_ndim, axis, naxis, NULL); - e2_ref = dpnp_max_c<_DataType>(q_ref, arr, max_arr, result_size, - input_shape, input_ndim, axis, naxis, NULL); - - shape_elem_type *_strides = reinterpret_cast( - sycl::malloc_shared(result_ndim * sizeof(shape_elem_type), q)); - get_shape_offsets_inkernel(result_shape, result_ndim, _strides); - - e3_ref = dpnp_subtract_c<_DataType, _DataType, _DataType>( - q_ref, result, result_size, result_ndim, result_shape, result_strides, - max_arr, result_size, result_ndim, result_shape, _strides, min_arr, - result_size, result_ndim, result_shape, _strides, NULL, NULL); - - DPCTLEvent_Wait(e1_ref); - DPCTLEvent_Wait(e2_ref); - DPCTLEvent_Wait(e3_ref); - DPCTLEvent_Delete(e1_ref); - DPCTLEvent_Delete(e2_ref); - DPCTLEvent_Delete(e3_ref); - - sycl::free(min_arr, q); - sycl::free(max_arr, q); - sycl::free(_strides, q); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_ptp_c(void *result1_out, - const size_t result_size, - const size_t result_ndim, - const shape_elem_type *result_shape, - const shape_elem_type *result_strides, - const void *input1_in, - const size_t input_size, - const size_t input_ndim, - const shape_elem_type *input_shape, - const shape_elem_type *input_strides, - const shape_elem_type *axis, - const size_t naxis) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_ptp_c<_DataType>( - q_ref, result1_out, result_size, result_ndim, result_shape, - result_strides, input1_in, input_size, input_ndim, input_shape, - input_strides, axis, naxis, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_ptp_default_c)(void *, - const size_t, - const size_t, - const shape_elem_type *, - const shape_elem_type *, - const void *, - const size_t, - const size_t, - const shape_elem_type *, - const shape_elem_type *, - const shape_elem_type *, - const size_t) = dpnp_ptp_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_vander_c(DPCTLSyclQueueRef q_ref, - const void *array1_in, - void *result1, - const size_t size_in, - const size_t N, - const int increasing, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - DPCTLSyclEventRef event_ref = nullptr; - - if ((array1_in == nullptr) || (result1 == nullptr)) - return event_ref; - - if (!size_in || !N) - return event_ref; - - sycl::queue q = *(reinterpret_cast(q_ref)); - - validate_type_for_device<_DataType_input>(q); - validate_type_for_device<_DataType_output>(q); - - DPNPC_ptr_adapter<_DataType_input> input1_ptr(q_ref, array1_in, size_in, - true); - DPNPC_ptr_adapter<_DataType_output> result_ptr(q_ref, result1, size_in * N, - true, true); - const _DataType_input *array_in = input1_ptr.get_ptr(); - _DataType_output *result = result_ptr.get_ptr(); - - if (N == 1) { - return dpnp_ones_c<_DataType_output>(q_ref, result, size_in, - dep_event_vec_ref); - } - - if (increasing) { - for (size_t i = 0; i < size_in; ++i) { - result[i * N] = 1; - } - for (size_t i = 1; i < N; ++i) { - for (size_t j = 0; j < size_in; ++j) { - result[j * N + i] = result[j * N + i - 1] * array_in[j]; - } - } - } - else { - for (size_t i = 0; i < size_in; ++i) { - result[i * N + N - 1] = 1; - } - for (size_t i = N - 2; i > 0; --i) { - for (size_t j = 0; j < size_in; ++j) { - result[j * N + i] = result[j * N + i + 1] * array_in[j]; - } - } - - for (size_t i = 0; i < size_in; ++i) { - result[i * N] = result[i * N + 1] * array_in[i]; - } - } - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_vander_c(const void *array1_in, - void *result1, - const size_t size_in, - const size_t N, - const int increasing) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_vander_c<_DataType_input, _DataType_output>( - q_ref, array1_in, result1, size_in, N, increasing, - dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_vander_default_c)(const void *, - void *, - const size_t, - const size_t, - const int) = - dpnp_vander_c<_DataType_input, _DataType_output>; - -template -class dpnp_trace_c_kernel; - -template -DPCTLSyclEventRef dpnp_trace_c(DPCTLSyclQueueRef q_ref, - const void *array1_in, - void *result_in, - const shape_elem_type *shape_, - const size_t ndim, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!array1_in || !result_in || !shape_ || !ndim) { - return event_ref; - } - - const size_t last_dim = shape_[ndim - 1]; - const size_t size = std::accumulate(shape_, shape_ + (ndim - 1), 1, - std::multiplies()); - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - validate_type_for_device<_DataType>(q); - validate_type_for_device<_ResultType>(q); - - const _DataType *input = static_cast(array1_in); - _ResultType *result = static_cast<_ResultType *>(result_in); - - sycl::range<1> gws(size); - auto kernel_parallel_for_func = [=](auto index) { - size_t i = index[0]; - _ResultType acc = _ResultType(0); - - for (size_t j = 0; j < last_dim; ++j) { - acc += input[i * last_dim + j]; - } - - result[i] = acc; - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - auto event = q.submit(kernel_func); - event_ref = reinterpret_cast(&event); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_trace_c(const void *array1_in, - void *result_in, - const shape_elem_type *shape_, - const size_t ndim) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_trace_c<_DataType, _ResultType>( - q_ref, array1_in, result_in, shape_, ndim, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_trace_default_c)(const void *, - void *, - const shape_elem_type *, - const size_t) = - dpnp_trace_c<_DataType, _ResultType>; - -template -class dpnp_tri_c_kernel; - -template -DPCTLSyclEventRef dpnp_tri_c(DPCTLSyclQueueRef q_ref, - void *result1, - const size_t N, - const size_t M, - const int k, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - sycl::event event; - - if (!result1 || !N || !M) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - validate_type_for_device<_DataType>(q); - - _DataType *result = static_cast<_DataType *>(result1); - - size_t idx = N * M; - sycl::range<1> gws(idx); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - size_t ind = global_id[0]; - size_t i = ind / M; - size_t j = ind % M; - - int val = i + k + 1; - size_t diag_idx_ = (val > 0) ? (size_t)val : 0; - size_t diag_idx = (M < diag_idx_) ? M : diag_idx_; - - if (j < diag_idx) { - result[ind] = 1; - } - else { - result[ind] = 0; - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - event_ref = reinterpret_cast(&event); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_tri_c(void *result1, const size_t N, const size_t M, const int k) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_tri_c<_DataType>(q_ref, result1, N, M, k, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_tri_default_c)(void *, const size_t, const size_t, const int) = - dpnp_tri_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_tril_c(DPCTLSyclQueueRef q_ref, - void *array_in, - void *result1, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if ((array_in == nullptr) || (result1 == nullptr)) { - return event_ref; - } - - if ((shape == nullptr) || (res_shape == nullptr)) { - return event_ref; - } - - if ((ndim == 0) || (res_ndim == 0)) { - return event_ref; - } - - const size_t res_size = std::accumulate(res_shape, res_shape + res_ndim, 1, - std::multiplies()); - if (res_size == 0) { - return event_ref; - } - - const size_t input_size = std::accumulate( - shape, shape + ndim, 1, std::multiplies()); - if (input_size == 0) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - validate_type_for_device<_DataType>(q); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array_in, input_size, true); - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result1, res_size, true, - true); - _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; - for (size_t j = 0; j < res_ndim; ++j) { - n /= res_shape[j]; - size_t p = val / n; - ids[j] = p; - if (p != 0) { - val = val - p * n; - } - } - - int diag_idx_ = - (ids[res_ndim - 2] + k > -1) ? (ids[res_ndim - 2] + k) : -1; - int values = res_shape[res_ndim - 1]; - int diag_idx = (values < diag_idx_) ? values : diag_idx_; - - if (ids[res_ndim - 1] <= diag_idx) { - result[i] = array_m[ids[res_ndim - 1]]; - } - else { - result[i] = 0; - } - } - } - else { - for (size_t i = 0; i < res_size; ++i) { - size_t n = res_size; - size_t val = i; - for (size_t j = 0; j < res_ndim; ++j) { - n /= res_shape[j]; - size_t p = val / n; - ids[j] = p; - if (p != 0) { - val = val - p * n; - } - } - - int diag_idx_ = - (ids[res_ndim - 2] + k > -1) ? (ids[res_ndim - 2] + k) : -1; - int values = res_shape[res_ndim - 1]; - int diag_idx = (values < diag_idx_) ? values : diag_idx_; - - if (ids[res_ndim - 1] <= diag_idx) { - result[i] = array_m[i]; - } - else { - result[i] = 0; - } - } - } - - delete[] ids; - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_tril_c(void *array_in, - void *result1, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_tril_c<_DataType>(q_ref, array_in, result1, k, shape, res_shape, - ndim, res_ndim, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_tril_default_c)(void *, - void *, - const int, - shape_elem_type *, - shape_elem_type *, - const size_t, - const size_t) = dpnp_tril_c<_DataType>; - -template -DPCTLSyclEventRef dpnp_triu_c(DPCTLSyclQueueRef q_ref, - void *array_in, - void *result1, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if ((array_in == nullptr) || (result1 == nullptr)) { - return event_ref; - } - - if ((shape == nullptr) || (res_shape == nullptr)) { - return event_ref; - } - - if ((ndim == 0) || (res_ndim == 0)) { - return event_ref; - } - - const size_t res_size = std::accumulate(res_shape, res_shape + res_ndim, 1, - std::multiplies()); - if (res_size == 0) { - return event_ref; - } - - const size_t input_size = std::accumulate( - shape, shape + ndim, 1, std::multiplies()); - if (input_size == 0) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - validate_type_for_device<_DataType>(q); - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array_in, input_size, true); - DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result1, res_size, true, - true); - _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; - for (size_t j = 0; j < res_ndim; ++j) { - n /= res_shape[j]; - size_t p = val / n; - ids[j] = p; - if (p != 0) { - val = val - p * n; - } - } - - int diag_idx_ = - (ids[res_ndim - 2] + k > -1) ? (ids[res_ndim - 2] + k) : -1; - int values = res_shape[res_ndim - 1]; - int diag_idx = (values < diag_idx_) ? values : diag_idx_; - - if (ids[res_ndim - 1] >= diag_idx) { - result[i] = array_m[ids[res_ndim - 1]]; - } - else { - result[i] = 0; - } - } - } - else { - for (size_t i = 0; i < res_size; ++i) { - size_t n = res_size; - size_t val = i; - for (size_t j = 0; j < res_ndim; ++j) { - n /= res_shape[j]; - size_t p = val / n; - ids[j] = p; - if (p != 0) { - val = val - p * n; - } - } - - int diag_idx_ = - (ids[res_ndim - 2] + k > -1) ? (ids[res_ndim - 2] + k) : -1; - int values = res_shape[res_ndim - 1]; - int diag_idx = (values < diag_idx_) ? values : diag_idx_; - - if (ids[res_ndim - 1] >= diag_idx) { - result[i] = array_m[i]; - } - else { - result[i] = 0; - } - } - } - - delete[] ids; - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_triu_c(void *array_in, - void *result1, - const int k, - shape_elem_type *shape, - shape_elem_type *res_shape, - const size_t ndim, - const size_t res_ndim) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_triu_c<_DataType>(q_ref, array_in, result1, k, shape, res_shape, - ndim, res_ndim, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_triu_default_c)(void *, - void *, - const int, - shape_elem_type *, - shape_elem_type *, - const size_t, - const size_t) = dpnp_triu_c<_DataType>; - template DPCTLSyclEventRef dpnp_zeros_c(DPCTLSyclQueueRef q_ref, void *result, @@ -1130,72 +155,7 @@ void (*dpnp_zeros_like_default_c)(void *, void func_map_init_arraycreation(func_map_t &fmap) { - fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_arange_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_arange_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_arange_default_c}; - fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_arange_default_c}; - - fmap[DPNPFuncName::DPNP_FN_DIAG][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_diag_default_c}; - fmap[DPNPFuncName::DPNP_FN_DIAG][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_diag_default_c}; - fmap[DPNPFuncName::DPNP_FN_DIAG][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_diag_default_c}; - fmap[DPNPFuncName::DPNP_FN_DIAG][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_diag_default_c}; - - fmap[DPNPFuncName::DPNP_FN_EYE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_eye_default_c}; - fmap[DPNPFuncName::DPNP_FN_EYE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_eye_default_c}; - fmap[DPNPFuncName::DPNP_FN_EYE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_eye_default_c}; - fmap[DPNPFuncName::DPNP_FN_EYE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_eye_default_c}; - - fmap[DPNPFuncName::DPNP_FN_FULL][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_full_default_c}; - fmap[DPNPFuncName::DPNP_FN_FULL][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_full_default_c}; - fmap[DPNPFuncName::DPNP_FN_FULL][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_full_default_c}; - fmap[DPNPFuncName::DPNP_FN_FULL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_full_default_c}; - fmap[DPNPFuncName::DPNP_FN_FULL][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_full_default_c}; - fmap[DPNPFuncName::DPNP_FN_FULL][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_full_default_c>}; - - fmap[DPNPFuncName::DPNP_FN_FULL_LIKE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_full_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_FULL_LIKE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_full_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_FULL_LIKE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_full_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_FULL_LIKE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_full_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_FULL_LIKE][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_full_like_default_c}; - fmap[DPNPFuncName::DPNP_FN_FULL_LIKE][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_full_like_default_c>}; - - fmap[DPNPFuncName::DPNP_FN_IDENTITY][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_identity_default_c}; - fmap[DPNPFuncName::DPNP_FN_IDENTITY][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_identity_default_c}; - fmap[DPNPFuncName::DPNP_FN_IDENTITY][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_identity_default_c}; - fmap[DPNPFuncName::DPNP_FN_IDENTITY][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_identity_default_c}; - fmap[DPNPFuncName::DPNP_FN_IDENTITY][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_identity_default_c}; - fmap[DPNPFuncName::DPNP_FN_IDENTITY][eft_C128][eft_C128] = { - eft_C128, (void *)dpnp_identity_default_c>}; - + // Used in dpnp_rng_geometric_c fmap[DPNPFuncName::DPNP_FN_ONES][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_ones_default_c}; fmap[DPNPFuncName::DPNP_FN_ONES][eft_LNG][eft_LNG] = { @@ -1222,90 +182,8 @@ void func_map_init_arraycreation(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ONES_LIKE][eft_C128][eft_C128] = { eft_C128, (void *)dpnp_ones_like_default_c>}; - fmap[DPNPFuncName::DPNP_FN_PTP][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_ptp_default_c}; - fmap[DPNPFuncName::DPNP_FN_PTP][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_ptp_default_c}; - fmap[DPNPFuncName::DPNP_FN_PTP][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_ptp_default_c}; - fmap[DPNPFuncName::DPNP_FN_PTP][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_ptp_default_c}; - - fmap[DPNPFuncName::DPNP_FN_VANDER][eft_INT][eft_INT] = { - eft_LNG, (void *)dpnp_vander_default_c}; - fmap[DPNPFuncName::DPNP_FN_VANDER][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_vander_default_c}; - fmap[DPNPFuncName::DPNP_FN_VANDER][eft_FLT][eft_FLT] = { - eft_DBL, (void *)dpnp_vander_default_c}; - fmap[DPNPFuncName::DPNP_FN_VANDER][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_vander_default_c}; - fmap[DPNPFuncName::DPNP_FN_VANDER][eft_BLN][eft_BLN] = { - eft_LNG, (void *)dpnp_vander_default_c}; - fmap[DPNPFuncName::DPNP_FN_VANDER][eft_C128][eft_C128] = { - eft_C128, - (void *) - dpnp_vander_default_c, std::complex>}; - - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_LNG][eft_INT] = { - eft_INT, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_FLT][eft_INT] = { - eft_INT, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_DBL][eft_INT] = { - eft_INT, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_FLT][eft_LNG] = { - eft_LNG, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_DBL][eft_LNG] = { - eft_LNG, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_INT][eft_FLT] = { - eft_FLT, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_LNG][eft_FLT] = { - eft_FLT, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_DBL][eft_FLT] = { - eft_FLT, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_trace_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRACE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_trace_default_c}; - - fmap[DPNPFuncName::DPNP_FN_TRI][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_tri_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRI][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_tri_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRI][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_tri_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRI][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_tri_default_c}; - - fmap[DPNPFuncName::DPNP_FN_TRIL][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_tril_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRIL][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_tril_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRIL][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_tril_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRIL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_tril_default_c}; - - fmap[DPNPFuncName::DPNP_FN_TRIU][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_triu_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRIU][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_triu_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRIU][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_triu_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRIU][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_triu_default_c}; - + // Used in dpnp_rng_binomial_c, dpnp_rng_gamma_c, dpnp_rng_hypergeometric_c + // dpnp_rng_laplace_c, dpnp_rng_multinomial_c, dpnp_rng_weibull_c fmap[DPNPFuncName::DPNP_FN_ZEROS][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_zeros_default_c}; fmap[DPNPFuncName::DPNP_FN_ZEROS][eft_LNG][eft_LNG] = { diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 20be65f53cab..e3797bd22e6e 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -1026,19 +1026,6 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) #include -template -static void func_map_elemwise_2arg_3type_core(func_map_t &fmap) -{ - // dpnp_subtract_c_ext is implicitly used by dpnp_ptp_c - ((fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][FT1][FTs] = - {populate_func_types(), - (void *)dpnp_subtract_c_ext< - func_type_map_t::find_type()>, - func_type_map_t::find_type, - func_type_map_t::find_type>}), - ...); -} - template static void func_map_elemwise_2arg_3type_short_core(func_map_t &fmap) { @@ -1072,12 +1059,6 @@ static void func_map_elemwise_2arg_3type_short_core(func_map_t &fmap) ...); } -template -static void func_map_elemwise_2arg_3type_helper(func_map_t &fmap) -{ - ((func_map_elemwise_2arg_3type_core(fmap)), ...); -} - template static void func_map_elemwise_2arg_3type_short_helper(func_map_t &fmap) { @@ -1189,9 +1170,6 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) (void *)dpnp_multiply_c_default< std::complex, std::complex, std::complex>}; - func_map_elemwise_2arg_3type_helper(fmap); - func_map_elemwise_2arg_3type_short_helper(fmap); diff --git a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp index dcbf6ca906c7..523acd447c64 100644 --- a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp @@ -125,31 +125,6 @@ DPCTLSyclEventRef (*dpnp_choose_ext_c)(DPCTLSyclQueueRef, const DPCTLEventVectorRef) = dpnp_choose_c<_DataType1, _DataType2>; -template -DPCTLSyclEventRef - dpnp_diag_indices_c(DPCTLSyclQueueRef q_ref, - void *result1, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - return dpnp_arange_c<_DataType>(q_ref, 0, 1, result1, size, - dep_event_vec_ref); -} - -template -void dpnp_diag_indices_c(void *result1, size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_diag_indices_c<_DataType>(q_ref, result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_diag_indices_default_c)(void *, - size_t) = dpnp_diag_indices_c<_DataType>; - template DPCTLSyclEventRef dpnp_diagonal_c(DPCTLSyclQueueRef q_ref, void *array1_in, @@ -873,15 +848,6 @@ void func_map_init_indexing_func(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_LNG][eft_DBL] = { eft_DBL, (void *)dpnp_choose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_DIAG_INDICES][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_diag_indices_default_c}; - fmap[DPNPFuncName::DPNP_FN_DIAG_INDICES][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_diag_indices_default_c}; - fmap[DPNPFuncName::DPNP_FN_DIAG_INDICES][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_diag_indices_default_c}; - fmap[DPNPFuncName::DPNP_FN_DIAG_INDICES][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_diag_indices_default_c}; - fmap[DPNPFuncName::DPNP_FN_DIAGONAL][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_diagonal_default_c}; fmap[DPNPFuncName::DPNP_FN_DIAGONAL][eft_LNG][eft_LNG] = { diff --git a/dpnp/backend/kernels/dpnp_krnl_manipulation.cpp b/dpnp/backend/kernels/dpnp_krnl_manipulation.cpp deleted file mode 100644 index aaaa5a179dd7..000000000000 --- a/dpnp/backend/kernels/dpnp_krnl_manipulation.cpp +++ /dev/null @@ -1,235 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include -#include -#include - -#include - -#include "dpnp_fptr.hpp" -#include "dpnp_utils.hpp" -#include "dpnpc_memory_adapter.hpp" -#include "queue_sycl.hpp" - -template -class dpnp_repeat_c_kernel; - -template -DPCTLSyclEventRef dpnp_repeat_c(DPCTLSyclQueueRef q_ref, - const void *array1_in, - void *result1, - const size_t repeats, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!array1_in || !result1) { - return event_ref; - } - - if (!size || !repeats) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size); - const _DataType *array_in = input1_ptr.get_ptr(); - _DataType *result = reinterpret_cast<_DataType *>(result1); - - sycl::range<2> gws(size, repeats); - auto kernel_parallel_for_func = [=](sycl::id<2> global_id) { - size_t idx1 = global_id[0]; - size_t idx2 = global_id[1]; - result[(idx1 * repeats) + idx2] = array_in[idx1]; - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - - event_ref = reinterpret_cast(&event); - - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_repeat_c(const void *array1_in, - void *result1, - const size_t repeats, - const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_repeat_c<_DataType>( - q_ref, array1_in, result1, repeats, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_repeat_default_c)(const void *, - void *, - const size_t, - const size_t) = dpnp_repeat_c<_DataType>; - -template -class dpnp_elemwise_transpose_c_kernel; - -template -DPCTLSyclEventRef - dpnp_elemwise_transpose_c(DPCTLSyclQueueRef q_ref, - void *array1_in, - const shape_elem_type *input_shape, - const shape_elem_type *result_shape, - const shape_elem_type *permute_axes, - size_t ndim, - void *result1, - size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!size) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; - - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size); - _DataType *array1 = input1_ptr.get_ptr(); - _DataType *result = reinterpret_cast<_DataType *>(result1); - - shape_elem_type *input_offset_shape = reinterpret_cast( - sycl::malloc_shared(ndim * sizeof(shape_elem_type), q)); - get_shape_offsets_inkernel(input_shape, ndim, input_offset_shape); - - shape_elem_type *temp_result_offset_shape = - reinterpret_cast( - sycl::malloc_shared(ndim * sizeof(shape_elem_type), q)); - get_shape_offsets_inkernel(result_shape, ndim, temp_result_offset_shape); - - shape_elem_type *result_offset_shape = reinterpret_cast( - sycl::malloc_shared(ndim * sizeof(shape_elem_type), q)); - for (size_t axis = 0; axis < ndim; ++axis) { - result_offset_shape[permute_axes[axis]] = - temp_result_offset_shape[axis]; - } - - sycl::range<1> gws(size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - const size_t idx = global_id[0]; - - size_t output_index = 0; - size_t reminder = idx; - for (size_t axis = 0; axis < ndim; ++axis) { - /* reconstruct [x][y][z] from given linear idx */ - size_t xyz_id = reminder / input_offset_shape[axis]; - reminder = reminder % input_offset_shape[axis]; - - /* calculate destination index based on reconstructed [x][y][z] */ - output_index += (xyz_id * result_offset_shape[axis]); - } - - result[output_index] = array1[idx]; - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); - - event.wait(); - - sycl::free(input_offset_shape, q); - sycl::free(temp_result_offset_shape, q); - sycl::free(result_offset_shape, q); - - return event_ref; -} - -template -void dpnp_elemwise_transpose_c(void *array1_in, - const shape_elem_type *input_shape, - const shape_elem_type *result_shape, - const shape_elem_type *permute_axes, - size_t ndim, - void *result1, - size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_elemwise_transpose_c<_DataType>( - q_ref, array1_in, input_shape, result_shape, permute_axes, ndim, - result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_elemwise_transpose_default_c)(void *, - const shape_elem_type *, - const shape_elem_type *, - const shape_elem_type *, - size_t, - void *, - size_t) = - dpnp_elemwise_transpose_c<_DataType>; - -void func_map_init_manipulation(func_map_t &fmap) -{ - fmap[DPNPFuncName::DPNP_FN_REPEAT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_repeat_default_c}; - fmap[DPNPFuncName::DPNP_FN_REPEAT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_repeat_default_c}; - fmap[DPNPFuncName::DPNP_FN_REPEAT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_repeat_default_c}; - fmap[DPNPFuncName::DPNP_FN_REPEAT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_repeat_default_c}; - - fmap[DPNPFuncName::DPNP_FN_TRANSPOSE][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_elemwise_transpose_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRANSPOSE][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_elemwise_transpose_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRANSPOSE][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_elemwise_transpose_default_c}; - fmap[DPNPFuncName::DPNP_FN_TRANSPOSE][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_elemwise_transpose_default_c}; - return; -} diff --git a/dpnp/backend/src/dpnp_fptr.hpp b/dpnp/backend/src/dpnp_fptr.hpp index 2a9c42eb1720..73d627812a5c 100644 --- a/dpnp/backend/src/dpnp_fptr.hpp +++ b/dpnp/backend/src/dpnp_fptr.hpp @@ -331,7 +331,6 @@ void func_map_init_fft_func(func_map_t &fmap); void func_map_init_indexing_func(func_map_t &fmap); void func_map_init_linalg(func_map_t &fmap); void func_map_init_logic(func_map_t &fmap); -void func_map_init_manipulation(func_map_t &fmap); void func_map_init_mathematical(func_map_t &fmap); void func_map_init_random(func_map_t &fmap); void func_map_init_reduction(func_map_t &fmap); diff --git a/dpnp/backend/src/dpnp_iface_fptr.cpp b/dpnp/backend/src/dpnp_iface_fptr.cpp index f8214212728d..f80c5b358639 100644 --- a/dpnp/backend/src/dpnp_iface_fptr.cpp +++ b/dpnp/backend/src/dpnp_iface_fptr.cpp @@ -96,45 +96,6 @@ void (*dpnp_dot_default_c)(void *, const shape_elem_type *) = dpnp_dot_c<_DataType_output, _DataType_input1, _DataType_input2>; -void *get_backend_function_name(const char *func_name, const char *type_name) -{ - /** Implement it in this way to allow easier play with it */ - const char *supported_func_name = "dpnp_dot"; - const char *supported_type1_name = "double"; - const char *supported_type2_name = "float"; - const char *supported_type3_name = "long"; - const char *supported_type4_name = "int"; - - /** of coerce it will be converted into std::map later */ - if (!strncmp(func_name, supported_func_name, strlen(supported_func_name))) { - if (!strncmp(type_name, supported_type1_name, - strlen(supported_type1_name))) { - return reinterpret_cast( - dpnp_dot_default_c); - } - else if (!strncmp(type_name, supported_type2_name, - strlen(supported_type2_name))) - { - return reinterpret_cast( - dpnp_dot_default_c); - } - else if (!strncmp(type_name, supported_type3_name, - strlen(supported_type3_name))) - { - return reinterpret_cast( - dpnp_dot_default_c); - } - else if (!strncmp(type_name, supported_type4_name, - strlen(supported_type4_name))) - { - return reinterpret_cast( - dpnp_dot_default_c); - } - } - - throw std::runtime_error("DPNP Error: Unsupported function call"); -} - /** * This operator is needed for compatibility with Cython 0.29 which has a bug in * Enum handling @@ -172,7 +133,6 @@ static func_map_t func_map_init() func_map_init_indexing_func(fmap); func_map_init_linalg(fmap); func_map_init_logic(fmap); - func_map_init_manipulation(fmap); func_map_init_mathematical(fmap); func_map_init_random(fmap); func_map_init_reduction(fmap); diff --git a/dpnp/backend/src/queue_sycl.cpp b/dpnp/backend/src/queue_sycl.cpp index 5e6df29d21d2..786752facd60 100644 --- a/dpnp/backend/src/queue_sycl.cpp +++ b/dpnp/backend/src/queue_sycl.cpp @@ -80,36 +80,6 @@ } #endif -#if defined(DPNPC_TOUCH_KERNEL_TO_LINK) -/** - * Function push the SYCL kernels to be linked (final stage of the compilation) - * for the current queue - * - * TODO it is not the best idea to just a call some kernel. Needs better - * solution. - */ -static long dpnp_kernels_link() -{ - /* must use memory pre-allocated at the current queue */ - long *value_ptr = - reinterpret_cast(dpnp_memory_alloc_c(1 * sizeof(long))); - long *result_ptr = - reinterpret_cast(dpnp_memory_alloc_c(1 * sizeof(long))); - long result = 1; - - *value_ptr = 2; - - dpnp_square_c(value_ptr, result_ptr, 1); - - result = *result_ptr; - - dpnp_memory_free_c(result_ptr); - dpnp_memory_free_c(value_ptr); - - return result; -} -#endif - size_t dpnp_queue_is_cpu_c() { const auto &be = backend_sycl::get(); diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 37663bee8343..0c8bd1134a78 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -84,7 +84,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_RNG_WALD_EXT DPNP_FN_RNG_WEIBULL_EXT DPNP_FN_RNG_ZIPF_EXT - DPNP_FN_TRAPZ_EXT cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncType": # need this namespace for Enum import cdef enum DPNPFuncType "DPNPFuncType": diff --git a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi index fca1e6dc3036..28b89ce60a1a 100644 --- a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi @@ -40,16 +40,12 @@ __all__ += [ "dpnp_fmax", "dpnp_fmin", "dpnp_modf", - "dpnp_trapz", ] ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_1in_2out_t)(c_dpctl.DPCTLSyclQueueRef, void * , void * , void * , size_t, const c_dpctl.DPCTLEventVectorRef) -ctypedef c_dpctl.DPCTLSyclEventRef(*ftpr_custom_trapz_2in_1out_with_2size_t)(c_dpctl.DPCTLSyclQueueRef, - void *, void * , void * , double, size_t, size_t, - const c_dpctl.DPCTLEventVectorRef) cpdef utils.dpnp_descriptor dpnp_ediff1d(utils.dpnp_descriptor x1): @@ -166,41 +162,3 @@ cpdef tuple dpnp_modf(utils.dpnp_descriptor x1): c_dpctl.DPCTLEvent_Delete(event_ref) return (result1.get_pyobj(), result2.get_pyobj()) - - -cpdef utils.dpnp_descriptor dpnp_trapz(utils.dpnp_descriptor y1, utils.dpnp_descriptor x1, double dx): - - cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(y1.dtype) - cdef DPNPFuncType param2_type = dpnp_dtype_to_DPNPFuncType(x1.dtype) - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_TRAPZ_EXT, param1_type, param2_type) - - result_sycl_device, result_usm_type, result_sycl_queue = utils.get_common_usm_allocation(y1, x1) - - # create result array with type given by FPTR data - cdef shape_type_c result_shape = (1,) - cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, - kernel_data.return_type, - None, - device=result_sycl_device, - usm_type=result_usm_type, - sycl_queue=result_sycl_queue) - - result_sycl_queue = result.get_array().sycl_queue - - cdef c_dpctl.SyclQueue q = result_sycl_queue - cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() - - cdef ftpr_custom_trapz_2in_1out_with_2size_t func = kernel_data.ptr - cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, - y1.get_data(), - x1.get_data(), - result.get_data(), - dx, - y1.size, - x1.size, - NULL) # dep_events_ref - - with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) - c_dpctl.DPCTLEvent_Delete(event_ref) - - return result diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 1fe7839f5967..1caf1359be3e 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -64,7 +64,6 @@ dpnp_fmax, dpnp_fmin, dpnp_modf, - dpnp_trapz, ) from .dpnp_algo.dpnp_elementwise_common import ( DPNPAngle, @@ -3287,36 +3286,6 @@ def trapz(y1, x1=None, dx=1.0, axis=-1): """ - y_desc = dpnp.get_dpnp_descriptor(y1, copy_when_nondefault_queue=False) - if y_desc: - if y_desc.ndim > 1: - pass - else: - y_obj = y_desc.get_array() - if x1 is None: - x_obj = dpnp.empty( - y_desc.shape, - dtype=y_desc.dtype, - device=y_obj.sycl_device, - usm_type=y_obj.usm_type, - sycl_queue=y_obj.sycl_queue, - ) - else: - x_obj = x1 - - x_desc = dpnp.get_dpnp_descriptor( - x_obj, copy_when_nondefault_queue=False - ) - # TODO: change to "not x_desc" - if x_desc: - pass - elif y_desc.size != x_desc.size: - pass - elif y_desc.shape != x_desc.shape: - pass - else: - return dpnp_trapz(y_desc, x_desc, dx).get_pyobj() - return call_origin(numpy.trapz, y1, x1, dx, axis)