diff options
Diffstat (limited to 'silx/opencl/processing.py')
| -rw-r--r-- | silx/opencl/processing.py | 61 |
1 files changed, 51 insertions, 10 deletions
diff --git a/silx/opencl/processing.py b/silx/opencl/processing.py index ca46701..1997a55 100644 --- a/silx/opencl/processing.py +++ b/silx/opencl/processing.py @@ -4,7 +4,7 @@ # Project: S I L X project # https://github.com/silx-kit/silx # -# Copyright (C) European Synchrotron Radiation Facility, Grenoble, France +# Copyright (C) 2012-2017 European Synchrotron Radiation Facility, Grenoble, France # # Principal author: Jérôme Kieffer (Jerome.Kieffer@ESRF.eu) # @@ -41,7 +41,7 @@ __author__ = "Jerome Kieffer" __contact__ = "Jerome.Kieffer@ESRF.eu" __license__ = "MIT" __copyright__ = "European Synchrotron Radiation Facility, Grenoble, France" -__date__ = "03/02/2017" +__date__ = "03/10/2017" __status__ = "stable" @@ -51,14 +51,34 @@ import gc from collections import namedtuple import numpy import threading -from .common import ocl, pyopencl, release_cl_buffers +from .common import ocl, pyopencl, release_cl_buffers, kernel_workgroup_size from .utils import concatenate_cl_kernel BufferDescription = namedtuple("BufferDescription", ["name", "size", "dtype", "flags"]) EventDescription = namedtuple("EventDescription", ["name", "event"]) -logger = logging.getLogger("pyFAI.opencl.processing") +logger = logging.getLogger(__name__) + + +class KernelContainer(object): + """Those object holds a copy of all kernels accessible as attributes""" + + def __init__(self, program): + """Constructor of the class + + :param program: the OpenCL program as generated by PyOpenCL + """ + for kernel in program.all_kernels(): + self.__setattr__(kernel.function_name, kernel) + + def get_kernels(self): + "return the dictionary with all kernels" + return self.__dict__.copy() + + def get_kernel(self, name): + "get a kernel from its name" + return self.__dict__.get(name) class OpenclProcessing(object): @@ -85,9 +105,10 @@ class OpenclProcessing(object): :param devicetype: type of device, can be "CPU", "GPU", "ACC" or "ALL" :param platformid: integer with the platform_identifier, as given by clinfo :param deviceid: Integer with the device identifier, as given by clinfo - :param block_size: preferred workgroup size, may vary depending on the outpcome of the compilation - :param profile: switch on profiling to be able to profile at the kernel level, - store profiling elements (makes code slightly slower) + :param block_size: preferred workgroup size, may vary depending on the + out come of the compilation + :param profile: switch on profiling to be able to profile at the kernel + level, store profiling elements (makes code slightly slower) """ self.sem = threading.Semaphore() self.profile = None @@ -108,6 +129,7 @@ class OpenclProcessing(object): self.set_profiling(profile) self.block_size = block_size self.program = None + self.kernels = None def __del__(self): """Destructor: release all buffers and programs @@ -145,8 +167,8 @@ class OpenclProcessing(object): ualloc = 0 for buf in buffers: ualloc += numpy.dtype(buf.dtype).itemsize * buf.size - logger.info("%.3fMB are needed on device which has %.3fMB", - ualloc / 1.0e6, self.device.memory / 1.0e6) + logger.info("%.3fMB are needed on device: %s, which has %.3fMB", + ualloc / 1.0e6, self.device, self.device.memory / 1.0e6) if ualloc >= self.device.memory: raise MemoryError("Fatal error in allocate_buffers. Not enough " @@ -157,13 +179,29 @@ class OpenclProcessing(object): try: for buf in buffers: size = numpy.dtype(buf.dtype).itemsize * buf.size - mem[buf.name] = pyopencl.Buffer(self.ctx, buf.flags, size) + mem[buf.name] = pyopencl.Buffer(self.ctx, buf.flags, int(size)) except pyopencl.MemoryError as error: release_cl_buffers(mem) raise MemoryError(error) self.cl_mem.update(mem) + def add_to_cl_mem(self, parrays): + """ + Add pyopencl.array, which are allocated by pyopencl, to self.cl_mem. + This should be used before calling allocate_buffers(). + + :param parrays: a dictionary of `pyopencl.array.Array` or `pyopencl.Buffer` + """ + mem = self.cl_mem + for name, parr in parrays.items(): + mem[name] = parr + self.cl_mem.update(mem) + + def check_workgroup_size(self, kernel_name): + kernel = self.kernels.get_kernel(kernel_name) + self.compiletime_workgroup_size = kernel_workgroup_size(self.program, kernel) + def free_buffers(self): """free all device.memory allocated on the device """ @@ -199,12 +237,15 @@ class OpenclProcessing(object): self.program = pyopencl.Program(self.ctx, kernel_src).build(options=compile_options) except (pyopencl.MemoryError, pyopencl.LogicError) as error: raise MemoryError(error) + else: + self.kernels = KernelContainer(self.program) def free_kernels(self): """Free all kernels """ for kernel in self.cl_kernel_args: self.cl_kernel_args[kernel] = [] + self.kernels = None self.program = None def set_profiling(self, value=True): |