diff --git a/image_analysis/image_preprocessing/ImagePreprocessorCPU.cpp b/image_analysis/image_preprocessing/ImagePreprocessorCPU.cpp
index dd71c20d..8803a73d 100644
--- a/image_analysis/image_preprocessing/ImagePreprocessorCPU.cpp
+++ b/image_analysis/image_preprocessing/ImagePreprocessorCPU.cpp
@@ -46,13 +46,15 @@ ImageStatistics ImagePreprocessorCPU::Analyze(ImagePreprocessorBuffer &processed
         if (mask_1bit[i] != 0) {
             processed_image[i] = INT32_MIN;
             ++ret.masked_pixel_count;
+        } else if (image[i] == err_pixel_val) {
+            // Error/invalid marker = the pixel type's extreme value (0xFFFFFFFF for EIGER uint32).
+            // Tested before saturation, since for unsigned types the marker also exceeds sat_pixel_val
+            // (which is clipped above to the HDF5 saturation_value).
+            processed_image[i] = INT32_MIN;
+            ++ret.error_pixel_count;
         } else if (image[i] >= sat_pixel_val) {
             processed_image[i] = INT32_MAX;
             ++ret.saturated_pixel_count;
-        } else if (std::is_signed<T>::value && (image[i] == err_pixel_val)) {
-            // Error pixels are possible only for signed types
-            processed_image[i] = INT32_MIN;
-            ++ret.error_pixel_count;
         } else {
             processed_image[i] = static_cast<int32_t>(image[i]);
             if (image[i] > ret.max_value)
diff --git a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu
index 5f13339a..1a1f36cb 100644
--- a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu
+++ b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu
@@ -35,27 +35,28 @@ __global__ void preprocess_kernel(
     long long local_max = INT64_MIN;
     long long local_min = INT64_MAX;
 
-    constexpr bool is_signed = std::is_signed<T>::value;
-
     for (int i = blockIdx.x * blockDim.x + threadIdx.x;
          i < npixels;
          i += blockDim.x * gridDim.x) {
         T v = input[i];
 
         bool is_masked = mask[i];
-        bool is_sat = v >= saturation_limit;
-        bool is_err = is_signed && (v == err_value);
+        // Error/invalid marker = the pixel type's extreme value (0xFFFFFFFF for EIGER uint32); tested
+        // before saturation, since for unsigned types the marker also exceeds saturation_limit (which is
+        // clipped to the HDF5 saturation_value). Priority: masked > error > saturated.
+        bool is_err = (v == err_value);
+        bool is_sat = !is_err && (v >= saturation_limit);
 
         bool valid = !(is_masked || is_sat || is_err);
 
         // Output
         output[i] =
-                is_masked ? INT32_MIN : is_sat ? INT32_MAX : is_err ? INT32_MIN : (int32_t) v;
+                is_masked ? INT32_MIN : is_err ? INT32_MIN : is_sat ? INT32_MAX : (int32_t) v;
 
         // Counters
         local_masked += is_masked;
-        local_saturated += (!is_masked && is_sat);
-        local_error += (!is_masked && !is_sat && is_err);
+        local_error += (!is_masked && is_err);
+        local_saturated += (!is_masked && !is_err && is_sat);
 
         // Min/max only for valid
         if (valid) {