Review comments fixes

Author: AlexanderKalistratov

.pre-commit-config.yaml

@@ -52,7 +52,7 @@ repos:
     rev: 24.4.2
     hooks:
     -   id: black
-        args: ["--check", "--diff", "--color"]
+        args: ["--diff", "--color"]
 -   repo: https://github.com/pycqa/isort
     rev: 5.13.2
     hooks:

dpnp/backend/extensions/sycl_ext/CMakeLists.txt

@@ -34,6 +34,20 @@ set(_module_src
 pybind11_add_module(${python_module_name} MODULE ${_module_src})
 add_sycl_to_target(TARGET ${python_module_name} SOURCES ${_module_src})
 
+if(_dpnp_sycl_targets)
+    # make fat binary
+    target_compile_options(
+        ${python_module_name}
+        PRIVATE
+        -fsycl-targets=${_dpnp_sycl_targets}
+    )
+    target_link_options(
+        ${python_module_name}
+        PRIVATE
+        -fsycl-targets=${_dpnp_sycl_targets}
+    )
+endif()
+
 if (WIN32)
     if (${CMAKE_VERSION} VERSION_LESS "3.27")
         # this is a work-around for target_link_options inserting option after -link option, cause

dpnp/backend/extensions/sycl_ext/histogram.cpp

@@ -50,7 +50,7 @@ namespace
 {
 
 template <typename T, typename BinsT, typename HistType = size_t>
-static sycl::event histogram_impl(sycl::queue exec_q,
+static sycl::event histogram_impl(sycl::queue &exec_q,
                                   const void *vin,
                                   const void *vbins_edges,
                                   const void *vweights,
@@ -71,15 +71,15 @@ static sycl::event histogram_impl(sycl::queue exec_q,
             ? 256
             : device.get_info<sycl::info::device::max_work_group_size>();
 
-    uint32_t WorkPI = 128; // empirically found number
+    constexpr uint32_t WorkPI = 128; // empirically found number
     auto global_size = Align(CeilDiv(size, WorkPI), local_size);
 
     auto nd_range =
         sycl::nd_range(sycl::range<1>(global_size), sycl::range<1>(local_size));
 
     return exec_q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(depends);
-        uint32_t dims = 1;
+        constexpr uint32_t dims = 1;
 
         auto dispatch_edges = [&](uint32_t local_mem, auto &weights,
                                   auto &hist) {
@@ -239,7 +239,7 @@ std::tuple<sycl::event, sycl::event>
             exec_q, {sample, bins, histogram}, {ev});
     }
 
-    return {ev, args_ev};
+    return {args_ev, ev};
 }
 
 std::unique_ptr<Histogram> hist;

dpnp/backend/extensions/sycl_ext/histogram.hpp

@@ -28,15 +28,14 @@
 #include <sycl/sycl.hpp>
 
 namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
-// using dpctl::tensor::type_dispatch::num_types;
 
 namespace sycl_ext
 {
 namespace histogram
 {
 struct Histogram
 {
-    using FnT = sycl::event (*)(sycl::queue,
+    using FnT = sycl::event (*)(sycl::queue &,
                                 const void *,
                                 const void *,
                                 const void *,

dpnp/backend/extensions/sycl_ext/histogram_common.hpp

@@ -43,13 +43,13 @@ namespace histogram
 {
 
 template <typename N, typename D>
-N CeilDiv(N n, D d)
+constexpr auto CeilDiv(N n, D d)
 {
     return (n + d - 1) / d;
 }
 
 template <typename N, typename D>
-N Align(N n, D d)
+constexpr auto Align(N n, D d)
 {
     return CeilDiv(n, d) * d;
 }

dpnp/dpnp_iface_histograms.py

@@ -44,7 +44,12 @@
 import numpy
 
 import dpnp
-import dpnp.backend.extensions.sycl_ext._sycl_ext_impl as sycl_ext  # pylint: disable=C0301,E0611
+
+# pylint: disable=no-name-in-module
+import dpnp.backend.extensions.sycl_ext._sycl_ext_impl as sycl_ext
+
+# pylint: disable=no-name-in-module
+from .dpnp_utils import map_dtype_to_device
 
 __all__ = [
     "digitize",
@@ -318,19 +323,14 @@ def _find_supported_dtype(dt, supported):
     return None
 
 
-def _result_type(dtype1, dtype2, has_fp64):
+def _result_type_for_device(dtype1, dtype2, device):
     rt = dpnp.result_type(dtype1, dtype2)
-    if rt == dpnp.float64 and not has_fp64:
-        return dpnp.float32
-
-    if rt == dpnp.complex128 and not has_fp64:
-        return dpnp.complex64
+    return map_dtype_to_device(rt, device)
 
-    return rt
 
-
-def _align_dtypes(a_dtype, bins_dtype, ntype, has_fp64):
-    a_bin_dtype = _result_type(a_dtype, bins_dtype, has_fp64)
+def _align_dtypes(a_dtype, bins_dtype, ntype, device):
+    has_fp64 = device.has_aspect_fp64
+    a_bin_dtype = _result_type_for_device(a_dtype, bins_dtype, device)
 
     supported_types = (dpnp.float32, dpnp.int64, numpy.uint64, dpnp.complex64)
     if has_fp64:
@@ -347,7 +347,7 @@ def _align_dtypes(a_dtype, bins_dtype, ntype, has_fp64):
     if (a_bin_dtype in float_types and hist_dtype in float_types) or (
         a_bin_dtype in complex_types and hist_dtype in complex_types
     ):
-        common_type = _result_type(a_bin_dtype, hist_dtype, has_fp64)
+        common_type = _result_type_for_device(a_bin_dtype, hist_dtype, device)
         a_bin_dtype = common_type
         hist_dtype = common_type
 
@@ -463,10 +463,10 @@ def histogram(a, bins=10, range=None, density=None, weights=None):
         ntype = weights.dtype
 
     queue = a.sycl_queue
-    has_fp64 = queue.sycl_device.has_aspect_fp64
+    device = queue.sycl_device
 
     a_bin_dtype, hist_dtype = _align_dtypes(
-        a.dtype, bin_edges.dtype, ntype, has_fp64
+        a.dtype, bin_edges.dtype, ntype, device
     )
 
     if a_bin_dtype is None or hist_dtype is None:
@@ -487,7 +487,10 @@ def histogram(a, bins=10, range=None, density=None, weights=None):
         else None
     )
 
+    # histogram implementation uses atomics, but atomics doesn't work with
+    # host usm memory
     n_usm_type = "device" if usm_type == "host" else usm_type
+
     n_casted = dpnp.zeros(
         bin_edges.size - 1,
         dtype=hist_dtype,
@@ -521,7 +524,6 @@ def histogram(a, bins=10, range=None, density=None, weights=None):
         n = dpnp.astype(n_casted, ntype, copy=False)
 
     if density:
-        # pylint: disable=possibly-used-before-assignment
         db = dpnp.diff(bin_edges).astype(
             dpnp.default_float_type(sycl_queue=queue)
         )