summaryrefslogtreecommitdiff
path: root/silx/resources/opencl/preprocess.cl
diff options
context:
space:
mode:
Diffstat (limited to 'silx/resources/opencl/preprocess.cl')
-rw-r--r--silx/resources/opencl/preprocess.cl567
1 files changed, 567 insertions, 0 deletions
diff --git a/silx/resources/opencl/preprocess.cl b/silx/resources/opencl/preprocess.cl
new file mode 100644
index 0000000..de35c48
--- /dev/null
+++ b/silx/resources/opencl/preprocess.cl
@@ -0,0 +1,567 @@
+/*
+ * Project: Azimuthal regroupping OpenCL kernel for PyFAI.
+ * Preprocessing program
+ *
+ *
+ * Copyright (C) 2012-2017 European Synchrotron Radiation Facility
+ * Grenoble, France
+ *
+ * Principal authors: J. Kieffer (kieffer@esrf.fr)
+ * Last revision: 19/01/2017
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ */
+
+/**
+ * \file
+ *
+ * \brief OpenCL kernels for image array casting, array mem-setting and normalizing
+ *
+ * Constant to be provided at build time:
+ * NIMAGE: size of the image
+ *
+ */
+
+#include "for_eclipse.h"
+
+/**
+ * \brief cast values of an array of int8 into a float output array.
+ *
+ * - array_s8: Pointer to global memory with the input data as signed8 array
+ * - array_float: Pointer to global memory with the output data as float array
+ */
+__kernel void
+s8_to_float(__global char *array_s8,
+ __global float *array_float
+)
+{
+ int i = get_global_id(0);
+ //Global memory guard for padding
+ if (i < NIMAGE)
+ array_float[i] = (float)array_s8[i];
+}
+
+
+/**
+ * \brief cast values of an array of uint8 into a float output array.
+ *
+ * - array_u8: Pointer to global memory with the input data as unsigned8 array
+ * - array_float: Pointer to global memory with the output data as float array
+ */
+__kernel void
+u8_to_float(__global unsigned char *array_u8,
+ __global float *array_float
+)
+{
+ int i = get_global_id(0);
+ //Global memory guard for padding
+ if (i < NIMAGE)
+ array_float[i] = (float)array_u8[i];
+}
+
+
+/**
+ * \brief cast values of an array of int16 into a float output array.
+ *
+ * - array_s16: Pointer to global memory with the input data as signed16 array
+ * - array_float: Pointer to global memory with the output data as float array
+ */
+__kernel void
+s16_to_float(__global short *array_s16,
+ __global float *array_float
+)
+{
+ int i = get_global_id(0);
+ //Global memory guard for padding
+ if (i < NIMAGE)
+ array_float[i] = (float)array_s16[i];
+}
+
+
+/**
+ * \brief cast values of an array of uint16 into a float output array.
+ *
+ * - array_u16: Pointer to global memory with the input data as unsigned16 array
+ * - array_float: Pointer to global memory with the output data as float array
+ */
+__kernel void
+u16_to_float(__global unsigned short *array_u16,
+ __global float *array_float
+)
+{
+ int i = get_global_id(0);
+ //Global memory guard for padding
+ if (i < NIMAGE)
+ array_float[i] = (float)array_u16[i];
+}
+
+/**
+ * \brief cast values of an array of uint32 into a float output array.
+ *
+ * - array_u32: Pointer to global memory with the input data as unsigned32 array
+ * - array_float: Pointer to global memory with the output data as float array
+ */
+__kernel void
+u32_to_float(__global unsigned int *array_u32,
+ __global float *array_float
+)
+{
+ int i = get_global_id(0);
+ //Global memory guard for padding
+ if (i < NIMAGE)
+ array_float[i] = (float)array_u32[i];
+}
+
+/**
+ * \brief convert values of an array of int32 into a float output array.
+ *
+ * - array_int: Pointer to global memory with the data as unsigned32 array
+ * - array_float: Pointer to global memory with the data float array
+ */
+__kernel void
+s32_to_float(__global int *array_int,
+ __global float *array_float
+)
+{
+ int i = get_global_id(0);
+ //Global memory guard for padding
+ if (i < NIMAGE)
+ array_float[i] = (float)(array_int[i]);
+}
+
+
+/**
+ * Functions to be called from an actual kernel.
+ * \brief Performs the normalization of input image by dark subtraction,
+ * variance is propagated to second member of the float3
+ * flatfield, solid angle, polarization and absorption are stored in
+ * third member of the float3 returned.
+ *
+ * Invalid/Dummy pixels will all have the third-member set to 0, i.e. no weight
+ *
+ * - image Float pointer to global memory storing the input image.
+ * - do_dark Bool/int: shall dark-current correction be applied ?
+ * - dark Float pointer to global memory storing the dark image.
+ * - do_flat Bool/int: shall flat-field correction be applied ?
+ * - flat Float pointer to global memory storing the flat image.
+ * - do_solidangle Bool/int: shall flat-field correction be applied ?
+ * - solidangle Float pointer to global memory storing the solid angle of each pixel.
+ * - do_polarization Bool/int: shall polarization correction be applied ?
+ * - polarization Float pointer to global memory storing the polarization of each pixel.
+ * - do_absorption Bool/int: shall absorption correction be applied ?
+ * - absorption Float pointer to global memory storing the effective absoption of each pixel.
+ * - do_mask perform mask correction ?
+ * - mask Bool/char pointer to mask array
+ * - do_dummy Bool/int: shall the dummy pixel be checked. Dummy pixel are pixels marked as bad and ignored
+ * - dummy Float: value for bad pixels
+ * - delta_dummy Float: precision for bad pixel value
+ * - normalization_factor : divide the input by this value
+ *
+**/
+
+static float3 _preproc3(const __global float *image,
+ const __global float *variance,
+ const char do_dark,
+ const __global float *dark,
+ const char do_dark_variance,
+ const __global float *dark_variance,
+ const char do_flat,
+ const __global float *flat,
+ const char do_solidangle,
+ const __global float *solidangle,
+ const char do_polarization,
+ const __global float *polarization,
+ const char do_absorption,
+ const __global float *absorption,
+ const char do_mask,
+ const __global char *mask,
+ const char do_dummy,
+ const float dummy,
+ const float delta_dummy,
+ const float normalization_factor)
+{
+ size_t i= get_global_id(0);
+ float3 result = (float3)(0.0, 0.0, 0.0);
+ if (i < NIMAGE)
+ {
+ if ((!do_mask) || (!mask[i]))
+ {
+ result.s0 = image[i];
+ if (variance != 0)
+ result.s1 = variance[i];
+ result.s2 = normalization_factor;
+ if ( (!do_dummy)
+ ||((delta_dummy != 0.0f) && (fabs(result.s0-dummy) > delta_dummy))
+ ||((delta_dummy == 0.0f) && (result.s0 != dummy)))
+ {
+ if (do_dark)
+ result.s0 -= dark[i];
+ if (do_dark_variance)
+ result.s1 += dark_variance[i];
+ if (do_flat)
+ {
+ float one_flat = flat[i];
+ if ( (!do_dummy)
+ ||((delta_dummy != 0.0f) && (fabs(one_flat-dummy) > delta_dummy))
+ ||((delta_dummy == 0.0f) && (one_flat != dummy)))
+ result.s2 *= one_flat;
+ else
+ result.s2 = 0.0f;
+ }
+ if (do_solidangle)
+ result.s2 *= solidangle[i];
+ if (do_polarization)
+ result.s2 *= polarization[i];
+ if (do_absorption)
+ result.s2 *= absorption[i];
+ if (isnan(result.s0) || isnan(result.s1) || isnan(result.s2) || (result.s2 == 0.0f))
+ result = (float3)(0.0, 0.0, 0.0);
+ }
+ else
+ {
+ result = (float3)(0.0, 0.0, 0.0);
+ }//end if do_dummy
+ } // end if mask
+ };//end if NIMAGE
+ return result;
+};//end function
+
+
+/**
+ * \brief Performs the normalization of input image by dark subtraction,
+ * flatfield, solid angle, polarization and absorption division.
+ *
+ * Intensities of images are corrected by:
+ * - dark (read-out) noise subtraction
+ * - Solid angle correction (division)
+ * - polarization correction (division)
+ * - flat fiels correction (division)
+ * Corrections are made in place unless the pixel is dummy.
+ * Dummy pixels are left untouched so that they remain dummy
+ *
+ * - image Float pointer to global memory storing the input image.
+ * - do_dark Bool/int: shall dark-current correction be applied ?
+ * - dark Float pointer to global memory storing the dark image.
+ * - do_flat Bool/int: shall flat-field correction be applied ?
+ * - flat Float pointer to global memory storing the flat image.
+ * - do_solidangle Bool/int: shall flat-field correction be applied ?
+ * - solidangle Float pointer to global memory storing the solid angle of each pixel.
+ * - do_polarization Bool/int: shall polarization correction be applied ?
+ * - polarization Float pointer to global memory storing the polarization of each pixel.
+ * - do_absorption Bool/int: shall absorption correction be applied ?
+ * - absorption Float pointer to global memory storing the effective absoption of each pixel.
+ * - do_mask perform mask correction ?
+ * - mask Bool/char pointer to mask array
+ * - do_dummy Bool/int: shall the dummy pixel be checked. Dummy pixel are pixels marked as bad and ignored
+ * - dummy Float: value for bad pixels
+ * - delta_dummy Float: precision for bad pixel value
+ * - normalization_factor : divide the input by this value
+ *
+**/
+
+__kernel void
+corrections(const __global float *image,
+ const char do_dark,
+ const __global float *dark,
+ const char do_flat,
+ const __global float *flat,
+ const char do_solidangle,
+ const __global float *solidangle,
+ const char do_polarization,
+ const __global float *polarization,
+ const char do_absorption,
+ const __global float *absorption,
+ const char do_mask,
+ const __global char *mask,
+ const char do_dummy,
+ const float dummy,
+ const float delta_dummy,
+ const float normalization_factor,
+ __global float *output
+ )
+{
+ size_t i= get_global_id(0);
+ float3 result = (float3)(0.0, 0.0, 0.0);
+ if (i < NIMAGE)
+ {
+ result = _preproc3(image,
+ 0,
+ do_dark,
+ dark,
+ 0,
+ 0,
+ do_flat,
+ flat,
+ do_solidangle,
+ solidangle,
+ do_polarization,
+ polarization,
+ do_absorption,
+ absorption,
+ do_mask,
+ mask,
+ do_dummy,
+ dummy,
+ delta_dummy,
+ normalization_factor);
+ if (result.s2 != 0.0f)
+ output[i] = result.s0 / result.s2;
+ else
+ output[i] = dummy;
+ };//end if NIMAGE
+
+};//end kernel
+
+
+/**
+ * \brief Performs Normalization of input image with float2 output (num,denom)
+ *
+ * Intensities of images are corrected by:
+ * - dark (read-out) noise subtraction for the data
+ * - Solid angle correction (denominator)
+ * - polarization correction (denominator)
+ * - flat fiels correction (denominator)
+ *
+ * Corrections are made out of place.
+ * Dummy pixels set both the numerator and denominator to 0
+ *
+ * - image Float pointer to global memory storing the input image.
+ * - do_dark Bool/int: shall dark-current correction be applied ?
+ * - dark Float pointer to global memory storing the dark image.
+ * - do_flat Bool/int: shall flat-field correction be applied ?
+ * - flat Float pointer to global memory storing the flat image.
+ * - do_solidangle Bool/int: shall flat-field correction be applied ?
+ * - solidangle Float pointer to global memory storing the solid angle of each pixel.
+ * - do_polarization Bool/int: shall flat-field correction be applied ?
+ * - polarization Float pointer to global memory storing the polarization of each pixel.
+ * - do_dummy Bool/int: shall the dummy pixel be checked. Dummy pixel are pixels marked as bad and ignored
+ * - dummy Float: value for bad pixels
+ * - delta_dummy Float: precision for bad pixel value
+ * - normalization_factor : divide the input by this value
+ *
+ *
+**/
+__kernel void
+corrections2(const __global float *image,
+ const char do_dark,
+ const __global float *dark,
+ const char do_flat,
+ const __global float *flat,
+ const char do_solidangle,
+ const __global float *solidangle,
+ const char do_polarization,
+ const __global float *polarization,
+ const char do_absorption,
+ const __global float *absorption,
+ const char do_mask,
+ const __global char *mask,
+ const char do_dummy,
+ const float dummy,
+ const float delta_dummy,
+ const float normalization_factor,
+ __global float2 *output
+ )
+{
+ size_t i= get_global_id(0);
+ float3 result = (float3)(0.0, 0.0, 0.0);
+ if (i < NIMAGE)
+ {
+ result = _preproc3(image,
+ 0,
+ do_dark,
+ dark,
+ 0,
+ 0,
+ do_flat,
+ flat,
+ do_solidangle,
+ solidangle,
+ do_polarization,
+ polarization,
+ do_absorption,
+ absorption,
+ do_mask,
+ mask,
+ do_dummy,
+ dummy,
+ delta_dummy,
+ normalization_factor);
+ output[i] = (float2)(result.s0, result.s2);
+ };//end if NIMAGE
+};//end kernel
+
+/**
+ * \brief Performs Normalization of input image with float3 output (signal, variance, normalization) assuming poissonian signal
+ *
+ * Intensities of images are corrected by:
+ * - dark (read-out) noise subtraction for the data
+ * - Solid angle correction (denominator)
+ * - polarization correction (denominator)
+ * - flat fiels correction (denominator)
+ *
+ * Corrections are made out of place.
+ * Dummy pixels set both the numerator and denominator to 0
+ *
+ * - image Float pointer to global memory storing the input image.
+ * - do_dark Bool/int: shall dark-current correction be applied ?
+ * - dark Float pointer to global memory storing the dark image.
+ * - do_flat Bool/int: shall flat-field correction be applied ?
+ * - flat Float pointer to global memory storing the flat image.
+ * - do_solidangle Bool/int: shall flat-field correction be applied ?
+ * - solidangle Float pointer to global memory storing the solid angle of each pixel.
+ * - do_polarization Bool/int: shall flat-field correction be applied ?
+ * - polarization Float pointer to global memory storing the polarization of each pixel.
+ * - do_dummy Bool/int: shall the dummy pixel be checked. Dummy pixel are pixels marked as bad and ignored
+ * - dummy Float: value for bad pixels
+ * - delta_dummy Float: precision for bad pixel value
+ * - normalization_factor : divide the input by this value
+ *
+ *
+**/
+__kernel void
+corrections3Poisson( const __global float *image,
+ const char do_dark,
+ const __global float *dark,
+ const char do_flat,
+ const __global float *flat,
+ const char do_solidangle,
+ const __global float *solidangle,
+ const char do_polarization,
+ const __global float *polarization,
+ const char do_absorption,
+ const __global float *absorption,
+ const char do_mask,
+ const __global char *mask,
+ const char do_dummy,
+ const float dummy,
+ const float delta_dummy,
+ const float normalization_factor,
+ __global float3 *output
+ )
+{
+ size_t i= get_global_id(0);
+ float3 result = (float3)(0.0, 0.0, 0.0);
+ if (i < NIMAGE)
+ {
+ result = _preproc3(image,
+ image,
+ do_dark,
+ dark,
+ do_dark,
+ dark,
+ do_flat,
+ flat,
+ do_solidangle,
+ solidangle,
+ do_polarization,
+ polarization,
+ do_absorption,
+ absorption,
+ do_mask,
+ mask,
+ do_dummy,
+ dummy,
+ delta_dummy,
+ normalization_factor);
+ output[i] = result;
+ };//end if NIMAGE
+};//end kernel
+
+
+/**
+ * \brief Performs Normalization of input image with float3 output (signal, variance, normalization)
+ *
+ * Intensities of images are corrected by:
+ * - dark (read-out) noise subtraction for the data
+ * - Solid angle correction (division)
+ * - polarization correction (division)
+ * - flat fiels correction (division)
+ * Corrections are made in place unless the pixel is dummy.
+ * Dummy pixels are left untouched so that they remain dummy
+ *
+ * - image Float pointer to global memory storing the input image.
+ * - do_dark Bool/int: shall dark-current correction be applied ?
+ * - dark Float pointer to global memory storing the dark image.
+ * - do_flat Bool/int: shall flat-field correction be applied ?
+ * - flat Float pointer to global memory storing the flat image.
+ * - do_solidangle Bool/int: shall flat-field correction be applied ?
+ * - solidangle Float pointer to global memory storing the solid angle of each pixel.
+ * - do_polarization Bool/int: shall flat-field correction be applied ?
+ * - polarization Float pointer to global memory storing the polarization of each pixel.
+ * - do_dummy Bool/int: shall the dummy pixel be checked. Dummy pixel are pixels marked as bad and ignored
+ * - dummy Float: value for bad pixels
+ * - delta_dummy Float: precision for bad pixel value
+ * - normalization_factor : divide the input by this value
+ *
+ *
+**/
+
+__kernel void
+corrections3(const __global float *image,
+ const __global float *variance,
+ const char do_dark,
+ const __global float *dark,
+ const char do_dark_variance,
+ const __global float *dark_variance,
+ const char do_flat,
+ const __global float *flat,
+ const char do_solidangle,
+ const __global float *solidangle,
+ const char do_polarization,
+ const __global float *polarization,
+ const char do_absorption,
+ const __global float *absorption,
+ const char do_mask,
+ const __global char *mask,
+ const char do_dummy,
+ const float dummy,
+ const float delta_dummy,
+ const float normalization_factor,
+ __global float3 *output
+ )
+{
+ size_t i= get_global_id(0);
+ float3 result = (float3)(0.0, 0.0, 0.0);
+ if (i < NIMAGE)
+ {
+ result = _preproc3( image,
+ variance,
+ do_dark,
+ dark,
+ do_dark_variance,
+ dark_variance,
+ do_flat,
+ flat,
+ do_solidangle,
+ solidangle,
+ do_polarization,
+ polarization,
+ do_absorption,
+ absorption,
+ do_mask,
+ mask,
+ do_dummy,
+ dummy,
+ delta_dummy,
+ normalization_factor);
+ output[i] = result;
+ };//end if NIMAGE
+};//end kernel
+
+
generated by cgit v1.2.3 (git 2.25.1) at 2024-04-26 18:30:06 +0000