Function "linalg.qr" raised an exception with mode='reduced' (#1148)

* DPNP function "linalg.qr" raised an exception when mode='reduced'.

Author: antonwolfy

dpnp/backend/kernels/dpnp_krnl_linalg.cpp

@@ -662,35 +662,36 @@ DPCTLSyclEventRef dpnp_qr_c(DPCTLSyclQueueRef q_ref,
         }
     }
 
-    DPNPC_ptr_adapter<_ComputeDT> result1_ptr(q_ref, result1, size_m * size_m, true, true);
-    DPNPC_ptr_adapter<_ComputeDT> result2_ptr(q_ref, result2, size_m * size_n, true, true);
-    DPNPC_ptr_adapter<_ComputeDT> result3_ptr(q_ref, result3, std::min(size_m, size_n), true, true);
+    const size_t min_size_m_n = std::min<size_t>(size_m, size_n);
+    DPNPC_ptr_adapter<_ComputeDT> result1_ptr(q_ref, result1, size_m * min_size_m_n, true, true);
+    DPNPC_ptr_adapter<_ComputeDT> result2_ptr(q_ref, result2, min_size_m_n * size_n, true, true);
+    DPNPC_ptr_adapter<_ComputeDT> result3_ptr(q_ref, result3, min_size_m_n, true, true);
     _ComputeDT* res_q = result1_ptr.get_ptr();
     _ComputeDT* res_r = result2_ptr.get_ptr();
     _ComputeDT* tau = result3_ptr.get_ptr();
 
     const std::int64_t lda = size_m;
 
-    const std::int64_t geqrf_scratchpad_size =
-        mkl_lapack::geqrf_scratchpad_size<_ComputeDT>(q, size_m, size_n, lda);
+    const std::int64_t geqrf_scratchpad_size = mkl_lapack::geqrf_scratchpad_size<_ComputeDT>(q, size_m, size_n, lda);
 
     _ComputeDT* geqrf_scratchpad =
         reinterpret_cast<_ComputeDT*>(sycl::malloc_shared(geqrf_scratchpad_size * sizeof(_ComputeDT), q));
 
     std::vector<sycl::event> depends(1);
     set_barrier_event(q, depends);
 
-    event =
-        mkl_lapack::geqrf(q, size_m, size_n, in_a, lda, tau, geqrf_scratchpad, geqrf_scratchpad_size, depends);
-
+    event = mkl_lapack::geqrf(q, size_m, size_n, in_a, lda, tau, geqrf_scratchpad, geqrf_scratchpad_size, depends);
     event.wait();
 
-    verbose_print("oneapi::mkl::lapack::geqrf", depends.front(), event);
+    if (!depends.empty()) {
+        verbose_print("oneapi::mkl::lapack::geqrf", depends.front(), event);
+    }
 
     sycl::free(geqrf_scratchpad, q);
 
     // R
-    for (size_t i = 0; i < size_m; ++i)
+    size_t mrefl = min_size_m_n;
+    for (size_t i = 0; i < mrefl; ++i)
     {
         for (size_t j = 0; j < size_n; ++j)
         {
@@ -706,37 +707,30 @@ DPCTLSyclEventRef dpnp_qr_c(DPCTLSyclQueueRef q_ref,
     }
 
     // Q
-    const size_t nrefl = std::min<size_t>(size_m, size_n);
+    const size_t nrefl = min_size_m_n;
     const std::int64_t orgqr_scratchpad_size =
-        mkl_lapack::orgqr_scratchpad_size<_ComputeDT>(q, size_m, size_m, nrefl, lda);
+        mkl_lapack::orgqr_scratchpad_size<_ComputeDT>(q, size_m, nrefl, nrefl, lda);
 
     _ComputeDT* orgqr_scratchpad =
         reinterpret_cast<_ComputeDT*>(sycl::malloc_shared(orgqr_scratchpad_size * sizeof(_ComputeDT), q));
 
-    depends.clear();
     set_barrier_event(q, depends);
 
-    event = mkl_lapack::orgqr(
-        q, size_m, size_m, nrefl, in_a, lda, tau, orgqr_scratchpad, orgqr_scratchpad_size, depends);
-
+    event =
+        mkl_lapack::orgqr(q, size_m, nrefl, nrefl, in_a, lda, tau, orgqr_scratchpad, orgqr_scratchpad_size, depends);
     event.wait();
 
-    verbose_print("oneapi::mkl::lapack::orgqr", depends.front(), event);
+    if (!depends.empty()) {
+        verbose_print("oneapi::mkl::lapack::orgqr", depends.front(), event);
+    }
 
     sycl::free(orgqr_scratchpad, q);
 
     for (size_t i = 0; i < size_m; ++i)
     {
-        for (size_t j = 0; j < size_m; ++j)
+        for (size_t j = 0; j < nrefl; ++j)
         {
-            if (j < nrefl)
-            {
-                res_q[i * size_m + j] = in_a[j * size_m + i];
-            }
-            else
-            {
-                res_q[i * size_m + j] = _ComputeDT(0);
-            }
+            res_q[i * nrefl + j] = in_a[j * size_m + i];
         }
     }
 

dpnp/linalg/dpnp_algo_linalg.pyx

@@ -308,13 +308,14 @@ cpdef object dpnp_norm(object input, ord=None, axis=None):
 cpdef tuple dpnp_qr(utils.dpnp_descriptor x1, str mode):
     cdef size_t size_m = x1.shape[0]
     cdef size_t size_n = x1.shape[1]
-    cdef size_t size_tau = min(size_m, size_n)
+    cdef size_t min_m_n = min(size_m, size_n)
+    cdef size_t size_tau = min_m_n
 
     cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype)
     cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_QR, param1_type, param1_type)
 
-    cdef utils.dpnp_descriptor res_q = utils.create_output_descriptor((size_m, size_m), kernel_data.return_type, None)
-    cdef utils.dpnp_descriptor res_r = utils.create_output_descriptor((size_m, size_n), kernel_data.return_type, None)
+    cdef utils.dpnp_descriptor res_q = utils.create_output_descriptor((size_m, min_m_n), kernel_data.return_type, None)
+    cdef utils.dpnp_descriptor res_r = utils.create_output_descriptor((min_m_n, size_n), kernel_data.return_type, None)
     cdef utils.dpnp_descriptor tau = utils.create_output_descriptor((size_tau, ), kernel_data.return_type, None)
 
     cdef custom_linalg_1in_3out_shape_t func = < custom_linalg_1in_3out_shape_t > kernel_data.ptr

dpnp/linalg/dpnp_iface_linalg.py

@@ -391,7 +391,7 @@ def qr(x1, mode='reduced'):
     Limitations
     -----------
     Input array is supported as :obj:`dpnp.ndarray`.
-    Parameter mode='complete' is supported.
+    Parameter mode='reduced' is supported.
 
     """
 
@@ -400,7 +400,7 @@ def qr(x1, mode='reduced'):
         if mode != 'reduced':
             pass
         else:
-            result_tup = dpnp_qr(x1, mode)
+            result_tup = dpnp_qr(x1_desc, mode)
 
             return result_tup
 

tests/skipped_tests_gpu.tbl

@@ -263,10 +263,6 @@ tests/test_linalg.py::test_eig_arange[16-float64]
 tests/test_linalg.py::test_eig_arange[16-float32]
 tests/test_linalg.py::test_eig_arange[16-int64]
 tests/test_linalg.py::test_eig_arange[16-int32]
-tests/test_linalg.py::test_qr[(16,16)-float64]
-tests/test_linalg.py::test_qr[(16,16)-float32]
-tests/test_linalg.py::test_qr[(16,16)-int64]
-tests/test_linalg.py::test_qr[(16,16)-int32]
 tests/test_random.py::test_randn_normal_distribution
 tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_multidim_outer
 tests/third_party/cupy/random_tests/test_sample.py::TestRandintDtype::test_dtype

tests/test_linalg.py

@@ -219,12 +219,15 @@ def test_norm3(array, ord, axis):
 @pytest.mark.parametrize("shape",
                          [(2, 2), (3, 4), (5, 3), (16, 16)],
                          ids=['(2,2)', '(3,4)', '(5,3)', '(16,16)'])
-def test_qr(type, shape):
+@pytest.mark.parametrize("mode",
+                         ['complete', 'reduced'],
+                         ids=['complete', 'reduced'])
+def test_qr(type, shape, mode):
     a = numpy.arange(shape[0] * shape[1], dtype=type).reshape(shape)
     ia = inp.array(a)
 
-    np_q, np_r = numpy.linalg.qr(a, "complete")
-    dpnp_q, dpnp_r = inp.linalg.qr(ia, "complete")
+    np_q, np_r = numpy.linalg.qr(a, mode)
+    dpnp_q, dpnp_r = inp.linalg.qr(ia, mode)
 
     assert (dpnp_q.dtype == np_q.dtype)
     assert (dpnp_r.dtype == np_r.dtype)