pyFAI.opencl package

pyFAI.opencl.OCLFullSplit module

class pyFAI.opencl.OCLFullSplit.OCLFullSplit1d(pos, bins=100, pos0Range=None, pos1Range=None, mask=None, mask_checksum=None, allow_pos0_neg=False, unit='undefined', workgroup_size=256, devicetype='all', platformid=None, deviceid=None, profile=False)

Bases: object

__init__(pos, bins=100, pos0Range=None, pos1Range=None, mask=None, mask_checksum=None, allow_pos0_neg=False, unit='undefined', workgroup_size=256, devicetype='all', platformid=None, deviceid=None, profile=False)

Initialize self. See help(type(self)) for accurate signature.

get_platform()
get_queue()

pyFAI.opencl.azim_csr module

class pyFAI.opencl.azim_csr.OCL_CSR_Integrator(lut, image_size, checksum=None, empty=None, unit=None, bin_centers=None, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=None, profile=False)

Bases: silx.opencl.processing.OpenclProcessing

Class in charge of doing a sparse-matrix multiplication in OpenCL using the CSR representation of the matrix.

It also performs the preprocessing using the preproc kernel

BLOCK_SIZE = 64
__init__(lut, image_size, checksum=None, empty=None, unit=None, bin_centers=None, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=None, profile=False)
Parameters

  • lut – 3-tuple of arrays data: coefficient of the matrix in a 1D vector of float32 - size of nnz indices: Column index position for the data (same size as data) indptr: row pointer indicates the start of a given row. len nbin+1

  • image_size – Expected image size: image.size

  • checksum – pre-calculated checksum of the LUT to prevent re-calculating it :)

  • empty – value to be assigned to bins without contribution from any pixel

  • unit – Storage for the unit related to the LUT

  • bin_centers – the radial position of the bin_center, place_holder

  • ctx – actual working context, left to None for automatic initialization from device type or platformid/deviceid

  • devicetype – type of device, can be “CPU”, “GPU”, “ACC” or “ALL”

  • platformid – integer with the platform_identifier, as given by clinfo

  • deviceid – Integer with the device identifier, as given by clinfo

  • block_size – preferred workgroup size, may vary depending on the outpcome of the compilation

  • profile – switch on profiling to be able to profile at the kernel level, store profiling elements (makes code slightly slower)

buffers = [BufferDescription(name='output', size=1, dtype=<class 'numpy.float32'>, flags=2), BufferDescription(name='output4', size=4, dtype=<class 'numpy.float32'>, flags=2), BufferDescription(name='image_raw', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='image', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='variance', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='dark', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='dark_variance', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='flat', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='polarization', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='solidangle', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='absorption', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='mask', size=1, dtype=<class 'numpy.int8'>, flags=4)]
compile_kernels(kernel_file=None)

Call the OpenCL compiler :param kernel_file: path to the kernel (by default use the one in the resources directory)

integrate(data, dummy=None, delta_dummy=None, dark=None, flat=None, solidangle=None, polarization=None, absorption=None, dark_checksum=None, flat_checksum=None, solidangle_checksum=None, polarization_checksum=None, absorption_checksum=None, preprocess_only=False, safe=True, normalization_factor=1.0, coef_power=1, out_merged=None, out_sum_data=None, out_sum_count=None)

Before performing azimuthal integration, the preprocessing is:

\[data = (data - dark) / (flat * solidangle * polarization)\]

Integration is performed using the CSR representation of the look-up table

Parameters

  • dark – array of same shape as data for pre-processing

  • flat – array of same shape as data for pre-processing

  • solidangle – array of same shape as data for pre-processing

  • polarization – array of same shape as data for pre-processing

  • dark_checksum – CRC32 checksum of the given array

  • flat_checksum – CRC32 checksum of the given array

  • solidangle_checksum – CRC32 checksum of the given array

  • polarization_checksum – CRC32 checksum of the given array

  • safe – if True (default) compares arrays on GPU according to their checksum, unless, use the buffer location is used

  • preprocess_only – return the dark subtracted; flat field & solidangle & polarization corrected image, else

  • normalization_factor – divide raw signal by this value

  • coef_power – set to 2 for variance propagation, leave to 1 for mean calculation

  • out_merged – destination array or pyopencl array for averaged data

  • out_sum_data – destination array or pyopencl array for sum of all data

  • out_sum_count – destination array or pyopencl array for sum of the number of pixels

Returns

averaged data, weighted histogram, unweighted histogram

integrate_legacy(data, dummy=None, delta_dummy=None, dark=None, flat=None, solidangle=None, polarization=None, absorption=None, dark_checksum=None, flat_checksum=None, solidangle_checksum=None, polarization_checksum=None, absorption_checksum=None, preprocess_only=False, safe=True, normalization_factor=1.0, coef_power=1, out_merged=None, out_sum_data=None, out_sum_count=None)

Before performing azimuthal integration, the preprocessing is:

\[data = (data - dark) / (flat * solidangle * polarization)\]

Integration is performed using the CSR representation of the look-up table

Parameters

  • dark – array of same shape as data for pre-processing

  • flat – array of same shape as data for pre-processing

  • solidangle – array of same shape as data for pre-processing

  • polarization – array of same shape as data for pre-processing

  • dark_checksum – CRC32 checksum of the given array

  • flat_checksum – CRC32 checksum of the given array

  • solidangle_checksum – CRC32 checksum of the given array

  • polarization_checksum – CRC32 checksum of the given array

  • safe – if True (default) compares arrays on GPU according to their checksum, unless, use the buffer location is used

  • preprocess_only – return the dark subtracted; flat field & solidangle & polarization corrected image, else

  • normalization_factor – divide raw signal by this value

  • coef_power – set to 2 for variance propagation, leave to 1 for mean calculation

  • out_merged – destination array or pyopencl array for averaged data

  • out_sum_data – destination array or pyopencl array for sum of all data

  • out_sum_count – destination array or pyopencl array for sum of the number of pixels

Returns

averaged data, weighted histogram, unweighted histogram

integrate_ng(data, dark=None, dummy=None, delta_dummy=None, poissonian=None, variance=None, dark_variance=None, flat=None, solidangle=None, polarization=None, absorption=None, dark_checksum=None, flat_checksum=None, solidangle_checksum=None, polarization_checksum=None, absorption_checksum=None, dark_variance_checksum=None, safe=True, normalization_factor=1.0, out_avgint=None, out_stderr=None, out_merged=None)

Before performing azimuthal integration with proper variance propagation, the preprocessing is:

\[signal = (raw - dark) variance = variance + dark_variance normalization = normalization_factor*(flat * solidangle * polarization * absortoption) count = number of pixel contributing\]

Integration is performed using the CSR representation of the look-up table on all arrays: signal, variance, normalization and count

Parameters

  • dark – array of same shape as data for pre-processing

  • dummy – value for invalid data

  • delta_dummy – precesion for dummy assessement

  • poissonian – set to use signal as variance (minimum 1)

  • variance – array of same shape as data for pre-processing

  • dark_variance – array of same shape as data for pre-processing

  • flat – array of same shape as data for pre-processing

  • solidangle – array of same shape as data for pre-processing

  • polarization – array of same shape as data for pre-processing

  • dark_checksum – CRC32 checksum of the given array

  • flat_checksum – CRC32 checksum of the given array

  • solidangle_checksum – CRC32 checksum of the given array

  • polarization_checksum – CRC32 checksum of the given array

  • safe – if True (default) compares arrays on GPU according to their checksum, unless, use the buffer location is used

  • preprocess_only – return the dark subtracted; flat field & solidangle & polarization corrected image, else

  • normalization_factor – divide raw signal by this value

  • out_avgint – destination array or pyopencl array for sum of all data

  • out_stderr – destination array or pyopencl array for sum of the number of pixels

  • out_merged – destination array or pyopencl array for averaged data (float8!)

Returns

out_avgint, out_stderr, out_merged

kernel_files = ['pyfai:openCL/kahan.cl', 'pyfai:openCL/preprocess.cl', 'pyfai:openCL/memset.cl', 'pyfai:openCL/ocl_azim_CSR.cl']
mapping = {<class 'numpy.int8'>: 's8_to_float', <class 'numpy.uint8'>: 'u8_to_float', <class 'numpy.int16'>: 's16_to_float', <class 'numpy.uint16'>: 'u16_to_float', <class 'numpy.uint32'>: 'u32_to_float', <class 'numpy.int32'>: 's32_to_float'}
send_buffer(data, dest, checksum=None)

Send a numpy array to the device, including the cast on the device if possible

Parameters

  • data – numpy array with data

  • dest – name of the buffer as registered in the class

set_kernel_arguments()

Tie arguments of OpenCL kernel-functions to the actual kernels

sigma_clip(data, dark=None, dummy=None, delta_dummy=None, variance=None, dark_variance=None, flat=None, solidangle=None, polarization=None, absorption=None, dark_checksum=None, flat_checksum=None, solidangle_checksum=None, polarization_checksum=None, absorption_checksum=None, dark_variance_checksum=None, safe=True, error_model=None, normalization_factor=1.0, cutoff=4.0, cycle=5, out_avgint=None, out_stderr=None, out_merged=None)

Perform a sigma-clipping iterative filter within each along each row. see the doc of scipy.stats.sigmaclip for more descriptions.

If the error model is “azimuthal”: the variance is the variance within a bin, which is refined at each iteration, can be costly !

Else, the error is propagated according to:

\[signal = (raw - dark) variance = variance + dark_variance normalization = normalization_factor*(flat * solidangle * polarization * absortoption) count = number of pixel contributing\]

Integration is performed using the CSR representation of the look-up table on all arrays: signal, variance, normalization and count

Parameters

  • dark – array of same shape as data for pre-processing

  • dummy – value for invalid data

  • delta_dummy – precesion for dummy assessement

  • variance – array of same shape as data for pre-processing

  • dark_variance – array of same shape as data for pre-processing

  • flat – array of same shape as data for pre-processing

  • solidangle – array of same shape as data for pre-processing

  • polarization – array of same shape as data for pre-processing

  • dark_checksum – CRC32 checksum of the given array

  • flat_checksum – CRC32 checksum of the given array

  • solidangle_checksum – CRC32 checksum of the given array

  • polarization_checksum – CRC32 checksum of the given array

  • safe – if True (default) compares arrays on GPU according to their checksum, unless, use the buffer location is used

  • preprocess_only – return the dark subtracted; flat field & solidangle & polarization corrected image, else

  • normalization_factor – divide raw signal by this value

  • cutoff – discard all points with |value - avg| > cutoff * sigma. 3-4 is quite common

  • cycle – perform at maximum this number of cycles. 5 is common.

  • out_avgint – destination array or pyopencl array for sum of all data

  • out_stderr – destination array or pyopencl array for sum of the number of pixels

  • out_merged – destination array or pyopencl array for averaged data (float8!)

Returns

namedtuple with “position intensity error signal variance normalization count”

pyFAI.opencl.azim_hist module

class pyFAI.opencl.azim_hist.Integrator1d(filename=None)

Bases: object

Attempt to implements ocl_azim using pyopencl

Implementation based on Dimitris’ work. Deprecated, marked for removal in version 1.0

BLOCK_SIZE = 128
__init__(filename=None)
Parameters

filename – file in which profiling information are saved

clean(preserve_context=False)

Free OpenCL related resources allocated by the library.

clean() is used to reinitiate the library back in a vanilla state. It may be asked to preserve the context created by init or completely clean up OpenCL. Guard/Status flags that are set will be reset.

Parameters

preserve_context (bool) – preserves or destroys all OpenCL resources

configure(kernel=None)

The method configure() allocates the OpenCL resources required and compiled the OpenCL kernels. An active context must exist before a call to configure() and getConfiguration() must have been called at least once. Since the compiled OpenCL kernels carry some information on the integration parameters, a change to any of the parameters of getConfiguration() requires a subsequent call to configure() for them to take effect.

If a configuration exists and configure() is called, the configuration is cleaned up first to avoid OpenCL memory leaks

Parameters

kernel_path – is the path to the actual kernel

execute(image)

Perform a 1D azimuthal integration

execute() may be called only after an OpenCL device is configured and a Tth array has been loaded (at least once) It takes the input image and based on the configuration provided earlier it performs the 1D integration. Notice that if the provided image is bigger than N then only N points will be taked into account, while if the image is smaller than N the result may be catastrophic. set/unset and loadTth methods have a direct impact on the execute() method. All the rest of the methods will require at least a new configuration via configure().

Takes an image, integrate and return the histogram and weights

Parameters

image – image to be processed as a numpy array

Returns

tth_out, histogram, bins

TODO: to improve performances, the image should be casted to float32 in an optimal way: currently using numpy machinery but would be better if done in OpenCL

getConfiguration(Nimage, Nbins, useFp64=None)

getConfiguration gets the description of the integrations to be performed and keeps an internal copy

Parameters

  • Nimage – number of pixel in image

  • Nbins – number of bins in regrouped histogram

  • useFp64 – use double precision. By default the same as init!

get_status()

return a dictionnary with the status of the integrator: for compatibilty with former implementation

init(devicetype='GPU', useFp64=True, platformid=None, deviceid=None)

Initial configuration: Choose a device and initiate a context. Devicetypes can be GPU, gpu, CPU, cpu, DEF, ACC, ALL. Suggested are GPU,CPU. For each setting to work there must be such an OpenCL device and properly installed. E.g.: If Nvidia driver is installed, GPU will succeed but CPU will fail. The AMD SDK kit (AMD APP) is required for CPU via OpenCL.

Parameters

  • devicetype – string in [“cpu”,”gpu”, “all”, “acc”]

  • useFp64 – boolean specifying if double precision will be used

  • platformid – integer

  • devid – integer

loadTth(tth, dtth, tth_min=None, tth_max=None)

Load the 2th arrays along with the min and max value.

loadTth maybe be recalled at any time of the execution in order to update the 2th arrays.

loadTth is required and must be called at least once after a configure()

log(**kwarg)

log in a file all opencl events

setDummyValue(dummy, delta_dummy)

Enables dummy value functionality and uploads the value to the OpenCL device.

Image values that are similar to the dummy value are set to 0.

Parameters

  • dummy – value in image of missing values (masked pixels?)

  • delta_dummy – precision for dummy values

setMask(mask)

Enables the use of a Mask during integration. The Mask can be updated by recalling setMask at any point.

The Mask must be a PyFAI Mask. Pixels with 0 are masked out. TODO: check and invert!

Parameters

mask – numpy.ndarray of integer.

setRange(lowerBound, upperBound)

Instructs the program to use a user - defined range for 2th values

setRange is optional. By default the integration will use the tth_min and tth_max given by loadTth() as integration range. When setRange is called it sets a new integration range without affecting the 2th array. All values outside that range will then be discarded when interpolating. Currently, if the interval of 2th (2th + -d2th) is not all inside the range specified, it is discarded. The bins of the histogram are RESCALED to the defined range and not the original tth_max - tth_min range.

setRange can be called at any point and as many times required after a valid configuration is created.

Parameters

  • lowerBound (float) – lower bound of the integration range

  • upperBound (float) – upper bound of the integration range

setSolidAngle(solidAngle)

Enables SolidAngle correction and uploads the suitable array to the OpenCL device.

By default the program will assume no solidangle correction unless setSolidAngle() is called. From then on, all integrations will be corrected via the SolidAngle array.

If the SolidAngle array needs to be changes, one may just call setSolidAngle() again with that array

Parameters

solidAngle (ndarray) – the solid angle of the given pixel

unsetDummyValue()

Disable a dummy value. May be re-enabled at any time by setDummyValue

unsetMask()

Disables the use of a Mask from that point. It may be re-enabled at any point via setMask

unsetRange()

Disable the use of a user-defined 2th range and revert to tth_min,tth_max range

unsetRange instructs the program to revert to its default integration range. If the method is called when no user-defined range had been previously specified, no action will be performed

unsetSolidAngle()

Instructs the program to not perform solidangle correction from now on.

SolidAngle correction may be turned back on at any point

class pyFAI.opencl.azim_hist.OCL_Histogram1d(radial, bins, radial_checksum=None, empty=None, unit=None, azimuthal=None, azimuthal_checksum=None, mask=None, mask_checksum=None, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=None, profile=False)

Bases: silx.opencl.processing.OpenclProcessing

Class in charge of performing histogram calculation in OpenCL using atomic_add

It also performs the preprocessing using the preproc kernel

BLOCK_SIZE = 32
__init__(radial, bins, radial_checksum=None, empty=None, unit=None, azimuthal=None, azimuthal_checksum=None, mask=None, mask_checksum=None, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=None, profile=False)
Parameters

  • radial – array with the radial position of every single pixel. Same as image size

  • bins – number of bins on which to histogram

  • checksum – pre-calculated checksum of the radial array to prevent re-calculating it :)

  • empty – value to be assigned to bins without contribution from any pixel

  • unit – just a place_holder for the units of radial array.

  • azimuthal – array with the azimuthal position, same size as radial

  • azimuthal_checksum – Checksum of the azimuthal array

  • mask – Array with the mask, 0 for valid values, anything for masked pixels, same size as radial

  • ctx – actual working context, left to None for automatic initialization from device type or platformid/deviceid

  • devicetype – type of device, can be “CPU”, “GPU”, “ACC” or “ALL”

  • platformid – integer with the platform_identifier, as given by clinfo

  • deviceid – Integer with the device identifier, as given by clinfo

  • block_size – preferred workgroup size, may vary depending on the outpcome of the compilation

  • profile – switch on profiling to be able to profile at the kernel level, store profiling elements (makes code slightly slower)

buffers = [BufferDescription(name='output4', size=4, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='radial', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='azimuthal', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='image_raw', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='image', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='variance', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='dark', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='dark_variance', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='flat', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='polarization', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='solidangle', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='absorption', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='mask', size=1, dtype=<class 'numpy.int8'>, flags=4)]
compile_kernels(kernel_file=None)

Call the OpenCL compiler

Parameters

kernel_file – path to the kernel (by default use the one in the resources directory)

integrate(data, dark=None, dummy=None, delta_dummy=None, poissonian=False, variance=None, dark_variance=None, flat=None, solidangle=None, polarization=None, absorption=None, dark_checksum=None, flat_checksum=None, solidangle_checksum=None, polarization_checksum=None, absorption_checksum=None, preprocess_only=False, safe=True, normalization_factor=1.0, radial_range=None, azimuth_range=None, histo_signal=None, histo_variance=None, histo_normalization=None, histo_count=None, intensity=None, error=None)

Performing azimuthal integration, the preprocessing is:

\[Signal= (data - dark) Variance= (variance + dark_variance) Normalization= (normalization_factor * flat * solidangle * polarization * absorption) Count= 1 per valid pixel\]

Integration is performed using the histograms (based on atomic adds

Parameters

  • dark – array of same shape as data for pre-processing

  • dummy – value for invalid data

  • delta_dummy – precesion for dummy assessement

  • poissonian – set to assume variance is data (minimum 1)

  • variance – array of same shape as data for pre-processing

  • dark_variance – array of same shape as data for pre-processing

  • flat – array of same shape as data for pre-processing

  • solidangle – array of same shape as data for pre-processing

  • polarization – array of same shape as data for pre-processing

  • dark_checksum – CRC32 checksum of the given array

  • flat_checksum – CRC32 checksum of the given array

  • solidangle_checksum – CRC32 checksum of the given array

  • polarization_checksum – CRC32 checksum of the given array

  • safe – if True (default) compares arrays on GPU according to their checksum, unless, use the buffer location is used

  • preprocess_only – return the dark subtracted; flat field & solidangle & polarization corrected image, else

  • normalization_factor – divide raw signal by this value

  • radial_range – provide lower and upper bound for radial array

  • azimuth_range – provide lower and upper bound for azimuthal array

  • histo_signal – destination array or pyopencl array for sum of signals

  • histo_normalization – destination array or pyopencl array for sum of normalization

  • histo_count – destination array or pyopencl array for counting pixels

  • intensity – destination PyOpenCL array for integrated intensity

  • error – destination PyOpenCL array for standart deviation

Returns

bin_positions, averaged data, histogram of signal, histogram of variance, histogram of normalization, count of pixels

kernel_files = ['pyfai:openCL/kahan.cl', 'pyfai:openCL/preprocess.cl', 'pyfai:openCL/ocl_histo.cl']
mapping = {<class 'numpy.int8'>: 's8_to_float', <class 'numpy.uint8'>: 'u8_to_float', <class 'numpy.int16'>: 's16_to_float', <class 'numpy.uint16'>: 'u16_to_float', <class 'numpy.uint32'>: 'u32_to_float', <class 'numpy.int32'>: 's32_to_float'}
send_buffer(data, dest, checksum=None)

Send a numpy array to the device, including the cast on the device if possible

Parameters

  • data – numpy array with data

  • dest – name of the buffer as registered in the class

set_kernel_arguments()

Tie arguments of OpenCL kernel-functions to the actual kernels

class pyFAI.opencl.azim_hist.OCL_Histogram2d(radial, azimuthal, bins_radial, bins_azimuthal, radial_checksum=None, azimuthal_checksum=None, empty=None, unit=None, mask=None, mask_checksum=None, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=None, profile=False)

Bases: pyFAI.opencl.azim_hist.OCL_Histogram1d

Class in charge of performing histogram calculation in OpenCL using atomic_add

It also performs the preprocessing using the preproc kernel

BLOCK_SIZE = 32
__init__(radial, azimuthal, bins_radial, bins_azimuthal, radial_checksum=None, azimuthal_checksum=None, empty=None, unit=None, mask=None, mask_checksum=None, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=None, profile=False)
Parameters

  • radial – array with the radial position of every single pixel. Same as image size

  • azimuthal – array with the azimuthal position of every single pixel. Same as image size

  • bins_radial – number of bins on which to histogram is calculated in radial direction

  • bins_azimuthal – number of bins on which to histogram is calculated in azimuthal direction

  • radial_checksum – pre-calculated checksum of the position array to prevent re-calculating it :)

  • azimuthal_checksum – pre-calculated checksum of the position array to prevent re-calculating it :)

  • empty – value to be assigned to bins without contribution from any pixel

  • ctx – actual working context, left to None for automatic initialization from device type or platformid/deviceid

  • devicetype – type of device, can be “CPU”, “GPU”, “ACC” or “ALL”

  • platformid – integer with the platform_identifier, as given by clinfo

  • deviceid – Integer with the device identifier, as given by clinfo

  • block_size – preferred workgroup size, may vary depending on the outpcome of the compilation

  • profile – switch on profiling to be able to profile at the kernel level, store profiling elements (makes code slightly slower)

buffers = [BufferDescription(name='output4', size=4, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='radial', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='azimuthal', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='image_raw', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='image', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='variance', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='dark', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='dark_variance', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='flat', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='polarization', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='solidangle', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='absorption', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='mask', size=1, dtype=<class 'numpy.int8'>, flags=4)]
integrate(data, dark=None, dummy=None, delta_dummy=None, variance=None, dark_variance=None, flat=None, solidangle=None, polarization=None, absorption=None, dark_checksum=None, flat_checksum=None, solidangle_checksum=None, polarization_checksum=None, absorption_checksum=None, preprocess_only=False, safe=True, normalization_factor=1.0, radial_range=None, azimuthal_range=None, histo_signal=None, histo_variance=None, histo_normalization=None, histo_count=None, intensity=None, error=None)

Performing azimuthal integration, the preprocessing is:

\[Signal= (data - dark) Variance= (variance + dark_variance) Normalization= (normalization_factor * flat * solidangle * polarization * absorption) Count= 1 per valid pixel\]

Integration is performed using the histograms (based on atomic adds

Parameters

  • dark – array of same shape as data for pre-processing

  • dummy – value for invalid data

  • delta_dummy – precesion for dummy assessement

  • variance – array of same shape as data for pre-processing

  • dark_variance – array of same shape as data for pre-processing

  • flat – array of same shape as data for pre-processing

  • solidangle – array of same shape as data for pre-processing

  • polarization – array of same shape as data for pre-processing

  • dark_checksum – CRC32 checksum of the given array

  • flat_checksum – CRC32 checksum of the given array

  • solidangle_checksum – CRC32 checksum of the given array

  • polarization_checksum – CRC32 checksum of the given array

  • safe – if True (default) compares arrays on GPU according to their checksum, unless, use the buffer location is used

  • preprocess_only – return the dark subtracted; flat field & solidangle & polarization corrected image, else

  • normalization_factor – divide raw signal by this value

  • radial_range – provide lower and upper bound for radial array

  • azimuth_range – provide lower and upper bound for azimuthal array

  • histo_signal – destination array or pyopencl array for sum of signals

  • histo_normalization – destination array or pyopencl array for sum of normalization

  • histo_count – destination array or pyopencl array for counting pixels

  • intensity – destination PyOpenCL array for integrated intensity

  • error – destination PyOpenCL array for standart deviation

Returns

bin_positions, averaged data, histogram of signal, histogram of variance, histogram of normalization, count of pixels

kernel_files = ['pyfai:openCL/kahan.cl', 'pyfai:openCL/preprocess.cl', 'pyfai:openCL/ocl_histo.cl']
mapping = {<class 'numpy.int8'>: 's8_to_float', <class 'numpy.uint8'>: 'u8_to_float', <class 'numpy.int16'>: 's16_to_float', <class 'numpy.uint16'>: 'u16_to_float', <class 'numpy.uint32'>: 'u32_to_float', <class 'numpy.int32'>: 's32_to_float'}
set_kernel_arguments()

Tie arguments of OpenCL kernel-functions to the actual kernels

pyFAI.opencl.azim_lut module

class pyFAI.opencl.azim_lut.OCL_LUT_Integrator(lut, image_size, checksum=None, empty=None, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=None, profile=False)

Bases: silx.opencl.processing.OpenclProcessing

Class in charge of doing a sparse-matrix multiplication in OpenCL using the LUT representation of the matrix. It also performs the preprocessing using the preproc kernel

BLOCK_SIZE = 16
__init__(lut, image_size, checksum=None, empty=None, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=None, profile=False)

Constructor of the OCL_LUT_Integrator class

Parameters

  • lut – array of int32 - float32 with shape (nbins, lut_size) with indexes and coefficients

  • image_size – Expected image size: image.size

  • checksum – pre-calculated checksum of the LUT to prevent re-calculating it :)

  • empty – value to be assigned to bins without contribution from any pixel

  • ctx – actual working context, left to None for automatic initialization from device type or platformid/deviceid

  • devicetype – type of device, can be “CPU”, “GPU”, “ACC” or “ALL”

  • platformid – integer with the platform_identifier, as given by clinfo

  • deviceid – Integer with the device identifier, as given by clinfo

  • block_size – preferred workgroup size, may vary depending on the outpcome of the compilation

  • profile – switch on profiling to be able to profile at the kernel level, store profiling elements (makes code slightly slower)

buffers = [BufferDescription(name='output', size=1, dtype=<class 'numpy.float32'>, flags=2), BufferDescription(name='image_raw', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='image', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='variance', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='dark', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='dark_variance', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='flat', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='polarization', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='solidangle', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='absorption', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='mask', size=1, dtype=<class 'numpy.int8'>, flags=4)]
compile_kernels(kernel_file=None)

Call the OpenCL compiler :param kernel_file: path to the kernel (by default use the one in the resources directory)

integrate(data, dummy=None, delta_dummy=None, dark=None, flat=None, solidangle=None, polarization=None, absorption=None, dark_checksum=None, flat_checksum=None, solidangle_checksum=None, polarization_checksum=None, absorption_checksum=None, preprocess_only=False, safe=True, normalization_factor=1.0, coef_power=1, out_merged=None, out_sum_data=None, out_sum_count=None)

Before performing azimuthal integration, the preprocessing is:

\[data = (data - dark) / (flat * solidangle * polarization * absorption)\]

Integration is performed using the LUT representation of the look-up table

Parameters

  • dark – array of same shape as data for pre-processing

  • flat – array of same shape as data for pre-processing

  • solidangle – array of same shape as data for pre-processing

  • polarization – array of same shape as data for pre-processing

  • dark_checksum – CRC32 checksum of the given array

  • flat_checksum – CRC32 checksum of the given array

  • solidangle_checksum – CRC32 checksum of the given array

  • polarization_checksum – CRC32 checksum of the given array

  • safe – if True (default) compares arrays on GPU according to their checksum, unless, use the buffer location is used

  • preprocess_only – return the dark subtracted; flat field & solidangle & polarization corrected image, else

  • normalization_factor – divide raw signal by this value

  • coef_power – set to 2 for variance propagation, leave to 1 for mean calculation

  • out_merged – destination array or pyopencl array for averaged data

  • out_sum_data – destination array or pyopencl array for sum of all data

  • out_sum_count – destination array or pyopencl array for sum of the number of pixels

Returns

averaged data, weighted histogram, unweighted histogram

integrate_legacy(data, dummy=None, delta_dummy=None, dark=None, flat=None, solidangle=None, polarization=None, absorption=None, dark_checksum=None, flat_checksum=None, solidangle_checksum=None, polarization_checksum=None, absorption_checksum=None, preprocess_only=False, safe=True, normalization_factor=1.0, coef_power=1, out_merged=None, out_sum_data=None, out_sum_count=None)

Before performing azimuthal integration, the preprocessing is:

\[data = (data - dark) / (flat * solidangle * polarization * absorption)\]

Integration is performed using the LUT representation of the look-up table

Parameters

  • dark – array of same shape as data for pre-processing

  • flat – array of same shape as data for pre-processing

  • solidangle – array of same shape as data for pre-processing

  • polarization – array of same shape as data for pre-processing

  • dark_checksum – CRC32 checksum of the given array

  • flat_checksum – CRC32 checksum of the given array

  • solidangle_checksum – CRC32 checksum of the given array

  • polarization_checksum – CRC32 checksum of the given array

  • safe – if True (default) compares arrays on GPU according to their checksum, unless, use the buffer location is used

  • preprocess_only – return the dark subtracted; flat field & solidangle & polarization corrected image, else

  • normalization_factor – divide raw signal by this value

  • coef_power – set to 2 for variance propagation, leave to 1 for mean calculation

  • out_merged – destination array or pyopencl array for averaged data

  • out_sum_data – destination array or pyopencl array for sum of all data

  • out_sum_count – destination array or pyopencl array for sum of the number of pixels

Returns

averaged data, weighted histogram, unweighted histogram

integrate_ng(*arg, **kwargs)
kernel_files = ['pyfai:openCL/kahan.cl', 'pyfai:openCL/preprocess.cl', 'pyfai:openCL/memset.cl', 'pyfai:openCL/ocl_azim_LUT.cl']
mapping = {<class 'numpy.int8'>: 's8_to_float', <class 'numpy.uint8'>: 'u8_to_float', <class 'numpy.int16'>: 's16_to_float', <class 'numpy.uint16'>: 'u16_to_float', <class 'numpy.uint32'>: 'u32_to_float', <class 'numpy.int32'>: 's32_to_float'}
send_buffer(data, dest, checksum=None)

Send a numpy array to the device, including the cast on the device if possible

Parameters

  • data – numpy array with data

  • dest – name of the buffer as registered in the class

set_kernel_arguments()

Tie arguments of OpenCL kernel-functions to the actual kernels

set_kernel_arguments() is a private method, called by configure(). It uses the dictionary _cl_kernel_args. Note that by default, since TthRange is disabled, the integration kernels have tth_min_max tied to the tthRange argument slot. When setRange is called it replaces that argument with tthRange low and upper bounds. When unsetRange is called, the argument slot is reset to tth_min_max.

pyFAI.opencl.ocl_hist_pixelsplit module

class pyFAI.opencl.ocl_hist_pixelsplit.OCL_Hist_Pixelsplit(pos, bins, image_size, pos0Range=None, pos1Range=None, devicetype='all', padded=False, block_size=32, platformid=None, deviceid=None, checksum=None, profile=False)

Bases: object

__init__(pos, bins, image_size, pos0Range=None, pos1Range=None, devicetype='all', padded=False, block_size=32, platformid=None, deviceid=None, checksum=None, profile=False)
Parameters

  • lut – 3-tuple of arrays data: coefficient of the matrix in a 1D vector of float32 - size of nnz indices: Column index position for the data (same size as data) indptr: row pointer indicates the start of a given row. len nbin+1

  • image_size – size of the image (for pre-processing)

  • devicetype – can be “cpu”,”gpu”,”acc” or “all”

  • platformid (int) – number of the platform as given by clinfo

  • deviceid (int) – number of the device as given by clinfo

  • checksum – pre - calculated checksum to prevent re - calculating it :)

  • profile – store profiling elements

integrate(data, dummy=None, delta_dummy=None, dark=None, flat=None, solidAngle=None, polarization=None, dark_checksum=None, flat_checksum=None, solidAngle_checksum=None, polarization_checksum=None)
log_profile()

If we are in profiling mode, prints out all timing for every single OpenCL call

pyFAI.opencl.preproc module

OpenCL implementation of the preproc module

class pyFAI.opencl.preproc.OCL_Preproc(image_size=None, image_dtype=None, image=None, dark=None, flat=None, solidangle=None, polarization=None, absorption=None, mask=None, dummy=None, delta_dummy=None, empty=None, split_result=False, calc_variance=False, poissonian=False, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=32, profile=False)

Bases: silx.opencl.processing.OpenclProcessing

OpenCL class for pre-processing … mainly for demonstration

__init__(image_size=None, image_dtype=None, image=None, dark=None, flat=None, solidangle=None, polarization=None, absorption=None, mask=None, dummy=None, delta_dummy=None, empty=None, split_result=False, calc_variance=False, poissonian=False, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=32, profile=False)
Parameters

  • image_size – (int) number of element of the input image

  • image_dtype – dtype of the input image

  • image – retrieve image_size and image_dtype from template

  • dark – dark current image as numpy array

  • flat – flat field image as numpy array

  • solidangle – solid angle image as numpy array

  • absorption – absorption image as numpy array

  • mask – array of int8 with 0 where the data are valid

  • dummy – value of impossible values: dynamic mask

  • delta_dummy – precision for dummy values

  • empty – value to be assigned to pixel without contribution (i.e masked)

  • split_result – return the result a tuple: data, [variance], normalization, so the last dim becomes 2 or 3

  • calc_variance – report the result as data, variance, normalization

  • poissonian – assumes poisson law for data and dark,

  • ctx – actual working context, left to None for automatic initialization from device type or platformid/deviceid

  • devicetype – type of device, can be “CPU”, “GPU”, “ACC” or “ALL”

  • platformid – integer with the platform_identifier, as given by clinfo

  • deviceid – Integer with the device identifier, as given by clinfo

  • block_size – preferred workgroup size, may vary depending on the outpcome of the compilation

  • profile – switch on profiling to be able to profile at the kernel level, store profiling elements (makes code slower)

buffers = [BufferDescription(name='output', size=3, dtype=<class 'numpy.float32'>, flags=2), BufferDescription(name='image_raw', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='image', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='variance', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='dark', size=1, dtype=<class 'numpy.float32'>, flags=1), BufferDescription(name='dark_variance', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='flat', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='polarization', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='solidangle', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='absorption', size=1, dtype=<class 'numpy.float32'>, flags=4), BufferDescription(name='mask', size=1, dtype=<class 'numpy.int8'>, flags=4)]
compile_kernels(kernel_files=None, compile_options=None)

Call the OpenCL compiler

Parameters

kernel_files – list of path to the kernel (by default use the one declared in the class)

property delta_dummy
property dummy
property empty
kernel_files = ['pyfai:openCL/preprocess.cl']
mapping = {<class 'numpy.int8'>: 's8_to_float', <class 'numpy.uint8'>: 'u8_to_float', <class 'numpy.int16'>: 's16_to_float', <class 'numpy.uint16'>: 'u16_to_float', <class 'numpy.uint32'>: 'u32_to_float', <class 'numpy.int32'>: 's32_to_float'}
process(image, dark=None, variance=None, dark_variance=None, normalization_factor=1.0, poissonian=None, split_result=None)

Perform the pixel-wise operation of the array

Parameters

  • raw – numpy array with the input image

  • dark – numpy array with the dark-current image

  • variance – numpy array with the variance of input image

  • dark_variance – numpy array with the variance of dark-current image

  • normalization_factor – divide the result by this

  • poissonian – set to true to set variance=signal (minimum 1). None uses the default from constructor

Returns

array with processed data, may be an array of (data,variance,normalization) depending on class initialization

send_buffer(data, dest)

Send a numpy array to the device

Parameters

  • data – numpy array with data

  • dest – name of the buffer as registered in the class

set_kernel_arguments()

Tie arguments of OpenCL kernel-functions to the actual kernels

pyFAI.opencl.preproc.preproc(raw, dark=None, flat=None, solidangle=None, polarization=None, absorption=None, mask=None, dummy=None, delta_dummy=None, normalization_factor=1.0, empty=None, split_result=False, variance=None, dark_variance=None, poissonian=False, dtype=<class 'numpy.float32'>)

Common preprocessing step, implemented using OpenCL. May be inefficient

Parameters

  • data – raw value, as a numpy array, 1D or 2D

  • mask – array non null where data should be ignored

  • dummy – value of invalid data

  • delta_dummy – precision for invalid data

  • dark – array containing the value of the dark noise, to be subtracted

  • flat – Array containing the flatfield image. It is also checked for dummies if relevant.

  • solidangle – the value of the solid_angle. This processing may be performed during the rebinning instead. left for compatibility

  • polarization – Correction for polarization of the incident beam

  • absorption – Correction for absorption in the sensor volume

  • normalization_factor – final value is divided by this

  • empty – value to be given for empty pixels

  • split_result – set to true to separate numerator from denominator and return an array of float2 or float3 (with variance)

  • variance – provide an estimation of the variance, enforce split_result=True and return an float3 array with variance in second position.

  • poissonian – set to “True” for assuming the detector is poissonian and variance = raw + dark

  • dtype – dtype for all processing

All calculation are performed in single precision floating point (32 bits).

NaN are always considered as invalid values

if neither empty nor dummy is provided, empty pixels are 0. Empty pixels are always zero in “split_result” mode

Split result:

  • When set to False, i.e the default, the pixel-wise operation is: I = (raw - dark)/(flat * solidangle * polarization * absorption) Invalid pixels are set to the dummy or empty value.

  • When split_ressult is set to True, each result result is a float2 or a float3 (with an additional value for the variance) as such:

    I = [(raw - dark), (variance), (flat * solidangle * polarization * absorption)]

    Empty pixels will have all their 2 or 3 values to 0 (and not to dummy or empty value)

  • If poissonian is set to True, the variance is evaluated as (raw + dark) minimum (1)

pyFAI.opencl.setup module

pyFAI.opencl.setup.configuration(parent_package='', top_path=None)

pyFAI.opencl.sort module

Module for 2D sort based on OpenCL for median filtering and Bragg/amorphous separation on GPU.

class pyFAI.opencl.sort.Separator(npt_height=512, npt_width=1024, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=None, profile=False)

Bases: silx.opencl.processing.OpenclProcessing

Implementation of sort, median filter and trimmed-mean in pyopencl

DUMMY = -3.4028235e+38
__init__(npt_height=512, npt_width=1024, ctx=None, devicetype='all', platformid=None, deviceid=None, block_size=None, profile=False)
Parameters

  • ctx – context

  • block_size – 1 on macOSX on CPU

  • profile – turn on profiling

allocate_buffers(*arg, **kwarg)

Allocate OpenCL buffers required for a specific configuration

Note that an OpenCL context also requires some memory, as well as Event and other OpenCL functionalities which cannot and are not taken into account here. The memory required by a context varies depending on the device. Typical for GTX580 is 65Mb but for a 9300m is ~15Mb In addition, a GPU will always have at least 3-5Mb of memory in use. Unfortunately, OpenCL does NOT have a built-in way to check the actual free memory on a device, only the total memory.

filter_horizontal(data, dummy=None, quantile=0.5)

Sort the data along the vertical axis (azimuthal)

Parameters

  • data – numpy or pyopencl array

  • dummy – dummy value

  • quantile –

Returns

pyopencl array

filter_vertical(data, dummy=None, quantile=0.5)

Sort the data along the vertical axis (azimuthal)

Parameters

  • data – numpy or pyopencl array

  • dummy – dummy value

  • quantile –

Returns

pyopencl array

kernel_files = ['pyfai:openCL/kahan.cl', 'pyfai:openCL/bitonic.cl', 'pyfai:openCL/separate.cl', 'pyfai:openCL/sigma_clip.cl']
mean_std_horizontal(data, dummy=None)

calculates the mean and std along a row

mean_std_vertical(data, dummy=None)

calculates the mean and std along a column, column size has to be multiple of 8 and <8192

set_kernel_arguments()

Tie arguments of OpenCL kernel-functions to the actual kernels

set_kernel_arguments() is a private method, called by configure(). It uses the dictionary _cl_kernel_args.

Note that by default, since TthRange is disabled, the integration kernels have tth_min_max tied to the tthRange argument slot.

When setRange is called it replaces that argument with tthRange low and upper bounds. When unsetRange is called, the argument slot is reset to tth_min_max.

sigma_clip_horizontal(data, sigma_lo=3, sigma_hi=None, max_iter=5, dummy=None)

calculates iterative sigma-clipped mean and std per row. column size has to be multiple of 8 and <8192

sigma_clip_vertical(data, sigma_lo=3, sigma_hi=None, max_iter=5, dummy=None)

calculates iterative sigma-clipped mean and std per column. column size has to be multiple of 8 and <8192

sort_horizontal(data, dummy=None)

Sort the data along the horizontal axis (radial)

Parameters

  • data – numpy or pyopencl array

  • dummy – dummy value

Returns

pyopencl array

sort_vertical(data, dummy=None)

Sort the data along the vertical axis (azimuthal)

Parameters

  • data – numpy or pyopencl array

  • dummy – dummy value

Returns

pyopencl array

trimmed_mean_horizontal(data, dummy=None, quantiles=(0.5, 0.5))

Perform a trimmed mean (mean without the extremes) After sorting the data along the vertical axis (azimuthal)

Parameters

  • data – numpy or pyopencl array

  • dummy – dummy value

  • quantile –

Returns

pyopencl array

trimmed_mean_vertical(data, dummy=None, quantiles=(0.5, 0.5))

Perform a trimmed mean (mean without the extremes) After sorting the data along the vertical axis (azimuthal)

Parameters

  • data – numpy or pyopencl array

  • dummy – dummy value

  • quantile –

Returns

pyopencl array

Module contents

Contains all OpenCL implementation.

pyFAI.opencl.get_x87_volatile_option(ctx)