[OCL] Explicitely disable fp64

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 <https://doi.org/10.48550/arXiv.2411.09515>`_ in **J. Applied Crystallography** (2024); 
+  `Accepted <https://doi.org/10.48550/arXiv.2411.09515>`_ 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 <https://github.com/silx-kit/pyFAI/network/dependents?dependent_type=PACKAGE>`_ 
+prestigious scientific journals (Nature, PNAS, ...) and
+40 other `applications <https://github.com/silx-kit/pyFAI/network/dependents?dependent_type=PACKAGE>`_
 using pyFAI as a library.

src/pyFAI/opencl/__init__.py

@@ -36,7 +36,7 @@
 __contact__ = "[email protected]"
 __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"

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__ = "[email protected]"
 
@@ -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

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__ = "[email protected]"
 
@@ -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:

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__ = "[email protected]"
 
@@ -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

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__ = "[email protected]"
 
@@ -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)

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 ([email protected])
 #
@@ -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__ = "[email protected]"
 
@@ -328,15 +328,17 @@ 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
                              (by default use the one declared in the class)
         """
         # 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):

src/pyFAI/opencl/test/test_preproc.py

@@ -33,7 +33,7 @@
 __contact__ = "[email protected]"
 __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}")
 
 

src/pyFAI/resources/openCL/preprocess.cl

@@ -7,7 +7,7 @@
  *                           Grenoble, France
  *
  *   Principal authors: J. Kieffer ([email protected])
- *   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
     }