diff --git a/src/calng/AgipdCorrection.py b/src/calng/AgipdCorrection.py
index 136ea595e587e184ef8498aa3027e3ef9c466ff6..b7d4d629909da9adc2c337301edcda627c02881a 100644
--- a/src/calng/AgipdCorrection.py
+++ b/src/calng/AgipdCorrection.py
@@ -1,13 +1,21 @@
import timeit
import numpy as np
-from karabo.bound import BOOL_ELEMENT, KARABO_CLASSINFO
+from karabo.bound import (
+ BOOL_ELEMENT,
+ KARABO_CLASSINFO,
+ STRING_ELEMENT,
+ TABLE_ELEMENT,
+ UINT32_ELEMENT,
+ Hash,
+ Schema,
+)
from karabo.common.states import State
from . import utils
from ._version import version as deviceVersion
+from .agipd_gpu import AgipdGpuRunner, BadPixelValues, CorrectionFlags
from .base_correction import BaseCorrection
-from .agipd_gpu import AgipdGpuRunner, CorrectionFlags
@KARABO_CLASSINFO("AgipdCorrection", deviceVersion)
@@ -80,6 +88,31 @@ class AgipdCorrection(BaseCorrection):
.commit(),
)
+ # TODO: hook this up to actual correction done
+ bad_pixel_selection_schema = Schema()
+ (
+ STRING_ELEMENT(bad_pixel_selection_schema).key("name").readOnly().commit(),
+ UINT32_ELEMENT(bad_pixel_selection_schema).key("value").readOnly().commit(),
+ BOOL_ELEMENT(bad_pixel_selection_schema)
+ .key("apply")
+ .assignmentOptional()
+ .defaultValue(True)
+ .reconfigurable()
+ .commit(),
+ TABLE_ELEMENT(expected)
+ .key("corrections.badPixelFlagsApplied")
+ .setNodeSchema(bad_pixel_selection_schema)
+ .assignmentOptional()
+ .defaultValue(
+ [
+ Hash("name", field.name, "value", field.value, "apply", True)
+ for field in BadPixelValues
+ ]
+ )
+ .reconfigurable()
+ .commit(),
+ )
+
@property
def input_data_shape(self):
return (
@@ -128,7 +161,7 @@ class AgipdCorrection(BaseCorrection):
# TODO: what are these empty things for?
if not data.has("image"):
- self.log.INFO("Ignoring hash without image node")
+ # self.log.INFO("Ignoring hash without image node")
return
time_start = timeit.default_timer()
@@ -147,7 +180,6 @@ class AgipdCorrection(BaseCorrection):
self.set(
"status", f"Updating input shapes based on received {image_data.shape}"
)
- # TODO: truncate if > 800
self.set("dataFormat.memoryCells", image_data.shape[0])
with self._buffer_lock:
# TODO: pulse filter update after reimplementation
@@ -235,7 +267,6 @@ class AgipdCorrection(BaseCorrection):
)
def _load_constant_to_gpu(self, constant_name, constant_data):
- # TODO: also hook flushConstants or whatever it is called
if constant_name == "ThresholdsDark":
self.gpu_runner.load_thresholds(constant_data)
# TODO: encode correction / constant dependencies in a clever way
@@ -263,7 +294,11 @@ class AgipdCorrection(BaseCorrection):
self.set("corrections.preview.relGainXray", True)
elif "BadPixels" in constant_name:
# TODO: implement loading bad pixels
- ...
+ self.gpu_runner.load_badpixels_map(constant_data)
+ if not self.get("corrections.available.badpixels"):
+ self.set("corrections.available.badpixels", True)
+ self.set("corrections.enabled.badpixels", True)
+ self.set("corrections.preview.badpixels", True)
def _update_pulse_filter(self, filter_string):
"""Called whenever the pulse filter changes, typically followed by
diff --git a/src/calng/agipd_gpu.py b/src/calng/agipd_gpu.py
index 196aa30acd9d46ac1559e1e0af4711a63d408788..4ec912cb48fe1f65a91753e288d53df07b848ef2 100644
--- a/src/calng/agipd_gpu.py
+++ b/src/calng/agipd_gpu.py
@@ -28,6 +28,7 @@ class AgipdGpuRunner(base_gpu.BaseGpuRunner):
output_transpose=(1, 2, 0), # default: memorycells-fast
input_data_dtype=np.uint16,
output_data_dtype=np.float32,
+ badpixel_mask_value=np.float32(np.nan),
):
self.input_shape = (memory_cells, 2, pixels_x, pixels_y)
self.processed_shape = (memory_cells, pixels_x, pixels_y)
@@ -51,9 +52,11 @@ class AgipdGpuRunner(base_gpu.BaseGpuRunner):
)
self.offset_map_gpu = cupy.zeros(self.gm_map_shape, dtype=np.float32)
self.rel_gain_pc_map_gpu = cupy.ones(self.gm_map_shape, dtype=np.float32)
+ # not gm_map_shape because it only applies to medium gain pixels
self.md_additional_offset_gpu = cupy.zeros(self.map_shape, dtype=np.float32)
self.rel_gain_xray_map_gpu = cupy.ones(self.map_shape, dtype=np.float32)
- self.badpixel_map_gpu = cupy.zeros(self.map_shape, dtype=np.uint32)
+ self.badpixel_map_gpu = cupy.zeros(self.gm_map_shape, dtype=np.uint32)
+ self.badpixel_mask_value = badpixel_mask_value # TODO: make this configurable
self.update_block_size((1, 1, 64))
@@ -112,7 +115,7 @@ class AgipdGpuRunner(base_gpu.BaseGpuRunner):
)
rel_gain_map[1] = rel_gain_map[0] * frac_hg_mg
rel_gain_map[2] = rel_gain_map[1] * 4.48
- # TODO: enable overriding this based on user input
+ # TODO: enable overriding md_additional_offset_gpu based on user input
self.rel_gain_pc_map_gpu.set(np.transpose(rel_gain_map, (1, 3, 2, 0)))
self.md_additional_offset_gpu.set(np.transpose(md_additional_offset, (0, 2, 1)))
@@ -131,11 +134,34 @@ class AgipdGpuRunner(base_gpu.BaseGpuRunner):
# TODO: maybe clamp and replace NaNs like slopes_pc
self.rel_gain_xray_map_gpu.set(np.transpose(slopes_ff_map).astype(np.float32))
+ def load_badpixels_map(self, badpixels_map):
+ print(f"Loading bad pixels with shape: {badpixels_map.shape}")
+ # will simply OR with already loaded, does not take into account which ones
+ # TODO: simple loading
+ # TODO: allow configuring subset of bad pixels to care about
+ # TODO: allow configuring value for masked pixels
+ # TODO: inquire what "mask for double size pixels" means
+ if len(badpixels_map.shape) == 3:
+ # BadPixelsFF is not per gain stage...
+ # TODO: broadcast?
+ ...
+ else:
+ self.badpixel_map_gpu |= cupy.asarray(
+ np.transpose(badpixels_map, (2, 1, 0, 3)), dtype=np.uint32
+ )
+
+ 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.badpixel_map_gpu.fill(0)
+
+ # TODO: baseline shift
+
def correct(self, flags):
if flags & CorrectionFlags.BLSHIFT:
raise NotImplementedError("Baseline shift not implemented yet")
- if flags & CorrectionFlags.BPMASK:
- raise NotImplementedError("Bad pixel masking not implemented yet")
self.correction_kernel(
self.full_grid,
self.full_block,
@@ -149,6 +175,7 @@ class AgipdGpuRunner(base_gpu.BaseGpuRunner):
self.md_additional_offset_gpu,
self.rel_gain_xray_map_gpu,
self.badpixel_map_gpu,
+ self.badpixel_mask_value,
self.gain_map_gpu,
self.processed_data_gpu,
),
@@ -168,3 +195,31 @@ class AgipdGpuRunner(base_gpu.BaseGpuRunner):
)
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 = 0b000000000000000000001 # bit 1
+ NOISE_OUT_OF_THRESHOLD = 0b000000000000000000010 # bit 2
+ OFFSET_NOISE_EVAL_ERROR = 0b000000000000000000100 # bit 3
+ NO_DARK_DATA = 0b000000000000000001000 # bit 4
+ CI_GAIN_OF_OF_THRESHOLD = 0b000000000000000010000 # bit 5
+ CI_LINEAR_DEVIATION = 0b000000000000000100000 # bit 6
+ CI_EVAL_ERROR = 0b000000000000001000000 # bit 7
+ FF_GAIN_EVAL_ERROR = 0b000000000000010000000 # bit 8
+ FF_GAIN_DEVIATION = 0b000000000000100000000 # bit 9
+ FF_NO_ENTRIES = 0b000000000001000000000 # bit 10
+ CI2_EVAL_ERROR = 0b000000000010000000000 # bit 11
+ VALUE_IS_NAN = 0b000000000100000000000 # bit 12
+ VALUE_OUT_OF_RANGE = 0b000000001000000000000 # bit 13
+ GAIN_THRESHOLDING_ERROR = 0b000000010000000000000 # bit 14
+ DATA_STD_IS_ZERO = 0b000000100000000000000 # bit 15
+ ASIC_STD_BELOW_NOISE = 0b000001000000000000000 # bit 16
+ INTERPOLATED = 0b000010000000000000000 # bit 17
+ NOISY_ADC = 0b000100000000000000000 # bit 18
+ OVERSCAN = 0b001000000000000000000 # bit 19
+ NON_SENSITIVE = 0b010000000000000000000 # bit 20
+ NON_LIN_RESPONSE_REGION = 0b100000000000000000000 # bit 21
diff --git a/src/calng/agipd_gpu_kernels.cpp b/src/calng/agipd_gpu_kernels.cpp
index 898e0b87866457a7d57ba4fdc379b6c3f52fde17..357c82a8b8270f9dfee94af3957c7a8d0a40f469 100644
--- a/src/calng/agipd_gpu_kernels.cpp
+++ b/src/calng/agipd_gpu_kernels.cpp
@@ -1,5 +1,4 @@
#include <cuda_fp16.h>
-#include <math_constants.h>
{{corr_enum}}
@@ -20,6 +19,7 @@ extern "C" {
const float* md_additional_offset,
const float* rel_gain_xray_map,
const unsigned int* bad_pixel_map,
+ const float bad_pixel_mask_value,
unsigned char* gain_map,
{{output_data_dtype}}* output) {
const size_t X = {{pixels_x}};
@@ -107,7 +107,7 @@ extern "C" {
x * gm_map_stride_x;
if ((corr_flags & BPMASK) && bad_pixel_map[gm_map_index]) {
- corrected = CUDART_NAN_F;
+ corrected = bad_pixel_mask_value;
} else {
if (corr_flags & OFFSET) {
corrected -= offset_map[gm_map_index];
diff --git a/src/calng/base_correction.py b/src/calng/base_correction.py
index cfc3c3a6dde280853ba3b2116c2e544239c7541b..1f87b97322fa94c981f346fcaa33beeea5b28f71 100644
--- a/src/calng/base_correction.py
+++ b/src/calng/base_correction.py
@@ -463,8 +463,6 @@ class BaseCorrection(calibrationBase.CalibrationReceiverBaseDevice):
def requestConstant(self, name, mostRecent=False, tryRemote=True):
"""constantLoaded hook would have gotten called without naming constant, so here
we go. Ugly hooking it."""
- # TODO: clear from device, too
- # TODO: update correction capability flag
if name in self._cached_constants:
del self._cached_constants[name]
super().requestConstant(name, mostRecent, tryRemote)
@@ -477,6 +475,14 @@ class BaseCorrection(calibrationBase.CalibrationReceiverBaseDevice):
self._load_constant_to_gpu(name, constant)
self._update_correction_flags()
+ def flush_constants(self):
+ """Override from CalibrationReceiverBaseDevice to also flush GPU buffers"""
+ super().flush_constants()
+ for correction_step, _ in self._correction_slot_names:
+ self.set(f"corrections.available.{correction_step}", False)
+ self.gpu_runner.flush_buffers()
+ self._update_correction_flags()
+
def _write_output(self, data, old_metadata):
metadata = ChannelMetaData(
old_metadata.get("source"),
diff --git a/src/calng/base_gpu.py b/src/calng/base_gpu.py
index 5347fc2a1b730c10cd3fd5663e9c16a73342a603..fc716ee3e3d4d87dd4655a8d537493f86c5a4dbb 100644
--- a/src/calng/base_gpu.py
+++ b/src/calng/base_gpu.py
@@ -78,14 +78,18 @@ class BaseGpuRunner:
# functions to get data from respective buffers to cell, x, y shape for preview computation
# TODO: handle shape juggling programmatically, removing need for these two helpers
- def _preview_preprocess_raw():
+ def _preview_preprocess_raw(self):
"""Should return view of self.input_data_gpu with shape (cell, x, y)"""
raise NotImplementedError()
- def _preview_preprocess_corr():
+ def _preview_preprocess_corr(self):
"""Should return view of self.processed_data_gpu with shape (cell, x, y)"""
raise NotImplementedError()
+ def flush_buffers(self):
+ """Optional reset GPU buffers (implement in appropriate subclasses"""
+ pass
+
def correct(self, flags):
"""Correct (already loaded) image data according to flags