diff options
Diffstat (limited to 'silx/opencl/statistics.py')
| -rw-r--r-- | silx/opencl/statistics.py | 44 |
1 files changed, 31 insertions, 13 deletions
diff --git a/silx/opencl/statistics.py b/silx/opencl/statistics.py index bd8e7b7..a96ee33 100644 --- a/silx/opencl/statistics.py +++ b/silx/opencl/statistics.py @@ -29,13 +29,10 @@ large data where numpy is not very efficient. """ -from __future__ import absolute_import, print_function, with_statement, division - - __author__ = "Jerome Kieffer" __license__ = "MIT" -__date__ = "11/01/2019" -__copyright__ = "2012-2017, ESRF, Grenoble" +__date__ = "19/05/2021" +__copyright__ = "2012-2019, ESRF, Grenoble" __contact__ = "jerome.kieffer@esrf.fr" import logging @@ -128,7 +125,7 @@ class Statistics(OpenclProcessing): self.kernel_files, "-D NIMAGE=%i" % self.size) compiler_options = self.get_compiler_options(x87_volatile=True) - src = concatenate_cl_kernel(("kahan.cl", "statistics.cl")) + src = concatenate_cl_kernel(("doubleword.cl", "statistics.cl")) self.reduction_comp = ReductionKernel(self.ctx, dtype_out=float8, neutral=zero8, @@ -146,6 +143,19 @@ class Statistics(OpenclProcessing): preamble=src, options=compiler_options) + if "cl_khr_fp64" in self.device.extensions: + self.reduction_double = ReductionKernel(self.ctx, + dtype_out=float8, + neutral=zero8, + map_expr="map_statistics(data, i)", + reduce_expr="reduce_statistics_double(a,b)", + arguments="__global float *data", + preamble=src, + options=compiler_options) + else: + logger.info("Device %s does not support double-precision arithmetics, fall-back on compensated one", self.device) + self.reduction_double = self.reduction_comp + def send_buffer(self, data, dest): """ Send a numpy array to the device, including the cast on the device if @@ -154,7 +164,7 @@ class Statistics(OpenclProcessing): :param numpy.ndarray data: numpy array with data :param dest: name of the buffer as registered in the class """ - + logger.info("send data to %s", dest) dest_type = numpy.dtype([i.dtype for i in self.buffers if i.name == dest][0]) events = [] if (data.dtype == dest_type) or (data.dtype.itemsize > dest_type.itemsize): @@ -173,8 +183,8 @@ class Statistics(OpenclProcessing): self.cl_mem["raw"].data, self.cl_mem[dest].data) events += [ - EventDescription("copy H->D %s" % dest, copy_image), - EventDescription("cast to float", cast_to_float) + EventDescription("copy H->D raw", copy_image), + EventDescription(f"cast to float {dest}", cast_to_float) ] if self.profile: self.events += events @@ -193,16 +203,24 @@ class Statistics(OpenclProcessing): size = data.size assert size <= self.size, "size is OK" events = [] + if comp is True: + comp = "comp" + elif comp is False: + comp = "single" + else: + comp = comp.lower() with self.sem: self.send_buffer(data, "converted") - if comp: - reduction = self.reduction_comp - else: + if comp in ("single", "fp32", "float32"): reduction = self.reduction_simple + elif comp in ("double", "fp64", "float64"): + reduction = self.reduction_double + else: + reduction = self.reduction_comp res_d, evt = reduction(self.cl_mem["converted"][:self.size], queue=self.queue, return_event=True) - events.append(EventDescription("statistical reduction %s" % ("comp"if comp else "simple"), evt)) + events.append(EventDescription(f"statistical reduction {comp}", evt)) if self.profile: self.events += events res_h = res_d.get() |