fix negative strides

Author: vtavana

dpnp/backend/extensions/blas/dot.cpp

@@ -193,6 +193,15 @@ std::pair<sycl::event, sycl::event> dot(sycl::queue exec_q,
     char *b_typeless_ptr = vectorB.get_data();
     char *r_typeless_ptr = result.get_data();
 
+    const int a_elemsize = vectorA.get_elemsize();
+    const int b_elemsize = vectorB.get_elemsize();
+    if (str_a < 0) {
+        a_typeless_ptr -= (n - 1) * std::abs(str_a) * a_elemsize;
+    }
+    if (str_b < 0) {
+        b_typeless_ptr -= (n - 1) * std::abs(str_b) * b_elemsize;
+    }
+
     sycl::event dot_ev = dot_fn(exec_q, n, a_typeless_ptr, str_a,
                                 b_typeless_ptr, str_b, r_typeless_ptr, depends);
 

dpnp/backend/extensions/blas/dotu.cpp

@@ -195,6 +195,15 @@ std::pair<sycl::event, sycl::event>
     char *b_typeless_ptr = vectorB.get_data();
     char *r_typeless_ptr = result.get_data();
 
+    const int a_elemsize = vectorA.get_elemsize();
+    const int b_elemsize = vectorB.get_elemsize();
+    if (str_a < 0) {
+        a_typeless_ptr -= (n - 1) * std::abs(str_a) * a_elemsize;
+    }
+    if (str_b < 0) {
+        b_typeless_ptr -= (n - 1) * std::abs(str_b) * b_elemsize;
+    }
+
     sycl::event dotu_ev =
         dotu_fn(exec_q, n, a_typeless_ptr, str_a, b_typeless_ptr, str_b,
                 r_typeless_ptr, depends);

tests/test_dot.py

@@ -181,21 +181,28 @@ def test_dot_ndarray(dtype, array_info):
     assert_dtype_allclose(result, expected)
 
 
-@pytest.mark.parametrize("dtype", get_all_dtypes())
+@pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True))
 def test_dot_strided(dtype):
-    a = numpy.ones(25, dtype=dtype)
-    b = numpy.ones(25, dtype=dtype)
+    a = numpy.arange(25, dtype=dtype)
+    b = numpy.arange(25, dtype=dtype)
     ia = dpnp.array(a)
     ib = dpnp.array(b)
 
     result = dpnp.dot(ia[::3], ib[::3])
     expected = numpy.dot(a[::3], b[::3])
     assert_dtype_allclose(result, expected)
 
-    # TODO: unmute when the problem with negative stride is fixed
-    # result = dpnp.dot(ia, ib[::-1])
-    # expected = numpy.dot(a, b[::-1])
-    # assert_dtype_allclose(result, expected)
+    result = dpnp.dot(ia, ib[::-1])
+    expected = numpy.dot(a, b[::-1])
+    assert_dtype_allclose(result, expected)
+
+    result = dpnp.dot(ia[::-2], ib[::-2])
+    expected = numpy.dot(a[::-2], b[::-2])
+    assert_dtype_allclose(result, expected)
+
+    result = dpnp.dot(ia[::-5], ib[::-5])
+    expected = numpy.dot(a[::-5], b[::-5])
+    assert_dtype_allclose(result, expected)
 
 
 @pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True))