Merge pull request #712 from IntelPython/master

Merge master to gold

Author: densmirn

doc/conf.py

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

dpnp/backend/CMakeLists.txt

@@ -27,8 +27,8 @@
 
 cmake_minimum_required(VERSION 3.10 FATAL_ERROR)
 
-# set(DPNP_VERSION 0.5.1)
-# set(DPNP_API_VERSION 0.5)
+# set(DPNP_VERSION 0.6.2)
+# set(DPNP_API_VERSION 0.6)
 
 # 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.5.1
+PROJECT_NUMBER         = 0.6.2
 
 # 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,6 +705,18 @@ 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,6 +199,7 @@ 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,6 +211,68 @@ 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;
 
@@ -539,6 +601,23 @@ 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,54 +475,72 @@ 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)
+    if (!size || !result)
     {
         return;
     }
     _DataType* result1 = reinterpret_cast<_DataType*>(result);
 
-    const _DataType displacement = _DataType(0.0);
-
-    const _DataType scalefactor = _DataType(1.0);
+    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);
 
-    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();
+        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;
 }
 
 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)
+    if (!size || !result)
     {
         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 (dpnp_queue_is_cpu_c())
+    if (ntrial == 0)
     {
-        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();
+        dpnp_zeros_c<_DataType>(result, size);
     }
     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)
+        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())
         {
-            throw std::runtime_error("DPNP RNG Error: dpnp_rng_multinomial_c() failed.");
+            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)
+            {
+                throw std::runtime_error("DPNP RNG Error: dpnp_rng_multinomial_c() failed.");
+            }
         }
     }
+    return;
 }
 
 template <typename _DataType>
@@ -946,17 +964,20 @@ 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 (!(size) || !(high_dim_size > 1))
+    if (!result)
     {
         return;
     }
 
-    char* result1 = reinterpret_cast<char*>(result);
+    if (!size || !ndim || !(high_dim_size > 1))
+    {
+        return;
+    }
 
-    double* Uvec = nullptr;
+    char* result1 = reinterpret_cast<char*>(result);
 
     size_t uvec_size = high_dim_size - 1;
-    Uvec = reinterpret_cast<double*>(dpnp_memory_alloc_c(uvec_size * sizeof(double)));
+    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();
@@ -966,42 +987,52 @@ 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).
-        // TODO
-        // kernel
-        char* buf = nullptr;
-        buf = reinterpret_cast<char*>(dpnp_memory_alloc_c(itemsize * sizeof(char)));
+        char* 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]));
-            memcpy(buf, result1 + j * itemsize, itemsize);
-            memcpy(result1 + j * itemsize, result1 + i * itemsize, itemsize);
-            memcpy(result1 + i * itemsize, buf, itemsize);
+            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();
+            }
         }
-
         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
-        buf = reinterpret_cast<char*>(dpnp_memory_alloc_c(step_size * sizeof(char)));
+        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)
             {
-                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);
+                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();
             }
         }
-
         dpnp_memory_free_c(buf);
     }
-
     dpnp_memory_free_c(Uvec);
 }
 

dpnp/backend/src/queue_sycl.cpp

@@ -97,6 +97,7 @@ 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
  *
@@ -120,6 +121,7 @@ static long dpnp_kernels_link()
 
     return result;
 }
+#endif
 
 #if defined(DPNP_LOCAL_QUEUE)
 // Catch asynchronous exceptions
@@ -177,7 +179,10 @@ 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,6 +171,7 @@ 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