Update build system/tests for extension to run on Iris Xe

Author: oleksandr-pavlyk

common_ext/cy_kde.pyx

@@ -57,7 +57,7 @@ def kde_eval(
         raise ValueError("Evaluation and observation data must be non-empty")
 
     if cython.floating is float:
-        f = np.empty(x.shape[0], dtype=np.float)
+        f = np.empty(x.shape[0], dtype=np.single)
     else:
         f = np.empty(x.shape[0], dtype=np.double)
 

common_src/kde.hpp

@@ -16,6 +16,7 @@
 
 #include <CL/sycl.hpp>
 #include <cstdint>
+#include <iostream>
 
 namespace example {
 
@@ -66,14 +67,20 @@ kernel_density_estimate(
     sycl::buffer<T, 1> buf_F(f, sycl::range<1>(n), buf_props);
 
     // initialize buffer of function values with zeros
+    try {
     q.submit(
 	[&](sycl::handler &cgh) {
 	    sycl::accessor acc_F(buf_F, cgh, sycl::write_only, sycl::no_init);
 	    cgh.fill(acc_F, T(0));
 	});
+    } catch (const std::exception &e){
+      std::cout << e.what() << std::endl;
+      std::rethrow_exception(std::current_exception());
+    }
 
     // populate function values buffer
     //   perform 2D loop in parallel
+    try{
     q.submit(
 	[&](sycl::handler &cgh) {
 	    sycl::accessor acc_X(buf_X, cgh, sycl::read_only);
@@ -135,6 +142,10 @@ kernel_density_estimate(
 		    }
 		});
 	});
+    } catch (const std::exception &e) {
+      std::cout << e.what() << std::endl;
+      std::rethrow_exception(std::current_exception());
+    }
 
     return;
 }

common_src/kmeans_draft.hpp

@@ -17,7 +17,7 @@
 
 #include <CL/sycl.hpp>
 #include <limits>
-#include "kmeans.hpp"
+#include <cstdint>
 
 namespace example {
 
@@ -49,17 +49,17 @@ kmeans_host_data(
     unsigned int *labels // optional pointer to array where to give centroid assignments
     )
 {
-    const sycl::property_list buf_props = {sycl::property::buffer::use_host_ptr()};
+    const sycl::property_list buf_props {sycl::property::buffer::use_host_ptr()};
     sycl::buffer<T, 2> buf_X(data, sycl::range<2>(n, dim), buf_props);
     sycl::buffer<T, 2> buf_C(centroids, sycl::range<2>(k, dim), buf_props);
 
-    // temporary space for new centroids
+    // temporary space for new centroid coordinaties
     sycl::buffer<T, 2> buf_Cn{ sycl::range<2>(k, dim) };
+    // temporary space for point labels (index of centroid it is closest to)
+    sycl::buffer<unsigned int, 1> buf_L {sycl::range<1>(n)};
     // temporary space for new centroid sizes
     sycl::buffer<size_t, 1> buf_Cs{ sycl::range<1>(k) };
 
-    sycl::buffer<unsigned int, 1> buf_L {sycl::range<1>(n)};
-
     if (labels != nullptr) {
 	buf_L = sycl::buffer<unsigned int, 1>(labels, sycl::range<1>(n), buf_props);
     }
@@ -71,7 +71,7 @@ kmeans_host_data(
 	[&](sycl::handler &cgh) {
 
 	    // this kernel only reads data-points and centroids
-	    //   hence no reason to copy data from device to host
+	    //   hence no reason to copy back data from device to host
 	    sycl::accessor acc_X(buf_X, cgh, sycl::read_only);
 	    sycl::accessor acc_C(buf_C, cgh, sycl::read_only);
 	    // it only assigns labels and overwrites the buffer completely

kde_setuptools/README.md

@@ -3,7 +3,7 @@
 Assuming `Cython` and `pybind11` are installed and DCP++ has been activated:
 
 ```bash
-CC=dpcpp LDSHARED="dpcpp --shared" python setup.py develop
+python setup.py develop
 pytest -m tests
 ```
 

kde_setuptools/setup.py

@@ -16,6 +16,15 @@
 from pybind11.setup_helpers import Pybind11Extension
 import dpctl
 
+from setuptools.command.build_ext import build_ext
+
+class custom_build_ext(build_ext):
+    def build_extensions(self):
+        self.compiler.set_executable("compiler_so", "dpcpp -fPIC")
+        self.compiler.set_executable("compiler_cxx", "dpcpp -fPIC")
+        self.compiler.set_executable("linker_so", "dpcpp -shared -fpic -fsycl-device-code-split=per_kernel")
+        build_ext.build_extensions(self)
+
 ext_modules = [
     Pybind11Extension(
         "kde_setuptools._pybind11_kde",
@@ -43,5 +52,6 @@
     """,
     license="Apache 2.0",
     url="https://intelpython.github.io/oneAPI-for-SciPy",
-    ext_modules=ext_modules
+    ext_modules=ext_modules,
+    cmdclass={"build_ext": custom_build_ext}
 )

kde_setuptools/tests/test_cy_kde.py

@@ -41,6 +41,10 @@ def test_1d():
     x = np.linspace(0.1, 0.9, num=15).reshape((-1, 1))
     data = np.random.rand(128,1)
 
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
+
     f_cy = kde_eval(q, x, data, 0.1)
     f_ref = ref_kde(x, data, 0.1)
 
@@ -62,6 +66,10 @@ def test_2d():
 
     data = np.random.rand(128*128, 2)
 
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
+
     f_cy = kde_eval(q, x, data, 0.05)
     f_ref = ref_kde(x, data, 0.05)
 
@@ -84,6 +92,10 @@ def test_3d():
 
     data = np.random.rand(16000,3)
 
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
+
     f_cy = kde_eval(q, x, data, 0.01)
     f_ref = ref_kde(x, data, 0.01)
 

kde_setuptools/tests/test_py_kde.py

@@ -41,6 +41,10 @@ def test_1d():
     x = np.linspace(0.1, 0.9, num=15).reshape((-1, 1))
     data = np.random.rand(128,1)
 
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
+
     f_cy = kde_eval(q, x, data, 0.1)
     f_ref = ref_kde(x, data, 0.1)
 
@@ -62,6 +66,10 @@ def test_2d():
 
     data = np.random.rand(128*128, 2)
 
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
+
     f_cy = kde_eval(q, x, data, 0.05)
     f_ref = ref_kde(x, data, 0.05)
 
@@ -84,6 +92,10 @@ def test_3d():
 
     data = np.random.rand(16000,3)
 
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
+
     f_cy = kde_eval(q, x, data, 0.01)
     f_ref = ref_kde(x, data, 0.01)
 

kde_skbuild/CMakeLists.txt

@@ -41,6 +41,7 @@ pybind11_add_module(${py_module_name}
     MODULE
     ${CMAKE_SOURCE_DIR}/src/pybind11_kde.cpp
 )
+target_link_options(${py_module_name} PRIVATE -fsycl-device-code-split=per_kernel)
 target_include_directories(${py_module_name}
     PUBLIC ${CMAKE_SOURCE_DIR}/src ${Dpctl_INCLUDE_DIR}
 )
@@ -57,6 +58,7 @@ add_cython_target(${cy_module_name} ${CMAKE_SOURCE_DIR}/src/_cython_kde.pyx CXX
 add_library(${cy_module_name} MODULE ${_generated_src})
 
 target_include_directories(${cy_module_name} PRIVATE ${NumPy_INCLUDE_DIR} ${Dpctl_INCLUDE_DIR} ${CMAKE_SOURCE_DIR}/src)
+target_link_options(${cy_module_name} PRIVATE -fsycl-device-code-split=per_kernel)
 python_extension_module(${cy_module_name})
 install(TARGETS ${cy_module_name} LIBRARY DESTINATION kde_skbuild)
 

kde_skbuild/tests/test_cy_kde.py

@@ -1,4 +1,4 @@
-#  Copyright 2020-2021 Intel Corporation
+#  Copyright 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.
@@ -41,6 +41,10 @@ def test_1d():
     x = np.linspace(0.1, 0.9, num=15).reshape((-1, 1))
     data = np.random.rand(128,1)
 
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
+    
     f_cy = kde_eval(q, x, data, 0.1)
     f_ref = ref_kde(x, data, 0.1)
 
@@ -62,6 +66,10 @@ def test_2d():
 
     data = np.random.rand(128*128, 2)
 
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
+
     f_cy = kde_eval(q, x, data, 0.05)
     f_ref = ref_kde(x, data, 0.05)
 
@@ -84,6 +92,10 @@ def test_3d():
 
     data = np.random.rand(16000,3)
 
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
+
     f_cy = kde_eval(q, x, data, 0.01)
     f_ref = ref_kde(x, data, 0.01)
 

kde_skbuild/tests/test_py_kde.py

@@ -41,6 +41,10 @@ def test_1d():
     x = np.linspace(0.1, 0.9, num=15).reshape((-1, 1))
     data = np.random.rand(128,1)
 
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
+
     f_cy = kde_eval(q, x, data, 0.1)
     f_ref = ref_kde(x, data, 0.1)
 
@@ -61,6 +65,10 @@ def test_2d():
     ).reshape(-1, 2)
 
     data = np.random.rand(128*128, 2)
+    
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
 
     f_cy = kde_eval(q, x, data, 0.05)
     f_ref = ref_kde(x, data, 0.05)
@@ -83,7 +91,11 @@ def test_3d():
     ).reshape(-1, 3)
 
     data = np.random.rand(16000,3)
-
+    
+    if not q.sycl_device.has_aspect_fp64:
+        x = np.asarray(x, dtype=np.float32)
+        data = np.asarray(data, dtype=np.float32)
+    
     f_cy = kde_eval(q, x, data, 0.01)
     f_ref = ref_kde(x, data, 0.01)