Added an example to exercise pybind11 bindings for dpctl.program.SyclKernel

Author: oleksandr-pavlyk

examples/pybind11/use_dpctl_sycl_kernel/CMakeLists.txt

@@ -0,0 +1,36 @@
+cmake_minimum_required(VERSION 3.21)
+
+project(use_queue_device LANGUAGES CXX)
+
+set(DPCTL_CMAKE_MODULES_PATH "${CMAKE_SOURCE_DIR}/../../../cmake")
+set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${DPCTL_CMAKE_MODULES_PATH})
+find_package(IntelDPCPP REQUIRED PATHS ${DPCTL_CMAKE_MODULES_PATH} NO_DEFAULT_PATH)
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED True)
+set(CMAKE_BUILD_TYPE Debug)
+
+# Fetch pybind11
+include(FetchContent)
+FetchContent_Declare(
+  pybind11
+  URL https://github.com/pybind/pybind11/archive/refs/tags/v2.10.0.tar.gz
+  URL_HASH SHA256=eacf582fa8f696227988d08cfc46121770823839fe9e301a20fbce67e7cd70ec
+)
+FetchContent_MakeAvailable(pybind11)
+
+find_package(PythonExtensions REQUIRED)
+find_package(Dpctl REQUIRED)
+find_package(NumPy REQUIRED)
+
+set(py_module_name _use_kernel)
+pybind11_add_module(${py_module_name}
+    MODULE
+    use_kernel/_example.cpp
+)
+target_include_directories(${py_module_name} PUBLIC ${Dpctl_INCLUDE_DIRS})
+install(TARGETS ${py_module_name}
+  DESTINATION use_kernel
+)
+
+set(ignoreMe "${SKBUILD}")

examples/pybind11/use_dpctl_sycl_kernel/README.md

@@ -0,0 +1,29 @@
+# Usage of dpctl Entities in Pybind11
+
+## Description
+
+This extension demonstrates how you can use dpctl Python types,
+such as ``dpctl.SyclQueue``, in Pybind11
+extensions.
+
+
+## Building
+
+To build the extension, run:
+```
+source /opt/intel/oneapi/compiler/latest/env/vars.sh
+CXX=icpx python setup.py build_ext --inplace
+python -m pytest tests
+python example.py
+```
+
+# Sample output
+
+```
+(idp) [17:25:27 ansatnuc04 use_dpctl_syclqueue]$ python example.py
+EU count returned by Pybind11 extension 24
+EU count computed by dpctl 24
+
+Computing modular reduction using SYCL on a NumPy array
+Offloaded result agrees with reference one computed by NumPy
+```

examples/pybind11/use_dpctl_sycl_kernel/example.py

@@ -0,0 +1,41 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2022 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+# coding: utf-8
+
+import use_kernel as eg
+
+import dpctl
+import dpctl.program as dppr
+import dpctl.tensor as dpt
+
+q = dpctl.SyclQueue()
+
+with open("resource/double_it.spv", "br") as fh:
+    il = fh.read()
+
+pr = dppr.create_program_from_spirv(q, il, "")
+assert pr.has_sycl_kernel("double_it")
+
+krn = pr.get_sycl_kernel("double_it")
+assert krn.num_args == 2
+
+x = dpt.arange(0, stop=13, step=1, dtype="i4", sycl_queue=q)
+y = dpt.empty_like(x)
+
+eg.submit_custom_kernel(q, krn, x, y)
+
+print(dpt.asnumpy(y))

examples/pybind11/use_dpctl_sycl_kernel/resource/README.md

@@ -0,0 +1,8 @@
+# Rebuilding SPIR-V file from source
+
+```bash
+export TOOLS_DIR=$(dirname $(dirname $(which icx)))/bin-llvm
+$TOOLS_DIR/clang -cc1 -triple spir double_it.cl -finclude-default-header -flto -emit -llvm-bc -o double_it.bc
+$TOOLS_DIR/llvm-spirv double_it.bc -o double_it.spv
+rm double_it.bc
+```

examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.cl

@@ -0,0 +1,30 @@
+/*
+                      Data Parallel Control (dpctl)
+
+   Copyright 2020-2022 Intel Corporation
+
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+
+      http://www.apache.org/licenses/LICENSE-2.0
+
+   Unless required by applicable law or agreed to in writing, software
+   distributed under the License is distributed on an "AS IS" BASIS,
+   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+   See the License for the specific language governing permissions and
+   limitations under the License.
+
+===----------------------------------------------------------------------===
+
+ \file
+ This file implements a sample OpenCL kernel to use in this example.
+
+===----------------------------------------------------------------------===
+*/
+
+__kernel void double_it(__global int *x, __global int *y) {
+    uint idx = get_global_id(0);
+
+    y[idx] = 2 * x[idx];
+}

examples/pybind11/use_dpctl_sycl_kernel/resource/double_it.spv

@@ -0,0 +1,26 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2022 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+from skbuild import setup
+
+setup(
+    name="use_kernel",
+    version="0.0.1",
+    description="an example of SYCL-powered Python package (with pybind11)",
+    author="Intel Scripting",
+    license="Apache 2.0",
+    packages=["use_kernel"],
+)

examples/pybind11/use_dpctl_sycl_kernel/setup.py

@@ -0,0 +1,69 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2022 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+# coding: utf-8
+
+import os.path
+
+import numpy as np
+import pytest
+import use_kernel as uk
+
+import dpctl
+import dpctl.program as dpm
+import dpctl.tensor as dpt
+
+
+def _get_spv_path():
+    uk_dir = os.path.dirname(os.path.abspath(uk.__file__))
+    proj_dir = os.path.dirname(uk_dir)
+    return os.path.join(proj_dir, "resource", "double_it.spv")
+
+
+def test_spv_file_exists():
+    assert os.path.exists(_get_spv_path())
+
+
+def test_kernel_can_be_found():
+    fn = _get_spv_path()
+    with open(fn, "br") as f:
+        il = f.read()
+    try:
+        q = dpctl.SyclQueue()
+    except dpctl.SyclQueueCreationError:
+        pytest.skip("Could not create default queue")
+    pr = dpm.create_program_from_spirv(q, il, "")
+    assert pr.has_sycl_kernel("double_it")
+
+
+def test_kernel_submit_through_extension():
+    fn = _get_spv_path()
+    with open(fn, "br") as f:
+        il = f.read()
+    try:
+        q = dpctl.SyclQueue()
+    except dpctl.SyclQueueCreationError:
+        pytest.skip("Could not create default queue")
+    pr = dpm.create_program_from_spirv(q, il, "")
+    krn = pr.get_sycl_kernel("double_it")
+    assert krn.num_args == 2
+
+    x = dpt.arange(0, stop=13, step=1, dtype="i4", sycl_queue=q)
+    y = dpt.zeros_like(x)
+
+    uk.submit_custom_kernel(q, krn, x, y, [])
+
+    assert np.array_equal(dpt.asnumpy(y), np.arange(0, 26, step=2, dtype="i4"))

examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py

@@ -0,0 +1,34 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2022 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+# coding: utf-8
+
+from ._use_kernel import submit_custom_kernel
+
+__all__ = [
+    "submit_custom_kernel",
+]
+
+__doc__ = """
+Example pybind11 extension demonstrating binding of dpctl entities to
+SYCL entities.
+
+dpctl provides type casters that bind ``sycl::kernel`` to
+`dpctl.program.SyclKernel`, ``sycl::device`` to `dpctl.SyclDevice`, etc.
+
+Use of these type casters simplifies writing of Python extensions and compile
+then using SYCL C++ compilers, such as Intel(R) oneAPI DPC++ compiler.
+"""

examples/pybind11/use_dpctl_sycl_kernel/use_kernel/__init__.py

@@ -0,0 +1,92 @@
+//==- _example.cpp - Example of Pybind11 extension working with  =---------===//
+//  dpctl Python objects.
+//
+//                      Data Parallel Control (dpctl)
+//
+// Copyright 2020-2022 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file implements Pybind11-generated extension exposing functions that
+/// take dpctl Python objects, such as dpctl.SyclQueue, dpctl.SyclDevice as
+/// arguments.
+///
+//===----------------------------------------------------------------------===//
+
+#include "dpctl4pybind11.hpp"
+#include <CL/sycl.hpp>
+#include <cstdint>
+#include <pybind11/pybind11.h>
+#include <pybind11/stl.h>
+#include <vector>
+
+namespace py = pybind11;
+
+void submit_custom_kernel(sycl::queue q,
+                          sycl::kernel krn,
+                          dpctl::tensor::usm_ndarray x,
+                          dpctl::tensor::usm_ndarray y,
+                          const std::vector<sycl::event> &depends = {})
+{
+    if (x.get_ndim() != 1 || !x.is_c_contiguous() || y.get_ndim() != 1 ||
+        !y.is_c_contiguous())
+    {
+        throw py::value_error(
+            "src and dst arguments must be 1D and contiguous.");
+    }
+
+    auto const &api = dpctl::detail::dpctl_capi::get();
+    if (x.get_typenum() != api.UAR_INT32_ || y.get_typenum() != api.UAR_INT32_)
+    {
+        throw py::value_error(
+            "src and dst arguments must have int32 element data types.");
+    }
+
+    size_t n_x = x.get_size();
+    size_t n_y = y.get_size();
+
+    if (n_x != n_y) {
+        throw py::value_error("src and dst arguments must have equal size.");
+    }
+
+    if (!dpctl::utils::queues_are_compatible(q, {x.get_queue(), y.get_queue()}))
+    {
+        throw std::runtime_error(
+            "Execution queue is not compatible with allocation queues");
+    }
+
+    void *x_data = x.get_data<void>();
+    void *y_data = y.get_data<void>();
+
+    sycl::event e = q.submit([&](sycl::handler &cgh) {
+        cgh.depends_on(depends);
+        cgh.set_arg(0, x_data);
+        cgh.set_arg(1, y_data);
+        cgh.parallel_for(sycl::range<1>(n_x), krn);
+    });
+
+    e.wait();
+
+    return;
+}
+
+PYBIND11_MODULE(_use_kernel, m)
+{
+    m.def("submit_custom_kernel", &submit_custom_kernel,
+          "Submit given kernel with arguments (int *, int *) to queue",
+          py::arg("queue"), py::arg("kernel"), py::arg("src"), py::arg("dst"),
+          py::arg("depends") = py::list());
+}