Skip to content

Commit

Permalink
Merge pull request #76 from clEsperanto/add-execute-function
Browse files Browse the repository at this point in the history
ADD: execute function for running custom `opencl.cl` files
  • Loading branch information
StRigaud authored May 4, 2023
2 parents ea1bd4b + 0a0f483 commit 3247ccc
Show file tree
Hide file tree
Showing 7 changed files with 518 additions and 178 deletions.
2 changes: 1 addition & 1 deletion CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -40,7 +40,7 @@ option(BUILD_BENCHMARK OFF)
FetchContent_Declare(
CLIc_lib
GIT_REPOSITORY https://github.com/clEsperanto/CLIc_prototype.git
GIT_TAG ${CLIC_VERSION}
GIT_TAG improve-custom-kernel-call #${CLIC_VERSION}
)
FetchContent_MakeAvailable(CLIc_lib)

Expand Down
14 changes: 14 additions & 0 deletions demos/add_image_and_scalar.cl
Original file line number Diff line number Diff line change
@@ -0,0 +1,14 @@
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

__kernel void add_image_and_scalar(IMAGE_src_TYPE src, IMAGE_dst_TYPE dst,
const float scalar) {
const int x = get_global_id(0);
const int y = get_global_id(1);
const int z = get_global_id(2);

const float value =
(float)READ_IMAGE(src, sampler, POS_src_INSTANCE(x, y, z, 0)).x;
WRITE_IMAGE(dst, POS_dst_INSTANCE(x, y, z, 0),
CONVERT_dst_PIXEL_TYPE(value + scalar));
}
296 changes: 296 additions & 0 deletions demos/run_custom_kernel.ipynb

Large diffs are not rendered by default.

1 change: 1 addition & 0 deletions pyclesperanto/__init__.py
Original file line number Diff line number Diff line change
Expand Up @@ -26,6 +26,7 @@
)
from ._functionalities import (
imshow,
execute,
)
from ._tier1 import (
add_image_and_scalar,
Expand Down
50 changes: 49 additions & 1 deletion pyclesperanto/_functionalities.py
Original file line number Diff line number Diff line change
@@ -1,6 +1,54 @@
from ._memory_operations import pull
from os import path
from typing import Optional, Union

from ._memory_operations import pull
from ._device import Device, get_device

def execute(anchor: str, opencl_filename: str, kernel_name: str, parameters: dict, global_range: Optional[tuple] = None, device: Optional[Device] = None):
"""Execute a custom OpenCL kernel.
Parameters
----------
anchor : str
Path to the directory where the kernel file is located.
opencl_kernel_filename : str
Name of the OpenCL kernel file, e.g. "my_kernel.cl".
kernel_name : str
Name of the kernel function to be executed in the OpenCL kernel file.
parameters : dict
Dictionary of parameters to be passed to the kernel function, e.g. {"src": src, "dst": dst}.
global_range : tuple, optional
Global size of the kernel execution, by default None.
device : Device, optional
Device to be used for execution, by default None.
"""
from ._pyclesperanto import _std_variant as std_variant
from ._pyclesperanto import _CustomKernel_Call as op

if global_range is None:
global_range = (0,0,0)
else:
if len(global_range) == 2:
global_range = (1, global_range[0], global_range[1])
if len(global_range) == 1:
global_range = (1, 1, global_range[0])

if device is None:
device = parameters['src'].device or parameters['src1'].device or get_device()

cpp_parameter_map = {key: std_variant(val) for key, val in parameters.items()}
op(
device,
file_name=str(path.join(anchor, opencl_filename)),
kernel_name=kernel_name,
dx=global_range[2],
dy=global_range[1],
dz=global_range[0],
parameters=cpp_parameter_map,
# constants=cpp_constant_map,
)



def imshow(
image,
Expand Down
1 change: 1 addition & 0 deletions pyproject.toml
Original file line number Diff line number Diff line change
Expand Up @@ -43,4 +43,5 @@ skip = ["cp27-*", "cp36-*", "pp*", "*-win32"]

[tool.cibuildwheel.macos]
before-all = "bash {package}/scripts/build-opencl-macos.sh"
environment = {MACOSX_DEPLOYMENT_TARGET = '10.14'}
skip = ["cp27-*", "cp36-*", "pp*", "*_universal2"]
332 changes: 156 additions & 176 deletions wrapper/cleTier1.cpp

Large diffs are not rendered by default.

0 comments on commit 3247ccc

Please sign in to comment.