Draft: add Jungfrau correction device

Following https://git.xfel.eu/karaboDevices/calng/-/merge_requests/3, leaving as WIP pointed against add-agipd-device for now. This MR would add a Jungfrau correction device supporting dynamic / fixed gain mode and single cell / burst mode. Initially, with same type of GPU kernel runner as DSSC and AGIPD devices (though I believe a well-implemented CPU kernel runner will be plenty here; something to try later).

src/calng/kernels/jungfrau_gpu.cu

+#include <cuda_fp16.h>
+
+{{corr_enum}}
+
+extern "C" {
+	__global__ void correct(const {{input_data_dtype}}* data, // shape: memory cell, y, x
+	                        const unsigned char* gain_stage, // same shape
+	                        const unsigned char* cell_table,
+	                        const unsigned char corr_flags,
+	                        const float* offset_map,
+	                        const float* rel_gain_map,
+	                        const unsigned int* bad_pixel_map,
+	                        const float bad_pixel_mask_value,
+	                        {{output_data_dtype}}* output) {
+		const size_t X = {{pixels_x}};
+		const size_t Y = {{pixels_y}};
+		const size_t memory_cells = {{data_memory_cells}};
+		const size_t map_memory_cells = {{constant_memory_cells}};
+
+		const size_t memory_cell = blockIdx.x * blockDim.x + threadIdx.x;
+		const size_t y = blockIdx.y * blockDim.y + threadIdx.y;
+		const size_t x = blockIdx.z * blockDim.z + threadIdx.z;
+
+		if (memory_cell >= memory_cells || y >= Y || x >= X) {
+			return;
+		}
+
+		const size_t data_stride_x = 1;
+		const size_t data_stride_y = X * data_stride_x;
+		const size_t data_stride_cell = Y * data_stride_y;
+		const size_t data_index = memory_cell * data_stride_cell +
+			y * data_stride_y +
+			x * data_stride_x;
+		float res = (float)data[data_index];
+
+		// gain mapped constant shape: cell, y, x, gain_level (dim size 3)
+		// note: in fixed gain mode, constant still provides data for three stages
+		const size_t map_stride_gain = 1;
+		const size_t map_stride_x = 3 * map_stride_gain;
+		const size_t map_stride_y = X * map_stride_x;
+		const size_t map_stride_cell = Y * map_stride_y;
+
+		// TODO: warn user about cell_table value of 255 in either mode
+		// note: cell table may contain 255 if data didn't arrive
+		{% if burst_mode %}
+		// burst mode: "cell 255" will get copied
+		// TODO: consider masking "cell 255"
+		const size_t map_cell = cell_table[memory_cell];
+		{% else %}
+		// single cell: "cell 255" will get "corrected"
+		const size_t map_cell = 0;
+		{% endif %}
+
+		if (map_cell < map_memory_cells) {
+			unsigned char gain = gain_stage[data_index];
+			if (gain == 2) {
+				// gain should be read as 0, 1, or 3; value of 2 indicates issue
+				res = bad_pixel_mask_value;
+			} else {
+				if (gain == 3) {
+					gain = 2;
+				}
+				const size_t map_index = map_cell * map_stride_cell +
+					y * map_stride_y +
+					x * map_stride_x +
+					gain * map_stride_gain;
+				if ((corr_flags & BPMASK) && bad_pixel_map[map_index]) {
+					res = bad_pixel_mask_value;
+				} else {
+					if (corr_flags & OFFSET) {
+						res -= offset_map[map_index];
+					}
+					if (corr_flags & REL_GAIN) {
+						res /= rel_gain_map[map_index];
+					}
+				}
+			}
+		}
+
+		{% if output_data_dtype == "half" %}
+		output[data_index] = __float2half(res);
+		{% else %}
+		output[data_index] = ({{output_data_dtype}})res;
+		{% endif %}
+	}
+}