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..4c0db511e 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.2+") + 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..163aed8aa 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.2+") + 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..4db6efe89 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.2+") + 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..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__ = "08/04/2024" +__date__ = "20/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.2+") + 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..b5dc5ec50 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 @@ -58,6 +58,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): + 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..6d805be49 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 @@ -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, @@ -199,11 +199,15 @@ 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)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 }