Skip to content
Snippets Groups Projects
Commit c892b65c authored by David Hammer's avatar David Hammer
Browse files

Reuse correction flag enum definition

parent c714edee
No related branches found
No related tags found
2 merge requests!12Snapshot: field test deployed version as of end of run 202201,!3Base correction device, CalCat interaction, DSSC and AGIPD devices
...@@ -145,6 +145,7 @@ class AgipdGpuRunner(base_gpu.BaseGpuRunner): ...@@ -145,6 +145,7 @@ class AgipdGpuRunner(base_gpu.BaseGpuRunner):
"constant_memory_cells": self.constant_memory_cells, "constant_memory_cells": self.constant_memory_cells,
"input_data_dtype": utils.np_dtype_to_c_type(self.input_data_dtype), "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), "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.source_module = cupy.RawModule(code=kernel_source)
......
#include <cuda_fp16.h> #include <cuda_fp16.h>
#include <math_constants.h> #include <math_constants.h>
const unsigned char CORR_THRESHOLD = 1; {{corr_enum}}
const unsigned char CORR_OFFSET = 2;
const unsigned char CORR_BLSHIFT = 4;
const unsigned char CORR_REL_GAIN_PC = 8;
const unsigned char CORR_REL_GAIN_XRAY = 16;
const unsigned char CORR_BPMASK = 32;
extern "C" { extern "C" {
/* /*
...@@ -82,7 +77,7 @@ extern "C" { ...@@ -82,7 +77,7 @@ extern "C" {
if (map_cell < map_cells) { if (map_cell < map_cells) {
unsigned char gain = 0; unsigned char gain = 0;
if (corr_flags & CORR_THRESHOLD) { if (corr_flags & THRESHOLD) {
const float threshold_0 = threshold_map[0 * threshold_map_stride_threshold + const float threshold_0 = threshold_map[0 * threshold_map_stride_threshold +
map_cell * threshold_map_stride_cell + map_cell * threshold_map_stride_cell +
y * threshold_map_stride_y + y * threshold_map_stride_y +
...@@ -110,20 +105,20 @@ extern "C" { ...@@ -110,20 +105,20 @@ extern "C" {
y * gm_map_stride_y + y * gm_map_stride_y +
x * gm_map_stride_x; x * gm_map_stride_x;
if ((corr_flags & CORR_BPMASK) && bad_pixel_map[gm_map_index]) { if ((corr_flags & BPMASK) && bad_pixel_map[gm_map_index]) {
corrected = CUDART_NAN_F; corrected = CUDART_NAN_F;
} else { } else {
if (corr_flags & CORR_OFFSET) { if (corr_flags & OFFSET) {
corrected -= offset_map[gm_map_index]; corrected -= offset_map[gm_map_index];
} }
// TODO: baseline shift // TODO: baseline shift
if (corr_flags & CORR_REL_GAIN_PC) { if (corr_flags & REL_GAIN_PC) {
corrected *= rel_gain_pc_map[gm_map_index]; corrected *= rel_gain_pc_map[gm_map_index];
if (gain == 1) { if (gain == 1) {
corrected += md_additional_offset[map_index]; corrected += md_additional_offset[map_index];
} }
} }
if (corr_flags & CORR_REL_GAIN_XRAY) { if (corr_flags & REL_GAIN_XRAY) {
// TODO // TODO
//corrected *= rel_gain_xray_map[map_index]; //corrected *= rel_gain_xray_map[map_index];
} }
......
...@@ -31,6 +31,14 @@ def np_dtype_to_c_type(dtype): ...@@ -31,6 +31,14 @@ def np_dtype_to_c_type(dtype):
return _np_typechar_to_c_typestring[as_char] return _np_typechar_to_c_typestring[as_char]
def enum_to_c_template(enum_class):
res = [f"enum {enum_class.__name__} {{"]
for field in enum_class:
res.append(f"\t{field.name} = {field.value},")
res.append("};")
return "\n".join(res)
def ceil_div(num, denom): def ceil_div(num, denom):
return (num + denom - 1) // denom return (num + denom - 1) // denom
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment