Merge pull request #2341 from kif/2339_fix_buggy_Apple_compiler

[OCL] Explicitely disable fp64

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
     }