diff --git a/pyclesperanto/__init__.py b/pyclesperanto/__init__.py index b06d13c1..adcebdb6 100644 --- a/pyclesperanto/__init__.py +++ b/pyclesperanto/__init__.py @@ -1,5 +1,6 @@ from ._core import ( gpu_info, + execute, select_backend, select_device, get_device, diff --git a/pyclesperanto/_core.py b/pyclesperanto/_core.py index c84d1b6f..eadccbab 100644 --- a/pyclesperanto/_core.py +++ b/pyclesperanto/_core.py @@ -1,8 +1,11 @@ from typing import Optional, Union +from pathlib import Path +import numpy as np import warnings from ._pyclesperanto import _Device as Device from ._pyclesperanto import _BackendManager as BackendManager +from ._pyclesperanto import _execute class _current_device: @@ -156,6 +159,63 @@ def default_initialisation(): ) +def execute(anchor = '__file__', kernel_source: str = '', kernel_name: str = '', global_size: tuple = (1, 1, 1), parameters: dict = {}, constants: dict = {}, device: Device = None): + """Execute a kernel from a file or a string + + Call, build, and execute a kernel compatible with CLIj framework. + The kernel can be called from a file or a string. + + Parameters + ---------- + anchor : str, default = '__file__' + Enter __file__ when calling this method and the corresponding open.cl + file lies in the same folder as the python file calling it. + Ignored if kernel_source is a string. + kernel_source : str + Filename of the open.cl file to be called or string containing the open.cl source code + kernel_name : str + Kernel method inside the open.cl file to be called + most clij/clesperanto kernel functions have the same name as the file they are in + global_size : tuple (z,y,x), default = (1, 1, 1) + Global_size according to OpenCL definition (usually shape of the destination image). + parameters : dict(str, [Array, float, int]) + Dictionary containing parameters. Take care: They must be of the + right type and in the right order as specified in the open.cl file. + constants: dict(str, int), optional + Dictionary with names/values which will be added to the define + statements. They are necessary, e.g. to create arrays of a given + maximum size in OpenCL as variable array lengths are not supported. + device : Device, default = None + The device to execute the kernel on. If None, use the current device + """ + + # load the kernel file + def load_file(anchor, filename): + """Load the opencl kernel file as a string""" + if anchor is None: + kernel = Path(filename).read_text() + else: + kernel = (Path(anchor).parent / filename).read_text() + return kernel + + # test if kernel_source ends with .cl or .cu + if kernel_source.endswith('.cl') or kernel_source.endswith('.cu'): + kernel_source = load_file(anchor, kernel_source) + + # manage the device if not given + if not device: + device = get_device() + + # manage global range + if not isinstance(global_size, tuple): + if isinstance(global_size, list) or isinstance(global_size, np.ndarray): + global_size = tuple(global_size) + else: + global_size = (global_size,) + + _execute(device, kernel_name, kernel_source, parameters, global_size, constants) + + def gpu_info(): device_list = list_available_devices() info = [] diff --git a/src/wrapper/execute_.cpp b/src/wrapper/execute_.cpp new file mode 100644 index 00000000..b50d1c94 --- /dev/null +++ b/src/wrapper/execute_.cpp @@ -0,0 +1,84 @@ + +#include "pycle_wrapper.hpp" + +#include "device.hpp" +#include "array.hpp" +#include "utils.hpp" +#include "execution.hpp" + +#include +#include +#include +#include + +namespace py = pybind11; + +auto py_execute(const cle::Device::Pointer &device, const std::string &kernel_name, const std::string &kernel_source, const py::dict ¶meters, const py::tuple &range, const py::dict &constants) -> void +{ + cle::RangeArray global_range = {1, 1, 1}; + switch (range.size()) + { + case 1: + global_range[0] = range[0].cast(); + break; + case 2: + global_range[0] = range[1].cast(); + global_range[1] = range[0].cast(); + break; + case 3: + global_range[0] = range[2].cast(); + global_range[1] = range[1].cast(); + global_range[2] = range[0].cast(); + break; + default: + throw std::invalid_argument("Error: range tuple must have 3 elements or less. Received " + std::to_string(range.size()) + " elements."); + break; + } + + // manage kernel name and code + const cle::KernelInfo kernel_info = {kernel_name, kernel_source}; + + // convert py::dict paramter to vector> + cle::ParameterList clic_parameters; + for (auto item : parameters) + { + // check if item.second is cle::Array::Pointer + if (py::isinstance(item.second)) + { + clic_parameters.push_back({item.first.cast(), item.second.cast()}); + } + else if (py::isinstance(item.second)) + { + // convert py::int to int + clic_parameters.push_back({item.first.cast(), item.second.cast()}); + } + else if (py::isinstance(item.second)) + { + // convert py::float to float + clic_parameters.push_back({item.first.cast(), item.second.cast()}); + } + else + { + throw std::invalid_argument("Error: parameter type not supported. Received " + std::string(py::str(item.second.get_type()).cast())); + } + } + + // convert py::dict constant to vector> + cle::ConstantList clic_constants; + if (!constants.empty()) + { + for (auto item : constants) + { + clic_constants.push_back({item.first.cast(), item.second.cast()}); + } + } + + // execute + cle::execute(device, kernel_info, clic_parameters, global_range, clic_constants); +} + +auto execute_(py::module_ &m) -> void +{ + m.def("_execute", &py_execute, "Call execute function from C++.", + py::arg("device"), py::arg("kernel_name"), py::arg("kernel_source"), py::arg("parameters"), py::arg("range"), py::arg("constants")); +} diff --git a/src/wrapper/pycle_wrapper.cpp b/src/wrapper/pycle_wrapper.cpp index b2d58f1a..9fdfe309 100644 --- a/src/wrapper/pycle_wrapper.cpp +++ b/src/wrapper/pycle_wrapper.cpp @@ -8,6 +8,7 @@ PYBIND11_MODULE(_pyclesperanto, m) types_(m); core_(m); array_(m); + execute_(m); tier1_(m); tier2_(m); diff --git a/src/wrapper/pycle_wrapper.hpp b/src/wrapper/pycle_wrapper.hpp index 83e68873..2845b54f 100644 --- a/src/wrapper/pycle_wrapper.hpp +++ b/src/wrapper/pycle_wrapper.hpp @@ -9,6 +9,7 @@ auto types_(pybind11::module_ &module) -> void; auto core_(pybind11::module_ &module) -> void; auto array_(pybind11::module_ &module) -> void; +auto execute_(pybind11::module_ &module) -> void; auto tier1_(pybind11::module_ &module) -> void; auto tier2_(pybind11::module_ &module) -> void; diff --git a/tests/test_execute.py b/tests/test_execute.py new file mode 100644 index 00000000..1dc3d64d --- /dev/null +++ b/tests/test_execute.py @@ -0,0 +1,40 @@ +import pyclesperanto as cle +import numpy as np + +absolute_ocl = """ +__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST; + +__kernel void absolute( + IMAGE_src_TYPE src, + IMAGE_dst_TYPE dst +) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2); + + IMAGE_src_PIXEL_TYPE value = READ_IMAGE(src, sampler, POS_src_INSTANCE(x,y,z,0)).x; + if ( value < 0 ) { + value = -1 * value; + } + + WRITE_IMAGE(dst, POS_dst_INSTANCE(x,y,z,0), CONVERT_dst_PIXEL_TYPE(value)); +} +""" + +def test_execute_absolute(): + input = cle.push(np.asarray([ + [1, -1], + [1, -1] + ])) + output = cle.create(input) + + param = {'src': input, 'dst': output} + cle.execute(kernel_source=absolute_ocl, kernel_name="absolute", global_size=input.shape, parameters=param) + + print(output) + + a = cle.pull(output) + assert (np.min(a) == 1) + assert (np.max(a) == 1) + assert (np.mean(a) == 1)