Merge changes for nightly tests in CI (#1173)

* dpnp_take_c uses SYCL kernel, no need to use no_sycl parameter in adapter

The reason this crashed with CPU device and gave incorrect results on
Windows was deeper.

   1. Adapter call allocated USM-shared buffer and copies data into it
   2. Kernel is submitted to work on USM-shared pointer
   3. dpnp_take_c returns kernel submission even
   4. Adapter class goes out of scope and frees USM allocation
      without making sure that the kernel that works on it has completed
      its execution
   5. If kernel execution was in progress we got a crash on CPU, or
      incorrect result on GPU

      If kernel execution was complete it worked as expected.

This change fixes the problem because it removes creation of unprotected
USM-shared temporary.

* Change to DPNPC_adapter to set/use events upon which deallocation must depend

The deallocation routine simply calls sycl::event::wait on the stored vector.

* Used DPNPC_ptr_adapter::depends_on

Also applied DPCTLEvent_Delete in legacy interfaces to avoid memory leak.

* Get rid of "Improper Null Termination" issue

Add a null-terminated symbol at the end of char array to avoid "Improper Null Termination" issue reported by Checkmarx scan.

* implemented PR feedback

* Reworked solution with a pointer on void

* Update dpnp/backend/kernels/dpnp_krnl_random.cpp

Co-authored-by: Oleksandr Pavlyk <[email protected]>

* Update dpnp/backend/kernels/dpnp_krnl_random.cpp

Co-authored-by: Oleksandr Pavlyk <[email protected]>

* Skip for two more tests till waiting fix (#1171)

* Skip for two more tests till waiting fix

tests/third_party/cupy/statistics_tests/test_order.py::TestOrder::test_ptp_all_nan
tests/third_party/cupy/statistics_tests/test_order.py::TestOrder::test_ptp_nan

Need to skip them because CI does not work due to this.

* The same tests skip for gpu

* dpnp_take failed on Windows due to memory corruption (#1172)

* dpnp_take failed on Windows due to memory corruption

* Add more tests

* Integer indexes types with different types of input data

* Add trailing empty line to .gitignore

Co-authored-by: Oleksandr Pavlyk <[email protected]>
Co-authored-by: Alexander Rybkin <[email protected]>

Author: 3 people

.gitignore

@@ -1,7 +1,15 @@
+# CMake build and local install directory
 build
 build_cython
+
+# Byte-compiled / optimized / DLL files
 __pycache__/
+
+# Code project files
+.vscode
+
 *dpnp_backend*
 dpnp/**/*.cpython*.so
 dpnp/**/*.pyd
-*~
+*~
+core

dpnp/backend/kernels/dpnp_krnl_indexing.cpp

@@ -901,10 +901,8 @@ DPCTLSyclEventRef dpnp_take_c(DPCTLSyclQueueRef q_ref,
     DPCTLSyclEventRef event_ref = nullptr;
     sycl::queue q = *(reinterpret_cast<sycl::queue*>(q_ref));
 
-    DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, array1_size);
-    DPNPC_ptr_adapter<_IndecesType> input2_ptr(q_ref, indices1, size);
-    _DataType* array_1 = input1_ptr.get_ptr();
-    _IndecesType* indices = input2_ptr.get_ptr();
+    _DataType* array_1 = reinterpret_cast<_DataType*>(array1_in);
+    _IndecesType* indices = reinterpret_cast<_IndecesType*>(indices1);
     _DataType* result = reinterpret_cast<_DataType*>(result1);
 
     sycl::range<1> gws(size);
@@ -920,7 +918,6 @@ DPCTLSyclEventRef dpnp_take_c(DPCTLSyclQueueRef q_ref,
     sycl::event event = q.submit(kernel_func);
 
     event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
-
     return DPCTLEvent_Copy(event_ref);
 }
 
@@ -937,6 +934,7 @@ void dpnp_take_c(void* array1_in, const size_t array1_size, void* indices1, void
                                                                        size,
                                                                        dep_event_vec_ref);
     DPCTLEvent_WaitAndThrow(event_ref);
+    DPCTLEvent_Delete(event_ref);
 }
 
 template <typename _DataType, typename _IndecesType>
@@ -1073,21 +1071,36 @@ void func_map_init_indexing_func(func_map_t& fmap)
     fmap[DPNPFuncName::DPNP_FN_PUT_ALONG_AXIS_EXT][eft_DBL][eft_DBL] = {eft_DBL,
                                                                         (void*)dpnp_put_along_axis_ext_c<double>};
 
-    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_BLN][eft_BLN] = {eft_BLN, (void*)dpnp_take_default_c<bool, int64_t>};
-    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_take_default_c<int32_t, int64_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_BLN][eft_INT] = {eft_BLN, (void*)dpnp_take_default_c<bool, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_take_default_c<int32_t, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_LNG][eft_INT] = {eft_LNG, (void*)dpnp_take_default_c<int64_t, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_FLT][eft_INT] = {eft_FLT, (void*)dpnp_take_default_c<float, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_DBL][eft_INT] = {eft_DBL, (void*)dpnp_take_default_c<double, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_C128][eft_INT] = {eft_C128,
+                                                           (void*)dpnp_take_default_c<std::complex<double>, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_BLN][eft_LNG] = {eft_BLN, (void*)dpnp_take_default_c<bool, int64_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_INT][eft_LNG] = {eft_INT, (void*)dpnp_take_default_c<int32_t, int32_t>};
     fmap[DPNPFuncName::DPNP_FN_TAKE][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_take_default_c<int64_t, int64_t>};
-    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_take_default_c<float, int64_t>};
-    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_take_default_c<double, int64_t>};
-    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_C128][eft_C128] = {eft_C128,
-                                                            (void*)dpnp_take_default_c<std::complex<double>, int64_t>};
-
-    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_BLN][eft_BLN] = {eft_BLN, (void*)dpnp_take_ext_c<bool, int64_t>};
-    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_take_ext_c<int32_t, int64_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_FLT][eft_LNG] = {eft_FLT, (void*)dpnp_take_default_c<float, int64_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_DBL][eft_LNG] = {eft_DBL, (void*)dpnp_take_default_c<double, int64_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE][eft_C128][eft_LNG] = {eft_C128,
+                                                           (void*)dpnp_take_default_c<std::complex<double>, int64_t>};
+
+    // TODO: add a handling of other indexes types once DPCtl implementation of data copy is ready
+    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_BLN][eft_INT] = {eft_BLN, (void*)dpnp_take_ext_c<bool, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_take_ext_c<int32_t, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_LNG][eft_INT] = {eft_LNG, (void*)dpnp_take_ext_c<int64_t, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_FLT][eft_INT] = {eft_FLT, (void*)dpnp_take_ext_c<float, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_DBL][eft_INT] = {eft_DBL, (void*)dpnp_take_ext_c<double, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_C128][eft_INT] = {eft_C128,
+                                                               (void*)dpnp_take_ext_c<std::complex<double>, int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_BLN][eft_LNG] = {eft_BLN, (void*)dpnp_take_ext_c<bool, int64_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_INT][eft_LNG] = {eft_INT, (void*)dpnp_take_ext_c<int32_t, int64_t>};
     fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_take_ext_c<int64_t, int64_t>};
-    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_take_ext_c<float, int64_t>};
-    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_take_ext_c<double, int64_t>};
-    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_C128][eft_C128] = {eft_C128,
-                                                                (void*)dpnp_take_ext_c<std::complex<double>, int64_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_FLT][eft_LNG] = {eft_FLT, (void*)dpnp_take_ext_c<float, int64_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_DBL][eft_LNG] = {eft_DBL, (void*)dpnp_take_ext_c<double, int64_t>};
+    fmap[DPNPFuncName::DPNP_FN_TAKE_EXT][eft_C128][eft_LNG] = {eft_C128,
+                                                               (void*)dpnp_take_ext_c<std::complex<double>, int64_t>};
 
     return;
 }

tests/skipped_tests_gpu.tbl

@@ -138,19 +138,6 @@ tests/test_indexing.py::test_nonzero[[[0, 1, 2], [3, 0, 5], [6, 7, 0]]]
 tests/test_indexing.py::test_nonzero[[[0, 1, 0, 3, 0], [5, 0, 7, 0, 9]]]
 tests/test_indexing.py::test_nonzero[[[[1, 2], [0, 4]], [[0, 2], [0, 1]], [[0, 0], [3, 1]]]]
 tests/test_indexing.py::test_nonzero[[[[[1, 2, 3], [3, 4, 5]], [[1, 2, 3], [2, 1, 0]]], [[[1, 3, 5], [3, 1, 0]], [[0, 1, 2], [1, 3, 4]]]]]
-tests/test_indexing.py::test_take[[[0, 1, 2], [3, 4, 5], [6, 7, 8]]-[[0, 0], [0, 0]]]
-tests/test_indexing.py::test_take[[[0, 1, 2], [3, 4, 5], [6, 7, 8]]-[[1, 2], [1, 2]]]
-tests/test_indexing.py::test_take[[[0, 1, 2], [3, 4, 5], [6, 7, 8]]-[[1, 2], [3, 4]]]
-tests/test_indexing.py::test_take[[[0, 1, 2, 3, 4], [5, 6, 7, 8, 9]]-[[1, 2], [1, 2]]]
-tests/test_indexing.py::test_take[[[0, 1, 2, 3, 4], [5, 6, 7, 8, 9]]-[[1, 2], [3, 4]]]
-tests/test_indexing.py::test_take[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-[[1, 2], [1, 2]]]
-tests/test_indexing.py::test_take[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-[[1, 2], [3, 4]]]
-tests/test_indexing.py::test_take[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-[[1, 2], [1, 2]]]
-tests/test_indexing.py::test_take[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-[[1, 2], [3, 4]]]
-tests/test_indexing.py::test_take[[[[[1, 2, 3], [3, 4, 5]], [[1, 2, 3], [2, 1, 0]]], [[[1, 3, 5], [3, 1, 0]], [[0, 1, 2], [1, 3, 4]]]]-[[1, 2], [1, 2]]]
-tests/test_indexing.py::test_take[[[[[1, 2, 3], [3, 4, 5]], [[1, 2, 3], [2, 1, 0]]], [[[1, 3, 5], [3, 1, 0]], [[0, 1, 2], [1, 3, 4]]]]-[[1, 2], [3, 4]]]
-tests/test_indexing.py::test_take[[[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]], [[[13, 14, 15], [16, 17, 18]], [[19, 20, 21], [22, 23, 24]]]]-[[1, 2], [1, 2]]]
-tests/test_indexing.py::test_take[[[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]], [[[13, 14, 15], [16, 17, 18]], [[19, 20, 21], [22, 23, 24]]]]-[[1, 2], [3, 4]]]
 tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_arange_no_dtype_int
 tests/third_party/cupy/indexing_tests/test_indexing.py::TestIndexing::test_take_no_axis
 tests/third_party/cupy/indexing_tests/test_insert.py::TestPlace_param_3_{n_vals=1, shape=(7,)}::test_place

tests/test_indexing.py

@@ -374,6 +374,12 @@ def test_select():
     numpy.testing.assert_array_equal(expected, result)
 
 
+@pytest.mark.parametrize("array_type",
+                         [numpy.bool8, numpy.int32, numpy.int64, numpy.float32, numpy.float64, numpy.complex128],
+                         ids=['bool8', 'int32', 'int64', 'float32', 'float64', 'complex128'])
+@pytest.mark.parametrize("indices_type",
+                         [numpy.int32, numpy.int64],
+                         ids=['int32', 'int64'])
 @pytest.mark.parametrize("indices",
                          [[[0, 0], [0, 0]],
                           [[1, 2], [1, 2]],
@@ -395,9 +401,9 @@ def test_select():
                               '[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]',
                               '[[[[1, 2, 3], [3, 4, 5]], [[1, 2, 3], [2, 1, 0]]], [[[1, 3, 5], [3, 1, 0]], [[0, 1, 2], [1, 3, 4]]]]',
                               '[[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]], [[[13, 14, 15], [16, 17, 18]], [[19, 20, 21], [22, 23, 24]]]]'])
-def test_take(array, indices):
-    a = numpy.array(array)
-    ind = numpy.array(indices)
+def test_take(array, indices, array_type, indices_type):
+    a = numpy.array(array, dtype=array_type)
+    ind = numpy.array(indices, dtype=indices_type)
     ia = dpnp.array(a)
     iind = dpnp.array(ind)
     expected = numpy.take(a, ind)