From e77e31d11588b1149e85ae42de1598c831da1d2f Mon Sep 17 00:00:00 2001 From: Edgar Gutierrez Fernandez Date: Thu, 14 Nov 2024 17:03:33 +0100 Subject: [PATCH 1/9] simplify fiber integrator --- src/pyFAI/integrator/fiber.py | 173 +++++----------------------------- 1 file changed, 25 insertions(+), 148 deletions(-) diff --git a/src/pyFAI/integrator/fiber.py b/src/pyFAI/integrator/fiber.py index d92eb022e..ac9442ba8 100644 --- a/src/pyFAI/integrator/fiber.py +++ b/src/pyFAI/integrator/fiber.py @@ -30,7 +30,7 @@ __contact__ = "edgar.gutierrez-fernandez@esr.fr" __license__ = "MIT" __copyright__ = "European Synchrotron Radiation Facility, Grenoble, France" -__date__ = "13/11/2024" +__date__ = "14/11/2024" __status__ = "stable" __docformat__ = 'restructuredtext' @@ -155,6 +155,8 @@ def integrate_fiber(self, data, :param pyFAI.units.UnitFiber/str unit_ip: unit to describe the in-plane axis. If not provided, it takes qip_nm^-1 :param list ip_range: The lower and upper range of the in-plane unit. If not provided, range is simply (data.min(), data.max()). Values outside the range are ignored. Optional. :param bool vertical_integration: If True, integrates along unit_ip; if False, integrates along unit_oop + :param incident_angle: tilting of the sample towards the beam (analog to rot2): in radians + :param tilt_angle: tilting of the sample orthogonal to the beam direction (analog to rot3): in radians :param int sample_orientation: 1-4, four different orientation of the fiber axis regarding the detector main axis, from 1 to 4 is +90º :param str filename: output filename in 2/3 column ascii format :param bool correctSolidAngle: correct for solid angle of each pixel if True @@ -183,8 +185,14 @@ def integrate_fiber(self, data, unit_ip = unit_ip or 'qip_nm^-1' unit_oop = unit_oop or 'qoop_nm^-1' - unit_ip = parse_fiber_unit(unit=unit_ip, sample_orientation=sample_orientation) - unit_oop = parse_fiber_unit(unit=unit_oop, sample_orientation=unit_ip.sample_orientation) + unit_ip = parse_fiber_unit(unit=unit_ip, + incident_angle=kwargs.get('incident_angle', None), + tilt_angle=kwargs.get('tilt_angle', None), + sample_orientation=sample_orientation) + unit_oop = parse_fiber_unit(unit=unit_oop, + incident_angle=unit_ip.incident_angle, + tilt_angle=unit_ip.tilt_angle, + sample_orientation=unit_ip.sample_orientation) self.reset_integrator(incident_angle=unit_ip.incident_angle, tilt_angle=unit_ip.tilt_angle, @@ -262,78 +270,7 @@ def integrate_fiber(self, data, return result - def integrate_grazing_incidence(self, data, - npt_oop=None, unit_oop=None, oop_range=None, - npt_ip=None, unit_ip=None, ip_range=None, - vertical_integration = True, - incident_angle=None, tilt_angle=None, sample_orientation=None, - filename=None, - correctSolidAngle=True, - mask=None, dummy=None, delta_dummy=None, - polarization_factor=None, dark=None, flat=None, - method=("no", "histogram", "cython"), - normalization_factor=1.0, - **kwargs): - """Calculate the integrated profile curve along a specific FiberUnit, additional inputs for incident angle, tilt angle and sample_orientation - - :param ndarray data: 2D array from the Detector/CCD camera - :param int npt_oop: number of points to be used along the out-of-plane axis - :param pyFAI.units.UnitFiber/str unit_oop: unit to describe the out-of-plane axis. If not provided, it takes qoop_nm^-1 - :param list oop_range: The lower and upper range of the out-of-plane unit. If not provided, range is simply (data.min(), data.max()). Values outside the range are ignored. Optional. - :param int npt_ip: number of points to be used along the in-plane axis - :param pyFAI.units.UnitFiber/str unit_ip: unit to describe the in-plane axis. If not provided, it takes qip_nm^-1 - :param list ip_range: The lower and upper range of the in-plane unit. If not provided, range is simply (data.min(), data.max()). Values outside the range are ignored. Optional. - :param bool vertical_integration: If True, integrates along unit_ip; if False, integrates along unit_oop - :param incident_angle: tilting of the sample towards the beam (analog to rot2): in radians - :param tilt_angle: tilting of the sample orthogonal to the beam direction (analog to rot3): in radians - :param int sample_orientation: 1-4, four different orientation of the fiber axis regarding the detector main axis, from 1 to 4 is +90º - :param str filename: output filename in 2/3 column ascii format - :param bool correctSolidAngle: correct for solid angle of each pixel if True - :param ndarray mask: array (same size as image) with 1 for masked pixels, and 0 for valid pixels - :param float dummy: value for dead/masked pixels - :param float delta_dummy: precision for dummy value - :param float polarization_factor: polarization factor between -1 (vertical) and +1 (horizontal). - * 0 for circular polarization or random, - * None for no correction, - * True for using the former correction - :param ndarray dark: dark noise image - :param ndarray flat: flat field image - :param IntegrationMethod method: IntegrationMethod instance or 3-tuple with (splitting, algorithm, implementation) - :param float normalization_factor: Value of a normalization monitor - :return: chi bins center positions and regrouped intensity - :rtype: Integrate1dResult - """ - deprecated_params = get_deprecated_params_1d(**kwargs) - npt_oop = deprecated_params.get('npt_oop', None) or npt_oop - npt_ip = deprecated_params.get('npt_ip', None) or npt_ip - unit_oop = deprecated_params.get('unit_oop', None) or unit_oop - unit_ip = deprecated_params.get('unit_ip', None) or unit_ip - oop_range = deprecated_params.get('oop_range', None) or oop_range - ip_range = deprecated_params.get('ip_range', None) or ip_range - vertical_integration = deprecated_params.get('vertical_integration', None) or vertical_integration - - unit_ip = unit_ip or 'qip_nm^-1' - unit_oop = unit_oop or 'qoop_nm^-1' - unit_ip = parse_fiber_unit(unit=unit_ip, incident_angle=incident_angle, tilt_angle=tilt_angle, sample_orientation=sample_orientation) - unit_oop = parse_fiber_unit(unit=unit_oop, incident_angle=unit_ip.incident_angle, tilt_angle=unit_ip.tilt_angle, sample_orientation=unit_ip.sample_orientation) - - self.reset_integrator(incident_angle=unit_ip.incident_angle, - tilt_angle=unit_ip.tilt_angle, - sample_orientation=unit_ip.sample_orientation) - - return self.integrate_fiber(data=data, - npt_oop=npt_oop, unit_oop=unit_oop, oop_range=oop_range, - npt_ip=npt_ip, unit_ip=unit_ip, ip_range=ip_range, - vertical_integration=vertical_integration, - sample_orientation=sample_orientation, - filename=filename, - correctSolidAngle=correctSolidAngle, - mask=mask, dummy=dummy, delta_dummy=delta_dummy, - polarization_factor=polarization_factor, dark=dark, flat=flat, - method=method, - normalization_factor=normalization_factor, - ) - + integrate_grazing_incidence = integrate_fiber integrate1d_grazing_incidence = integrate_grazing_incidence integrate1d_fiber = integrate_fiber @@ -356,6 +293,8 @@ def integrate2d_fiber(self, data, :param int npt_oop: number of points to be used along the out-of-plane axis :param pyFAI.units.UnitFiber/str unit_oop: unit to describe the out-of-plane axis. If not provided, it takes qoop_nm^-1 :param list oop_range: The lower and upper range of the out-of-plane unit. If not provided, range is simply (data.min(), data.max()). Values outside the range are ignored. Optional. + :param incident_angle: tilting of the sample towards the beam (analog to rot2): in radians + :param tilt_angle: tilting of the sample orthogonal to the beam direction (analog to rot3): in radians :param int sample_orientation: 1-4, four different orientation of the fiber axis regarding the detector main axis, from 1 to 4 is +90º :param str filename: output filename in 2/3 column ascii format :param bool correctSolidAngle: correct for solid angle of each pixel if True @@ -383,8 +322,15 @@ def integrate2d_fiber(self, data, unit_ip = unit_ip or 'qip_nm^-1' unit_oop = unit_oop or 'qoop_nm^-1' - unit_ip = parse_fiber_unit(unit=unit_ip, sample_orientation=sample_orientation) - unit_oop = parse_fiber_unit(unit=unit_oop, sample_orientation=unit_ip.sample_orientation) + unit_ip = parse_fiber_unit(unit=unit_ip, + sample_orientation=sample_orientation, + incident_angle=kwargs.get('incident_angle', None), + tilt_angle=kwargs.get('tilt_angle', None), + ) + unit_oop = parse_fiber_unit(unit=unit_oop, + incident_angle=unit_ip.incident_angle, + tilt_angle=unit_ip.tilt_angle, + sample_orientation=unit_ip.sample_orientation) self.reset_integrator(incident_angle=unit_ip.incident_angle, tilt_angle=unit_ip.tilt_angle, @@ -403,74 +349,5 @@ def integrate2d_fiber(self, data, azimuth_range=oop_range, unit=(unit_ip, unit_oop), filename=filename) - - def integrate2d_grazing_incidence(self, data, - npt_ip=1000, unit_ip=None, ip_range=None, - npt_oop=1000, unit_oop=None, oop_range=None, - incident_angle=None, tilt_angle=None, sample_orientation=None, - filename=None, - correctSolidAngle=True, - mask=None, dummy=None, delta_dummy=None, - polarization_factor=None, dark=None, flat=None, - method=("no", "histogram", "cython"), - normalization_factor=1.0, **kwargs): - """Reshapes the data pattern as a function of two FiberUnits, additional inputs for incident angle, tilt angle and sample_orientation - - :param ndarray data: 2D array from the Detector/CCD camera - :param int npt_ip: number of points to be used along the in-plane axis - :param pyFAI.units.UnitFiber/str unit_ip: unit to describe the in-plane axis. If not provided, it takes qip_nm^-1 - :param list ip_range: The lower and upper range of the in-plane unit. If not provided, range is simply (data.min(), data.max()). Values outside the range are ignored. Optional. - :param int npt_oop: number of points to be used along the out-of-plane axis - :param pyFAI.units.UnitFiber/str unit_oop: unit to describe the out-of-plane axis. If not provided, it takes qoop_nm^-1 - :param list oop_range: The lower and upper range of the out-of-plane unit. If not provided, range is simply (data.min(), data.max()). Values outside the range are ignored. Optional. - :param incident_angle: tilting of the sample towards the beam (analog to rot2): in radians - :param tilt_angle: tilting of the sample orthogonal to the beam direction (analog to rot3): in radians - :param int sample_orientation: 1-4, four different orientation of the fiber axis regarding the detector main axis, from 1 to 4 is +90º - :param str filename: output filename in 2/3 column ascii format - :param bool correctSolidAngle: correct for solid angle of each pixel if True - :param ndarray mask: array (same size as image) with 1 for masked pixels, and 0 for valid pixels - :param float dummy: value for dead/masked pixels - :param float delta_dummy: precision for dummy value - :param float polarization_factor: polarization factor between -1 (vertical) and +1 (horizontal). - * 0 for circular polarization or random, - * None for no correction, - * True for using the former correction - :param ndarray dark: dark noise image - :param ndarray flat: flat field image - :param IntegrationMethod method: IntegrationMethod instance or 3-tuple with (splitting, algorithm, implementation) - :param float normalization_factor: Value of a normalization monitor - :return: regrouped intensity and unit arrays - :rtype: Integrate2dResult - """ - deprecated_params = get_deprecated_params_2d(**kwargs) - npt_oop = deprecated_params.get('npt_oop', None) or npt_oop - npt_ip = deprecated_params.get('npt_ip', None) or npt_ip - unit_oop = deprecated_params.get('unit_oop', None) or unit_oop - unit_ip = deprecated_params.get('unit_ip', None) or unit_ip - oop_range = deprecated_params.get('oop_range', None) or oop_range - ip_range = deprecated_params.get('ip_range', None) or ip_range - - unit_ip = unit_ip or 'qip_nm^-1' - unit_oop = unit_oop or 'qoop_nm^-1' - unit_ip = parse_fiber_unit(unit=unit_ip, incident_angle=incident_angle, tilt_angle=tilt_angle, sample_orientation=sample_orientation) - unit_oop = parse_fiber_unit(unit=unit_oop, incident_angle=unit_ip.incident_angle, tilt_angle=unit_ip.tilt_angle, sample_orientation=unit_ip.sample_orientation) - - self.reset_integrator(incident_angle=unit_ip.incident_angle, - tilt_angle=unit_ip.tilt_angle, - sample_orientation=unit_ip.sample_orientation) - - if (isinstance(method, (tuple, list)) and method[0] != "no") or (isinstance(method, IntegrationMethod) and method.split != "no"): - logger.warning(f"Method {method} is using a pixel-splitting scheme. GI integration should be use WITHOUT PIXEL-SPLITTING! The results could be wrong!") - - return self.integrate2d_fiber(data=data, npt_ip=npt_ip, npt_oop=npt_oop, - unit_ip=unit_ip, unit_oop=unit_oop, - ip_range=ip_range, - oop_range=oop_range, - sample_orientation=sample_orientation, - filename=filename, - correctSolidAngle=correctSolidAngle, - mask=mask, dummy=dummy, delta_dummy=delta_dummy, - polarization_factor=polarization_factor, dark=dark, flat=flat, - method=method, - normalization_factor=normalization_factor, - ) + + integrate2d_grazing_incidence = integrate2d_fiber \ No newline at end of file From a249f163743bdd6a444379fc6bfa6e0a7350136b Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Thu, 14 Nov 2024 16:04:15 +0000 Subject: [PATCH 2/9] [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --- src/pyFAI/integrator/fiber.py | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/src/pyFAI/integrator/fiber.py b/src/pyFAI/integrator/fiber.py index ac9442ba8..4ac92f916 100644 --- a/src/pyFAI/integrator/fiber.py +++ b/src/pyFAI/integrator/fiber.py @@ -185,7 +185,7 @@ def integrate_fiber(self, data, unit_ip = unit_ip or 'qip_nm^-1' unit_oop = unit_oop or 'qoop_nm^-1' - unit_ip = parse_fiber_unit(unit=unit_ip, + unit_ip = parse_fiber_unit(unit=unit_ip, incident_angle=kwargs.get('incident_angle', None), tilt_angle=kwargs.get('tilt_angle', None), sample_orientation=sample_orientation) @@ -322,7 +322,7 @@ def integrate2d_fiber(self, data, unit_ip = unit_ip or 'qip_nm^-1' unit_oop = unit_oop or 'qoop_nm^-1' - unit_ip = parse_fiber_unit(unit=unit_ip, + unit_ip = parse_fiber_unit(unit=unit_ip, sample_orientation=sample_orientation, incident_angle=kwargs.get('incident_angle', None), tilt_angle=kwargs.get('tilt_angle', None), @@ -349,5 +349,5 @@ def integrate2d_fiber(self, data, azimuth_range=oop_range, unit=(unit_ip, unit_oop), filename=filename) - - integrate2d_grazing_incidence = integrate2d_fiber \ No newline at end of file + + integrate2d_grazing_incidence = integrate2d_fiber From de78dcf6e3b32564def4e5ef3c8db869e2fb23e3 Mon Sep 17 00:00:00 2001 From: Jerome Kieffer Date: Tue, 19 Nov 2024 18:02:09 +0100 Subject: [PATCH 3/9] Explicitely Disable fp64 Work aroung a bug in Apple Silicon GPU driver --- src/pyFAI/opencl/__init__.py | 20 +++++++++++++++++++- src/pyFAI/opencl/azim_csr.py | 16 ++++++++++------ src/pyFAI/opencl/azim_hist.py | 20 ++++++++++++-------- src/pyFAI/opencl/azim_lut.py | 19 ++++++++++--------- src/pyFAI/opencl/peak_finder.py | 17 +++++++---------- src/pyFAI/opencl/preproc.py | 10 ++++++---- src/pyFAI/opencl/test/test_preproc.py | 5 +++-- src/pyFAI/resources/openCL/preprocess.cl | 4 ++-- 8 files changed, 69 insertions(+), 42 deletions(-) diff --git a/src/pyFAI/opencl/__init__.py b/src/pyFAI/opencl/__init__.py index 15c16ff37..49bdccf52 100644 --- a/src/pyFAI/opencl/__init__.py +++ b/src/pyFAI/opencl/__init__.py @@ -36,7 +36,7 @@ __contact__ = "Jerome.Kieffer@ESRF.eu" __license__ = "MIT" __copyright__ = "2012-2024 European Synchrotron Radiation Facility, Grenoble, France" -__date__ = "15/04/2024" +__date__ = "19/11/2024" __status__ = "stable" import os @@ -79,6 +79,24 @@ def get_x87_volatile_option(ctx): else: return "" +def get_compiler_options(ctx, x87_volatile=False, apple_gpu=False): + """Provide a set of common compiler options to work around known bugs: + + :x87_volatile: set to true to declare all x87 operation as volatile, needed on PoCL x86 32bits + :apple_gpu: redefine the cl_khr_fp64 to zero when the device is Apple GPU + which wrongly declares fp64 compatibility. See #2339 + :return: compilation directive as string. + """ + + if x87_volatile: + options = get_x87_volatile_option(ctx) + else: + options = "" + if apple_gpu: + fp64_support = 1 if "cl_khr_fp64" in ctx.devices[0].extensions else 0 + options += f" -D cl_khr_fp64={fp64_support}" + return options.strip() + def dtype_converter(dtype): "convert a numpy dtype as a int8" diff --git a/src/pyFAI/opencl/azim_csr.py b/src/pyFAI/opencl/azim_csr.py index 62c24b8de..5a8c61ff5 100644 --- a/src/pyFAI/opencl/azim_csr.py +++ b/src/pyFAI/opencl/azim_csr.py @@ -28,7 +28,7 @@ __authors__ = ["Jérôme Kieffer", "Giannis Ashiotis"] __license__ = "MIT" -__date__ = "06/09/2024" +__date__ = "19/11/2024" __copyright__ = "ESRF, Grenoble" __contact__ = "jerome.kieffer@esrf.fr" @@ -275,11 +275,15 @@ def compile_kernels(self, kernel_file=None): kernel_file = kernel_file or self.kernel_files[-1] kernels = self.kernel_files[:-1] + [kernel_file] - compile_options = f"-D NBINS={self.bins} -D NIMAGE={self.size}" - default_compiler_options = self.get_compiler_options(x87_volatile=True) - if default_compiler_options: - compile_options += " " + default_compiler_options - OpenclProcessing.compile_kernels(self, kernels, compile_options) + try: + compile_options = self.get_compiler_options(x87_volatile=True, apple_gpu=True) + except (AttributeError, TypeError): # Silx version too old + logger.warning("Please upgrade to silx v2.1+") + from . import get_compiler_options + compile_options = get_compiler_options(self.ctx, x87_volatile=True, apple_gpu=True) + + compile_options += f" -D NBINS={self.bins} -D NIMAGE={self.size}" + OpenclProcessing.compile_kernels(self, kernels, compile_options.strip()) for kernel_name in self.kernels.__dict__: if kernel_name.startswith("_"): continue diff --git a/src/pyFAI/opencl/azim_hist.py b/src/pyFAI/opencl/azim_hist.py index b8a1dcd91..cdaaf1079 100644 --- a/src/pyFAI/opencl/azim_hist.py +++ b/src/pyFAI/opencl/azim_hist.py @@ -32,7 +32,7 @@ """ __author__ = "Jérôme Kieffer" __license__ = "MIT" -__date__ = "25/04/2024" +__date__ = "19/11/2024" __copyright__ = "2012-2021, ESRF, Grenoble" __contact__ = "jerome.kieffer@esrf.fr" @@ -47,7 +47,7 @@ raise ImportError("pyopencl is not installed") from . import allocate_cl_buffers, release_cl_buffers, kernel_workgroup_size -from . import concatenate_cl_kernel, get_x87_volatile_option, processing, OpenclProcessing +from . import concatenate_cl_kernel, processing, OpenclProcessing from ..containers import Integrate1dtpl, Integrate2dtpl, ErrorModel from ..utils.decorators import deprecated EventDescription = processing.EventDescription @@ -257,13 +257,17 @@ def compile_kernels(self, kernel_file=None): # concatenate all needed source files into a single openCL module kernel_file = kernel_file or self.kernel_files[-1] kernels = self.kernel_files[:-1] + [kernel_file] - default_compiler_options = get_x87_volatile_option(self.ctx) - compile_options = "-D NBINS=%i -D NIMAGE=%i -D WORKGROUP_SIZE=%i" % \ - (self.bins, self.size, self.BLOCK_SIZE) - if default_compiler_options: - compile_options += " " + default_compiler_options try: - OpenclProcessing.compile_kernels(self, kernels, compile_options) + compile_options = self.get_compiler_options(x87_volatile=True, apple_gpu=True) + except (AttributeError, TypeError): # Silx version too old + logger.warning("Please upgrade to silx v2.1+") + from . import get_compiler_options + compile_options = get_compiler_options(self.ctx, x87_volatile=True, apple_gpu=True) + + + compile_options += f" -D NBINS={self.bins} -D NIMAGE={self.size} -D WORKGROUP_SIZE={self.BLOCK_SIZE}" + try: + OpenclProcessing.compile_kernels(self, kernels, compile_options.strip()) except Exception as error: # This error may be related to issue #1219. Provides an ugly work around. if "cl_khr_int64_base_atomics" in self.ctx.devices[0].extensions: diff --git a/src/pyFAI/opencl/azim_lut.py b/src/pyFAI/opencl/azim_lut.py index b590b0359..f93f92f07 100644 --- a/src/pyFAI/opencl/azim_lut.py +++ b/src/pyFAI/opencl/azim_lut.py @@ -27,7 +27,7 @@ __author__ = "Jérôme Kieffer" __license__ = "MIT" -__date__ = "24/04/2024" +__date__ = "19/11/2024" __copyright__ = "2012-2024, ESRF, Grenoble" __contact__ = "jerome.kieffer@esrf.fr" @@ -205,14 +205,15 @@ def compile_kernels(self, kernel_file=None): # concatenate all needed source files into a single openCL module kernel_file = kernel_file or self.kernel_files[-1] kernels = self.kernel_files[:-1] + [kernel_file] - - compile_options = "-D NBINS=%i -D NIMAGE=%i -D NLUT=%i -D ON_CPU=%i" % \ - (self.bins, self.size, self.lut_size, int(self.device.type == "CPU")) - - default_compiler_options = self.get_compiler_options(x87_volatile=True) - if default_compiler_options: - compile_options += " " + default_compiler_options - OpenclProcessing.compile_kernels(self, kernels, compile_options) + try: + compile_options = self.get_compiler_options(x87_volatile=True, apple_gpu=True) + except (AttributeError, TypeError): # Silx version too old + logger.warning("Please upgrade to silx v2.1+") + from . import get_compiler_options + compile_options = get_compiler_options(self.ctx, x87_volatile=True, apple_gpu=True) + + compile_options += f" -D NBINS={self.bins} -D NIMAGE={self.size} -D NLUT={self.lut_size} -D ON_CPU={int(self.device.type == 'CPU')}" + OpenclProcessing.compile_kernels(self, kernels, compile_options.strip()) def set_kernel_arguments(self): """Tie arguments of OpenCL kernel-functions to the actual kernels diff --git a/src/pyFAI/opencl/peak_finder.py b/src/pyFAI/opencl/peak_finder.py index fc5c87850..761646def 100644 --- a/src/pyFAI/opencl/peak_finder.py +++ b/src/pyFAI/opencl/peak_finder.py @@ -29,7 +29,7 @@ __authors__ = ["Jérôme Kieffer"] __license__ = "MIT" -__date__ = "08/04/2024" +__date__ = "19/11/2024" __copyright__ = "2014-2023, ESRF, Grenoble" __contact__ = "jerome.kieffer@esrf.fr" @@ -40,7 +40,7 @@ from ..containers import SparseFrame, ErrorModel from ..utils import EPS32 from .azim_csr import OCL_CSR_Integrator, BufferDescription, EventDescription, mf, calc_checksum, pyopencl, OpenclProcessing -from . import get_x87_volatile_option, kernel_workgroup_size, dtype_converter +from . import kernel_workgroup_size, dtype_converter logger = logging.getLogger(__name__) @@ -908,15 +908,12 @@ def compile_kernels(self, kernel_file=None): kernels = self.kernel_files[:-1] + [kernel_file] try: - default_compiler_options = self.get_compiler_options(x87_volatile=True) - except AttributeError: # Silx version too old - logger.warning("Please upgrade to silx v0.10+") - default_compiler_options = get_x87_volatile_option(self.ctx) + compile_options = self.get_compiler_options(x87_volatile=True, apple_gpu=True) + except (AttributeError, TypeError): # Silx version too old + logger.warning("Please upgrade to silx v2.1+") + from . import get_compiler_options + compile_options = get_compiler_options(self.ctx, x87_volatile=True, apple_gpu=True) - if default_compiler_options: - compile_options = default_compiler_options - else: - compile_options = "" OpenclProcessing.compile_kernels(self, kernels, compile_options) for kernel_name, kernel in self.kernels.get_kernels().items(): wg = kernel_workgroup_size(self.program, kernel) diff --git a/src/pyFAI/opencl/preproc.py b/src/pyFAI/opencl/preproc.py index f6bc09681..bca8091c6 100644 --- a/src/pyFAI/opencl/preproc.py +++ b/src/pyFAI/opencl/preproc.py @@ -3,7 +3,7 @@ # Project: Azimuthal integration # https://github.com/silx-kit/pyFAI # -# Copyright (C) 2015-2018 European Synchrotron Radiation Facility, Grenoble, France +# Copyright (C) 2015-2024 European Synchrotron Radiation Facility, Grenoble, France # # Principal author: Jérôme Kieffer (Jerome.Kieffer@ESRF.eu) # @@ -31,7 +31,7 @@ __author__ = "Jérôme Kieffer" __license__ = "MIT" -__date__ = "23/10/2024" +__date__ = "19/11/2024" __copyright__ = "2015-2017, ESRF, Grenoble" __contact__ = "jerome.kieffer@esrf.fr" @@ -328,7 +328,7 @@ def set_kernel_arguments(self): ("output", self.cl_mem["output"]))) - def compile_kernels(self, kernel_files=None, compile_options=None): + def compile_kernels(self, kernel_files=None): """Call the OpenCL compiler :param kernel_files: list of path to the kernel @@ -336,7 +336,9 @@ def compile_kernels(self, kernel_files=None, compile_options=None): """ # concatenate all needed source files into a single openCL module kernel_files = kernel_files or self.kernel_files - compile_options = "-D NIMAGE=%i" % (self.size) + # Explicit handling of fp64 since Apple silicon compiler wrongly clams fp64 support see issue #2339 + fp64_support = 1 if "cl_khr_fp64" in self.ctx.devices[0].extensions else 0 + compile_options = f"-D NIMAGE={self.size} -D cl_khr_fp64={fp64_support}" OpenclProcessing.compile_kernels(self, kernel_files, compile_options) def send_buffer(self, data, dest, convert=True): diff --git a/src/pyFAI/opencl/test/test_preproc.py b/src/pyFAI/opencl/test/test_preproc.py index 0e4661f52..d4fbdb7e8 100644 --- a/src/pyFAI/opencl/test/test_preproc.py +++ b/src/pyFAI/opencl/test/test_preproc.py @@ -33,7 +33,7 @@ __contact__ = "jerome.kieffer@esrf.eu" __license__ = "MIT" __copyright__ = "European Synchrotron Radiation Facility, Grenoble, France" -__date__ = "04/10/2023" +__date__ = "19/11/2024" import logging import numpy @@ -57,7 +57,8 @@ def test_preproc(self): """ from ..preproc import preproc ary = numpy.arange(12).reshape(4,3) - for dtype in (numpy.uint8, numpy.int8, numpy.int16, numpy.uint16, numpy.uint32, numpy.int32, numpy.uint64, numpy.int64, numpy.float32): + for dtype in (numpy.uint8, numpy.int8, numpy.int16, numpy.uint16, numpy.uint32, numpy.int32):#, numpy.uint64, numpy.int64, numpy.float32): + import sys; sys.stderr.write(f"test {dtype}\n") self.assertEqual(abs(preproc(ary.astype(dtype),split_result=4)[..., 0]-ary).max(), 0, "Result OK for dtype {dtype}") diff --git a/src/pyFAI/resources/openCL/preprocess.cl b/src/pyFAI/resources/openCL/preprocess.cl index cf7400c12..a078be026 100644 --- a/src/pyFAI/resources/openCL/preprocess.cl +++ b/src/pyFAI/resources/openCL/preprocess.cl @@ -7,7 +7,7 @@ * Grenoble, France * * Principal authors: J. Kieffer (kieffer@esrf.fr) - * Last revision: 23/04/2024 + * Last revision: 19/11/2024 * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -203,7 +203,7 @@ static float _any2float(const global uchar* input, input[8*position+4],input[8*position+5], input[8*position+6],input[8*position+7]); value = convert_float(as_double(rval)); #else - if (get_global_id==0)printf("Doubleprecision arithmetics is not supported on this device !\n"); + if (get_global_id(0)==0)printf("Double precision arithmetics is not supported on this device !\n"); #endif } From d522f2bf86cf5fa035e24a83fb9084ae196551cc Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Tue, 19 Nov 2024 17:04:08 +0000 Subject: [PATCH 4/9] [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --- doc/source/publications.rst | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/doc/source/publications.rst b/doc/source/publications.rst index e77f2464e..9931a202c 100644 --- a/doc/source/publications.rst +++ b/doc/source/publications.rst @@ -34,11 +34,11 @@ Publications about pyFAI * *Application of signal separation to diffraction image compression and serial crystallography*; Jérôme Kieffer, Julien Orlans, Nicolas Coquelle, Samuel Debionne, Shibom Basu, Alejandro Homs, Gianluca Santonia and Daniele De Sanctis; - `Accepted `_ in **J. Applied Crystallography** (2024); + `Accepted `_ in **J. Applied Crystallography** (2024); In depth explainaion of sigma-clipping background assessment and error models. The latest paper should be the cited in publications using pyFAI. There are already 1400 publications referring to pyFAI, some of them in the most -prestigious scientific journals (Nature, PNAS, ...) and -40 other `applications `_ +prestigious scientific journals (Nature, PNAS, ...) and +40 other `applications `_ using pyFAI as a library. From 3327bedc71c2e5adc71414460c3ddcf26a34086d Mon Sep 17 00:00:00 2001 From: Jerome Kieffer Date: Tue, 19 Nov 2024 18:04:48 +0100 Subject: [PATCH 5/9] Typo --- src/pyFAI/resources/openCL/preprocess.cl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/pyFAI/resources/openCL/preprocess.cl b/src/pyFAI/resources/openCL/preprocess.cl index a078be026..a4ca007f0 100644 --- a/src/pyFAI/resources/openCL/preprocess.cl +++ b/src/pyFAI/resources/openCL/preprocess.cl @@ -147,7 +147,7 @@ s32_to_float(global int *array_int, } /* Function reading at the given position. - * Dtype is 1/-1 for char/uchar .... 8/-4 for int64/uint64 and 32/64 for float/double. + * Dtype is 1/-1 for char/uchar .... 8/-8 for int64/uint64 and 32/64 for float/double. */ static float _any2float(const global uchar* input, size_t position, From 8b232b904408f0dd60e8d0bcb16474385c8d979f Mon Sep 17 00:00:00 2001 From: Jerome Kieffer Date: Tue, 19 Nov 2024 18:08:25 +0100 Subject: [PATCH 6/9] Print warning message --- src/pyFAI/resources/openCL/preprocess.cl | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/src/pyFAI/resources/openCL/preprocess.cl b/src/pyFAI/resources/openCL/preprocess.cl index a4ca007f0..6d805be49 100644 --- a/src/pyFAI/resources/openCL/preprocess.cl +++ b/src/pyFAI/resources/openCL/preprocess.cl @@ -199,9 +199,13 @@ static float _any2float(const global uchar* input, } else if (dtype == 64){ #ifdef cl_khr_fp64 + #if cl_khr_fp64 uchar8 rval = (uchar8) (input[8*position],input[8*position+1], input[8*position+2],input[8*position+3], input[8*position+4],input[8*position+5], input[8*position+6],input[8*position+7]); value = convert_float(as_double(rval)); + #else + if (get_global_id(0)==0)printf("Double precision arithmetics is not supported on this device !\n"); + #endif #else if (get_global_id(0)==0)printf("Double precision arithmetics is not supported on this device !\n"); #endif From 870b4c12597c275a79a4fbb3648020349a317b34 Mon Sep 17 00:00:00 2001 From: Jerome Kieffer Date: Tue, 19 Nov 2024 18:12:14 +0100 Subject: [PATCH 7/9] restore full test --- src/pyFAI/opencl/test/test_preproc.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/pyFAI/opencl/test/test_preproc.py b/src/pyFAI/opencl/test/test_preproc.py index d4fbdb7e8..b5dc5ec50 100644 --- a/src/pyFAI/opencl/test/test_preproc.py +++ b/src/pyFAI/opencl/test/test_preproc.py @@ -57,7 +57,7 @@ def test_preproc(self): """ from ..preproc import preproc ary = numpy.arange(12).reshape(4,3) - for dtype in (numpy.uint8, numpy.int8, numpy.int16, numpy.uint16, numpy.uint32, numpy.int32):#, numpy.uint64, numpy.int64, numpy.float32): + for dtype in (numpy.uint8, numpy.int8, numpy.int16, numpy.uint16, numpy.uint32, numpy.int32, numpy.uint64, numpy.int64, numpy.float32): import sys; sys.stderr.write(f"test {dtype}\n") self.assertEqual(abs(preproc(ary.astype(dtype),split_result=4)[..., 0]-ary).max(), 0, "Result OK for dtype {dtype}") From 300fbe1911f459028c5851ed93dfa7c39cf20755 Mon Sep 17 00:00:00 2001 From: Jerome Kieffer Date: Wed, 20 Nov 2024 15:35:21 +0100 Subject: [PATCH 8/9] Retrigger CI --- src/pyFAI/integrator/fiber.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/pyFAI/integrator/fiber.py b/src/pyFAI/integrator/fiber.py index 4ac92f916..267bd5b38 100644 --- a/src/pyFAI/integrator/fiber.py +++ b/src/pyFAI/integrator/fiber.py @@ -30,7 +30,7 @@ __contact__ = "edgar.gutierrez-fernandez@esr.fr" __license__ = "MIT" __copyright__ = "European Synchrotron Radiation Facility, Grenoble, France" -__date__ = "14/11/2024" +__date__ = "20/11/2024" __status__ = "stable" __docformat__ = 'restructuredtext' From 0089ba9c0f2567c107be5eb36a23ded067d3747b Mon Sep 17 00:00:00 2001 From: Jerome Kieffer Date: Wed, 20 Nov 2024 16:53:59 +0100 Subject: [PATCH 9/9] Change version of silx --- src/pyFAI/opencl/azim_csr.py | 2 +- src/pyFAI/opencl/azim_hist.py | 2 +- src/pyFAI/opencl/azim_lut.py | 2 +- src/pyFAI/opencl/peak_finder.py | 4 ++-- 4 files changed, 5 insertions(+), 5 deletions(-) diff --git a/src/pyFAI/opencl/azim_csr.py b/src/pyFAI/opencl/azim_csr.py index 5a8c61ff5..4c0db511e 100644 --- a/src/pyFAI/opencl/azim_csr.py +++ b/src/pyFAI/opencl/azim_csr.py @@ -278,7 +278,7 @@ def compile_kernels(self, kernel_file=None): try: compile_options = self.get_compiler_options(x87_volatile=True, apple_gpu=True) except (AttributeError, TypeError): # Silx version too old - logger.warning("Please upgrade to silx v2.1+") + logger.warning("Please upgrade to silx v2.2+") from . import get_compiler_options compile_options = get_compiler_options(self.ctx, x87_volatile=True, apple_gpu=True) diff --git a/src/pyFAI/opencl/azim_hist.py b/src/pyFAI/opencl/azim_hist.py index cdaaf1079..163aed8aa 100644 --- a/src/pyFAI/opencl/azim_hist.py +++ b/src/pyFAI/opencl/azim_hist.py @@ -260,7 +260,7 @@ def compile_kernels(self, kernel_file=None): try: compile_options = self.get_compiler_options(x87_volatile=True, apple_gpu=True) except (AttributeError, TypeError): # Silx version too old - logger.warning("Please upgrade to silx v2.1+") + logger.warning("Please upgrade to silx v2.2+") from . import get_compiler_options compile_options = get_compiler_options(self.ctx, x87_volatile=True, apple_gpu=True) diff --git a/src/pyFAI/opencl/azim_lut.py b/src/pyFAI/opencl/azim_lut.py index f93f92f07..4db6efe89 100644 --- a/src/pyFAI/opencl/azim_lut.py +++ b/src/pyFAI/opencl/azim_lut.py @@ -208,7 +208,7 @@ def compile_kernels(self, kernel_file=None): try: compile_options = self.get_compiler_options(x87_volatile=True, apple_gpu=True) except (AttributeError, TypeError): # Silx version too old - logger.warning("Please upgrade to silx v2.1+") + logger.warning("Please upgrade to silx v2.2+") from . import get_compiler_options compile_options = get_compiler_options(self.ctx, x87_volatile=True, apple_gpu=True) diff --git a/src/pyFAI/opencl/peak_finder.py b/src/pyFAI/opencl/peak_finder.py index 761646def..1c5e4adb2 100644 --- a/src/pyFAI/opencl/peak_finder.py +++ b/src/pyFAI/opencl/peak_finder.py @@ -29,7 +29,7 @@ __authors__ = ["Jérôme Kieffer"] __license__ = "MIT" -__date__ = "19/11/2024" +__date__ = "20/11/2024" __copyright__ = "2014-2023, ESRF, Grenoble" __contact__ = "jerome.kieffer@esrf.fr" @@ -910,7 +910,7 @@ def compile_kernels(self, kernel_file=None): try: compile_options = self.get_compiler_options(x87_volatile=True, apple_gpu=True) except (AttributeError, TypeError): # Silx version too old - logger.warning("Please upgrade to silx v2.1+") + logger.warning("Please upgrade to silx v2.2+") from . import get_compiler_options compile_options = get_compiler_options(self.ctx, x87_volatile=True, apple_gpu=True)