Debug parallel implementation on CPU/Intel driver

src/pyFAI/opencl/azim_csr.py

@@ -32,6 +32,7 @@
 __copyright__ = "ESRF, Grenoble"
 __contact__ = "jerome.kieffer@esrf.fr"
 
+import math
 import logging
 from collections import OrderedDict
 import numpy
@@ -301,7 +302,16 @@ def compile_kernels(self, kernel_file=None):
             else:
                 wg_max = self.kernels.max_workgroup_size(kernel_name)
                 wg_min = self.kernels.min_workgroup_size(kernel_name)
-                self.workgroup_size[kernel_name] = (wg_min, wg_max)
+                if kernel_name=="csr_medfilt":
+                    # limit the wg size due to
+                    device = self.ctx.devices[0]
+                    maxthreads = device.local_mem_size/12/4
+                    self.workgroup_size[kernel_name] = (wg_min,
+                                                        min(wg_max, 2**(int(math.log2(maxthreads)))))
+                else:
+                    self.workgroup_size[kernel_name] = (wg_min, wg_max)
+
+
 
     def set_kernel_arguments(self):
         """Tie arguments of OpenCL kernel-functions to the actual kernels
@@ -1243,8 +1253,7 @@ def medfilt(self, data, dark=None, dummy=None, delta_dummy=None,
             events.append(EventDescription(kernel_correction_name, ev))
 
             kw_int["quant_min"] = numpy.float32(quant_min)
-            kw_int["quant_max"] = numpy.float32(quant_max)#*EPS32
-
+            kw_int["quant_max"] = numpy.float32(quant_max)
             wg_min = max(self.workgroup_size["csr_medfilt"])
             kw_int["shared_int"] = pyopencl.LocalMemory(4 * wg_min)
             kw_int["shared_float"] = pyopencl.LocalMemory(8 * wg_min)
@@ -1292,6 +1301,8 @@ def medfilt(self, data, dark=None, dummy=None, delta_dummy=None,
 
         res = Integrate1dtpl(self.bin_centers, avgint, sem, merged[:, 0], merged[:, 2], merged[:, 4], merged[:, 6],
                              std, sem, merged[:, 7])
+        # print(self.get_buffer("merged8")[0])
+        # raise RuntimeError()
         return res
 
 

src/pyFAI/opencl/test/test_ocl_azim_csr.py

@@ -33,7 +33,7 @@
 __contact__ = "jerome.kieffer@esrf.eu"
 __license__ = "MIT"
 __copyright__ = "2019-2021 European Synchrotron Radiation Facility, Grenoble, France"
-__date__ = "05/12/2024"
+__date__ = "06/12/2024"
 
 import logging
 import numpy
@@ -120,20 +120,25 @@ def integrate_ng(self, block_size=None, method_called="integrate_ng", extra=None
                 pix = csr[2][1:]-csr[2][:-1]
                 self.assertTrue(numpy.allclose(res.count, pix), "all pixels have been counted")
 
-            # Intensities are not that different:
-            delta = ref.intensity - res.intensity
-            self.assertLessEqual(abs(delta).max(), 1e-5, "intensity is almost the same")
 
             # histogram of normalization
-            err = abs((res.normalization - ref.sum_normalization)).max()
-            self.assertLess(err, 5e-4, "normalization content is the same: %s<5e-5" % (err))
+            print(ref.sum_normalization)
+            print(res.normalization)
+            err = abs((res.normalization - ref.sum_normalization))
+            print(err)
+            self.assertLess(err.max(), 5e-4, "normalization content is the same: %s<5e-5" % (err.max))
 
             # histogram of signal
             self.assertLess(abs((res.signal - ref.sum_signal)).max(), 5e-5, "signal content is the same")
 
             # histogram of variance
             self.assertLess(abs((res.variance - ref.sum_variance)).max(), 5e-5, "signal content is the same")
 
+            # Intensities are not that different:
+            delta = ref.intensity - res.intensity
+            print(delta)
+            self.assertLessEqual(abs(delta).max(), 1e-5, "intensity is almost the same")
+
 
     @unittest.skipUnless(ocl, "pyopencl is missing")
     def test_integrate_ng(self):