Merge pull request #171 from clEsperanto/add-execute

Add-execute

pyclesperanto/__init__.py

@@ -1,5 +1,6 @@
 from ._core import (
     gpu_info,
+    execute,
     select_backend,
     select_device,
     get_device,

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 = []

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 <pybind11/functional.h>
+#include <pybind11/numpy.h>
+#include <pybind11/pybind11.h>
+#include <pybind11/stl.h>
+
+namespace py = pybind11;
+
+auto py_execute(const cle::Device::Pointer &device, const std::string &kernel_name, const std::string &kernel_source, const py::dict &parameters, 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<size_t>();
+          break;
+     case 2:
+          global_range[0] = range[1].cast<size_t>();
+          global_range[1] = range[0].cast<size_t>();
+          break;
+     case 3:
+          global_range[0] = range[2].cast<size_t>();
+          global_range[1] = range[1].cast<size_t>();
+          global_range[2] = range[0].cast<size_t>();
+          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<pair<string, array>>
+     cle::ParameterList clic_parameters;
+     for (auto item : parameters)
+     {
+          // check if item.second is cle::Array::Pointer
+          if (py::isinstance<cle::Array>(item.second))
+          {
+               clic_parameters.push_back({item.first.cast<std::string>(), item.second.cast<cle::Array::Pointer>()});
+          }
+          else if (py::isinstance<py::int_>(item.second))
+          {
+               // convert py::int to int
+               clic_parameters.push_back({item.first.cast<std::string>(), item.second.cast<int>()});
+          }
+          else if (py::isinstance<py::float_>(item.second))
+          {
+               // convert py::float to float
+               clic_parameters.push_back({item.first.cast<std::string>(), item.second.cast<float>()});
+          }
+          else
+          {
+               throw std::invalid_argument("Error: parameter type not supported. Received " + std::string(py::str(item.second.get_type()).cast<std::string>()));
+          }
+     }
+
+     // convert py::dict constant to vector<pair<string, int>>
+     cle::ConstantList clic_constants;
+     if (!constants.empty())
+     {
+          for (auto item : constants)
+          {
+               clic_constants.push_back({item.first.cast<std::string>(), item.second.cast<int>()});
+          }
+     }
+
+     // 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"));
+}

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);

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;

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)