diff options
Diffstat (limited to 'silx/resources')
| -rw-r--r-- | silx/resources/__init__.py | 200 | ||||
| -rw-r--r-- | silx/resources/gui/icons/compare-mode-a-minus-b.png | bin | 0 -> 3862 bytes | |||
| -rw-r--r-- | silx/resources/gui/icons/compare-mode-a-minus-b.svg | 25 | ||||
| -rw-r--r-- | silx/resources/gui/icons/eye.png | bin | 0 -> 755 bytes | |||
| -rw-r--r-- | silx/resources/gui/icons/eye.svg | 23 | ||||
| -rw-r--r-- | silx/resources/gui/icons/pointing-hand.png | bin | 0 -> 680 bytes | |||
| -rw-r--r-- | silx/resources/gui/icons/pointing-hand.svg | 2 | ||||
| -rw-r--r-- | silx/resources/opencl/convolution.cl | 312 | ||||
| -rw-r--r-- | silx/resources/opencl/convolution_textures.cl | 374 | ||||
| -rw-r--r-- | silx/resources/opencl/sparse.cl | 84 |
10 files changed, 823 insertions, 197 deletions
diff --git a/silx/resources/__init__.py b/silx/resources/__init__.py index ba390d6..5346f48 100644 --- a/silx/resources/__init__.py +++ b/silx/resources/__init__.py @@ -56,20 +56,14 @@ of this modules to ensure access across different distribution schemes: __authors__ = ["V.A. Sole", "Thomas Vincent", "J. Kieffer"] __license__ = "MIT" -__date__ = "13/02/2019" +__date__ = "08/03/2019" import os import sys -import threading -import json import logging -import tempfile -import unittest import importlib -import six - logger = logging.getLogger(__name__) @@ -288,193 +282,5 @@ def _resource_filename(resource, default_directory=None): return pkg_resources.resource_filename(package_name, resource_name) -class ExternalResources(object): - """Utility class which allows to download test-data from www.silx.org - and manage the temporary data during the tests. - - """ - - def __init__(self, project, - url_base, - env_key=None, - timeout=60): - """Constructor of the class - - :param str project: name of the project, like "silx" - :param str url_base: base URL for the data, like "http://www.silx.org/pub" - :param str env_key: name of the environment variable which contains the - test_data directory, like "SILX_DATA". - If None (default), then the name of the - environment variable is built from the project argument: - "<PROJECT>_DATA". - The environment variable is optional: in case it is not set, - a directory in the temporary folder is used. - :param timeout: time in seconds before it breaks - """ - self.project = project - self._initialized = False - self.sem = threading.Semaphore() - - self.env_key = env_key or (self.project.upper() + "_TESTDATA") - self.url_base = url_base - self.all_data = set() - self.timeout = timeout - self._data_home = None - - @property - def data_home(self): - """Returns the data_home path and make sure it exists in the file - system.""" - if self._data_home is not None: - return self._data_home - - data_home = os.environ.get(self.env_key) - if data_home is None: - try: - import getpass - name = getpass.getuser() - except Exception: - if "getlogin" in dir(os): - name = os.getlogin() - elif "USER" in os.environ: - name = os.environ["USER"] - elif "USERNAME" in os.environ: - name = os.environ["USERNAME"] - else: - name = "uid" + str(os.getuid()) - - basename = "%s_testdata_%s" % (self.project, name) - data_home = os.path.join(tempfile.gettempdir(), basename) - if not os.path.exists(data_home): - os.makedirs(data_home) - self._data_home = data_home - return data_home - - def _initialize_data(self): - """Initialize for downloading test data""" - if not self._initialized: - with self.sem: - if not self._initialized: - self.testdata = os.path.join(self.data_home, "all_testdata.json") - if os.path.exists(self.testdata): - with open(self.testdata) as f: - self.all_data = set(json.load(f)) - self._initialized = True - - def clean_up(self): - pass - - def getfile(self, filename): - """Downloads the requested file from web-server available - at https://www.silx.org/pub/silx/ - - :param: relative name of the image. - :return: full path of the locally saved file. - """ - logger.debug("ExternalResources.getfile('%s')", filename) - - if not self._initialized: - self._initialize_data() - - fullfilename = os.path.abspath(os.path.join(self.data_home, filename)) - - if not os.path.isfile(fullfilename): - logger.debug("Trying to download image %s, timeout set to %ss", - filename, self.timeout) - dictProxies = {} - if "http_proxy" in os.environ: - dictProxies['http'] = os.environ["http_proxy"] - dictProxies['https'] = os.environ["http_proxy"] - if "https_proxy" in os.environ: - dictProxies['https'] = os.environ["https_proxy"] - if dictProxies: - proxy_handler = six.moves.urllib.request.ProxyHandler(dictProxies) - opener = six.moves.urllib.request.build_opener(proxy_handler).open - else: - opener = six.moves.urllib.request.urlopen - - logger.debug("wget %s/%s", self.url_base, filename) - try: - data = opener("%s/%s" % (self.url_base, filename), - data=None, timeout=self.timeout).read() - logger.info("Image %s successfully downloaded.", filename) - except six.moves.urllib.error.URLError: - raise unittest.SkipTest("network unreachable.") - - if not os.path.isdir(os.path.dirname(fullfilename)): - # Create sub-directory if needed - os.makedirs(os.path.dirname(fullfilename)) - - try: - with open(fullfilename, "wb") as outfile: - outfile.write(data) - except IOError: - raise IOError("unable to write downloaded \ - data to disk at %s" % self.data_home) - - if not os.path.isfile(fullfilename): - raise RuntimeError( - "Could not automatically \ - download test images %s!\n \ If you are behind a firewall, \ - please set both environment variable http_proxy and https_proxy.\ - This even works under windows ! \n \ - Otherwise please try to download the images manually from \n%s/%s" - % (filename, self.url_base, filename)) - - if filename not in self.all_data: - self.all_data.add(filename) - image_list = list(self.all_data) - image_list.sort() - try: - with open(self.testdata, "w") as fp: - json.dump(image_list, fp, indent=4) - except IOError: - logger.debug("Unable to save JSON list") - - return fullfilename - - def getdir(self, dirname): - """Downloads the requested tarball from the server - https://www.silx.org/pub/silx/ - and unzips it into the data directory - - :param: relative name of the image. - :return: list of files with their full path. - """ - lodn = dirname.lower() - if (lodn.endswith("tar") or lodn.endswith("tgz") or - lodn.endswith("tbz2") or lodn.endswith("tar.gz") or - lodn.endswith("tar.bz2")): - import tarfile - engine = tarfile.TarFile.open - elif lodn.endswith("zip"): - import zipfile - engine = zipfile.ZipFile - else: - raise RuntimeError("Unsupported archive format. Only tar and zip " - "are currently supported") - full_path = self.getfile(dirname) - root = os.path.dirname(full_path) - with engine(full_path, mode="r") as fd: - fd.extractall(self.data_home) - if lodn.endswith("zip"): - result = [os.path.join(root, i) for i in fd.namelist()] - else: - result = [os.path.join(root, i) for i in fd.getnames()] - return result - - def download_all(self, imgs=None): - """Download all data needed for the test/benchmarks - - :param imgs: list of files to download, by default all - :return: list of path with all files - """ - if not self._initialized: - self._initialize_data() - if not imgs: - imgs = self.all_data - res = [] - for fn in imgs: - logger.info("Downloading from silx.org: %s", fn) - res.append(self.getfile(fn)) - return res +# Expose ExternalResources for compatibility (since silx 0.11) +from ..utils.ExternalResources import ExternalResources diff --git a/silx/resources/gui/icons/compare-mode-a-minus-b.png b/silx/resources/gui/icons/compare-mode-a-minus-b.png Binary files differnew file mode 100644 index 0000000..75f8db3 --- /dev/null +++ b/silx/resources/gui/icons/compare-mode-a-minus-b.png diff --git a/silx/resources/gui/icons/compare-mode-a-minus-b.svg b/silx/resources/gui/icons/compare-mode-a-minus-b.svg new file mode 100644 index 0000000..bd4afbf --- /dev/null +++ b/silx/resources/gui/icons/compare-mode-a-minus-b.svg @@ -0,0 +1,25 @@ +<?xml version="1.0" encoding="UTF-8"?> +<!-- Created with Inkscape (http://www.inkscape.org/) --> +<svg id="svg2985" width="48px" height="48px" version="1.1" xmlns="http://www.w3.org/2000/svg" xmlns:cc="http://creativecommons.org/ns#" xmlns:dc="http://purl.org/dc/elements/1.1/" xmlns:rdf="http://www.w3.org/1999/02/22-rdf-syntax-ns#"> + <metadata id="metadata2990"> + <rdf:RDF> + <cc:Work rdf:about=""> + <dc:format>image/svg+xml</dc:format> + <dc:type rdf:resource="http://purl.org/dc/dcmitype/StillImage"/> + <dc:title/> + </cc:Work> + </rdf:RDF> + </metadata> + <g id="layer1"> + <path id="path2992-3" d="m33.031 9.9688c-0.99857 0.014569-2.0178 0.11579-3.0312 0.34375h-0.0625c-2.1183 0.48558-4.0182 1.4422-5.6562 2.6875 2.8649 2.1382 4.9928 5.2987 5.8125 9.125 1.3092 6.1116-1.1148 12.153-5.7188 15.656 3.3562 2.5072 7.6866 3.6093 12.062 2.625 8.1077-1.8236 13.248-10.036 11.469-18.344-1.5572-7.2694-7.885-12.196-14.875-12.094zm-8.6562 27.812c-2.8622-2.1382-5.0246-5.3011-5.8438-9.125-1.3093-6.1119 1.1417-12.153 5.75-15.656-3.3555-2.5044-7.689-3.6087-12.062-2.625l-0.0625 0.03125c-8.0776 1.8516-13.181 10.028-11.406 18.312 1.7797 8.3079 9.7986 13.574 17.906 11.75 2.1433-0.4821 4.0643-1.4286 5.7188-2.6875z"/> + <g id="text3834" fill="#fff"> + <path id="path3870" d="m5.8982 29.014c1.8906-4.8626 3.2728-8.4075 4.1465-10.635 0.05728-0.10741 0.18261-0.23631 0.37598-0.38672 0.20052-0.15754 0.34732-0.23631 0.44043-0.23633 0.02148 1.5e-5 0.04296 0.01076 0.06445 0.03223 0.02864 0.01434 0.05012 0.02508 0.06445 0.03223l4.3721 11.473c0.35806 0.92383 0.54068 1.3965 0.54785 1.418 0.09309 0.22201 0.31509 0.40104 0.66602 0.53711 0.35806 0.12891 0.62303 0.19336 0.79492 0.19336 0.04295 0.06445 0.06444 0.17188 0.06445 0.32227-1.4e-5 0.14323-0.0072 0.24349-0.02148 0.30078-1.4323-0.07161-2.2201-0.10742-2.3633-0.10742-0.29363 0-0.56935 0.0072-0.82715 0.02148-0.25066 0.0072-0.53354 0.02148-0.84863 0.04297-0.31511 0.02148-0.58725 0.03581-0.81641 0.04297-0.05014-0.05729-0.0752-0.16113-0.0752-0.31152 0.0072-0.15039 0.03938-0.25423 0.09668-0.31152 0.1647 1e-6 0.38313-0.02506 0.65527-0.0752 0.27929-0.05729 0.4619-0.12532 0.54785-0.2041 0.05728-0.05729 0.08593-0.15755 0.08594-0.30078-1e-5 -0.19336-0.12176-0.63737-0.36524-1.332-0.2435-0.70182-0.48699-1.3535-0.73047-1.9551l-0.36524-0.90234c-0.04298-0.05729-0.08595-0.08593-0.12891-0.08594-0.18621-0.01432-0.89519-0.02148-2.127-0.02148-0.90235 6e-6 -1.6758 0.02149-2.3203 0.06445l-0.10742 0.06445c-0.40105 1.0599-0.72331 1.9336-0.9668 2.6211-0.20052 0.58724-0.30078 1.0957-0.30078 1.5254-2.6e-6 0.12175 0.064451 0.22559 0.19336 0.31152 0.13606 0.08594 0.2972 0.15039 0.4834 0.19336 0.19336 0.03581 0.35807 0.06087 0.49414 0.0752 0.13606 0.01432 0.2399 0.02149 0.31152 0.02148 0.028642 0.02865 0.046545 0.12533 0.053711 0.29004 0.014326 0.15756 0.010745 0.26856-0.010735 0.33301-0.20769-0.007161-0.43327-0.021484-0.67676-0.042969-0.24349-0.021484-0.46908-0.035807-0.67676-0.042968-0.20768-0.01432-0.45117-0.02148-0.73047-0.02148-0.29362 0-0.55501 0.0072-0.78418 0.02148-0.22201 0.0072-0.46908 0.02148-0.74121 0.04297-0.27214 0.02148-0.52995 0.03581-0.77344 0.04297-0.028646-0.02865-0.042968-0.12174-0.042969-0.2793 3e-7 -0.15755 0.014323-0.27214 0.042969-0.34375 0.17904 1e-6 0.41536-0.03581 0.70898-0.10742 0.29362-0.07878 0.48698-0.17904 0.58008-0.30078 0.30078-0.39388 0.63737-1.0671 1.0098-2.0195m4.2109-8.6582-1.9766 5.2422c-4.3e-6 0.0573 0.014319 0.08594 0.042969 0.08594 0.74479 0.02865 1.3965 0.04298 1.9551 0.04297 0.5371 7e-6 1.1387-0.02148 1.8047-0.06445 0.01431 7e-6 0.02864-0.0072 0.04297-0.02148 0.01431-0.02864 0.02148-0.05012 0.02148-0.06445v-0.02148l-1.8906-5.1992"/> + </g> + <g id="text3834-3" fill="#fff"> + <path id="path3867" d="m35.655 25.1v4.6621c-0.0072 0.20768 0.0072 0.48698 0.04297 0.83789 0.0358 0.35091 0.08593 0.55859 0.15039 0.62305 0.15039 0.15755 0.35449 0.26139 0.6123 0.31152 0.26497 0.05013 0.62662 0.0752 1.085 0.07519 2.2773 1e-6 3.416-1.0527 3.416-3.1582-1e-5 -2.2344-1.332-3.3516-3.9961-3.3516h-1.3105m0-0.81641h1.7939c0.89518 8e-6 1.5684-0.24706 2.0195-0.74121 0.45116-0.49413 0.67675-1.085 0.67676-1.7725-9e-6 -0.9453-0.30079-1.6829-0.90234-2.2129-0.59441-0.52994-1.3965-0.79491-2.4062-0.79492-0.6517 1.4e-5 -1.0098 0.07163-1.0742 0.21484-0.08594 0.2077-0.12175 0.7448-0.10742 1.6113v3.6953m2.127-6.209c0.41536 1.4e-5 0.8164 0.03224 1.2031 0.09668 0.38671 0.06447 0.77343 0.17547 1.1602 0.33301 0.38671 0.1504 0.71972 0.34376 0.99902 0.58008 0.28645 0.23634 0.51561 0.54428 0.6875 0.92383 0.17902 0.37241 0.26854 0.79493 0.26855 1.2676-1e-5 0.73764-0.2435 1.3965-0.73047 1.9766-0.47983 0.58009-0.95964 0.95964-1.4395 1.1387-0.06446 0.02865-0.06446 0.06446 0 0.10742 0.27929 0.01433 0.59439 0.10385 0.94531 0.26856 0.3509 0.16472 0.69107 0.39031 1.0205 0.67676 0.33658 0.2793 0.61946 0.64454 0.84863 1.0957 0.22916 0.44402 0.34374 0.91667 0.34375 1.418-1.2e-5 1.3105-0.48341 2.3704-1.4502 3.1797-0.89519 0.74479-2.3955 1.1315-4.501 1.1602h-4.9414c-0.05013-0.05729-0.0752-0.16113-0.0752-0.31152 0.0072-0.15039 0.03223-0.25423 0.0752-0.31152 0.21484 1e-6 0.50846-0.03223 0.88086-0.09668 0.37239-0.06445 0.57291-0.14681 0.60156-0.24707 0.07161-0.22916 0.10742-0.75911 0.10742-1.5898v-9.1738c-2e-6 -0.83072-0.03581-1.3607-0.10742-1.5898-0.02865-0.10025-0.22917-0.1826-0.60156-0.24707-0.3724-0.06444-0.66602-0.09667-0.88086-0.09668-0.04297-0.05728-0.06804-0.16112-0.0752-0.31152-1e-6 -0.15038 0.02506-0.25422 0.0752-0.31152 1.4681 0.07163 2.3131 0.10744 2.5352 0.10742 0.24349 1.4e-5 0.69824-0.0071 1.3643-0.02149 0.66601-0.01431 1.2282-0.02147 1.6865-0.02148"/> + </g> + <g id="text3859"> + <path id="path3864" d="m22.387 24.902c2.7083-0.14322 4.2773-0.22135 4.707-0.23438 0.09114 5e-6 0.13671 0.0586 0.13672 0.17578-6e-6 0.34506-0.08464 0.69662-0.25391 1.0547-0.26693 0.01954-1.8359 0.09115-4.707 0.21484-0.09115 4e-6 -0.13672-0.07161-0.13672-0.21484-1e-6 -0.29947 0.08463-0.6315 0.25391-0.99609"/> + </g> + </g> +</svg> diff --git a/silx/resources/gui/icons/eye.png b/silx/resources/gui/icons/eye.png Binary files differnew file mode 100644 index 0000000..a2d1c23 --- /dev/null +++ b/silx/resources/gui/icons/eye.png diff --git a/silx/resources/gui/icons/eye.svg b/silx/resources/gui/icons/eye.svg new file mode 100644 index 0000000..7658d86 --- /dev/null +++ b/silx/resources/gui/icons/eye.svg @@ -0,0 +1,23 @@ +<?xml version="1.0" encoding="UTF-8"?> +<svg id="svg2" version="1.1" viewBox="0 0 32 32" xmlns="http://www.w3.org/2000/svg" xmlns:cc="http://creativecommons.org/ns#" xmlns:dc="http://purl.org/dc/elements/1.1/" xmlns:rdf="http://www.w3.org/1999/02/22-rdf-syntax-ns#"> + <metadata id="metadata16"> + <rdf:RDF> + <cc:Work rdf:about=""> + <dc:format>image/svg+xml</dc:format> + <dc:type rdf:resource="http://purl.org/dc/dcmitype/StillImage"/> + <dc:title/> + </cc:Work> + </rdf:RDF> + </metadata> + <g id="g3801" transform="translate(-.98262 -.54146)"> + <path id="path3789" transform="matrix(1.1637 0 0 1.1637 -1.7349 -11.417)" d="m20.738 24.026a4.6539 4.6539 0 1 1-0.0819-0.86928l-4.572 0.86928z" stroke="#000"/> + <path id="path3795" d="m12.628 15.997s0.08165-3.4564 3.4836-3.8374" fill="none" stroke="#FFF" stroke-width="1px"/> + </g> + <path id="path3816" d="m2.534 15.997c13.578 16.114 26.932 0 26.932 0" fill="none" stroke="#000" stroke-linecap="round" stroke-width="1px"/> + <path id="path3818" d="m2.534 15.997c13.578-16.114 26.932 0 26.932 0" fill="none" stroke="#000" stroke-linecap="round" stroke-width="1px"/> + <path id="path3822" d="m16 8.6956v-2.0658" fill="none" stroke="#000" stroke-width="1px"/> + <path id="path3824" d="m10.772 9.5672-0.8732-1.8723" fill="none" stroke="#000" stroke-width="1px"/> + <path id="path3826" d="m5.3345 12.715-1.3278-1.583" fill="none" stroke="#000" stroke-width="1px"/> + <path id="path3828" d="m26.665 12.715 1.328-1.583" fill="none" stroke="#000" stroke-width="1px"/> + <path id="path3830" d="m21.228 9.5672 0.873-1.8723" fill="none" stroke="#000" stroke-width="1px"/> +</svg> diff --git a/silx/resources/gui/icons/pointing-hand.png b/silx/resources/gui/icons/pointing-hand.png Binary files differnew file mode 100644 index 0000000..abc9cd1 --- /dev/null +++ b/silx/resources/gui/icons/pointing-hand.png diff --git a/silx/resources/gui/icons/pointing-hand.svg b/silx/resources/gui/icons/pointing-hand.svg new file mode 100644 index 0000000..2a755da --- /dev/null +++ b/silx/resources/gui/icons/pointing-hand.svg @@ -0,0 +1,2 @@ +<?xml version="1.0" encoding="UTF-8"?> +<svg id="svg3092" version="1.1" viewBox="0 0 32 32" xml:space="preserve" xmlns="http://www.w3.org/2000/svg" xmlns:cc="http://creativecommons.org/ns#" xmlns:dc="http://purl.org/dc/elements/1.1/" xmlns:rdf="http://www.w3.org/1999/02/22-rdf-syntax-ns#"><metadata id="metadata3123"><rdf:RDF><cc:Work rdf:about=""><dc:format>image/svg+xml</dc:format><dc:type rdf:resource="http://purl.org/dc/dcmitype/StillImage"/><dc:title/></cc:Work></rdf:RDF></metadata><g id="g3946" transform="translate(-31.092 .90251)"><path id="path3922" d="m43.172 16.218v-12.218" fill="none" stroke="#000" stroke-width="1px"/><path id="path3924" d="m46.736 15.2v-10.997" fill="none" stroke="#000" stroke-width="1px"/><path id="path3926" d="m50.299 15.2v-3.8692" fill="none" stroke="#000" stroke-width="1px"/><path id="path3930" d="m43.172 3.9994c0.10182-2.7491 3.4619-2.7491 3.5637 0.20364" fill="none" stroke="#000" stroke-linecap="round" stroke-width="1px"/><path id="path3932" d="m46.738 11.229c0.10182-2.7491 3.4619-2.7491 3.5637 0.20364" fill="none" stroke="#000" stroke-width="1px"/><path id="path3934" d="m50.305 12.254c0.10182-2.7491 3.4619-2.7491 3.5637 0.20364" fill="none" stroke="#000" stroke-linecap="round" stroke-width="1px"/><path id="path3936" d="m53.892 14.792v-2.7491" fill="none" stroke="#000" stroke-width="1px"/><path id="path3938" d="m53.888 13.061c0.10182-1.8328 2.7242-2.0364 2.826 0.20364v6.7201c-0.4936 3.3365-2.6441 1.5754-2.24 8.2474" fill="none" stroke="#000" stroke-width="1px"/><path id="path3940" d="m43.019 15.565c-3.3601-4.1746-6.7558-1.7904-5.1267 0.34784l3.3601 3.7674c1.945 2.7929 1.9311 1.7043 1.9346 8.5529" fill="none" stroke="#000" stroke-width="1px"/><rect id="rect3944" x="42.677" y="25.484" width="12.285" height="3.0344" ry="0"/></g></svg> diff --git a/silx/resources/opencl/convolution.cl b/silx/resources/opencl/convolution.cl new file mode 100644 index 0000000..629b8fc --- /dev/null +++ b/silx/resources/opencl/convolution.cl @@ -0,0 +1,312 @@ +#define MAX_CONST_SIZE 16384 + +/******************************************************************************/ +/**************************** Macros ******************************************/ +/******************************************************************************/ + +// Get the center index of the filter, +// and the "half-Left" and "half-Right" lengths. +// In the case of an even-sized filter, the center is shifted to the left. +#define GET_CENTER_HL(hlen){\ + if (hlen & 1) {\ + c = hlen/2;\ + hL = c;\ + hR = c;\ + }\ + else {\ + c = hlen/2 - 1;\ + hL = c;\ + hR = c+1;\ + }\ +}\ + +// Boundary handling modes +#define CONV_MODE_REFLECT 0 // cba|abcd|dcb +#define CONV_MODE_NEAREST 1 // aaa|abcd|ddd +#define CONV_MODE_WRAP 2 // bcd|abcd|abc +#define CONV_MODE_CONSTANT 3 // 000|abcd|000 +#ifndef USED_CONV_MODE + #define USED_CONV_MODE CONV_MODE_NEAREST +#endif + +#define CONV_PERIODIC_IDX_X int idx_x = gidx - c + jx; if (idx_x < 0) idx_x += Nx; if (idx_x >= Nx) idx_x -= Nx; +#define CONV_PERIODIC_IDX_Y int idx_y = gidy - c + jy; if (idx_y < 0) idx_y += Ny; if (idx_y >= Ny) idx_y -= Ny; +#define CONV_PERIODIC_IDX_Z int idx_z = gidz - c + jz; if (idx_z < 0) idx_z += Nz; if (idx_z >= Nz) idx_z -= Nz; + +#define CONV_NEAREST_IDX_X int idx_x = clamp((int) (gidx - c + jx), 0, Nx-1); +#define CONV_NEAREST_IDX_Y int idx_y = clamp((int) (gidy - c + jy), 0, Ny-1); +#define CONV_NEAREST_IDX_Z int idx_z = clamp((int) (gidz - c + jz), 0, Nz-1); + +#define CONV_REFLECT_IDX_X int idx_x = gidx - c + jx; if (idx_x < 0) idx_x = -idx_x-1; if (idx_x >= Nx) idx_x = Nx-(idx_x-(Nx-1)); +#define CONV_REFLECT_IDX_Y int idx_y = gidy - c + jy; if (idx_y < 0) idx_y = -idx_y-1; if (idx_y >= Ny) idx_y = Ny-(idx_y-(Ny-1)); +#define CONV_REFLECT_IDX_Z int idx_z = gidz - c + jz; if (idx_z < 0) idx_z = -idx_z-1; if (idx_z >= Nz) idx_z = Nz-(idx_z-(Nz-1)); + + +#if USED_CONV_MODE == CONV_MODE_REFLECT + #define CONV_IDX_X CONV_REFLECT_IDX_X + #define CONV_IDX_Y CONV_REFLECT_IDX_Y + #define CONV_IDX_Z CONV_REFLECT_IDX_Z +#elif USED_CONV_MODE == CONV_MODE_NEAREST + #define CONV_IDX_X CONV_NEAREST_IDX_X + #define CONV_IDX_Y CONV_NEAREST_IDX_Y + #define CONV_IDX_Z CONV_NEAREST_IDX_Z +#elif USED_CONV_MODE == CONV_MODE_WRAP + #define CONV_IDX_X CONV_PERIODIC_IDX_X + #define CONV_IDX_Y CONV_PERIODIC_IDX_Y + #define CONV_IDX_Z CONV_PERIODIC_IDX_Z +#elif USED_CONV_MODE == CONV_MODE_CONSTANT + #error "constant not implemented yet" +#else + #error "Unknown convolution mode" +#endif + + + +// Image access patterns +#define READ_IMAGE_1D_X input[(gidz*Ny + gidy)*Nx + idx_x] +#define READ_IMAGE_1D_Y input[(gidz*Ny + idx_y)*Nx + gidx] +#define READ_IMAGE_1D_Z input[(idx_z*Ny + gidy)*Nx + gidx] + +#define READ_IMAGE_2D_XY input[(gidz*Ny + idx_y)*Nx + idx_x] +#define READ_IMAGE_2D_XZ input[(idx_z*Ny + gidy)*Nx + idx_x] +#define READ_IMAGE_2D_YZ input[(idx_z*Ny + idx_y)*Nx + gidx] + +#define READ_IMAGE_3D_XYZ input[(idx_z*Ny + idx_y)*Nx + idx_x] + + + +/******************************************************************************/ +/**************************** 1D Convolution **********************************/ +/******************************************************************************/ + + +// Convolution with 1D kernel along axis "X" (fast dimension) +// Works for batched 1D on 2D and batched 2D on 3D, along axis "X". +__kernel void convol_1D_X( + const __global float * input, + __global float * output, + __global float * filter, + int L, // filter size + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(L); + float sum = 0.0f; + + for (int jx = 0; jx <= hR+hL; jx++) { + CONV_IDX_X; // Get index "x" + sum += READ_IMAGE_1D_X * filter[L-1 - jx]; + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} + + +// Convolution with 1D kernel along axis "Y" +// Works for batched 1D on 2D and batched 2D on 3D, along axis "Y". +__kernel void convol_1D_Y( + const __global float * input, + __global float * output, + __global float * filter, + int L, // filter size + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(L); + float sum = 0.0f; + + for (int jy = 0; jy <= hR+hL; jy++) { + CONV_IDX_Y; // Get index "y" + sum += READ_IMAGE_1D_Y * filter[L-1 - jy]; + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} + + +// Convolution with 1D kernel along axis "Z" +// Works for batched 1D on 2D and batched 2D on 3D, along axis "Z". +__kernel void convol_1D_Z( + const __global float * input, + __global float * output, + __global float * filter, + int L, // filter size + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(L); + float sum = 0.0f; + + for (int jz = 0; jz <= hR+hL; jz++) { + CONV_IDX_Z; // Get index "z" + sum += READ_IMAGE_1D_Z * filter[L-1 - jz]; + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} + + +/******************************************************************************/ +/**************************** 2D Convolution **********************************/ +/******************************************************************************/ + +// Convolution with 2D kernel +// Works for batched 2D on 3D. +__kernel void convol_2D_XY( + const __global float * input, + __global float * output, + __global float * filter, + int Lx, // filter number of columns, + int Ly, // filter number of rows, + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(Lx); + float sum = 0.0f; + + for (int jy = 0; jy <= hR+hL; jy++) { + CONV_IDX_Y; // Get index "y" + for (int jx = 0; jx <= hR+hL; jx++) { + CONV_IDX_X; // Get index "x" + sum += READ_IMAGE_2D_XY * filter[(Ly-1-jy)*Lx + (Lx-1 - jx)]; + } + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} + + +// Convolution with 2D kernel +// Works for batched 2D on 3D. +__kernel void convol_2D_XZ( + const __global float * input, + __global float * output, + __global float * filter, + int Lx, // filter number of columns, + int Lz, // filter number of rows, + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(Lx); + float sum = 0.0f; + + for (int jz = 0; jz <= hR+hL; jz++) { + CONV_IDX_Z; // Get index "z" + for (int jx = 0; jx <= hR+hL; jx++) { + CONV_IDX_X; // Get index "x" + sum += READ_IMAGE_2D_XZ * filter[(Lz-1-jz)*Lx + (Lx-1 - jx)]; + } + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} + + +// Convolution with 2D kernel +// Works for batched 2D on 3D. +__kernel void convol_2D_YZ( + const __global float * input, + __global float * output, + __global float * filter, + int Ly, // filter number of columns, + int Lz, // filter number of rows, + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(Ly); + float sum = 0.0f; + + for (int jz = 0; jz <= hR+hL; jz++) { + CONV_IDX_Z; // Get index "z" + for (int jy = 0; jy <= hR+hL; jy++) { + CONV_IDX_Y; // Get index "y" + sum += READ_IMAGE_2D_YZ * filter[(Lz-1-jz)*Ly + (Ly-1 - jy)]; + } + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} + + + +/******************************************************************************/ +/**************************** 3D Convolution **********************************/ +/******************************************************************************/ + +// Convolution with 3D kernel +__kernel void convol_3D_XYZ( + const __global float * input, + __global float * output, + __global float * filter, + int Lx, // filter number of columns, + int Ly, // filter number of rows, + int Lz, // filter number of rows, + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(Lx); + float sum = 0.0f; + + for (int jz = 0; jz <= hR+hL; jz++) { + CONV_IDX_Z; // Get index "z" + for (int jy = 0; jy <= hR+hL; jy++) { + CONV_IDX_Y; // Get index "y" + for (int jx = 0; jx <= hR+hL; jx++) { + CONV_IDX_X; // Get index "x" + sum += READ_IMAGE_3D_XYZ * filter[((Lz-1-jz)*Ly + (Ly-1-jy))*Lx + (Lx-1 - jx)]; + } + } + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} + diff --git a/silx/resources/opencl/convolution_textures.cl b/silx/resources/opencl/convolution_textures.cl new file mode 100644 index 0000000..517a67c --- /dev/null +++ b/silx/resources/opencl/convolution_textures.cl @@ -0,0 +1,374 @@ +/******************************************************************************/ +/**************************** Macros ******************************************/ +/******************************************************************************/ + +// Error handling +#ifndef IMAGE_DIMS + #error "IMAGE_DIMS must be defined" +#endif +#ifndef FILTER_DIMS + #error "FILTER_DIMS must be defined" +#endif +#if FILTER_DIMS > IMAGE_DIMS + #error "Filter cannot have more dimensions than image" +#endif + +// Boundary handling modes +#define CONV_MODE_REFLECT 0 // CLK_ADDRESS_MIRRORED_REPEAT : cba|abcd|dcb +#define CONV_MODE_NEAREST 1 // CLK_ADDRESS_CLAMP_TO_EDGE : aaa|abcd|ddd +#define CONV_MODE_WRAP 2 // CLK_ADDRESS_REPEAT : bcd|abcd|abc +#define CONV_MODE_CONSTANT 3 // CLK_ADDRESS_CLAMP : 000|abcd|000 +#ifndef USED_CONV_MODE + #define USED_CONV_MODE CONV_MODE_NEAREST +#endif +#if USED_CONV_MODE == CONV_MODE_REFLECT + #define CLK_BOUNDARY CLK_ADDRESS_MIRRORED_REPEAT + #define CLK_COORDS CLK_NORMALIZED_COORDS_TRUE + #define USE_NORM_COORDS +#elif USED_CONV_MODE == CONV_MODE_NEAREST + #define CLK_BOUNDARY CLK_ADDRESS_CLAMP_TO_EDGE + #define CLK_COORDS CLK_NORMALIZED_COORDS_FALSE +#elif USED_CONV_MODE == CONV_MODE_WRAP + #define CLK_BOUNDARY CLK_ADDRESS_REPEAT + #define CLK_COORDS CLK_NORMALIZED_COORDS_TRUE + #define USE_NORM_COORDS +#elif USED_CONV_MODE == CONV_MODE_CONSTANT + #define CLK_BOUNDARY CLK_ADDRESS_CLAMP + #define CLK_COORDS CLK_NORMALIZED_COORDS_FALSE +#else + #error "Unknown convolution mode" +#endif + + + +// Convolution index for filter +#define FILTER_INDEX(j) (Lx - 1 - j) + +// Filter access patterns +#define READ_FILTER_1D(j) read_imagef(filter, (int2) (FILTER_INDEX(j), 0)).x; +#define READ_FILTER_2D(jx, jy) read_imagef(filter, (int2) (FILTER_INDEX(jx), FILTER_INDEX(jy))).x; +#define READ_FILTER_3D(jx, jy, jz) read_imagef(filter, (int4) (FILTER_INDEX(jx), FILTER_INDEX(jy), FILTER_INDEX(jz), 0)).x; + + +// Convolution index for image +#ifdef USE_NORM_COORDS + #define IMAGE_INDEX_X (gidx*1.0f +0.5f - c + jx)/Nx + #define IMAGE_INDEX_Y (gidy*1.0f +0.5f - c + jy)/Ny + #define IMAGE_INDEX_Z (gidz*1.0f +0.5f - c + jz)/Nz + #define RET_TYPE_1 float + #define RET_TYPE_2 float2 + #define RET_TYPE_4 float4 + #define C_ZERO 0.5f + #define GIDX (gidx*1.0f + 0.5f)/Nx + #define GIDY (gidy*1.0f + 0.5f)/Ny + #define GIDZ (gidz*1.0f + 0.5f)/Nz +#else + #define IMAGE_INDEX_X (gidx - c + jx) + #define IMAGE_INDEX_Y (gidy - c + jy) + #define IMAGE_INDEX_Z (gidz - c + jz) + #define RET_TYPE_1 int + #define RET_TYPE_2 int2 + #define RET_TYPE_4 int4 + #define C_ZERO 0 + #define GIDX gidx + #define GIDY gidy + #define GIDZ gidz +#endif + +static const sampler_t sampler = CLK_COORDS | CLK_BOUNDARY | CLK_FILTER_NEAREST; + +// Image access patterns +#define READ_IMAGE_1D read_imagef(input, sampler, (RET_TYPE_2) (IMAGE_INDEX_X, C_ZERO)).x + +#define READ_IMAGE_2D_X read_imagef(input, sampler, (RET_TYPE_2) (IMAGE_INDEX_X , GIDY)).x +#define READ_IMAGE_2D_Y read_imagef(input, sampler, (RET_TYPE_2) (GIDX, IMAGE_INDEX_Y)).x +#define READ_IMAGE_2D_XY read_imagef(input, sampler, (RET_TYPE_2) (IMAGE_INDEX_X, IMAGE_INDEX_Y)).x + +#define READ_IMAGE_3D_X read_imagef(input, sampler, (RET_TYPE_4) (IMAGE_INDEX_X, GIDY, GIDZ, C_ZERO)).x +#define READ_IMAGE_3D_Y read_imagef(input, sampler, (RET_TYPE_4) (GIDX, IMAGE_INDEX_Y, GIDZ, C_ZERO)).x +#define READ_IMAGE_3D_Z read_imagef(input, sampler, (RET_TYPE_4) (GIDX, GIDY, IMAGE_INDEX_Z, C_ZERO)).x +#define READ_IMAGE_3D_XY read_imagef(input, sampler, (RET_TYPE_4) (IMAGE_INDEX_X, IMAGE_INDEX_Y, GIDZ, C_ZERO)).x +#define READ_IMAGE_3D_XZ read_imagef(input, sampler, (RET_TYPE_4) (IMAGE_INDEX_X, GIDY, IMAGE_INDEX_Z, C_ZERO)).x +#define READ_IMAGE_3D_YZ read_imagef(input, sampler, (RET_TYPE_4) (GIDX, IMAGE_INDEX_Y, IMAGE_INDEX_Z, C_ZERO)).x +#define READ_IMAGE_3D_XYZ read_imagef(input, sampler, (RET_TYPE_4) (IMAGE_INDEX_X, IMAGE_INDEX_Y, IMAGE_INDEX_Z, C_ZERO)).x + + +// NOTE: pyopencl and OpenCL < 1.2 do not support image1d_t +#if FILTER_DIMS == 1 + #define FILTER_TYPE image2d_t + #define READ_FILTER_VAL(j) READ_FILTER_1D(j) +#elif FILTER_DIMS == 2 + #define FILTER_TYPE image2d_t + #define READ_FILTER_VAL(jx, jy) READ_FILTER_2D(jx, jy) +#elif FILTER_DIMS == 3 + #define FILTER_TYPE image3d_t + #define READ_FILTER_VAL(jx, jy, jz) READ_FILTER_3D(jx, jy, jz) +#endif + +#if IMAGE_DIMS == 1 + #define IMAGE_TYPE image2d_t + #define READ_IMAGE_X READ_IMAGE_1D +#elif IMAGE_DIMS == 2 + #define IMAGE_TYPE image2d_t + #define READ_IMAGE_X READ_IMAGE_2D_X + #define READ_IMAGE_Y READ_IMAGE_2D_Y + #define READ_IMAGE_XY READ_IMAGE_2D_XY +#elif IMAGE_DIMS == 3 + #define IMAGE_TYPE image3d_t + #define READ_IMAGE_X READ_IMAGE_3D_X + #define READ_IMAGE_Y READ_IMAGE_3D_Y + #define READ_IMAGE_Z READ_IMAGE_3D_Z + #define READ_IMAGE_XY READ_IMAGE_3D_XY + #define READ_IMAGE_XZ READ_IMAGE_3D_XZ + #define READ_IMAGE_YZ READ_IMAGE_3D_YZ + #define READ_IMAGE_XYZ READ_IMAGE_3D_XYZ +#endif + + +// Get the center index of the filter, +// and the "half-Left" and "half-Right" lengths. +// In the case of an even-sized filter, the center is shifted to the left. +#define GET_CENTER_HL(hlen){\ + if (hlen & 1) {\ + c = hlen/2;\ + hL = c;\ + hR = c;\ + }\ + else {\ + c = hlen/2 - 1;\ + hL = c;\ + hR = c+1;\ + }\ +}\ + + + +/******************************************************************************/ +/**************************** 1D Convolution **********************************/ +/******************************************************************************/ + +#if FILTER_DIMS == 1 +// Convolution with 1D kernel along axis "X" (fast dimension) +// Works for batched 1D on 2D and batched 2D on 3D, along axis "X". +__kernel void convol_1D_X_tex( + read_only IMAGE_TYPE input, + __global float * output, + read_only FILTER_TYPE filter, + int Lx, // filter size + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(Lx); + float sum = 0.0f; + + for (int jx = 0; jx <= hR+hL; jx++) { + sum += READ_IMAGE_X * READ_FILTER_VAL(jx); + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} + + +#if IMAGE_DIMS >= 2 +// Convolution with 1D kernel along axis "Y" +// Works for batched 1D on 2D and batched 2D on 3D, along axis "Y". +__kernel void convol_1D_Y_tex( + read_only IMAGE_TYPE input, + __global float * output, + read_only FILTER_TYPE filter, + int Lx, // filter size + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(Lx); + float sum = 0.0f; + + for (int jy = 0; jy <= hR+hL; jy++) { + sum += READ_IMAGE_Y * READ_FILTER_VAL(jy); + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} +#endif + +#if IMAGE_DIMS == 3 +// Convolution with 1D kernel along axis "Z" +// Works for batched 1D on 2D and batched 2D on 3D, along axis "Z". +__kernel void convol_1D_Z_tex( + read_only IMAGE_TYPE input, + __global float * output, + read_only FILTER_TYPE filter, + int Lx, // filter size + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(Lx); + float sum = 0.0f; + + for (int jz = 0; jz <= hR+hL; jz++) { + sum += READ_IMAGE_Z * READ_FILTER_VAL(jz); + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} +#endif +#endif + +/******************************************************************************/ +/**************************** 2D Convolution **********************************/ +/******************************************************************************/ + +#if IMAGE_DIMS >= 2 && FILTER_DIMS == 2 +// Convolution with 2D kernel +// Works for batched 2D on 3D. +__kernel void convol_2D_XY_tex( + read_only IMAGE_TYPE input, + __global float * output, + read_only FILTER_TYPE filter, + int Lx, // filter number of columns, + int Ly, // filter number of rows, + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(Lx); + float sum = 0.0f; + + for (int jy = 0; jy <= hR+hL; jy++) { + for (int jx = 0; jx <= hR+hL; jx++) { + sum += READ_IMAGE_XY * READ_FILTER_VAL(jx, jy); + } + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} +#endif + +#if IMAGE_DIMS == 3 && FILTER_DIMS == 2 +// Convolution with 2D kernel +// Works for batched 2D on 3D. +__kernel void convol_2D_XZ_tex( + read_only IMAGE_TYPE input, + __global float * output, + read_only FILTER_TYPE filter, + int Lx, // filter number of columns, + int Lz, // filter number of rows, + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(Lx); + float sum = 0.0f; + + for (int jz = 0; jz <= hR+hL; jz++) { + for (int jx = 0; jx <= hR+hL; jx++) { + sum += READ_IMAGE_XZ * READ_FILTER_VAL(jx, jz); + } + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} + + +// Convolution with 2D kernel +// Works for batched 2D on 3D. +__kernel void convol_2D_YZ_tex( + read_only IMAGE_TYPE input, + __global float * output, + read_only FILTER_TYPE filter, + int Lx, // filter number of columns, + int Lz, // filter number of rows, + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(Lx); + float sum = 0.0f; + + for (int jz = 0; jz <= hR+hL; jz++) { + for (int jy = 0; jy <= hR+hL; jy++) { + sum += READ_IMAGE_YZ * READ_FILTER_VAL(jy, jz); + } + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} +#endif + + +/******************************************************************************/ +/**************************** 3D Convolution **********************************/ +/******************************************************************************/ + +#if IMAGE_DIMS == 3 && FILTER_DIMS == 3 +// Convolution with 3D kernel +__kernel void convol_3D_XYZ_tex( + read_only IMAGE_TYPE input, + __global float * output, + read_only FILTER_TYPE filter, + int Lx, // filter number of columns, + int Ly, // filter number of rows, + int Lz, // filter number of rows, + int Nx, // input/output number of columns + int Ny, // input/output number of rows + int Nz // input/output depth +) +{ + uint gidx = get_global_id(0); + uint gidy = get_global_id(1); + uint gidz = get_global_id(2); + if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return; + + int c, hL, hR; + GET_CENTER_HL(Lx); + float sum = 0.0f; + + for (int jz = 0; jz <= hR+hL; jz++) { + for (int jy = 0; jy <= hR+hL; jy++) { + for (int jx = 0; jx <= hR+hL; jx++) { + sum += READ_IMAGE_XYZ * READ_FILTER_VAL(jx, jy, jz); + } + } + } + output[(gidz*Ny + gidy)*Nx + gidx] = sum; +} +#endif diff --git a/silx/resources/opencl/sparse.cl b/silx/resources/opencl/sparse.cl new file mode 100644 index 0000000..29e09ad --- /dev/null +++ b/silx/resources/opencl/sparse.cl @@ -0,0 +1,84 @@ +/* + * Copyright (C) 2016-2019 European Synchrotron Radiation Facility + * Grenoble, France + * + * Permission is hereby granted, free of charge, to any person + * obtaining a copy of this software and associated documentation + * files (the "Software"), to deal in the Software without + * restriction, including without limitation the rights to use, + * copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the + * Software is furnished to do so, subject to the following + * conditions: + * + * The above copyright notice and this permission notice shall be + * included in all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, + * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES + * OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND + * NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT + * HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, + * WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR + * OTHER DEALINGS IN THE SOFTWARE. + */ + +#ifndef IMAGE_WIDTH + #error "Please define IMAGE_WIDTH parameter" +#endif + +/** + * Densify a matric from CSR format to "dense" 2D format. + * The input CSR data consists in 3 arrays: (data, ind, iptr). + * The output array is a 2D array of dimensions IMAGE_WIDTH * image_height. + * + * data: 1D array containing the nonzero data items. + * ind: 1D array containing the column indices of the nonzero data items. + * iptr: 1D array containing indirection indices, such that range + * [iptr[i], iptr[i+1]-1] of "data" and "ind" contain the relevant data + * of output row "i". + * output: 2D array containing the densified data. + * image_height: height (number of rows) of the output data. +**/ + +kernel void densify_csr( + const global float* data, + const global int* ind, + const global int* iptr, + global float* output, + int image_height +) +{ + uint tid = get_local_id(0); + uint row_idx = get_global_id(1); + if ((tid >= IMAGE_WIDTH) || (row_idx >= image_height)) return; + + local float line[IMAGE_WIDTH]; + + // Memset + //~ #pragma unroll + for (int k = 0; tid+k < IMAGE_WIDTH; k += get_local_size(0)) { + if (tid+k >= IMAGE_WIDTH) break; + line[tid+k] = 0.0f; + } + barrier(CLK_LOCAL_MEM_FENCE); + + + uint start = iptr[row_idx], end = iptr[row_idx+1]; + //~ #pragma unroll + for (int k = start; k < end; k += get_local_size(0)) { + // Current work group handles one line of the final array + // on the current line, write data[start+tid] at column index ind[start+tid] + if (k+tid < end) + line[ind[k+tid]] = data[k+tid]; + } + barrier(CLK_LOCAL_MEM_FENCE); + + // write the current line (shared mem) into the output array (global mem) + //~ #pragma unroll + for (int k = 0; tid+k < IMAGE_WIDTH; k += get_local_size(0)) { + output[row_idx*IMAGE_WIDTH + tid+k] = line[tid+k]; + if (k+tid >= IMAGE_WIDTH) return; + } +} |
