You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
Copy file name to clipboardExpand all lines: content/en/docs/_index.md
+3-1Lines changed: 3 additions & 1 deletion
Original file line number
Diff line number
Diff line change
@@ -10,4 +10,6 @@ by [Nikita Grigorian](https://github.com/ndgrigorian) and [Oleksandr Pavlyk](htt
10
10
11
11
This poster is intended to introduce writing portable data-parallel Python extensions using oneAPI.
12
12
13
-
We present several examples, starting with the basics of initializing a USM (universal shared memory) array, then a KDE (kernel density estimation) with pure DPC++/Sycl, then a KDE Python extension, and finally how to write a portable Python extension which uses oneMKL.
13
+
We present several examples, starting with the basics of initializing a USM (unified shared memory) array, then a KDE (kernel density estimation) with pure DPC++/Sycl, then a KDE Python extension, and finally how to write a portable Python extension which uses oneMKL.
14
+
15
+
The examples can be found [here](https://github.com/IntelPython/example-portable-data-parallel-extensions).
Use of temporary allocation can be avoided if each work-item atomically adds the value of the local sum to the
67
67
appropriate zero-initialized location in the output array, as in implementation ``kernel_density_estimation_atomic_ref``
@@ -119,10 +119,10 @@ in the work-group without accessing the global memory. This could be done effici
119
119
```
120
120
121
121
Complete implementation can be found in ``kernel_density_estimation_work_group_reduce_and_atomic_ref`` function
122
-
in ``"steps/kernel_density_estimation_cpp/kde.hpp"``.
122
+
in [``"steps/kernel_density_estimation_cpp/kde.hpp"``](https://github.com/IntelPython/example-portable-data-parallel-extensions/blob/main/steps/kernel_density_estimation_cpp/kde.hpp).
123
123
124
-
These implementations are called from C++ application ``"steps/kernel_density_estimation_cpp/app.cpp"``, which
124
+
These implementations are called from C++ application [``"steps/kernel_density_estimation_cpp/app.cpp"``](https://github.com/IntelPython/example-portable-data-parallel-extensions/blob/main/steps/kernel_density_estimation_cpp/app.cpp), which
125
125
samples data uniformly distributed over unit cuboid, and estimates the density using Kernel Density Estimation
126
126
and spherically symmetric multivariate Gaussian probability density function as the kernel.
127
127
128
-
The application can be built using `CMake`, or `Meson`, please refer to [README](steps/kernel_density_estimation_cpp/README.md) document in that folder.
128
+
The application can be built using `CMake`, or `Meson`, please refer to [README](https://github.com/IntelPython/example-portable-data-parallel-extensions/blob/main/steps/kernel_density_estimation_cpp/README.md) document in that folder.
Copy file name to clipboardExpand all lines: content/en/docs/kde-python.md
+9-7Lines changed: 9 additions & 7 deletions
Original file line number
Diff line number
Diff line change
@@ -5,7 +5,7 @@ date: 2024-07-02
5
5
weight: 3
6
6
---
7
7
8
-
Since SYCL builds on C++, we are going to use `pybind11` project to generate Python extension.
8
+
Since SYCL builds on C++, we are going to use the `pybind11` project to generate a Python extension.
9
9
We also need Python objects to carry USM allocations of input and output data, such as `dpctl` ([Data Parallel Control](https://github.com/IntelPython/dpctl.git) Python package). The `dpctl` package also provides Python objects corresponding to DPC++ runtime objects:
10
10
11
11
| Python object | SYCL C++ object |
@@ -15,9 +15,9 @@ We also need Python objects to carry USM allocations of input and output data, s
15
15
|``dpctl.SyclContext``|``sycl::context``|
16
16
|``dpctl.SyclEvent``|``sycl::event``|
17
17
18
-
`dpctl` provides integration with `pybind11` supporting castings between `dpctl` Python objects and corresponding C++ SYCL classes listed in the table above. Furthermore, the integration provides C++ class ``dpctl::tensor::usm_ndarray`` which derives from ``pybind11::object``.
19
-
It stores `dpctl.tensor.usm_ndarray` object and provides methods to query its attributes, such as data pointer, dimensionality, shape, strides
20
-
and elemental type information.
18
+
`dpctl` provides integration with `pybind11` supporting castings between `dpctl` Python objects and corresponding C++ SYCL classes listed in the table above. Furthermore, the integration provides the C++ class ``dpctl::tensor::usm_ndarray`` which derives from ``pybind11::object``.
19
+
It stores the `dpctl.tensor.usm_ndarray` object and provides methods to query its attributes, such as data pointer, dimensionality, shape, strides
20
+
and elemental type information. Underlying `dpctl.tensor.usm_ndarray` is a SYCL unified shared memory (USM) allocation. See the [SYCL standard](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:usm) or [dpctl.memory documentation](https://intelpython.github.io/dpctl/latest/api_reference/dpctl/memory.html#dpctl-memory-pyapi) for more details.
21
21
22
22
For illustration purpose, here is a sample extension source code:
23
23
@@ -29,7 +29,9 @@ For illustration purpose, here is a sample extension source code:
// Execution queue is the queue associated with input arrays
@@ -98,12 +100,12 @@ of the host task a chance at execution.
98
100
Of course, if USM memory is not managed by Python, it may be possible to avoid using GIL altogether.
99
101
100
102
An example of Python extension `"kde_sycl_ext"` that exposes kernel density estimation code from previous
101
-
section can be found in `"steps/sycl_python_extension"` folder (see [README](steps/sycl_python_extension/README.md)).
103
+
section can be found in [`"steps/sycl_python_extension"`](https://github.com/IntelPython/example-portable-data-parallel-extensions/tree/main/steps/sycl_python_extension) folder (see [README](https://github.com/IntelPython/example-portable-data-parallel-extensions/blob/main/steps/sycl_python_extension/README.md)).
102
104
103
105
The folder contains comparison between `dpctl`-based implementation of the KDE implementation following the NumPy
104
106
implementation [above](#kde_numpy) and the dedicated C++ code:
105
107
106
-
```
108
+
```bash
107
109
KDE for n_sample = 1000000, n_est = 17, n_dim = 7, h = 0.05
Copy file name to clipboardExpand all lines: content/en/docs/oneMKL.md
+32-4Lines changed: 32 additions & 4 deletions
Original file line number
Diff line number
Diff line change
@@ -5,9 +5,18 @@ date: 2024-07-02
5
5
weight: 4
6
6
---
7
7
8
-
Since `dpctl.tensor.usm_ndarray` is a Python object carrying a USM allocation, it is possible to write extensions which wrap `oneAPI Math Kernel Library Interfaces` ([oneMKL Interfaces](https://github.com/oneapi-src/oneMKL)) routines and then call them on the USM data underlying the `usm_ndarray` container from Python.
8
+
Given a matrix \\(A\\), the QR decomposition of \\(A\\) is defined as the decomposition of \\(A\\) into the product of matrices \\(Q\\) and \\(R\\) such that \\(Q\\) is orthonormal and \\(R\\) is upper-triangular.
9
+
10
+
QR factorization is a common routine in more optimized LAPACK libraries, so rather than write and implement an algorithm ourselves, it would be preferable to find a suitable library routine.
11
+
12
+
Since `dpctl.tensor.usm_ndarray` is a Python object with an underlying USM allocation, it is possible to write extensions which wrap `oneAPI Math Kernel Library Interfaces` ([oneMKL Interfaces](https://github.com/oneapi-src/oneMKL)) USM routines and then call them on the `dpctl.tensor.usm_ndarray` from Python. These low-level routines can greatly improve the performance of an extension.
13
+
14
+
`oneMKL Interfaces` can be built to dispatch to a variety of backends including `cuBLAS` and `rocBLAS` (see [oneMKL Interfaces README](https://github.com/oneapi-src/oneMKL?tab=readme-ov-file#oneapi-math-kernel-library-onemkl-interfaces)). The [`portBLAS`](https://github.com/codeplaysoftware/portBLAS) backend is also notable as it is open-source and written in pure SYCL.
15
+
16
+
`oneMKL` routines are essentially wrappers for the same routine in an underlying backend library, depending on the targeted device. This means that the same code can be used for NVidia, AMD, and Intel devices, making it highly portable.
17
+
18
+
Looking to the `oneMKL` documentation on [`geqrf`](https://spec.oneapi.io/versions/latest/elements/oneMKL/source/domains/lapack/geqrf.html#geqrf-usm-version):
9
19
10
-
For an example routine from the `oneMKL` documentation, take [`geqrf`](https://spec.oneapi.io/versions/latest/elements/oneMKL/source/domains/lapack/geqrf.html#geqrf-usm-version):
The `pybind11` castings discussed in the previous section enable us to write a simple wrapper function for this routine with `dpctl::tensor::usm_ndarray` inputs and outputs, so long as we take the same precautions to avoid deadlocks. As a result, we can write the extension in much the same way as the `kde_sycl_ext` extension in the previous chapter.
34
+
This general format (``sycl::queue``, arguments, and a vector of ``sycl::event``s) is more or less the same throughout the `oneMKL` USM routines.
35
+
36
+
The `pybind11` castings discussed in the previous section enable us to write a simple wrapper function for this routine with ``dpctl::tensor::usm_ndarray`` inputs and outputs, so long as we take the same precautions to avoid deadlocks. As a result, we can write the extension in much the same way as the `"kde_sycl_ext"` extension in the previous chapter.
26
37
27
-
An example of a Python extension "mkl_interface_ext" that uses `oneMKL` calls to implement a QR decomposition can be found in "steps/mkl_interface" folder (see [README](steps/mkl_interface/README.md)).
38
+
An example of a Python extension `"mkl_interface_ext"` that uses `oneMKL` calls to implement a QR decomposition can be found in [`"steps/mkl_interface"`](https://github.com/IntelPython/example-portable-data-parallel-extensions/tree/main/steps/mkl_interface) folder (see [README](https://github.com/IntelPython/example-portable-data-parallel-extensions/blob/main/steps/mkl_interface/README.md)).
39
+
40
+
The folder executes the tests found in [`"steps/mkl_interface/tests"`](https://github.com/IntelPython/example-portable-data-parallel-extensions/tree/main/steps/mkl_interface/tests) as well as running a larger benchmark which compares Numpy's `linalg.qr` (for reference) to the extension's implementation:
41
+
42
+
```bash
43
+
$ python run.py
44
+
Using device NVIDIA GeForce GT 1030
45
+
================================================= test session starts ==================================================
46
+
collected 8 items
47
+
48
+
tests/test_qr.py ........ [100%]
49
+
50
+
================================================== 8 passed in 0.45s ===================================================
51
+
QR decomposition for matrix of size = (3000, 3000)
0 commit comments