diff --git a/src/calng/AgipdCorrection.py b/src/calng/AgipdCorrection.py
index 14012742a24d9870efc0d368425b23e3b01eadca..d0e46321b1b1e0b6dbdbdb0b049eb60c67fc376c 100644
--- a/src/calng/AgipdCorrection.py
+++ b/src/calng/AgipdCorrection.py
@@ -1,6 +1,10 @@
+import enum
+
+import cupy
import numpy as np
from karabo.bound import (
BOOL_ELEMENT,
+ DOUBLE_ELEMENT,
FLOAT_ELEMENT,
KARABO_CLASSINFO,
NODE_ELEMENT,
@@ -10,11 +14,405 @@ from karabo.bound import (
)
from karabo.common.states import State
-from . import utils
+from . import base_gpu, calcat_utils, utils
from ._version import version as deviceVersion
-from .agipd_gpu import AgipdGainMode, AgipdGpuRunner, BadPixelValues, CorrectionFlags
from .base_correction import BaseCorrection, add_correction_step_schema, preview_schema
-from .calcat_utils import AgipdCalcatFriend, AgipdConstants
+
+
+class AgipdConstants(enum.Enum):
+ SlopesFF = enum.auto()
+ ThresholdsDark = enum.auto()
+ Offset = enum.auto()
+ SlopesPC = enum.auto()
+ BadPixelsDark = enum.auto()
+ BadPixelsPC = enum.auto()
+ BadPixelsFF = enum.auto()
+
+
+# from pycalibration's enum.py
+class AgipdGainMode(enum.IntEnum):
+ ADAPTIVE_GAIN = 0
+ FIXED_HIGH_GAIN = 1
+ FIXED_MEDIUM_GAIN = 2
+ FIXED_LOW_GAIN = 3
+
+
+class BadPixelValues(enum.IntFlag):
+ """The European XFEL Bad Pixel Encoding
+
+ Straight from pycalibration's enum.py"""
+
+ OFFSET_OUT_OF_THRESHOLD = 2 ** 0
+ NOISE_OUT_OF_THRESHOLD = 2 ** 1
+ OFFSET_NOISE_EVAL_ERROR = 2 ** 2
+ NO_DARK_DATA = 2 ** 3
+ CI_GAIN_OUT_OF_THRESHOLD = 2 ** 4
+ CI_LINEAR_DEVIATION = 2 ** 5
+ CI_EVAL_ERROR = 2 ** 6
+ FF_GAIN_EVAL_ERROR = 2 ** 7
+ FF_GAIN_DEVIATION = 2 ** 8
+ FF_NO_ENTRIES = 2 ** 9
+ CI2_EVAL_ERROR = 2 ** 10
+ VALUE_IS_NAN = 2 ** 11
+ VALUE_OUT_OF_RANGE = 2 ** 12
+ GAIN_THRESHOLDING_ERROR = 2 ** 13
+ DATA_STD_IS_ZERO = 2 ** 14
+ ASIC_STD_BELOW_NOISE = 2 ** 15
+ INTERPOLATED = 2 ** 16
+ NOISY_ADC = 2 ** 17
+ OVERSCAN = 2 ** 18
+ NON_SENSITIVE = 2 ** 19
+ NON_LIN_RESPONSE_REGION = 2 ** 20
+
+
+class CorrectionFlags(enum.IntFlag):
+ NONE = 0
+ THRESHOLD = 1
+ OFFSET = 2
+ BLSHIFT = 4
+ REL_GAIN_PC = 8
+ REL_GAIN_XRAY = 16
+ BPMASK = 32
+
+
+class AgipdGpuRunner(base_gpu.BaseGpuRunner):
+ _kernel_source_filename = "agipd_gpu.cu"
+ _corrected_axis_order = "cxy"
+
+ def __init__(
+ self,
+ pixels_x,
+ pixels_y,
+ memory_cells,
+ constant_memory_cells,
+ input_data_dtype=cupy.uint16,
+ output_data_dtype=cupy.float32,
+ bad_pixel_mask_value=cupy.nan,
+ gain_mode=AgipdGainMode.ADAPTIVE_GAIN,
+ g_gain_value=1,
+ ):
+ self.gain_mode = gain_mode
+ if self.gain_mode is AgipdGainMode.ADAPTIVE_GAIN:
+ self.default_gain = cupy.uint8(gain_mode)
+ else:
+ self.default_gain = cupy.uint8(gain_mode - 1)
+ self.input_shape = (memory_cells, 2, pixels_x, pixels_y)
+ self.processed_shape = (memory_cells, pixels_x, pixels_y)
+ super().__init__(
+ pixels_x,
+ pixels_y,
+ memory_cells,
+ constant_memory_cells,
+ input_data_dtype,
+ output_data_dtype,
+ )
+ self.gain_map_gpu = cupy.empty(self.processed_shape, dtype=cupy.float32)
+ self.preview_buffer_getters.extend(
+ [self._get_raw_gain_for_preview, self._get_gain_map_for_preview]
+ )
+
+ self.map_shape = (self.constant_memory_cells, self.pixels_x, self.pixels_y)
+ self.gm_map_shape = self.map_shape + (3,) # for gain-mapped constants
+ self.threshold_map_shape = self.map_shape + (2,)
+ # constants
+ self.gain_thresholds_gpu = cupy.empty(
+ self.threshold_map_shape, dtype=cupy.float32
+ )
+ self.offset_map_gpu = cupy.zeros(self.gm_map_shape, dtype=cupy.float32)
+ self.rel_gain_pc_map_gpu = cupy.ones(self.gm_map_shape, dtype=cupy.float32)
+ # not gm_map_shape because it only applies to medium gain pixels
+ self.md_additional_offset_gpu = cupy.zeros(self.map_shape, dtype=cupy.float32)
+ self.rel_gain_xray_map_gpu = cupy.ones(self.map_shape, dtype=cupy.float32)
+ self.bad_pixel_map_gpu = cupy.zeros(self.gm_map_shape, dtype=cupy.uint32)
+ self.set_bad_pixel_mask_value(bad_pixel_mask_value)
+ self.set_g_gain_value(g_gain_value)
+
+ self.update_block_size((1, 1, 64))
+
+ def _get_raw_for_preview(self):
+ return self.input_data_gpu[:, 0]
+
+ def _get_corrected_for_preview(self):
+ return self.processed_data_gpu
+
+ # special to AGIPD
+ def _get_raw_gain_for_preview(self):
+ return self.input_data_gpu[:, 1]
+
+ def _get_gain_map_for_preview(self):
+ return self.gain_map_gpu
+
+ def load_thresholds(self, threshold_map):
+ # shape: y, x, memory cell, thresholds and gain values
+ # note: the gain values are something like means used to derive thresholds
+ self.gain_thresholds_gpu.set(
+ np.transpose(threshold_map[..., :2], (2, 1, 0, 3)).astype(np.float32)
+ )
+
+ def load_offset_map(self, offset_map):
+ # shape: y, x, memory cell, gain stage
+ self.offset_map_gpu.set(
+ np.transpose(offset_map, (2, 1, 0, 3)).astype(np.float32)
+ )
+
+ def load_rel_gain_pc_map(self, slopes_pc_map, override_md_additional_offset=None):
+ # pc has funny shape (11, 352, 128, 512) from file
+ # this is (fi, memory cell, y, x)
+ slopes_pc_map = slopes_pc_map.astype(np.float32)
+ # the following may contain NaNs, though...
+ hg_slope = slopes_pc_map[0]
+ hg_intercept = slopes_pc_map[1]
+ mg_slope = slopes_pc_map[3]
+ mg_intercept = slopes_pc_map[4]
+ # TODO: remove sanitization (should happen in constant preparation notebook)
+ # from agipdlib.py: replace NaN with median (per memory cell)
+ # note: suffixes in agipdlib are "_m" and "_l", should probably be "_I"
+ for naughty_array in (hg_slope, hg_intercept, mg_slope, mg_intercept):
+ medians = np.nanmedian(naughty_array, axis=(1, 2))
+ nan_bool = np.isnan(naughty_array)
+ nan_cell, _, _ = np.where(nan_bool)
+ naughty_array[nan_bool] = medians[nan_cell]
+
+ too_low_bool = naughty_array < 0.8 * medians[:, np.newaxis, np.newaxis]
+ too_low_cell, _, _ = np.where(too_low_bool)
+ naughty_array[too_low_bool] = medians[too_low_cell]
+
+ too_high_bool = naughty_array > 1.2 * medians[:, np.newaxis, np.newaxis]
+ too_high_cell, _, _ = np.where(too_high_bool)
+ naughty_array[too_high_bool] = medians[too_high_cell]
+
+ frac_hg_mg = hg_slope / mg_slope
+ rel_gain_map = np.ones(
+ (3, self.constant_memory_cells, self.pixels_y, self.pixels_x),
+ dtype=np.float32,
+ )
+ rel_gain_map[1] = rel_gain_map[0] * frac_hg_mg
+ rel_gain_map[2] = rel_gain_map[1] * 4.48
+ self.rel_gain_pc_map_gpu.set(np.transpose(rel_gain_map, (1, 3, 2, 0)))
+ if override_md_additional_offset is None:
+ md_additional_offset = (hg_intercept - mg_intercept * frac_hg_mg).astype(
+ np.float32
+ )
+ self.md_additional_offset_gpu.set(
+ np.transpose(md_additional_offset, (0, 2, 1))
+ )
+ else:
+ self.override_md_additional_offset(override_md_additional_offset)
+
+ def override_md_additional_offset(self, override_value):
+ self.md_additional_offset_gpu.fill(override_value)
+
+ def load_rel_gain_ff_map(self, slopes_ff_map):
+ # constant shape: y, x, memory cell
+ if slopes_ff_map.shape[2] == 2:
+ # TODO: remove support for old format
+ # old format, is per pixel only (shape is y, x, 2)
+ # note: we should not support this in online
+ slopes_ff_map = np.broadcast_to(
+ slopes_ff_map[..., 0][..., np.newaxis],
+ (self.pixels_y, self.pixels_x, self.constant_memory_cells),
+ )
+ self.rel_gain_xray_map_gpu.set(np.transpose(slopes_ff_map).astype(np.float32))
+
+ def set_g_gain_value(self, override_value):
+ self.g_gain_value = cupy.float32(override_value)
+
+ def load_bad_pixels_map(self, bad_pixels_map, override_flags_to_use=None):
+ print(f"Loading bad pixels with shape: {bad_pixels_map.shape}")
+ # will simply OR with already loaded, does not take into account which ones
+ # TODO: inquire what "mask for double size pixels" means
+ if len(bad_pixels_map.shape) == 3:
+ if bad_pixels_map.shape == (
+ self.pixels_y,
+ self.pixels_x,
+ self.constant_memory_cells,
+ ):
+ # BadPixelsFF is not per gain stage - broadcasting along gain dimension
+ self.bad_pixel_map_gpu |= cupy.asarray(
+ np.broadcast_to(
+ np.transpose(bad_pixels_map)[..., np.newaxis],
+ self.gm_map_shape,
+ ),
+ dtype=np.uint32,
+ )
+ elif bad_pixels_map.shape == (
+ self.constant_memory_cells,
+ self.pixels_y,
+ self.pixels_x,
+ ):
+ # oh, can also be old bad pixels pc?
+ self.bad_pixel_map_gpu |= cupy.asarray(
+ np.broadcast_to(
+ np.transpose(bad_pixels_map, (0, 2, 1))[..., np.newaxis],
+ self.gm_map_shape,
+ ),
+ dtype=np.uint32,
+ )
+ else:
+ raise ValueError(
+ f"What in the world is this shape? {bad_pixels_map.shape}"
+ )
+ else:
+ self.bad_pixel_map_gpu |= cupy.asarray(
+ np.transpose(bad_pixels_map, (2, 1, 0, 3)), dtype=np.uint32
+ )
+
+ if override_flags_to_use is not None:
+ self.override_bad_pixel_flags_to_use(override_flags_to_use)
+
+ def override_bad_pixel_flags_to_use(self, override_value):
+ self.bad_pixel_map_gpu &= cupy.uint32(override_value)
+
+ def set_bad_pixel_mask_value(self, mask_value):
+ self.bad_pixel_mask_value = cupy.float32(mask_value)
+
+ def flush_buffers(self):
+ self.offset_map_gpu.fill(0)
+ self.rel_gain_pc_map_gpu.fill(1)
+ self.md_additional_offset_gpu.fill(0)
+ self.rel_gain_xray_map_gpu.fill(1)
+ self.bad_pixel_map_gpu.fill(0)
+
+ # TODO: baseline shift
+
+ def correct(self, flags):
+ if flags & CorrectionFlags.BLSHIFT:
+ raise NotImplementedError("Baseline shift not implemented yet")
+ if self.gain_mode is not AgipdGainMode.ADAPTIVE_GAIN and (
+ flags & CorrectionFlags.THRESHOLD
+ ):
+ raise ValueError("Cannot do gain thresholding in fixed gain mode")
+
+ self.correction_kernel(
+ self.full_grid,
+ self.full_block,
+ (
+ self.input_data_gpu,
+ self.cell_table_gpu,
+ cupy.uint8(flags),
+ self.default_gain,
+ self.gain_thresholds_gpu,
+ self.offset_map_gpu,
+ self.rel_gain_pc_map_gpu,
+ self.md_additional_offset_gpu,
+ self.rel_gain_xray_map_gpu,
+ self.g_gain_value,
+ self.bad_pixel_map_gpu,
+ self.bad_pixel_mask_value,
+ self.gain_map_gpu,
+ self.processed_data_gpu,
+ ),
+ )
+
+ def _init_kernels(self):
+ kernel_source = self._kernel_template.render(
+ {
+ "pixels_x": self.pixels_x,
+ "pixels_y": self.pixels_y,
+ "data_memory_cells": self.memory_cells,
+ "constant_memory_cells": self.constant_memory_cells,
+ "input_data_dtype": utils.np_dtype_to_c_type(self.input_data_dtype),
+ "output_data_dtype": utils.np_dtype_to_c_type(self.output_data_dtype),
+ "corr_enum": utils.enum_to_c_template(CorrectionFlags),
+ }
+ )
+ self.source_module = cupy.RawModule(code=kernel_source)
+ self.correction_kernel = self.source_module.get_function("correct")
+
+
+class AgipdCalcatFriend(calcat_utils.BaseCalcatFriend):
+ _constant_enum_class = AgipdConstants
+
+ def __init__(self, device, *args, **kwargs):
+ super().__init__(device, *args, **kwargs)
+ self._constants_need_conditions = {
+ AgipdConstants.ThresholdsDark: self.dark_condition,
+ AgipdConstants.Offset: self.dark_condition,
+ AgipdConstants.SlopesPC: self.dark_condition,
+ AgipdConstants.SlopesFF: self.illuminated_condition,
+ AgipdConstants.BadPixelsDark: self.dark_condition,
+ AgipdConstants.BadPixelsPC: self.dark_condition,
+ AgipdConstants.BadPixelsFF: self.illuminated_condition,
+ }
+
+ @staticmethod
+ def add_schema(
+ schema,
+ managed_keys,
+ param_prefix="constantParameters",
+ status_prefix="foundConstants",
+ ):
+ super(AgipdCalcatFriend, AgipdCalcatFriend).add_schema(
+ schema, managed_keys, "AGIPD-Type", param_prefix, status_prefix
+ )
+ schema.setDefaultValue(f"{param_prefix}.memoryCells", 352)
+
+ (
+ DOUBLE_ELEMENT(schema)
+ .key(f"{param_prefix}.acquisitionRate")
+ .assignmentOptional()
+ .defaultValue(1.1)
+ .reconfigurable()
+ .commit(),
+ DOUBLE_ELEMENT(schema)
+ .key(f"{param_prefix}.gainSetting")
+ .assignmentOptional()
+ .defaultValue(0)
+ .reconfigurable()
+ .commit(),
+ DOUBLE_ELEMENT(schema)
+ .key(f"{param_prefix}.photonEnergy")
+ .assignmentOptional()
+ .defaultValue(9.2)
+ .reconfigurable()
+ .commit(),
+ DOUBLE_ELEMENT(schema)
+ .key(f"{param_prefix}.gainMode")
+ .assignmentOptional()
+ .defaultValue(0)
+ .reconfigurable()
+ .commit(),
+ DOUBLE_ELEMENT(schema)
+ .key(f"{param_prefix}.integrationTime")
+ .assignmentOptional()
+ .defaultValue(12)
+ .reconfigurable()
+ .commit(),
+ )
+ managed_keys.add(f"{param_prefix}.acquisitionRate")
+ managed_keys.add(f"{param_prefix}.gainSetting")
+ managed_keys.add(f"{param_prefix}.photonEnergy")
+ managed_keys.add(f"{param_prefix}.integrationTime")
+
+ calcat_utils.add_status_schema_from_enum(schema, status_prefix, AgipdConstants)
+
+ def dark_condition(self):
+ res = calcat_utils.OperatingConditions()
+ res["Memory cells"] = self._get_param("memoryCells")
+ res["Sensor Bias Voltage"] = self._get_param("biasVoltage")
+ res["Pixels X"] = self._get_param("pixelsX")
+ res["Pixels Y"] = self._get_param("pixelsY")
+ res["Acquisition rate"] = self._get_param("acquisitionRate")
+ # TODO: make configurable whether or not to include gain setting?
+ integration_time = self._get_param("integrationTime")
+ if integration_time != 12:
+ res["Integration Time"] = integration_time
+ res["Gain Setting"] = self._get_param("gainSetting")
+ return res
+
+ def illuminated_condition(self):
+ res = calcat_utils.OperatingConditions()
+ res["Memory cells"] = self._get_param("memoryCells")
+ res["Sensor Bias Voltage"] = self._get_param("biasVoltage")
+ res["Pixels X"] = self._get_param("pixelsX")
+ res["Pixels Y"] = self._get_param("pixelsY")
+ res["Source Energy"] = self._get_param("photonEnergy")
+ res["Acquisition rate"] = self._get_param("acquisitionRate")
+ integration_time = self._get_param("integrationTime")
+ if integration_time != 12:
+ res["Integration Time"] = integration_time
+ res["Gain Setting"] = self._get_param("gainSetting")
+ return res
@KARABO_CLASSINFO("AgipdCorrection", deviceVersion)
@@ -192,8 +590,6 @@ class AgipdCorrection(BaseCorrection):
"g_gain_value": config.get("corrections.relGainXray.gGainValue"),
}
- self._shmem_buffer_gain_map = None
-
# configurability: overriding md_additional_offset
if config.get("corrections.relGainPc.overrideMdAdditionalOffset"):
self._override_md_additional_offset = config.get(
diff --git a/src/calng/DsscCorrection.py b/src/calng/DsscCorrection.py
index 8698cf68ecf1bc9b9c9aab457c32ef7a2588e370..8c586b3df556abf734badd0e80677e5c5ad1a20d 100644
--- a/src/calng/DsscCorrection.py
+++ b/src/calng/DsscCorrection.py
@@ -1,12 +1,150 @@
+import enum
+
+import cupy
import numpy as np
-from karabo.bound import KARABO_CLASSINFO, VECTOR_STRING_ELEMENT
+from karabo.bound import DOUBLE_ELEMENT, KARABO_CLASSINFO, VECTOR_STRING_ELEMENT
from karabo.common.states import State
-from . import utils
+from . import base_gpu, calcat_utils, utils
from ._version import version as deviceVersion
from .base_correction import BaseCorrection, add_correction_step_schema
-from .calcat_utils import DsscCalcatFriend, DsscConstants
-from .dssc_gpu import CorrectionFlags, DsscGpuRunner
+
+
+class CorrectionFlags(enum.IntFlag):
+ NONE = 0
+ OFFSET = 1
+
+
+class DsscConstants(enum.Enum):
+ Offset = enum.auto()
+
+
+class DsscGpuRunner(base_gpu.BaseGpuRunner):
+ _kernel_source_filename = "dssc_gpu.cu"
+ _corrected_axis_order = "cyx"
+
+ def __init__(
+ self,
+ pixels_x,
+ pixels_y,
+ memory_cells,
+ constant_memory_cells,
+ input_data_dtype=np.uint16,
+ output_data_dtype=np.float32,
+ ):
+ self.input_shape = (memory_cells, pixels_y, pixels_x)
+ self.processed_shape = self.input_shape
+ super().__init__(
+ pixels_x,
+ pixels_y,
+ memory_cells,
+ constant_memory_cells,
+ input_data_dtype,
+ output_data_dtype,
+ )
+
+ self.map_shape = (self.constant_memory_cells, self.pixels_y, self.pixels_x)
+ self.offset_map_gpu = cupy.empty(self.map_shape, dtype=np.float32)
+
+ self._init_kernels()
+
+ self.offset_map_gpu = cupy.empty(self.map_shape, dtype=np.float32)
+
+ self.update_block_size((1, 1, 64))
+
+ def _get_raw_for_preview(self):
+ return self.input_data_gpu
+
+ def _get_corrected_for_preview(self):
+ return self.processed_data_gpu
+
+ def load_offset_map(self, offset_map):
+ # can have an extra dimension for some reason
+ if len(offset_map.shape) == 4: # old format (see offsetcorrection_dssc.py)?
+ offset_map = offset_map[..., 0]
+ # shape (now): x, y, memory cell
+ offset_map = np.transpose(offset_map).astype(np.float32)
+ self.offset_map_gpu.set(offset_map)
+
+ def correct(self, flags):
+ self.correction_kernel(
+ self.full_grid,
+ self.full_block,
+ (
+ self.input_data_gpu,
+ self.cell_table_gpu,
+ np.uint8(flags),
+ self.offset_map_gpu,
+ self.processed_data_gpu,
+ ),
+ )
+
+ def _init_kernels(self):
+ kernel_source = self._kernel_template.render(
+ {
+ "pixels_x": self.pixels_x,
+ "pixels_y": self.pixels_y,
+ "data_memory_cells": self.memory_cells,
+ "constant_memory_cells": self.constant_memory_cells,
+ "input_data_dtype": utils.np_dtype_to_c_type(self.input_data_dtype),
+ "output_data_dtype": utils.np_dtype_to_c_type(self.output_data_dtype),
+ "corr_enum": utils.enum_to_c_template(CorrectionFlags),
+ }
+ )
+ self.source_module = cupy.RawModule(code=kernel_source)
+ self.correction_kernel = self.source_module.get_function("correct")
+
+
+class DsscCalcatFriend(calcat_utils.BaseCalcatFriend):
+ _constant_enum_class = DsscConstants
+
+ def __init__(self, device, *args, **kwargs):
+ super().__init__(device, *args, **kwargs)
+ self._constants_need_conditions = {
+ DsscConstants.Offset: self.dark_condition,
+ }
+
+ @staticmethod
+ def add_schema(
+ schema,
+ managed_keys,
+ param_prefix="constantParameters",
+ status_prefix="foundConstants",
+ ):
+ super(DsscCalcatFriend, DsscCalcatFriend).add_schema(
+ schema, managed_keys, "DSSC-Type", param_prefix, status_prefix
+ )
+ schema.setDefaultValue(f"{param_prefix}.memoryCells", 400)
+ (
+ DOUBLE_ELEMENT(schema)
+ .key(f"{param_prefix}.pulseIdChecksum")
+ .assignmentOptional()
+ .defaultValue(2.8866323107820637e-36)
+ .commit(),
+ DOUBLE_ELEMENT(schema)
+ .key(f"{param_prefix}.acquisitionRate")
+ .assignmentOptional()
+ .defaultValue(4.5)
+ .commit(),
+ DOUBLE_ELEMENT(schema)
+ .key(f"{param_prefix}.encodedGain")
+ .assignmentOptional()
+ .defaultValue(67328)
+ .commit(),
+ )
+
+ calcat_utils.add_status_schema_from_enum(schema, status_prefix, DsscConstants)
+
+ def dark_condition(self):
+ res = calcat_utils.OperatingConditions()
+ res["Memory cells"] = self._get_param("memoryCells")
+ res["Sensor Bias Voltage"] = self._get_param("biasVoltage")
+ res["Pixels X"] = self._get_param("pixelsX")
+ res["Pixels Y"] = self._get_param("pixelsY")
+ # res["Pulse id checksum"] = self._get_param("pulseIdChecksum")
+ # res["Acquisition rate"] = self._get_param("acquisitionRate")
+ # res["Encoded gain"] = self._get_param("encodedGain")
+ return res
@KARABO_CLASSINFO("DsscCorrection", deviceVersion)
diff --git a/src/calng/agipd_gpu.py b/src/calng/agipd_gpu.py
deleted file mode 100644
index 98d26c547e49f4a6bf30d12b937b74d17601d517..0000000000000000000000000000000000000000
--- a/src/calng/agipd_gpu.py
+++ /dev/null
@@ -1,297 +0,0 @@
-import enum
-
-import cupy
-import numpy as np
-
-from . import base_gpu, utils
-
-
-class CorrectionFlags(enum.IntFlag):
- NONE = 0
- THRESHOLD = 1
- OFFSET = 2
- BLSHIFT = 4
- REL_GAIN_PC = 8
- REL_GAIN_XRAY = 16
- BPMASK = 32
-
-
-# from pycalibration's enum.py
-class AgipdGainMode(enum.IntEnum):
- ADAPTIVE_GAIN = 0
- FIXED_HIGH_GAIN = 1
- FIXED_MEDIUM_GAIN = 2
- FIXED_LOW_GAIN = 3
-
-
-class AgipdGpuRunner(base_gpu.BaseGpuRunner):
- _kernel_source_filename = "agipd_gpu.cu"
- _corrected_axis_order = "cxy"
-
- def __init__(
- self,
- pixels_x,
- pixels_y,
- memory_cells,
- constant_memory_cells,
- input_data_dtype=cupy.uint16,
- output_data_dtype=cupy.float32,
- bad_pixel_mask_value=cupy.nan,
- gain_mode=AgipdGainMode.ADAPTIVE_GAIN,
- g_gain_value=1,
- ):
- self.gain_mode = gain_mode
- if self.gain_mode is AgipdGainMode.ADAPTIVE_GAIN:
- self.default_gain = cupy.uint8(gain_mode)
- else:
- self.default_gain = cupy.uint8(gain_mode - 1)
- self.input_shape = (memory_cells, 2, pixels_x, pixels_y)
- self.processed_shape = (memory_cells, pixels_x, pixels_y)
- super().__init__(
- pixels_x,
- pixels_y,
- memory_cells,
- constant_memory_cells,
- input_data_dtype,
- output_data_dtype,
- )
- self.gain_map_gpu = cupy.empty(self.processed_shape, dtype=cupy.float32)
- self.preview_buffer_getters.extend(
- [self._get_raw_gain_for_preview, self._get_gain_map_for_preview]
- )
-
- self.map_shape = (self.constant_memory_cells, self.pixels_x, self.pixels_y)
- self.gm_map_shape = self.map_shape + (3,) # for gain-mapped constants
- self.threshold_map_shape = self.map_shape + (2,)
- # constants
- self.gain_thresholds_gpu = cupy.empty(
- self.threshold_map_shape, dtype=cupy.float32
- )
- self.offset_map_gpu = cupy.zeros(self.gm_map_shape, dtype=cupy.float32)
- self.rel_gain_pc_map_gpu = cupy.ones(self.gm_map_shape, dtype=cupy.float32)
- # not gm_map_shape because it only applies to medium gain pixels
- self.md_additional_offset_gpu = cupy.zeros(self.map_shape, dtype=cupy.float32)
- self.rel_gain_xray_map_gpu = cupy.ones(self.map_shape, dtype=cupy.float32)
- self.bad_pixel_map_gpu = cupy.zeros(self.gm_map_shape, dtype=cupy.uint32)
- self.set_bad_pixel_mask_value(bad_pixel_mask_value)
- self.set_g_gain_value(g_gain_value)
-
- self.update_block_size((1, 1, 64))
-
- def _get_raw_for_preview(self):
- return self.input_data_gpu[:, 0]
-
- def _get_corrected_for_preview(self):
- return self.processed_data_gpu
-
- # special to AGIPD
- def _get_raw_gain_for_preview(self):
- return self.input_data_gpu[:, 1]
-
- def _get_gain_map_for_preview(self):
- return self.gain_map_gpu
-
- def load_thresholds(self, threshold_map):
- # shape: y, x, memory cell, thresholds and gain values
- # note: the gain values are something like means used to derive thresholds
- self.gain_thresholds_gpu.set(
- np.transpose(threshold_map[..., :2], (2, 1, 0, 3)).astype(np.float32)
- )
-
- def load_offset_map(self, offset_map):
- # shape: y, x, memory cell, gain stage
- self.offset_map_gpu.set(
- np.transpose(offset_map, (2, 1, 0, 3)).astype(np.float32)
- )
-
- def load_rel_gain_pc_map(self, slopes_pc_map, override_md_additional_offset=None):
- # pc has funny shape (11, 352, 128, 512) from file
- # this is (fi, memory cell, y, x)
- slopes_pc_map = slopes_pc_map.astype(np.float32)
- # the following may contain NaNs, though...
- hg_slope = slopes_pc_map[0]
- hg_intercept = slopes_pc_map[1]
- mg_slope = slopes_pc_map[3]
- mg_intercept = slopes_pc_map[4]
- # TODO: remove sanitization (should happen in constant preparation notebook)
- # from agipdlib.py: replace NaN with median (per memory cell)
- # note: suffixes in agipdlib are "_m" and "_l", should probably be "_I"
- for naughty_array in (hg_slope, hg_intercept, mg_slope, mg_intercept):
- medians = np.nanmedian(naughty_array, axis=(1, 2))
- nan_bool = np.isnan(naughty_array)
- nan_cell, _, _ = np.where(nan_bool)
- naughty_array[nan_bool] = medians[nan_cell]
-
- too_low_bool = naughty_array < 0.8 * medians[:, np.newaxis, np.newaxis]
- too_low_cell, _, _ = np.where(too_low_bool)
- naughty_array[too_low_bool] = medians[too_low_cell]
-
- too_high_bool = naughty_array > 1.2 * medians[:, np.newaxis, np.newaxis]
- too_high_cell, _, _ = np.where(too_high_bool)
- naughty_array[too_high_bool] = medians[too_high_cell]
-
- frac_hg_mg = hg_slope / mg_slope
- rel_gain_map = np.ones(
- (3, self.constant_memory_cells, self.pixels_y, self.pixels_x),
- dtype=np.float32,
- )
- rel_gain_map[1] = rel_gain_map[0] * frac_hg_mg
- rel_gain_map[2] = rel_gain_map[1] * 4.48
- self.rel_gain_pc_map_gpu.set(np.transpose(rel_gain_map, (1, 3, 2, 0)))
- if override_md_additional_offset is None:
- md_additional_offset = (hg_intercept - mg_intercept * frac_hg_mg).astype(
- np.float32
- )
- self.md_additional_offset_gpu.set(
- np.transpose(md_additional_offset, (0, 2, 1))
- )
- else:
- self.override_md_additional_offset(override_md_additional_offset)
-
- def override_md_additional_offset(self, override_value):
- self.md_additional_offset_gpu.fill(override_value)
-
- def load_rel_gain_ff_map(self, slopes_ff_map):
- # constant shape: y, x, memory cell
- if slopes_ff_map.shape[2] == 2:
- # TODO: remove support for old format
- # old format, is per pixel only (shape is y, x, 2)
- # note: we should not support this in online
- slopes_ff_map = np.broadcast_to(
- slopes_ff_map[..., 0][..., np.newaxis],
- (self.pixels_y, self.pixels_x, self.constant_memory_cells),
- )
- self.rel_gain_xray_map_gpu.set(np.transpose(slopes_ff_map).astype(np.float32))
-
- def set_g_gain_value(self, override_value):
- self.g_gain_value = cupy.float32(override_value)
-
- def load_bad_pixels_map(self, bad_pixels_map, override_flags_to_use=None):
- print(f"Loading bad pixels with shape: {bad_pixels_map.shape}")
- # will simply OR with already loaded, does not take into account which ones
- # TODO: inquire what "mask for double size pixels" means
- if len(bad_pixels_map.shape) == 3:
- if bad_pixels_map.shape == (
- self.pixels_y,
- self.pixels_x,
- self.constant_memory_cells,
- ):
- # BadPixelsFF is not per gain stage - broadcasting along gain dimension
- self.bad_pixel_map_gpu |= cupy.asarray(
- np.broadcast_to(
- np.transpose(bad_pixels_map)[..., np.newaxis],
- self.gm_map_shape,
- ),
- dtype=np.uint32,
- )
- elif bad_pixels_map.shape == (
- self.constant_memory_cells,
- self.pixels_y,
- self.pixels_x,
- ):
- # oh, can also be old bad pixels pc?
- self.bad_pixel_map_gpu |= cupy.asarray(
- np.broadcast_to(
- np.transpose(bad_pixels_map, (0, 2, 1))[..., np.newaxis],
- self.gm_map_shape,
- ),
- dtype=np.uint32,
- )
- else:
- raise ValueError(
- f"What in the world is this shape? {bad_pixels_map.shape}"
- )
- else:
- self.bad_pixel_map_gpu |= cupy.asarray(
- np.transpose(bad_pixels_map, (2, 1, 0, 3)), dtype=np.uint32
- )
-
- if override_flags_to_use is not None:
- self.override_bad_pixel_flags_to_use(override_flags_to_use)
-
- def override_bad_pixel_flags_to_use(self, override_value):
- self.bad_pixel_map_gpu &= cupy.uint32(override_value)
-
- def set_bad_pixel_mask_value(self, mask_value):
- self.bad_pixel_mask_value = cupy.float32(mask_value)
-
- def flush_buffers(self):
- self.offset_map_gpu.fill(0)
- self.rel_gain_pc_map_gpu.fill(1)
- self.md_additional_offset_gpu.fill(0)
- self.rel_gain_xray_map_gpu.fill(1)
- self.bad_pixel_map_gpu.fill(0)
-
- # TODO: baseline shift
-
- def correct(self, flags):
- if flags & CorrectionFlags.BLSHIFT:
- raise NotImplementedError("Baseline shift not implemented yet")
- if self.gain_mode is not AgipdGainMode.ADAPTIVE_GAIN and (
- flags & CorrectionFlags.THRESHOLD
- ):
- raise ValueError("Cannot do gain thresholding in fixed gain mode")
-
- self.correction_kernel(
- self.full_grid,
- self.full_block,
- (
- self.input_data_gpu,
- self.cell_table_gpu,
- cupy.uint8(flags),
- self.default_gain,
- self.gain_thresholds_gpu,
- self.offset_map_gpu,
- self.rel_gain_pc_map_gpu,
- self.md_additional_offset_gpu,
- self.rel_gain_xray_map_gpu,
- self.g_gain_value,
- self.bad_pixel_map_gpu,
- self.bad_pixel_mask_value,
- self.gain_map_gpu,
- self.processed_data_gpu,
- ),
- )
-
- def _init_kernels(self):
- kernel_source = self._kernel_template.render(
- {
- "pixels_x": self.pixels_x,
- "pixels_y": self.pixels_y,
- "data_memory_cells": self.memory_cells,
- "constant_memory_cells": self.constant_memory_cells,
- "input_data_dtype": utils.np_dtype_to_c_type(self.input_data_dtype),
- "output_data_dtype": utils.np_dtype_to_c_type(self.output_data_dtype),
- "corr_enum": utils.enum_to_c_template(CorrectionFlags),
- }
- )
- self.source_module = cupy.RawModule(code=kernel_source)
- self.correction_kernel = self.source_module.get_function("correct")
-
-
-class BadPixelValues(enum.IntFlag):
- """The European XFEL Bad Pixel Encoding
-
- Straight from pycalibration's enum.py"""
-
- OFFSET_OUT_OF_THRESHOLD = 2 ** 0
- NOISE_OUT_OF_THRESHOLD = 2 ** 1
- OFFSET_NOISE_EVAL_ERROR = 2 ** 2
- NO_DARK_DATA = 2 ** 3
- CI_GAIN_OUT_OF_THRESHOLD = 2 ** 4
- CI_LINEAR_DEVIATION = 2 ** 5
- CI_EVAL_ERROR = 2 ** 6
- FF_GAIN_EVAL_ERROR = 2 ** 7
- FF_GAIN_DEVIATION = 2 ** 8
- FF_NO_ENTRIES = 2 ** 9
- CI2_EVAL_ERROR = 2 ** 10
- VALUE_IS_NAN = 2 ** 11
- VALUE_OUT_OF_RANGE = 2 ** 12
- GAIN_THRESHOLDING_ERROR = 2 ** 13
- DATA_STD_IS_ZERO = 2 ** 14
- ASIC_STD_BELOW_NOISE = 2 ** 15
- INTERPOLATED = 2 ** 16
- NOISY_ADC = 2 ** 17
- OVERSCAN = 2 ** 18
- NON_SENSITIVE = 2 ** 19
- NON_LIN_RESPONSE_REGION = 2 ** 20
diff --git a/src/calng/calcat_utils.py b/src/calng/calcat_utils.py
index 49403228e64b6279de8ade0c75c391709cc9d4c3..0b9099cc02b671543d665d81b23c73b57323af6b 100644
--- a/src/calng/calcat_utils.py
+++ b/src/calng/calcat_utils.py
@@ -1,10 +1,8 @@
import copy
-import enum
import functools
import json
import pathlib
import threading
-import typing
import calibration_client
import h5py
@@ -51,7 +49,7 @@ class CalibrationClientConfigError(Exception):
pass
-def _add_status_schema_from_enum(schema, prefix, enum_class):
+def add_status_schema_from_enum(schema, prefix, enum_class):
for constant in enum_class:
constant_node = f"{prefix}.{constant.name}"
(
@@ -532,164 +530,3 @@ class BaseCalcatFriend:
if warning is not None:
self.device.log_status_warn(warning)
raise to_raise
-
-
-class AgipdConstants(enum.Enum):
- SlopesFF = enum.auto()
- ThresholdsDark = enum.auto()
- Offset = enum.auto()
- SlopesPC = enum.auto()
- BadPixelsDark = enum.auto()
- BadPixelsPC = enum.auto()
- BadPixelsFF = enum.auto()
-
-
-class AgipdCalcatFriend(BaseCalcatFriend):
- _constant_enum_class = AgipdConstants
-
- def __init__(self, device, *args, **kwargs):
- super().__init__(device, *args, **kwargs)
- self._constants_need_conditions = {
- AgipdConstants.ThresholdsDark: self.dark_condition,
- AgipdConstants.Offset: self.dark_condition,
- AgipdConstants.SlopesPC: self.dark_condition,
- AgipdConstants.SlopesFF: self.illuminated_condition,
- AgipdConstants.BadPixelsDark: self.dark_condition,
- AgipdConstants.BadPixelsPC: self.dark_condition,
- AgipdConstants.BadPixelsFF: self.illuminated_condition,
- }
-
- @staticmethod
- def add_schema(
- schema,
- managed_keys,
- param_prefix="constantParameters",
- status_prefix="foundConstants",
- ):
- super(AgipdCalcatFriend, AgipdCalcatFriend).add_schema(
- schema, managed_keys, "AGIPD-Type", param_prefix, status_prefix
- )
- schema.setDefaultValue(f"{param_prefix}.memoryCells", 352)
-
- (
- DOUBLE_ELEMENT(schema)
- .key(f"{param_prefix}.acquisitionRate")
- .assignmentOptional()
- .defaultValue(1.1)
- .reconfigurable()
- .commit(),
- DOUBLE_ELEMENT(schema)
- .key(f"{param_prefix}.gainSetting")
- .assignmentOptional()
- .defaultValue(0)
- .reconfigurable()
- .commit(),
- DOUBLE_ELEMENT(schema)
- .key(f"{param_prefix}.photonEnergy")
- .assignmentOptional()
- .defaultValue(9.2)
- .reconfigurable()
- .commit(),
- DOUBLE_ELEMENT(schema)
- .key(f"{param_prefix}.gainMode")
- .assignmentOptional()
- .defaultValue(0)
- .reconfigurable()
- .commit(),
- DOUBLE_ELEMENT(schema)
- .key(f"{param_prefix}.integrationTime")
- .assignmentOptional()
- .defaultValue(12)
- .reconfigurable()
- .commit(),
- )
- managed_keys.add(f"{param_prefix}.acquisitionRate")
- managed_keys.add(f"{param_prefix}.gainSetting")
- managed_keys.add(f"{param_prefix}.photonEnergy")
- managed_keys.add(f"{param_prefix}.integrationTime")
-
- _add_status_schema_from_enum(schema, status_prefix, AgipdConstants)
-
- def dark_condition(self):
- res = OperatingConditions()
- res["Memory cells"] = self._get_param("memoryCells")
- res["Sensor Bias Voltage"] = self._get_param("biasVoltage")
- res["Pixels X"] = self._get_param("pixelsX")
- res["Pixels Y"] = self._get_param("pixelsY")
- res["Acquisition rate"] = self._get_param("acquisitionRate")
- # TODO: make configurable whether or not to include gain setting?
- integration_time = self._get_param("integrationTime")
- if integration_time != 12:
- res["Integration Time"] = integration_time
- res["Gain Setting"] = self._get_param("gainSetting")
- return res
-
- def illuminated_condition(self):
- res = OperatingConditions()
- res["Memory cells"] = self._get_param("memoryCells")
- res["Sensor Bias Voltage"] = self._get_param("biasVoltage")
- res["Pixels X"] = self._get_param("pixelsX")
- res["Pixels Y"] = self._get_param("pixelsY")
- res["Source Energy"] = self._get_param("photonEnergy")
- res["Acquisition rate"] = self._get_param("acquisitionRate")
- integration_time = self._get_param("integrationTime")
- if integration_time != 12:
- res["Integration Time"] = integration_time
- res["Gain Setting"] = self._get_param("gainSetting")
- return res
-
-
-class DsscConstants(enum.Enum):
- Offset = enum.auto()
-
-
-class DsscCalcatFriend(BaseCalcatFriend):
- _constant_enum_class = DsscConstants
-
- def __init__(self, device, *args, **kwargs):
- super().__init__(device, *args, **kwargs)
- self._constants_need_conditions = {
- DsscConstants.Offset: self.dark_condition,
- }
-
- @staticmethod
- def add_schema(
- schema,
- managed_keys,
- param_prefix="constantParameters",
- status_prefix="foundConstants",
- ):
- super(DsscCalcatFriend, DsscCalcatFriend).add_schema(
- schema, managed_keys, "DSSC-Type", param_prefix, status_prefix
- )
- schema.setDefaultValue(f"{param_prefix}.memoryCells", 400)
- (
- DOUBLE_ELEMENT(schema)
- .key(f"{param_prefix}.pulseIdChecksum")
- .assignmentOptional()
- .defaultValue(2.8866323107820637e-36)
- .commit(),
- DOUBLE_ELEMENT(schema)
- .key(f"{param_prefix}.acquisitionRate")
- .assignmentOptional()
- .defaultValue(4.5)
- .commit(),
- DOUBLE_ELEMENT(schema)
- .key(f"{param_prefix}.encodedGain")
- .assignmentOptional()
- .defaultValue(67328)
- .commit(),
- )
-
- _add_status_schema_from_enum(schema, status_prefix, DsscConstants)
-
- def dark_condition(self):
- res = OperatingConditions()
- res["Memory cells"] = self._get_param("memoryCells")
- res["Sensor Bias Voltage"] = self._get_param("biasVoltage")
- res["Pixels X"] = self._get_param("pixelsX")
- res["Pixels Y"] = self._get_param("pixelsY")
- # res["Pulse id checksum"] = self._get_param("pulseIdChecksum")
- # res["Acquisition rate"] = self._get_param("acquisitionRate")
- # res["Encoded gain"] = self._get_param("encodedGain")
- return res
diff --git a/src/calng/dssc_gpu.py b/src/calng/dssc_gpu.py
deleted file mode 100644
index b9dc3b091c6e14374adddbb08a9ad4d04340e9f3..0000000000000000000000000000000000000000
--- a/src/calng/dssc_gpu.py
+++ /dev/null
@@ -1,87 +0,0 @@
-import enum
-
-import cupy
-import numpy as np
-
-from . import base_gpu, utils
-
-
-class CorrectionFlags(enum.IntFlag):
- NONE = 0
- OFFSET = 1
-
-
-class DsscGpuRunner(base_gpu.BaseGpuRunner):
- _kernel_source_filename = "dssc_gpu.cu"
- _corrected_axis_order = "cyx"
-
- def __init__(
- self,
- pixels_x,
- pixels_y,
- memory_cells,
- constant_memory_cells,
- input_data_dtype=np.uint16,
- output_data_dtype=np.float32,
- ):
- self.input_shape = (memory_cells, pixels_y, pixels_x)
- self.processed_shape = self.input_shape
- super().__init__(
- pixels_x,
- pixels_y,
- memory_cells,
- constant_memory_cells,
- input_data_dtype,
- output_data_dtype,
- )
-
- self.map_shape = (self.constant_memory_cells, self.pixels_y, self.pixels_x)
- self.offset_map_gpu = cupy.empty(self.map_shape, dtype=np.float32)
-
- self._init_kernels()
-
- self.offset_map_gpu = cupy.empty(self.map_shape, dtype=np.float32)
-
- self.update_block_size((1, 1, 64))
-
- def _get_raw_for_preview(self):
- return self.input_data_gpu
-
- def _get_corrected_for_preview(self):
- return self.processed_data_gpu
-
- def load_offset_map(self, offset_map):
- # can have an extra dimension for some reason
- if len(offset_map.shape) == 4: # old format (see offsetcorrection_dssc.py)?
- offset_map = offset_map[..., 0]
- # shape (now): x, y, memory cell
- offset_map = np.transpose(offset_map).astype(np.float32)
- self.offset_map_gpu.set(offset_map)
-
- def correct(self, flags):
- self.correction_kernel(
- self.full_grid,
- self.full_block,
- (
- self.input_data_gpu,
- self.cell_table_gpu,
- np.uint8(flags),
- self.offset_map_gpu,
- self.processed_data_gpu,
- ),
- )
-
- def _init_kernels(self):
- kernel_source = self._kernel_template.render(
- {
- "pixels_x": self.pixels_x,
- "pixels_y": self.pixels_y,
- "data_memory_cells": self.memory_cells,
- "constant_memory_cells": self.constant_memory_cells,
- "input_data_dtype": utils.np_dtype_to_c_type(self.input_data_dtype),
- "output_data_dtype": utils.np_dtype_to_c_type(self.output_data_dtype),
- "corr_enum": utils.enum_to_c_template(CorrectionFlags),
- }
- )
- self.source_module = cupy.RawModule(code=kernel_source)
- self.correction_kernel = self.source_module.get_function("correct")