Skip to content

Commit 8464d9b

Browse files
AlexanderKalistratovantonwolfy
AlexanderKalistratov
and
antonwolfy
authored
Implementation of histogram with sycl kernel (#2027)
Implementation of histogram with sycl kernel --------- Co-authored-by: Anton <[email protected]>
1 parent 4816161 commit 8464d9b

13 files changed

+1785
-47
lines changed

‎.pre-commit-config.yaml

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -52,7 +52,6 @@ repos:
5252
rev: 24.4.2
5353
hooks:
5454
- id: black
55-
args: ["--check", "--diff", "--color"]
5655
- repo: https://github.com/pycqa/isort
5756
rev: 5.13.2
5857
hooks:

‎dpnp/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -58,6 +58,7 @@ add_subdirectory(backend/extensions/fft)
5858
add_subdirectory(backend/extensions/lapack)
5959
add_subdirectory(backend/extensions/vm)
6060
add_subdirectory(backend/extensions/ufunc)
61+
add_subdirectory(backend/extensions/statistics)
6162

6263
add_subdirectory(dpnp_algo)
6364
add_subdirectory(dpnp_utils)

‎dpnp/backend/extensions/statistics/CMakeLists.txt

Lines changed: 88 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,88 @@
1+
# *****************************************************************************
2+
# Copyright (c) 2016-2024, Intel Corporation
3+
# All rights reserved.
4+
#
5+
# Redistribution and use in source and binary forms, with or without
6+
# modification, are permitted provided that the following conditions are met:
7+
# - Redistributions of source code must retain the above copyright notice,
8+
# this list of conditions and the following disclaimer.
9+
# - Redistributions in binary form must reproduce the above copyright notice,
10+
# this list of conditions and the following disclaimer in the documentation
11+
# and/or other materials provided with the distribution.
12+
#
13+
# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
14+
# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
15+
# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
16+
# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
17+
# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18+
# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19+
# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20+
# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21+
# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22+
# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23+
# THE POSSIBILITY OF SUCH DAMAGE.
24+
# *****************************************************************************
25+
26+
27+
set(python_module_name _statistics_impl)
28+
set(_module_src
29+
${CMAKE_CURRENT_SOURCE_DIR}/common.cpp
30+
${CMAKE_CURRENT_SOURCE_DIR}/histogram.cpp
31+
${CMAKE_CURRENT_SOURCE_DIR}/histogram_common.cpp
32+
${CMAKE_CURRENT_SOURCE_DIR}/statistics_py.cpp
33+
)
34+
35+
pybind11_add_module(${python_module_name} MODULE ${_module_src})
36+
add_sycl_to_target(TARGET ${python_module_name} SOURCES ${_module_src})
37+
38+
if(_dpnp_sycl_targets)
39+
# make fat binary
40+
target_compile_options(
41+
${python_module_name}
42+
PRIVATE
43+
-fsycl-targets=${_dpnp_sycl_targets}
44+
)
45+
target_link_options(
46+
${python_module_name}
47+
PRIVATE
48+
-fsycl-targets=${_dpnp_sycl_targets}
49+
)
50+
endif()
51+
52+
if (WIN32)
53+
if (${CMAKE_VERSION} VERSION_LESS "3.27")
54+
# this is a work-around for target_link_options inserting option after -link option, cause
55+
# linker to ignore it.
56+
set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} -fsycl-device-code-split=per_kernel")
57+
endif()
58+
endif()
59+
60+
set_target_properties(${python_module_name} PROPERTIES CMAKE_POSITION_INDEPENDENT_CODE ON)
61+
62+
target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../include)
63+
target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../src)
64+
65+
target_include_directories(${python_module_name} PUBLIC ${Dpctl_INCLUDE_DIR})
66+
target_include_directories(${python_module_name} PUBLIC ${Dpctl_TENSOR_INCLUDE_DIR})
67+
68+
if (WIN32)
69+
target_compile_options(${python_module_name} PRIVATE
70+
/clang:-fno-approx-func
71+
/clang:-fno-finite-math-only
72+
)
73+
else()
74+
target_compile_options(${python_module_name} PRIVATE
75+
-fno-approx-func
76+
-fno-finite-math-only
77+
)
78+
endif()
79+
80+
target_link_options(${python_module_name} PUBLIC -fsycl-device-code-split=per_kernel)
81+
82+
if (DPNP_GENERATE_COVERAGE)
83+
target_link_options(${python_module_name} PRIVATE -fprofile-instr-generate -fcoverage-mapping)
84+
endif()
85+
86+
install(TARGETS ${python_module_name}
87+
DESTINATION "dpnp/backend/extensions/statistics"
88+
)

‎dpnp/backend/extensions/statistics/common.cpp

Lines changed: 124 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,124 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2024, Intel Corporation
3+
// All rights reserved.
4+
//
5+
// Redistribution and use in source and binary forms, with or without
6+
// modification, are permitted provided that the following conditions are met:
7+
// - Redistributions of source code must retain the above copyright notice,
8+
// this list of conditions and the following disclaimer.
9+
// - Redistributions in binary form must reproduce the above copyright notice,
10+
// this list of conditions and the following disclaimer in the documentation
11+
// and/or other materials provided with the distribution.
12+
//
13+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
14+
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
15+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
16+
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
17+
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18+
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19+
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20+
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21+
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22+
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23+
// THE POSSIBILITY OF SUCH DAMAGE.
24+
//*****************************************************************************
25+
26+
#include "common.hpp"
27+
#include "utils/type_dispatch.hpp"
28+
#include <pybind11/pybind11.h>
29+
30+
namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
31+
32+
namespace statistics
33+
{
34+
namespace common
35+
{
36+
37+
size_t get_max_local_size(const sycl::device &device)
38+
{
39+
constexpr const int default_max_cpu_local_size = 256;
40+
constexpr const int default_max_gpu_local_size = 0;
41+
42+
return get_max_local_size(device, default_max_cpu_local_size,
43+
default_max_gpu_local_size);
44+
}
45+
46+
size_t get_max_local_size(const sycl::device &device,
47+
int cpu_local_size_limit,
48+
int gpu_local_size_limit)
49+
{
50+
int max_work_group_size =
51+
device.get_info<sycl::info::device::max_work_group_size>();
52+
if (device.is_cpu() && cpu_local_size_limit > 0) {
53+
return std::min(cpu_local_size_limit, max_work_group_size);
54+
}
55+
else if (device.is_gpu() && gpu_local_size_limit > 0) {
56+
return std::min(gpu_local_size_limit, max_work_group_size);
57+
}
58+
59+
return max_work_group_size;
60+
}
61+
62+
sycl::nd_range<1>
63+
make_ndrange(size_t global_size, size_t local_range, size_t work_per_item)
64+
{
65+
return make_ndrange(sycl::range<1>(global_size),
66+
sycl::range<1>(local_range),
67+
sycl::range<1>(work_per_item));
68+
}
69+
70+
size_t get_local_mem_size_in_bytes(const sycl::device &device)
71+
{
72+
// Reserving 1kb for runtime needs
73+
constexpr const size_t reserve = 1024;
74+
75+
return get_local_mem_size_in_bytes(device, reserve);
76+
}
77+
78+
size_t get_local_mem_size_in_bytes(const sycl::device &device, size_t reserve)
79+
{
80+
size_t local_mem_size =
81+
device.get_info<sycl::info::device::local_mem_size>();
82+
return local_mem_size - reserve;
83+
}
84+
85+
pybind11::dtype dtype_from_typenum(int dst_typenum)
86+
{
87+
dpctl_td_ns::typenum_t dst_typenum_t =
88+
static_cast<dpctl_td_ns::typenum_t>(dst_typenum);
89+
switch (dst_typenum_t) {
90+
case dpctl_td_ns::typenum_t::BOOL:
91+
return py::dtype("?");
92+
case dpctl_td_ns::typenum_t::INT8:
93+
return py::dtype("i1");
94+
case dpctl_td_ns::typenum_t::UINT8:
95+
return py::dtype("u1");
96+
case dpctl_td_ns::typenum_t::INT16:
97+
return py::dtype("i2");
98+
case dpctl_td_ns::typenum_t::UINT16:
99+
return py::dtype("u2");
100+
case dpctl_td_ns::typenum_t::INT32:
101+
return py::dtype("i4");
102+
case dpctl_td_ns::typenum_t::UINT32:
103+
return py::dtype("u4");
104+
case dpctl_td_ns::typenum_t::INT64:
105+
return py::dtype("i8");
106+
case dpctl_td_ns::typenum_t::UINT64:
107+
return py::dtype("u8");
108+
case dpctl_td_ns::typenum_t::HALF:
109+
return py::dtype("f2");
110+
case dpctl_td_ns::typenum_t::FLOAT:
111+
return py::dtype("f4");
112+
case dpctl_td_ns::typenum_t::DOUBLE:
113+
return py::dtype("f8");
114+
case dpctl_td_ns::typenum_t::CFLOAT:
115+
return py::dtype("c8");
116+
case dpctl_td_ns::typenum_t::CDOUBLE:
117+
return py::dtype("c16");
118+
default:
119+
throw py::value_error("Unrecognized dst_typeid");
120+
}
121+
}
122+
123+
} // namespace common
124+
} // namespace statistics

0 commit comments

Comments
 (0)