Merge branch 'master' into add-md-local-accessor

Author: ndgrigorian

.github/CODEOWNERS

@@ -1 +1 @@
-*       @oleksandr-pavlyk @ndgrigorian
+*       @ndgrigorian @antonwolfy @vlad-perevezentsev @vtavana

.github/workflows/conda-package.yml

@@ -48,7 +48,7 @@ jobs:
         run: echo $CONDA/bin >> $GITHUB_PATH
       - name: Install conda-build
         # FIXME: unpin when conda-build fixes issues with lief
-        run: conda install conda-build -c conda-forge --override-channels
+        run: conda install conda-build"<25.1.2" -c conda-forge --override-channels
       - name: Store conda paths as envs
         shell: bash -l {0}
         run: |
@@ -100,7 +100,7 @@ jobs:
         # FIXME: unpin when conda-build fixes issues with lief
         run: |
           conda activate
-          conda install -y conda-build
+          conda install -y conda-build"<25.1.2"
           conda list -n base
 
       - name: Cache conda packages
@@ -167,7 +167,9 @@ jobs:
         run: echo $CONDA/bin >> $GITHUB_PATH
       - name: Install conda-index
         # Needed to be able to run conda index
-        run: conda install conda-index -c conda-forge --override-channels
+        run: |
+          conda update -n base --all
+          conda install conda-index -c conda-forge --override-channels
       - name: Create conda channel
         run: |
           mkdir -p $GITHUB_WORKSPACE/channel/linux-64
@@ -508,7 +510,9 @@ jobs:
           echo ${{ env.CHANNELS }}
       - name: Install conda-index
         # Needed to be able to run conda index
-        run: conda install conda-index -c conda-forge --override-channels
+        run: |
+          conda update -n base --all
+          conda install conda-index -c conda-forge --override-channels
       - name: Checkout dpctl repo
         uses: actions/checkout@v4.2.2
         with:
@@ -696,7 +700,9 @@ jobs:
         run: echo $CONDA/bin >> $GITHUB_PATH
       - name: Install conda-index
         # Needed to be able to run conda index
-        run: conda install conda-index -c conda-forge --override-channels
+        run: |
+          conda update -n base --all
+          conda install conda-index -c conda-forge --override-channels
       - name: Create conda channel
         run: |
           mkdir -p $GITHUB_WORKSPACE/channel/linux-64

conda-recipe/meta.yaml

@@ -80,7 +80,7 @@ about:
 
 extra:
     recipe-maintainers:
-        - diptorupd
-        - oleksandr-pavlyk
         - ndgrigorian
-        - ZzEeKkAa
+        - antonwolfy
+        - vtavana
+        - vlad-perevezentsev

dpctl/tensor/_copy_utils.py

@@ -126,7 +126,7 @@ def from_numpy(np_ary, /, *, device=None, usm_type="device", sycl_queue=None):
         arg:
             Input convertible to :class:`numpy.ndarray`
         device (object): array API specification of device where the
-            output array is created. Device can be specified by a
+            output array is created. Device can be specified by
             a filter selector string, an instance of
             :class:`dpctl.SyclDevice`, an instance of
             :class:`dpctl.SyclQueue`, or an instance of
@@ -620,7 +620,7 @@ def astype(
             If this keyword is set to `False`, a view of the input array
             may be returned when possible.
         device (object): array API specification of device where the
-            output array is created. Device can be specified by a
+            output array is created. Device can be specified by
             a filter selector string, an instance of
             :class:`dpctl.SyclDevice`, an instance of
             :class:`dpctl.SyclQueue`, or an instance of

dpctl/tensor/_usmarray.pyx

@@ -1777,8 +1777,10 @@ cdef api object UsmNDArray_MakeSimpleFromPtr(
     Returns:
         Created usm_ndarray instance
     """
-    cdef size_t itemsize = type_bytesize(typenum)
-    cdef size_t nbytes = itemsize * nelems
+    cdef int itemsize = type_bytesize(typenum)
+    if (itemsize < 1):
+        raise ValueError("dtype with typenum=" + str(typenum) + " is not supported.")
+    cdef size_t nbytes = (<size_t> itemsize) * nelems
     cdef c_dpmem._Memory mobj = c_dpmem._Memory.create_from_usm_pointer_size_qref(
         ptr, nbytes, QRef, memory_owner=owner
     )
@@ -1817,7 +1819,7 @@ cdef api object UsmNDArray_MakeFromPtr(
     Returns:
         Created usm_ndarray instance
     """
-    cdef size_t itemsize = type_bytesize(typenum)
+    cdef int itemsize = type_bytesize(typenum)
     cdef int err = 0
     cdef size_t nelems = 1
     cdef Py_ssize_t min_disp = 0
@@ -1830,6 +1832,8 @@ cdef api object UsmNDArray_MakeFromPtr(
     cdef object obj_shape
     cdef object obj_strides
 
+    if (itemsize < 1):
+        raise ValueError("dtype with typenum=" + str(typenum) + " is not supported.")
     if (nd < 0):
         raise ValueError("Dimensionality must be non-negative")
     if (ptr is NULL or QRef is NULL):
@@ -1881,7 +1885,7 @@ cdef api object UsmNDArray_MakeFromPtr(
         raise ValueError(
             "Given shape, strides and offset reference out-of-bound memory"
         )
-    nbytes = itemsize * (offset + max_disp + 1)
+    nbytes = (<size_t> itemsize) * (offset + max_disp + 1)
     mobj = c_dpmem._Memory.create_from_usm_pointer_size_qref(
         ptr, nbytes, QRef, memory_owner=owner
     )

dpctl/tensor/libtensor/source/accumulators.cpp

@@ -210,7 +210,7 @@ std::size_t py_mask_positions(const dpctl::tensor::usm_ndarray &mask,
             sycl::event::wait(host_task_events);
 
             // ensure deleter of smart pointer is invoked with GIL released
-            shape_strides_owner.release();
+            shape_strides_owner.reset(nullptr);
         }
         throw std::runtime_error("Unexpected error");
     }
@@ -231,7 +231,7 @@ std::size_t py_mask_positions(const dpctl::tensor::usm_ndarray &mask,
 
         sycl::event::wait(host_task_events);
         // ensure deleter of smart pointer is invoked with GIL released
-        shape_strides_owner.release();
+        shape_strides_owner.reset(nullptr);
     }
 
     return total_set;
@@ -367,7 +367,7 @@ std::size_t py_cumsum_1d(const dpctl::tensor::usm_ndarray &src,
             sycl::event::wait(host_task_events);
 
             // ensure USM deleter is called with GIL released
-            shape_strides_owner.release();
+            shape_strides_owner.reset(nullptr);
         }
         throw std::runtime_error("Unexpected error");
     }
@@ -387,7 +387,7 @@ std::size_t py_cumsum_1d(const dpctl::tensor::usm_ndarray &src,
         sycl::event::wait(host_task_events);
 
         // ensure USM deleter is called with GIL released
-        shape_strides_owner.release();
+        shape_strides_owner.reset(nullptr);
     }
 
     return total;

dpctl/tensor/libtensor/source/copy_numpy_ndarray_into_usm_ndarray.cpp

@@ -325,7 +325,7 @@ void copy_numpy_ndarray_into_usm_ndarray(
             dst_offset, depends, {copy_shape_ev});
 
         // invoke USM deleter in smart pointer while GIL is held
-        shape_strides_owner.release();
+        shape_strides_owner.reset(nullptr);
     }
 
     return;

dpctl/tests/test_usm_ndarray_ctor.py

@@ -963,6 +963,75 @@ def test_pyx_capi_make_general():
     assert zd_arr._pointer == mat._pointer
 
 
+def test_pyx_capi_make_fns_invalid_typenum():
+    q = get_queue_or_skip()
+    usm_ndarray = dpt.empty(tuple(), dtype="i4", sycl_queue=q)
+
+    make_simple_from_ptr = _pyx_capi_fnptr_to_callable(
+        usm_ndarray,
+        "UsmNDArray_MakeSimpleFromPtr",
+        b"PyObject *(size_t, int, DPCTLSyclUSMRef, "
+        b"DPCTLSyclQueueRef, PyObject *)",
+        fn_restype=ctypes.py_object,
+        fn_argtypes=(
+            ctypes.c_size_t,
+            ctypes.c_int,
+            ctypes.c_void_p,
+            ctypes.c_void_p,
+            ctypes.py_object,
+        ),
+    )
+
+    nelems = 10
+    dtype = dpt.int64
+    arr = dpt.arange(nelems, dtype=dtype, sycl_queue=q)
+
+    with pytest.raises(ValueError):
+        make_simple_from_ptr(
+            ctypes.c_size_t(nelems),
+            -1,
+            arr._pointer,
+            arr.sycl_queue.addressof_ref(),
+            arr,
+        )
+
+    make_from_ptr = _pyx_capi_fnptr_to_callable(
+        usm_ndarray,
+        "UsmNDArray_MakeFromPtr",
+        b"PyObject *(int, Py_ssize_t const *, int, Py_ssize_t const *, "
+        b"DPCTLSyclUSMRef, DPCTLSyclQueueRef, Py_ssize_t, PyObject *)",
+        fn_restype=ctypes.py_object,
+        fn_argtypes=(
+            ctypes.c_int,
+            ctypes.POINTER(ctypes.c_ssize_t),
+            ctypes.c_int,
+            ctypes.POINTER(ctypes.c_ssize_t),
+            ctypes.c_void_p,
+            ctypes.c_void_p,
+            ctypes.c_ssize_t,
+            ctypes.py_object,
+        ),
+    )
+    c_shape = (ctypes.c_ssize_t * 1)(
+        nelems,
+    )
+    c_strides = (ctypes.c_ssize_t * 1)(
+        1,
+    )
+    with pytest.raises(ValueError):
+        make_from_ptr(
+            ctypes.c_int(1),
+            c_shape,
+            -1,
+            c_strides,
+            arr._pointer,
+            arr.sycl_queue.addressof_ref(),
+            ctypes.c_ssize_t(0),
+            arr,
+        )
+    del arr
+
+
 def _pyx_capi_int(X, pyx_capi_name, caps_name=b"int", val_restype=ctypes.c_int):
     import sys
 

libsyclinterface/tests/test_sycl_device_interface.cpp

@@ -208,7 +208,7 @@ TEST_P(TestDPCTLSyclDeviceInterface, ChkGetSubGroupSizes)
     else
         EXPECT_TRUE(sg_sizes_len > 0);
     for (size_t i = 0; i < sg_sizes_len; ++i) {
-        EXPECT_TRUE(sg_sizes > 0);
+        EXPECT_TRUE(sg_sizes[i] > 0);
     }
     EXPECT_NO_FATAL_FAILURE(DPCTLSize_t_Array_Delete(sg_sizes));
 }

libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp

@@ -37,13 +37,16 @@
 
 #include <filesystem>
 #include <fstream>
+#include <utility>
+
 #include <gtest/gtest.h>
 #include <sycl/sycl.hpp>
-#include <utility>
 
 namespace
 {
-constexpr std::size_t SIZE = 100;
+constexpr std::size_t SIZE = 320;
+
+static_assert(SIZE % 10 == 0);
 
 using namespace dpctl::syclinterface;
 
@@ -69,11 +72,13 @@ void submit_kernel(DPCTLSyclQueueRef QRef,
         a_ptr[i] = 0;
     }
 
-    auto la1 = MDLocalAccessor{1, kernelArgTy, SIZE / 10, 1, 1};
+    std::size_t lws = SIZE / 10;
+
+    auto la1 = MDLocalAccessor{1, kernelArgTy, lws, 1, 1};
 
     // Create kernel args for vector_add
     std::size_t gRange[] = {SIZE};
-    std::size_t lRange[] = {SIZE / 10};
+    std::size_t lRange[] = {lws};
     void *args_1d[NARGS] = {unwrap<void>(a), (void *)&la1};
     DPCTLKernelArgType addKernelArgTypes[] = {DPCTL_VOID_PTR,
                                               DPCTL_LOCAL_ACCESSOR};
@@ -84,7 +89,7 @@ void submit_kernel(DPCTLSyclQueueRef QRef,
     ASSERT_TRUE(E1Ref != nullptr);
 
     DPCTLSyclEventRef DepEv1[] = {E1Ref};
-    auto la2 = MDLocalAccessor{2, kernelArgTy, SIZE / 10, 1, 1};
+    auto la2 = MDLocalAccessor{2, kernelArgTy, lws, 1, 1};
     void *args_2d[NARGS] = {unwrap<void>(a), (void *)&la2};
 
     DPCTLSyclEventRef E2Ref =
@@ -93,7 +98,7 @@ void submit_kernel(DPCTLSyclQueueRef QRef,
     ASSERT_TRUE(E2Ref != nullptr);
 
     DPCTLSyclEventRef DepEv2[] = {E1Ref, E2Ref};
-    auto la3 = MDLocalAccessor{3, kernelArgTy, SIZE / 10, 1, 1};
+    auto la3 = MDLocalAccessor{3, kernelArgTy, lws, 1, 1};
     void *args_3d[NARGS] = {unwrap<void>(a), (void *)&la3};
 
     DPCTLSyclEventRef E3Ref =
@@ -103,10 +108,7 @@ void submit_kernel(DPCTLSyclQueueRef QRef,
 
     DPCTLEvent_Wait(E3Ref);
 
-    if (kernelArgTy != DPCTL_FLOAT32_T && kernelArgTy != DPCTL_FLOAT64_T)
-        ASSERT_TRUE(a_ptr[0] == 20);
-    else
-        ASSERT_TRUE(a_ptr[0] == 20.0);
+    ASSERT_TRUE(a_ptr[0] == T(lws * 2));
 
     // clean ups
     DPCTLEvent_Delete(E1Ref);
@@ -133,9 +135,9 @@ void submit_kernel(DPCTLSyclQueueRef QRef,
 // figure which SPV file contains the kernels, use `spirv-dis` from the
 // spirv-tools package to translate the SPV binary format to a human-readable
 // textual format.
-#include <CL/sycl.hpp>
 #include <iostream>
 #include <sstream>
+#include <sycl/sycl.hpp>
 
 template <typename T>
 class SyclKernel_SLM
@@ -195,7 +197,7 @@ int main(int argc, const char **argv)
     driver<int8_t>(N);
     driver<uint8_t>(N);
     driver<int16_t>(N);
-    driver<int32_t>(N);
+    driver<uint16_t>(N);
     driver<int32_t>(N);
     driver<uint32_t>(N);
     driver<int64_t>(N);
@@ -220,11 +222,10 @@ struct TestQueueSubmitWithLocalAccessor : public ::testing::Test
     {
         DPCTLSyclDeviceSelectorRef DSRef = nullptr;
         DPCTLSyclDeviceRef DRef = nullptr;
+        const char *test_spv_fn = "./local_accessor_kernel_inttys_fp32.spv";
 
-        spirvFile.open("./local_accessor_kernel_inttys_fp32.spv",
-                       std::ios::binary | std::ios::ate);
-        spirvFileSize_ = std::filesystem::file_size(
-            "./local_accessor_kernel_inttys_fp32.spv");
+        spirvFile.open(test_spv_fn, std::ios::binary | std::ios::ate);
+        spirvFileSize_ = std::filesystem::file_size(test_spv_fn);
         spirvBuffer_.reserve(spirvFileSize_);
         spirvFile.seekg(0, std::ios::beg);
         spirvFile.read(spirvBuffer_.data(), spirvFileSize_);
@@ -261,11 +262,10 @@ struct TestQueueSubmitWithLocalAccessorFP64 : public ::testing::Test
     TestQueueSubmitWithLocalAccessorFP64()
     {
         DPCTLSyclDeviceSelectorRef DSRef = nullptr;
+        const char *test_spv_fn = "./local_accessor_kernel_fp64.spv";
 
-        spirvFile.open("./local_accessor_kernel_fp64.spv",
-                       std::ios::binary | std::ios::ate);
-        spirvFileSize_ =
-            std::filesystem::file_size("./local_accessor_kernel_fp64.spv");
+        spirvFile.open(test_spv_fn, std::ios::binary | std::ios::ate);
+        spirvFileSize_ = std::filesystem::file_size(test_spv_fn);
         spirvBuffer_.reserve(spirvFileSize_);
         spirvFile.seekg(0, std::ios::beg);
         spirvFile.read(spirvBuffer_.data(), spirvFileSize_);

pyproject.toml

@@ -27,6 +27,7 @@ classifiers = [
   "Programming Language :: Python :: 3.9",
   "Programming Language :: Python :: 3.10",
   "Programming Language :: Python :: 3.11",
+  "Programming Language :: Python :: 3.12",
   "Programming Language :: Python :: Implementation :: CPython",
   "Topic :: Software Development",
   "Topic :: Scientific/Engineering",