Revert "Merge master to gold (#705)"

This reverts commit bdc950a.

Author: densmirn

doc/conf.py

@@ -31,9 +31,9 @@
 author = 'Intel'
 
 # The short X.Y version
-version = '0.6'
+version = '0.5'
 # The full version, including alpha/beta/rc tags
-release = '0.6.2'
+release = '0.5.1'
 
 
 # -- General configuration ---------------------------------------------------

dpnp/backend/CMakeLists.txt

@@ -27,8 +27,8 @@
 
 cmake_minimum_required(VERSION 3.10 FATAL_ERROR)
 
-# set(DPNP_VERSION 0.6.2)
-# set(DPNP_API_VERSION 0.6)
+# set(DPNP_VERSION 0.5.1)
+# set(DPNP_API_VERSION 0.5)
 
 # set directory where the custom finders live
 set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_SOURCE_DIR}/cmake/Modules")

dpnp/backend/doc/Doxyfile

@@ -38,7 +38,7 @@ PROJECT_NAME           = "DPNP C++ backend kernel library"
 # could be handy for archiving the generated documentation or if some version
 # control system is used.
 
-PROJECT_NUMBER         = 0.6.2
+PROJECT_NUMBER         = 0.5.1
 
 # Using the PROJECT_BRIEF tag one can provide an optional one line description
 # for a project that appears at the top of each page and should give viewer a

dpnp/backend/include/dpnp_iface.hpp

@@ -705,18 +705,6 @@ INP_DLLEXPORT void dpnp_std_c(
 template <typename _DataType, typename _IndecesType>
 INP_DLLEXPORT void dpnp_take_c(void* array, void* indices, void* result, size_t size);
 
-/**
- * @ingroup BACKEND_API
- * @brief math library implementation of trace function
- *
- * @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.
- */
-template <typename _DataType, typename _ResultType>
-INP_DLLEXPORT void dpnp_trace_c(const void* array, void* result, const size_t* shape, const size_t ndim);
-
 /**
  * @ingroup BACKEND_API
  * @brief math library implementation of take function

dpnp/backend/include/dpnp_iface_fptr.hpp

@@ -199,7 +199,6 @@ enum class DPNPFuncName : size_t
     DPNP_FN_TAN,                      /**< Used in numpy.tan() implementation  */
     DPNP_FN_TANH,                     /**< Used in numpy.tanh() implementation  */
     DPNP_FN_TRANSPOSE,                /**< Used in numpy.transpose() implementation  */
-    DPNP_FN_TRACE,                    /**< Used in numpy.trace() implementation  */
     DPNP_FN_TRAPZ,                    /**< Used in numpy.trapz() implementation  */
     DPNP_FN_TRI,                      /**< Used in numpy.tri() implementation  */
     DPNP_FN_TRIL,                     /**< Used in numpy.tril() implementation  */

dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp

@@ -211,68 +211,6 @@ void dpnp_vander_c(const void* array1_in, void* result1, const size_t size_in, c
     }
 }
 
-template <typename _DataType, typename _ResultType>
-class dpnp_trace_c_kernel;
-
-template <typename _DataType, typename _ResultType>
-void dpnp_trace_c(const void* array1_in, void* result1, const size_t* shape_, const size_t ndim)
-{
-    cl::sycl::event event;
-
-    if ((array1_in == nullptr) || (result1 == nullptr))
-    {
-        return;
-    }
-
-    const _DataType* array_in = reinterpret_cast<const _DataType*>(array1_in);
-    _ResultType* result = reinterpret_cast<_ResultType*>(result1);
-
-    if (shape_ == nullptr)
-    {
-        return;
-    }
-
-    if (ndim == 0)
-    {
-        return;
-    }
-
-    size_t size = 1;
-    for (size_t i = 0; i < ndim - 1; ++i)
-    {
-        size *= shape_[i];
-    }
-
-    if (size == 0)
-    {
-        return;
-    }
-
-    size_t* shape = reinterpret_cast<size_t*>(dpnp_memory_alloc_c(ndim * sizeof(size_t)));
-    auto memcpy_event = DPNP_QUEUE.memcpy(shape, shape_, ndim * sizeof(size_t));
-
-    cl::sycl::range<1> gws(size);
-    auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
-        size_t i = global_id[0];
-        result[i] = 0;
-        for (size_t j = 0; j < shape[ndim - 1]; ++j)
-        {
-            result[i] += array_in[i * shape[ndim - 1] + j];
-        }
-    };
-
-    auto kernel_func = [&](cl::sycl::handler& cgh) {
-        cgh.depends_on({memcpy_event});
-        cgh.parallel_for<class dpnp_trace_c_kernel<_DataType, _ResultType>>(gws, kernel_parallel_for_func);
-    };
-
-    event = DPNP_QUEUE.submit(kernel_func);
-
-    event.wait();
-
-    dpnp_memory_free_c(shape);
-}
-
 template <typename _DataType>
 class dpnp_tri_c_kernel;
 
@@ -601,23 +539,6 @@ void func_map_init_arraycreation(func_map_t& fmap)
     fmap[DPNPFuncName::DPNP_FN_VANDER][eft_C128][eft_C128] = {
         eft_C128, (void*)dpnp_vander_c<std::complex<double>, std::complex<double>>};
 
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_trace_c<int, int>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_LNG][eft_INT] = {eft_INT, (void*)dpnp_trace_c<long, int>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_FLT][eft_INT] = {eft_INT, (void*)dpnp_trace_c<float, int>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_DBL][eft_INT] = {eft_INT, (void*)dpnp_trace_c<double, int>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_INT][eft_LNG] = {eft_LNG, (void*)dpnp_trace_c<int, long>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_trace_c<long, long>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_FLT][eft_LNG] = {eft_LNG, (void*)dpnp_trace_c<float, long>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_DBL][eft_LNG] = {eft_LNG, (void*)dpnp_trace_c<double, long>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_INT][eft_FLT] = {eft_FLT, (void*)dpnp_trace_c<int, float>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_LNG][eft_FLT] = {eft_FLT, (void*)dpnp_trace_c<long, float>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_trace_c<float, float>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_DBL][eft_FLT] = {eft_FLT, (void*)dpnp_trace_c<double, float>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_INT][eft_DBL] = {eft_DBL, (void*)dpnp_trace_c<int, double>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_LNG][eft_DBL] = {eft_DBL, (void*)dpnp_trace_c<long, double>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_FLT][eft_DBL] = {eft_DBL, (void*)dpnp_trace_c<float, double>};
-    fmap[DPNPFuncName::DPNP_FN_TRACE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_trace_c<double, double>};
-
     fmap[DPNPFuncName::DPNP_FN_TRI][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_tri_c<int>};
     fmap[DPNPFuncName::DPNP_FN_TRI][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_tri_c<long>};
     fmap[DPNPFuncName::DPNP_FN_TRI][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_tri_c<float>};

dpnp/backend/kernels/dpnp_krnl_random.cpp

@@ -475,72 +475,54 @@ void dpnp_rng_logistic_c(void* result, const double loc, const double scale, con
 template <typename _DataType>
 void dpnp_rng_lognormal_c(void* result, const _DataType mean, const _DataType stddev, const size_t size)
 {
-    if (!size || !result)
+    if (!size)
     {
         return;
     }
     _DataType* result1 = reinterpret_cast<_DataType*>(result);
 
-    if (stddev == 0.0)
-    {
-        _DataType* fill_value = reinterpret_cast<_DataType*>(dpnp_memory_alloc_c(sizeof(_DataType)));
-        fill_value[0] = static_cast<_DataType>(std::exp(mean + (stddev * stddev) / 2));
-        dpnp_initval_c<_DataType>(result, fill_value, size);
-        dpnp_memory_free_c(fill_value);
-    }
-    else
-    {
-        const _DataType displacement = _DataType(0.0);
-        const _DataType scalefactor = _DataType(1.0);
+    const _DataType displacement = _DataType(0.0);
 
-        mkl_rng::lognormal<_DataType> distribution(mean, stddev, displacement, scalefactor);
-        auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1);
-        event_out.wait();
-    }
-    return;
+    const _DataType scalefactor = _DataType(1.0);
+
+    mkl_rng::lognormal<_DataType> distribution(mean, stddev, displacement, scalefactor);
+    // perform generation
+    auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1);
+    event_out.wait();
 }
 
 template <typename _DataType>
 void dpnp_rng_multinomial_c(
     void* result, const int ntrial, const double* p_vector, const size_t p_vector_size, const size_t size)
 {
-    if (!size || !result)
+    if (!size)
     {
         return;
     }
+    std::int32_t* result1 = reinterpret_cast<std::int32_t*>(result);
+    std::vector<double> p(p_vector, p_vector + p_vector_size);
+    // size = size
+    // `result` is a array for random numbers
+    // `size` is a `result`'s len. `size = n * p.size()`
+    // `n` is a number of random values to be generated.
+    size_t n = size / p.size();
 
-    if (ntrial == 0)
+    if (dpnp_queue_is_cpu_c())
     {
-        dpnp_zeros_c<_DataType>(result, size);
+        mkl_rng::multinomial<std::int32_t> distribution(ntrial, p);
+        // perform generation
+        auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, n, result1);
+        event_out.wait();
     }
     else
     {
-        std::int32_t* result1 = reinterpret_cast<std::int32_t*>(result);
-        std::vector<double> p(p_vector, p_vector + p_vector_size);
-        // size = size
-        // `result` is a array for random numbers
-        // `size` is a `result`'s len. `size = n * p.size()`
-        // `n` is a number of random values to be generated.
-        size_t n = size / p.size();
-
-        if (dpnp_queue_is_cpu_c())
-        {
-            mkl_rng::multinomial<std::int32_t> distribution(ntrial, p);
-            // perform generation
-            auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, n, result1);
-            event_out.wait();
-        }
-        else
+        int errcode = viRngMultinomial(
+            VSL_RNG_METHOD_MULTINOMIAL_MULTPOISSON, get_rng_stream(), n, result1, ntrial, p_vector_size, p_vector);
+        if (errcode != VSL_STATUS_OK)
         {
-            int errcode = viRngMultinomial(
-                VSL_RNG_METHOD_MULTINOMIAL_MULTPOISSON, get_rng_stream(), n, result1, ntrial, p_vector_size, p_vector);
-            if (errcode != VSL_STATUS_OK)
-            {
-                throw std::runtime_error("DPNP RNG Error: dpnp_rng_multinomial_c() failed.");
-            }
+            throw std::runtime_error("DPNP RNG Error: dpnp_rng_multinomial_c() failed.");
         }
     }
-    return;
 }
 
 template <typename _DataType>
@@ -964,20 +946,17 @@ template <typename _DataType>
 void dpnp_rng_shuffle_c(
     void* result, const size_t itemsize, const size_t ndim, const size_t high_dim_size, const size_t size)
 {
-    if (!result)
-    {
-        return;
-    }
-
-    if (!size || !ndim || !(high_dim_size > 1))
+    if (!(size) || !(high_dim_size > 1))
     {
         return;
     }
 
     char* result1 = reinterpret_cast<char*>(result);
 
+    double* Uvec = nullptr;
+
     size_t uvec_size = high_dim_size - 1;
-    double* Uvec = reinterpret_cast<double*>(dpnp_memory_alloc_c(uvec_size * sizeof(double)));
+    Uvec = reinterpret_cast<double*>(dpnp_memory_alloc_c(uvec_size * sizeof(double)));
     mkl_rng::uniform<double> uniform_distribution(0.0, 1.0);
     auto uniform_event = mkl_rng::generate(uniform_distribution, DPNP_RNG_ENGINE, uvec_size, Uvec);
     uniform_event.wait();
@@ -987,52 +966,42 @@ void dpnp_rng_shuffle_c(
         // Fast, statically typed path: shuffle the underlying buffer.
         // Only for non-empty, 1d objects of class ndarray (subclasses such
         // as MaskedArrays may not support this approach).
-        char* buf = reinterpret_cast<char*>(dpnp_memory_alloc_c(itemsize * sizeof(char)));
+        // TODO
+        // kernel
+        char* buf = nullptr;
+        buf = reinterpret_cast<char*>(dpnp_memory_alloc_c(itemsize * sizeof(char)));
         for (size_t i = uvec_size; i > 0; i--)
         {
             size_t j = (size_t)(floor((i + 1) * Uvec[i - 1]));
-            if (i != j)
-            {
-                auto memcpy1 =
-                    DPNP_QUEUE.submit([&](cl::sycl::handler& h) { h.memcpy(buf, result1 + j * itemsize, itemsize); });
-                auto memcpy2 = DPNP_QUEUE.submit([&](cl::sycl::handler& h) {
-                    h.depends_on({memcpy1});
-                    h.memcpy(result1 + j * itemsize, result1 + i * itemsize, itemsize);
-                });
-                auto memcpy3 = DPNP_QUEUE.submit([&](cl::sycl::handler& h) {
-                    h.depends_on({memcpy2});
-                    h.memcpy(result1 + i * itemsize, buf, itemsize);
-                });
-                memcpy3.wait();
-            }
+            memcpy(buf, result1 + j * itemsize, itemsize);
+            memcpy(result1 + j * itemsize, result1 + i * itemsize, itemsize);
+            memcpy(result1 + i * itemsize, buf, itemsize);
         }
+
         dpnp_memory_free_c(buf);
     }
     else
     {
         // Multidimensional ndarrays require a bounce buffer.
+        // TODO
+        // kernel
+        char* buf = nullptr;
         size_t step_size = (size / high_dim_size) * itemsize; // size in bytes for x[i] element
-        char* buf = reinterpret_cast<char*>(dpnp_memory_alloc_c(step_size * sizeof(char)));
+        buf = reinterpret_cast<char*>(dpnp_memory_alloc_c(step_size * sizeof(char)));
         for (size_t i = uvec_size; i > 0; i--)
         {
             size_t j = (size_t)(floor((i + 1) * Uvec[i - 1]));
             if (j < i)
             {
-                auto memcpy1 =
-                    DPNP_QUEUE.submit([&](cl::sycl::handler& h) { h.memcpy(buf, result1 + j * step_size, step_size); });
-                auto memcpy2 = DPNP_QUEUE.submit([&](cl::sycl::handler& h) {
-                    h.depends_on({memcpy1});
-                    h.memcpy(result1 + j * step_size, result1 + i * step_size, step_size);
-                });
-                auto memcpy3 = DPNP_QUEUE.submit([&](cl::sycl::handler& h) {
-                    h.depends_on({memcpy2});
-                    h.memcpy(result1 + i * step_size, buf, step_size);
-                });
-                memcpy3.wait();
+                memcpy(buf, result1 + j * step_size, step_size);
+                memcpy(result1 + j * step_size, result1 + i * step_size, step_size);
+                memcpy(result1 + i * step_size, buf, step_size);
             }
         }
+
         dpnp_memory_free_c(buf);
     }
+
     dpnp_memory_free_c(Uvec);
 }
 

dpnp/backend/src/queue_sycl.cpp

@@ -97,7 +97,6 @@ static cl::sycl::device get_default_sycl_device()
 }
 #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
  *
@@ -121,7 +120,6 @@ static long dpnp_kernels_link()
 
     return result;
 }
-#endif
 
 #if defined(DPNP_LOCAL_QUEUE)
 // Catch asynchronous exceptions
@@ -179,10 +177,7 @@ void backend_sycl::backend_sycl_queue_init(QueueOptions selector)
 #endif
 
     std::chrono::high_resolution_clock::time_point t3 = std::chrono::high_resolution_clock::now();
-#if defined(DPNPC_TOUCH_KERNEL_TO_LINK)
-    // Remove pre-link kernel library at startup time
     dpnp_kernels_link();
-#endif
     std::chrono::high_resolution_clock::time_point t4 = std::chrono::high_resolution_clock::now();
     std::chrono::duration<double> time_kernels_link =
         std::chrono::duration_cast<std::chrono::duration<double>>(t4 - t3);

dpnp/dpnp_algo/dpnp_algo.pxd

@@ -171,7 +171,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName":  # need this na
         DPNP_FN_TAKE
         DPNP_FN_TAN
         DPNP_FN_TANH
-        DPNP_FN_TRACE
         DPNP_FN_TRANSPOSE
         DPNP_FN_TRAPZ
         DPNP_FN_TRI