Merge pull request #2124 from silx-kit/2099_numpy2

Fix some tests related to numpy2

src/pyFAI/opencl/OCLFullSplit.py

@@ -33,7 +33,7 @@
 
 __authors__ = ["Jérôme Kieffer", "Giannis Ashiotis"]
 __license__ = "MIT"
-__date__ = "06/10/2022"
+__date__ = "08/04/2024"
 __copyright__ = "2014, ESRF, Grenoble"
 __contact__ = "[email protected]"
 
@@ -126,7 +126,7 @@ def __init__(self,
         self._cl_kernel_args = {}
         self._cl_mem = {}
         self.events = []
-        self.workgroup_size = workgroup_size
+        self.workgroup_size = int(workgroup_size)
         if self.size < self.workgroup_size:
             raise RuntimeError("Fatal error in workgroup size selection. Size (%d) must be >= workgroup size (%d)\n", self.size, self.workgroup_size)
         if (platformid is None) and (deviceid is None):
@@ -237,8 +237,8 @@ def _calc_LUT(self):
         self._cl_kernel_args["memset_outMax"] = [self._cl_mem["outMax"]]
         self._cl_kernel_args["lut_1"] = [self._cl_mem["pos"], self._cl_mem["minmax"], self.pos0_range.data, self.pos1_range.data, self._cl_mem["outMax"]]
         self._cl_kernel_args["lut_2"] = [self._cl_mem["outMax"], self._cl_mem["idx_ptr"], self._cl_mem["lutsize"]]
-        memset_size = (self.bins + self.workgroup_size - 1) & ~(self.workgroup_size - 1),
-        global_size = (self.size + self.workgroup_size - 1) & ~(self.workgroup_size - 1),
+        memset_size = int(self.bins + self.workgroup_size - 1) // self.workgroup_size * self.workgroup_size,
+        global_size = int(self.size + self.workgroup_size - 1) // self.workgroup_size * self.workgroup_size,
         with self._sem:
             memset_outMax = self._program.memset_outMax(self._queue, memset_size, (self.workgroup_size,), *self._cl_kernel_args["memset_outMax"])
             self.profile_add("memset_outMax", memset_outMax)

src/pyFAI/opencl/azim_csr.py

@@ -28,7 +28,7 @@
 
 __authors__ = ["Jérôme Kieffer", "Giannis Ashiotis"]
 __license__ = "MIT"
-__date__ = "04/10/2023"
+__date__ = "08/04/2024"
 __copyright__ = "ESRF, Grenoble"
 __contact__ = "[email protected]"
 
@@ -46,7 +46,6 @@
 EventDescription = processing.EventDescription
 BufferDescription = processing.BufferDescription
 
-
 logger = logging.getLogger(__name__)
 
 
@@ -409,7 +408,7 @@ def send_buffer(self, data, dest, checksum=None, workgroup_size=None, convert=Tr
                 copy_image = pyopencl.enqueue_copy(self.queue, dest_buffer, data.data)
                 events.append(EventDescription(f"copy D->D {dest}", copy_image))
             elif convert:
-                tmp_buffer =  self.cl_mem["tmp"]
+                tmp_buffer = self.cl_mem["tmp"]
                 dest_buffer = self.cl_mem[dest]
                 copy_image = pyopencl.enqueue_copy(self.queue, tmp_buffer, data.data)
                 kernel_name = self.mapping[data.dtype.type]
@@ -418,7 +417,7 @@ def send_buffer(self, data, dest, checksum=None, workgroup_size=None, convert=Tr
                 convert_to_float = kernel(self.queue, ((self.size + wg - 1) // wg * wg,), (wg,), tmp_buffer, dest_buffer)
                 events += [EventDescription(f"copy raw D->D {dest}", copy_image),
                            EventDescription(f"convert {kernel_name}", convert_to_float)]
-            else: # no convert
+            else:  # no convert
                 actual_dest = f"{dest}_raw"
                 dest_buffer = self.cl_mem[actual_dest]
                 if data.dtype.itemsize > dest_type.itemsize:
@@ -434,7 +433,7 @@ def send_buffer(self, data, dest, checksum=None, workgroup_size=None, convert=Tr
                 copy_image = pyopencl.enqueue_copy(self.queue, dest_buffer, numpy.ascontiguousarray(data, dest_type))
                 events.append(EventDescription(f"copy H->D {dest}", copy_image))
             elif convert:
-                tmp_buffer =  self.cl_mem["tmp"]
+                tmp_buffer = self.cl_mem["tmp"]
                 dest_buffer = self.cl_mem[dest]
                 copy_image = pyopencl.enqueue_copy(self.queue, tmp_buffer, numpy.ascontiguousarray(data))
                 kernel_name = self.mapping[data.dtype.type]
@@ -494,7 +493,7 @@ def integrate_legacy(self, data, dummy=None, delta_dummy=None,
         with self.sem:
             self.send_buffer(data, "image")
             wg = max(self.workgroup_size["memset_out"])
-            wdim_bins = (self.bins + wg - 1) & ~(wg - 1),
+            wdim_bins = int(self.bins + wg - 1) // wg * wg,
             memset = self.kernels.memset_out(self.queue, wdim_bins, (wg,), *list(self.cl_kernel_args["memset_out"].values()))
             events.append(EventDescription("memset_out", memset))
             kw_corr = self.cl_kernel_args["corrections"]
@@ -573,7 +572,7 @@ def integrate_legacy(self, data, dummy=None, delta_dummy=None,
             kw_corr["do_absorption"] = do_absorption
 
             wg = max(self.workgroup_size["corrections"])
-            wdim_data = (self.size + wg - 1) & ~(wg - 1),
+            wdim_data = int(self.size + wg - 1) // wg * wg,
             ev = self.kernels.corrections(self.queue, wdim_data, (wg,), *list(kw_corr.values()))
             events.append(EventDescription("corrections", ev))
 
@@ -589,7 +588,7 @@ def integrate_legacy(self, data, dummy=None, delta_dummy=None,
             if wg_max == 1:
                 # thread-synchronization is probably not possible.
                 wg_max = max(self.workgroup_size["csr_integrate_single"])
-                wdim_bins = (self.bins + wg_max - 1) & ~(wg_max - 1),
+                wdim_bins = int(self.bins + wg_max - 1) // wg_max * wg_max,
                 integrate = self.kernels.csr_integrate_single(self.queue, wdim_bins, (wg_max,), *kw_int.values())
                 events.append(EventDescription("csr_integrate_single", integrate))
             else:
@@ -682,7 +681,7 @@ def integrate_ng(self, data, dark=None, dummy=None, delta_dummy=None,
             corrections4 = self.kernels.corrections4a
             kw_corr = self.cl_kernel_args[kernel_correction_name]
             kw_corr["image"] = self.send_buffer(data, "image", convert=False)
-            kw_corr["dtype"] = numpy.int8(32) if data.dtype.itemsize>4 else dtype_converter(data.dtype)
+            kw_corr["dtype"] = numpy.int8(32) if data.dtype.itemsize > 4 else dtype_converter(data.dtype)
             wg = workgroup_size if workgroup_size else max(self.workgroup_size["memset_ng"])
             wdim_bins = (self.bins + wg - 1) // wg * wg,
             memset = self.kernels.memset_out(self.queue, wdim_bins, (wg,), *list(self.cl_kernel_args["memset_ng"].values()))
@@ -911,9 +910,9 @@ def sigma_clip(self, data, dark=None, dummy=None, delta_dummy=None,
             corrections4 = self.kernels.corrections4a
             kw_corr = self.cl_kernel_args[kernel_correction_name]
             kw_corr["image"] = self.send_buffer(data, "image", convert=False)
-            kw_corr["dtype"] = numpy.int8(32) if data.dtype.itemsize>4 else dtype_converter(data.dtype)
+            kw_corr["dtype"] = numpy.int8(32) if data.dtype.itemsize > 4 else dtype_converter(data.dtype)
             wg = max(self.workgroup_size["memset_ng"])
-            wdim_bins = (self.bins + wg - 1) & ~(wg - 1),
+            wdim_bins = int(self.bins + wg - 1) // wg * wg,
             memset = self.kernels.memset_out(self.queue, wdim_bins, (wg,), *list(self.cl_kernel_args["memset_ng"].values()))
             events.append(EventDescription("memset_ng", memset))
             kw_int = self.cl_kernel_args["csr_sigma_clip4"]
@@ -1000,7 +999,7 @@ def sigma_clip(self, data, dark=None, dummy=None, delta_dummy=None,
             kw_corr["do_absorption"] = do_absorption
 
             wg = max(self.workgroup_size[kernel_correction_name])
-            wdim_data = (self.size + wg - 1) & ~(wg - 1),
+            wdim_data = int(self.size + wg - 1) // wg * wg,
             ev = corrections4(self.queue, wdim_data, (wg,), *list(kw_corr.values()))
             events.append(EventDescription(kernel_correction_name, ev))
 

src/pyFAI/opencl/azim_hist.py

@@ -32,7 +32,7 @@
 """
 __author__ = "Jérôme Kieffer"
 __license__ = "MIT"
-__date__ = "04/10/2023"
+__date__ = "08/04/2024"
 __copyright__ = "2012-2021, ESRF, Grenoble"
 __contact__ = "[email protected]"
 
@@ -152,8 +152,10 @@ def __init__(self, radial, bins, radial_checksum=None, empty=None, unit=None,
 
         self.BLOCK_SIZE = min(block_size, self.device.max_work_group_size)
         self.workgroup_size = {}
-        self.wdim_bins = (self.bins + self.BLOCK_SIZE - 1) & ~(self.BLOCK_SIZE - 1),
-        self.wdim_data = (self.size + self.BLOCK_SIZE - 1) & ~(self.BLOCK_SIZE - 1),
+        print("bins", self.bins, type(self.bins), "BLOCK_SIZE", self.BLOCK_SIZE, type(self.BLOCK_SIZE))
+        print(self.bins + self.BLOCK_SIZE - 1)
+        self.wdim_bins = int(self.bins + self.BLOCK_SIZE - 1) // self.BLOCK_SIZE * self.BLOCK_SIZE,
+        self.wdim_data = int(self.size + self.BLOCK_SIZE - 1) // self.BLOCK_SIZE * self.BLOCK_SIZE,
 
         self.buffers = [BufferDescription(i.name, i.size * self.size, i.dtype, i.flags)
                         for i in self.__class__.buffers]

src/pyFAI/opencl/azim_lut.py

@@ -27,7 +27,7 @@
 
 __author__ = "Jérôme Kieffer"
 __license__ = "MIT"
-__date__ = "04/10/2023"
+__date__ = "08/04/2024"
 __copyright__ = "2012-2021, ESRF, Grenoble"
 __contact__ = "[email protected]"
 
@@ -45,7 +45,6 @@
 EventDescription = processing.EventDescription
 BufferDescription = processing.BufferDescription
 
-
 logger = logging.getLogger(__name__)
 
 
@@ -134,8 +133,8 @@ def __init__(self, lut, image_size, checksum=None,
 
         self.BLOCK_SIZE = min(self.BLOCK_SIZE, self.device.max_work_group_size)
         self.workgroup_size = self.BLOCK_SIZE,  # Note this is a tuple
-        self.wdim_bins = (self.bins + self.BLOCK_SIZE - 1) & ~(self.BLOCK_SIZE - 1),
-        self.wdim_data = (self.size + self.BLOCK_SIZE - 1) & ~(self.BLOCK_SIZE - 1),
+        self.wdim_bins = int(self.bins + self.BLOCK_SIZE - 1) // self.BLOCK_SIZE * self.BLOCK_SIZE,
+        self.wdim_data = int(self.size + self.BLOCK_SIZE - 1) // self.BLOCK_SIZE * self.BLOCK_SIZE,
 
         self.buffers = [BufferDescription(i.name, i.size * self.size, i.dtype, i.flags)
                         for i in self.__class__.buffers]
@@ -520,7 +519,7 @@ def integrate_ng(self, data, dark=None, dummy=None, delta_dummy=None,
             self.send_buffer(data, "image")
 
             wg = self.workgroup_size
-            wdim_bins = (self.bins + wg[0] - 1) & ~(wg[0] - 1),
+            wdim_bins = int(self.bins + wg[0] - 1) // wg[0] * wg[0],
             memset = self.kernels.memset_out(self.queue, wdim_bins, wg, *list(self.cl_kernel_args["memset_ng"].values()))
             events.append(EventDescription("memset_ng", memset))
 

src/pyFAI/opencl/ocl_hist_pixelsplit.py

@@ -33,7 +33,7 @@
 
 __authors__ = ["Jérôme Kieffer", "Giannis Ashiotis"]
 __license__ = "MIT"
-__date__ = "20/01/2021"
+__date__ = "08/04/2024"
 __copyright__ = "2014, ESRF, Grenoble"
 __contact__ = "[email protected]"
 
@@ -113,8 +113,8 @@ def __init__(self, pos, bins, image_size, pos0_range=None, pos1_range=None, devi
         self.device_type = self.device.type
         self.BLOCK_SIZE = min(self.device.max_work_group_size, block_size)
         self.workgroup_size = self.BLOCK_SIZE,
-        self.wdim_bins = (self.bins * self.BLOCK_SIZE),
-        self.wdim_data = (self.size + self.BLOCK_SIZE - 1) & ~(self.BLOCK_SIZE - 1),
+        self.wdim_bins = int(self.bins * self.BLOCK_SIZE),
+        self.wdim_data = int(self.size + self.BLOCK_SIZE - 1) // self.BLOCK_SIZE * self.BLOCK_SIZE,
         try:
             # self._ctx = pyopencl.Context(devices=[pyopencl.get_platforms()[platformid].get_devices()[deviceid]])
             self._ctx = pyopencl.create_some_context()
@@ -334,7 +334,7 @@ def integrate(self, data, dummy=None, delta_dummy=None, dark=None, flat=None, so
             events.append(("copy D->H outData", ev))
             ev = pyopencl.enqueue_copy(self._queue, outCount, self._cl_mem["outCount"])
             events.append(("copy D->H outCount", ev))
-            global_size_integrate2 = (self.bins + self.BLOCK_SIZE - 1) & ~(self.BLOCK_SIZE - 1),
+            global_size_integrate2 = int(self.bins + self.BLOCK_SIZE - 1) // self.BLOCK_SIZE * self.BLOCK_SIZE,
             integrate2 = self._program.integrate2(self._queue, global_size_integrate2, self.workgroup_size, *self._cl_kernel_args["integrate2"])
             events.append(("integrate2", integrate2))
             ev = pyopencl.enqueue_copy(self._queue, outMerge, self._cl_mem["outMerge"])

src/pyFAI/opencl/peak_finder.py

@@ -29,7 +29,7 @@
 
 __authors__ = ["Jérôme Kieffer"]
 __license__ = "MIT"
-__date__ = "04/10/2023"
+__date__ = "08/04/2024"
 __copyright__ = "2014-2023, ESRF, Grenoble"
 __contact__ = "[email protected]"
 
@@ -230,10 +230,10 @@ def _sigma_clip(self, data, dark=None, dummy=None, delta_dummy=None,
         corrections4 = self.kernels.corrections4a
         kw_corr = self.cl_kernel_args[kernel_correction_name]
         kw_corr["image"] = self.send_buffer(data, "image", convert=False)
-        kw_corr["dtype"] = numpy.int8(32) if data.dtype.itemsize>4 else dtype_converter(data.dtype)
+        kw_corr["dtype"] = numpy.int8(32) if data.dtype.itemsize > 4 else dtype_converter(data.dtype)
 
         wg = max(self.workgroup_size["memset_ng"])
-        wdim_bins = (self.bins + wg - 1) & ~(wg - 1),
+        wdim_bins = int(self.bins + wg - 1) // wg * wg,
         memset = self.kernels.memset_out(self.queue, wdim_bins, (wg,), *list(self.cl_kernel_args["memset_ng"].values()))
         events.append(EventDescription("memset_ng", memset))
 
@@ -320,7 +320,7 @@ def _sigma_clip(self, data, dark=None, dummy=None, delta_dummy=None,
         kw_corr["error_model"] = kw_int["error_model"] = numpy.int8(error_model.value)
 
         wg = max(self.workgroup_size[kernel_correction_name])
-        wdim_data = (self.size + wg - 1) & ~(wg - 1),
+        wdim_data = int(self.size + wg - 1) // wg * wg,
         ev = corrections4(self.queue, wdim_data, (wg,), *list(kw_corr.values()))
         events.append(EventDescription(kernel_correction_name, ev))
 
@@ -369,7 +369,7 @@ def _count_intense(self, noise=1.0, cutoff_pick=3.0, radial_range=None, events=N
             kw_proj["radius_max"] = numpy.float32(numpy.finfo(numpy.float32).max)
 
         wg = max(self.workgroup_size["find_intense"])
-        wdim_data = (self.size + wg - 1) // wg * wg,  # & ~(wg - 1),
+        wdim_data = int(self.size + wg - 1) // wg * wg,
         kw_proj["shared"] = pyopencl.LocalMemory(wg * 4)  # stores int
         peak_search = self.program.find_intense(self.queue, wdim_data, (wg,), *list(kw_proj.values()))
         events.append(EventDescription("find_intense", peak_search))
@@ -479,7 +479,7 @@ def _peak_picking(self, data, noise, cutoff_peak, radial_range=None, patch_size=
                 wg1 = int(swg * sqrt2)
                 wg0 = wg // wg1
 
-        wdim_data = (data.shape[0] + wg0 - 1) & ~(wg0 - 1), (data.shape[1] + wg1 - 1) & ~(wg1 - 1)
+        wdim_data = int(data.shape[0] + wg0 - 1) // wg0 * wg0, int(data.shape[1] + wg1 - 1) // wg1 * wg
         # allocate local memory: we store 4 bytes but at most 1 pixel out of 4 can be a peak
 
         hw = patch_size // 2  # Half width of the patch
@@ -666,7 +666,7 @@ def sparsify(self, data, dark=None, dummy=None, delta_dummy=None,
                 # Call kernel to copy intensities
                 kw = self.cl_kernel_args["copy_intense"]
                 kw["counter"] = cnt_intense
-                size = (cnt_intense + self.BLOCK_SIZE - 1) & ~(self.BLOCK_SIZE - 1)
+                size = int(cnt_intense + self.BLOCK_SIZE - 1) // self.BLOCK_SIZE * self.BLOCK_SIZE
                 ev0 = kernel(self.queue, (size,), (self.BLOCK_SIZE,),
                                          *list(kw.values()))
                 ev1 = pyopencl.enqueue_copy(self.queue, indexes, self.cl_mem["position"])
@@ -847,7 +847,7 @@ def __init__(self, image_shape=None, mask=None,
         self.wg = self.size_to_doublet(self.BLOCK_SIZE)
         if sum(i < j for i, j in zip(self.ctx.devices[0].max_work_item_sizes, self.wg)):
             self.wg = self.ctx.devices[0].max_work_item_sizes[:2]
-        self.wdim = tuple((shape + BLOCK_SIZE - 1) & ~(BLOCK_SIZE - 1) for shape, BLOCK_SIZE in zip(self.shape[-1::-1], self.wg))
+        self.wdim = tuple(int(shape + BLOCK_SIZE - 1) // BLOCK_SIZE * BLOCK_SIZE for shape, BLOCK_SIZE in zip(self.shape[-1::-1], self.wg))
 
         self.buffers = [BufferDescription(i.name, i.size * numpy.prod(self.shape), i.dtype, i.flags)
                         for i in self.__class__.buffers]
@@ -1078,7 +1078,7 @@ def sparsify(self,
         with self.sem:
             count = self._count(image, window, cutoff, radius, noise)
             kw = self.cl_kernel_args["copy_intense"]
-            size = (count + self.BLOCK_SIZE - 1) & ~(self.BLOCK_SIZE - 1)
+            size = int(count + self.BLOCK_SIZE - 1) // self.BLOCK_SIZE * self.BLOCK_SIZE
 
             copy_intense = self.program.copy_intense(self.queue, (size,), (self.BLOCK_SIZE,),
                                            *list(kw.values()))

src/pyFAI/opencl/sort.py

@@ -33,7 +33,7 @@
 
 __author__ = "Jérôme Kieffer"
 __license__ = "MIT"
-__date__ = "06/10/2022"
+__date__ = "08/04/2024"
 __copyright__ = "2015, ESRF, Grenoble"
 __contact__ = "[email protected]"
 
@@ -236,7 +236,7 @@ def sort_vertical(self, data, dummy=None):
                 kargs["dummy"] = dummy
                 kargs["src_size"] = numpy.int32(data.size),
                 wg = min(32, self.block_size)
-                size = ((self.npt_height * self.npt_width) + wg - 1) & ~(wg - 1)
+                size = int(self.npt_height * self.npt_width + wg - 1) // wg * wg
                 evt = self.kernels.copy_pad(self.queue, (size,), (wg,), *kargs.values())
                 events.append(EventDescription("copy_pad", evt))
             else:
@@ -323,7 +323,7 @@ def filter_vertical(self, data, dummy=None, quantile=0.5):
             dummy = numpy.float32(dummy)
         _sorted = self.sort_vertical(data, dummy)
         wg = min(32, self.block_size)
-        ws = (self.npt_width + wg - 1) & ~(wg - 1)
+        ws = int(self.npt_width + wg - 1) // wg * wg
         with self.sem:
             kargs = self.cl_kernel_args["filter_vertical"]
             kargs["dummy"] = dummy
@@ -347,7 +347,7 @@ def filter_horizontal(self, data, dummy=None, quantile=0.5):
             dummy = numpy.float32(dummy)
         _sorted = self.sort_horizontal(data, dummy)
         wg = min(32, self.block_size)
-        ws = (self.npt_height + wg - 1) & ~(wg - 1)
+        ws = int(self.npt_height + wg - 1) // wg * wg
         with self.sem:
             kargs = self.cl_kernel_args["filter_horizontal"]
             kargs["dummy"] = dummy
@@ -372,7 +372,7 @@ def trimmed_mean_vertical(self, data, dummy=None, quantiles=(0.5, 0.5)):
             dummy = numpy.float32(dummy)
         _sorted = self.sort_vertical(data, dummy)
         wg = min(32, self.block_size)
-        ws = (self.npt_width + wg - 1) & ~(wg - 1)
+        ws = int(self.npt_width + wg - 1) // wg * wg
         with self.sem:
             kargs = self.cl_kernel_args["trimmed_mean_vertical"]
             kargs["dummy"] = dummy
@@ -398,7 +398,7 @@ def trimmed_mean_horizontal(self, data, dummy=None, quantiles=(0.5, 0.5)):
             dummy = numpy.float32(dummy)
         _sorted = self.sort_horizontal(data, dummy)
         wg = min(32, self.block_size)
-        ws = (self.npt_height + wg - 1) & ~(wg - 1)
+        ws = int(self.npt_height + wg - 1) // wg * wg
         with self.sem:
             kargs = self.cl_kernel_args["trimmed_mean_horizontal"]
             kargs["dummy"] = dummy