Skip to content

Commit

Permalink
Code is integrated, but wrong for now...
Browse files Browse the repository at this point in the history
  • Loading branch information
@kif
kif committed Dec 4, 2024
1 parent f66755a commit 06225a7
Show file tree
Hide file tree
Showing 2 changed files with 33 additions and 22 deletions.
28 changes: 19 additions & 9 deletions src/pyFAI/opencl/azim_csr.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -28,7 +28,7 @@

__authors__ = ["Jérôme Kieffer", "Giannis Ashiotis"]
__license__ = "MIT"
__date__ = "20/11/2024"
__date__ = "04/12/2024"
__copyright__ = "ESRF, Grenoble"
__contact__ = "[email protected]"

Expand Down Expand Up @@ -70,11 +70,16 @@ class OCL_CSR_Integrator(OpenclProcessing):
BufferDescription("solidangle", 1, numpy.float32, mf.READ_ONLY),
BufferDescription("absorption", 1, numpy.float32, mf.READ_ONLY),
BufferDescription("mask", 1, numpy.int8, mf.READ_ONLY),

]
kernel_files = ["silx:opencl/doubleword.cl",
"pyfai:openCL/preprocess.cl",
"pyfai:openCL/memset.cl",
"pyfai:openCL/ocl_azim_CSR.cl"
"pyfai:openCL/ocl_azim_CSR.cl",
"pyfai:openCL/collective/reduction.cl",
"pyfai:openCL/collective/scan.cl",
"pyfai:openCL/collective/comb_sort.cl",
"pyfai:openCL/medfilt.cl"
]
mapping = {numpy.int8: "s8_to_float",
numpy.uint8: "u8_to_float",
Expand Down Expand Up @@ -162,6 +167,7 @@ def __init__(self, lut, image_size, checksum=None,
BufferDescription("sem", self.bins, numpy.float32, mf.READ_WRITE),
BufferDescription("merged", self.bins, numpy.float32, mf.WRITE_ONLY),
BufferDescription("merged8", (self.bins, 8), numpy.float32, mf.WRITE_ONLY),
BufferDescription("work4", (self.data_size, 4), numpy.float32, mf.READ_WRITE),
]
try:
self.set_profiling(profile)
Expand Down Expand Up @@ -383,6 +389,7 @@ def set_kernel_arguments(self):
self.cl_kernel_args["csr_integrate_single"] = self.cl_kernel_args["csr_integrate"]
self.cl_kernel_args["csr_integrate4_single"] = self.cl_kernel_args["csr_integrate4"]
self.cl_kernel_args["csr_medfilt"] = OrderedDict((("output4", self.cl_mem["output4"]),
("work4", self.cl_mem["work4"]),
("data", self.cl_mem["data"]),
("indices", self.cl_mem["indices"]),
("indptr", self.cl_mem["indptr"]),
Expand All @@ -394,7 +401,8 @@ def set_kernel_arguments(self):
("averint", self.cl_mem["averint"]),
("std", self.cl_mem["std"]),
("sem", self.cl_mem["sem"]),
("shared", pyopencl.LocalMemory(32))
("shared_int", pyopencl.LocalMemory(128)),
("shared_float", pyopencl.LocalMemory(128)),
))
self.cl_kernel_args["memset_out"] = OrderedDict(((i, self.cl_mem[i]) for i in ("sum_data", "sum_count", "merged")))
self.cl_kernel_args["memset_ng"] = OrderedDict(((i, self.cl_mem[i]) for i in ("averint", "std", "merged8")))
Expand Down Expand Up @@ -607,6 +615,7 @@ def integrate_legacy(self, data, dummy=None, delta_dummy=None,
integrate = self.kernels.csr_integrate_single(self.queue, wdim_bins, (wg_max,), *kw_int.values())
events.append(EventDescription("csr_integrate_single", integrate))
else:
#TODO reshape this kernel with (wg, nbin)\(wg,1) rather than (self.bins * wg_min)\(wg_min) #2348
wdim_bins = (self.bins * wg_min),
kw_int["shared"] = pyopencl.LocalMemory(16 * wg_min) # sizeof(float4) == 16
integrate = self.kernels.csr_integrate(self.queue, wdim_bins, (wg_min,), *kw_int.values())
Expand Down Expand Up @@ -1121,7 +1130,7 @@ def medfilt(self, data, dark=None, dummy=None, delta_dummy=None,
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"]
kw_int = self.cl_kernel_args["csr_medfilt"]

if dummy is not None:
do_dummy = numpy.int8(1)
Expand Down Expand Up @@ -1212,11 +1221,12 @@ def medfilt(self, data, dark=None, dummy=None, delta_dummy=None,
kw_int["quant_min"] = numpy.float32(quant_min)
kw_int["quant_max"] = numpy.float32(quant_max)

wg_min = min(self.workgroup_size["csr_sigma_clip4"])
kw_int["shared"] = pyopencl.LocalMemory(32 * wg_min)
wdim_bins = (self.bins * wg_min),
integrate = self.kernels.csr_sigma_clip4(self.queue, wdim_bins, (wg_min,), *kw_int.values())
events.append(EventDescription("csr_sigma_clip4", integrate))
wg_min = min(self.workgroup_size["csr_medfilt"])
kw_int["shared_int"] = pyopencl.LocalMemory(4 * wg_min)
kw_int["shared_float"] = pyopencl.LocalMemory(8 * wg_min)
wdim_bins = (wg_min, self.bins)
integrate = self.kernels.csr_medfilt(self.queue, wdim_bins, (wg_min,1), *kw_int.values())
events.append(EventDescription("medfilt", integrate))

if out_merged is None:
merged = numpy.empty((self.bins, 8), dtype=numpy.float32)
Expand Down
27 changes: 14 additions & 13 deletions src/pyFAI/resources/openCL/medfilt.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -81,19 +81,20 @@ float2 inline sum_float2_reduction(local float* shared)
* dim 0: collaboarative working group size, probably optimal in the range 32-128
* dim 1: index of bin, size=1
*
* @param weights Float pointer to global memory storing the input image.
* @param coefs Float pointer to global memory holding the coeficient part of the LUT
* @param indices Integer pointer to global memory holding the corresponding index of the coeficient
* @param indptr Integer pointer to global memory holding the pointers to the coefs and indices for the CSR matrix
* @param cutoff Discard any value with |value - mean| > cutoff*sigma
* @param cycle number of cycle
* @param error_model 0:disable, 1:variance, 2:poisson, 3:azimuthal, 4:hybrid
* @param summed contains all the data
* @param averint Average signal
* @param stdevpix Float pointer to the output 1D array with the propagated error (std)
* @param stdevpix Float pointer to the output 1D array with the propagated error (sem)
* @param shared Buffer of shared memory of size WORKGROUP_SIZE * 8 * sizeof(float)
*/
* @param weights Float pointer to global memory storing the input image.
* @param coefs Float pointer to global memory holding the coeficient part of the LUT
* @param indices Integer pointer to global memory holding the corresponding index of the coeficient
* @param indptr Integer pointer to global memory holding the pointers to the coefs and indices for the CSR matrix
* @param quant_min start percentile/100 to use. Use 0.5 for the median
* @param quant_max stop percentile/100 to use. Use 0.5 for the median
* @param error_model 0:disable, 1:variance, 2:poisson, 3:azimuthal, 4:hybrid
* @param summed contains all the data
* @param averint Average signal
* @param stdevpix Float pointer to the output 1D array with the propagated error (std)
* @param stdevpix Float pointer to the output 1D array with the propagated error (sem)
* @param shared_int Buffer of shared memory of size WORKGROUP_SIZE * sizeof(int)
* @param shared_float Buffer of shared memory of size WORKGROUP_SIZE * 2 * sizeof(float)
* */



Expand Down

0 comments on commit 06225a7

Please sign in to comment.