pre/postprocessor: add state

MartinPulec · MartinPulec · commit 6a79a2295b5c · 2025-07-10T12:10:00.000+02:00
instead of keeping just the kernel pointer, use state struct

the metadata for kernel doesn't need to be recomputed with every image now
diff --git a/src/gpujpeg_common.c b/src/gpujpeg_common.c
@@ -560,7 +560,6 @@ gpujpeg_coder_init(struct gpujpeg_coder * coder)
     coder->param.segment_info = -1;
     coder->param.color_space_internal = GPUJPEG_NONE;
     coder->param_image.color_space = GPUJPEG_NONE;
-    coder->preprocessor = NULL;
     coder->component = NULL;
     coder->d_component = NULL;
     coder->component_allocated_size = 0;
diff --git a/src/gpujpeg_common_internal.h b/src/gpujpeg_common_internal.h
@@ -41,6 +41,7 @@
 #include <stdlib.h>
 #include "../libgpujpeg/gpujpeg_common.h"
 #include "../libgpujpeg/gpujpeg_type.h"
+#include "gpujpeg_preprocessor_common.cuh"
 #include "gpujpeg_util.h"
 
 // static_assert compat
@@ -328,8 +329,8 @@ struct gpujpeg_coder
     /// Number of allocated segments
     int segment_allocated_size;
 
-    /// Preprocessor data (kernel function pointer)
-    void* preprocessor;
+    /// Preprocessor or postprocessor data
+    struct gpujpeg_preprocessor preprocessor;
 
     /// Maximum sampling factor from components
     struct gpujpeg_component_sampling_factor sampling_factor;
diff --git a/src/gpujpeg_postprocessor.cu b/src/gpujpeg_postprocessor.cu
@@ -33,8 +33,10 @@
  * to raw image. It also does color space transformations.
  */
 
-#include "gpujpeg_colorspace.h"
+#define PREPROCESSOR_INTERNAL_API
 #include "gpujpeg_preprocessor_common.cuh"
+
+#include "gpujpeg_colorspace.h"
 #include "gpujpeg_postprocessor.h"
 #include "gpujpeg_util.h"
 
@@ -410,7 +412,18 @@ gpujpeg_preprocessor_decode_no_transform(struct gpujpeg_coder* coder)
 int
 gpujpeg_postprocessor_decoder_init(struct gpujpeg_coder* coder)
 {
-    coder->preprocessor = NULL;
+    coder->preprocessor.kernel = NULL;
+
+    struct gpujpeg_preprocessor_data *data = &coder->preprocessor.data;;
+    *data = {};
+    for ( int comp = 0; comp < coder->param.comp_count; comp++ ) {
+        assert(coder->sampling_factor.horizontal % coder->component[comp].sampling_factor.horizontal == 0);
+        assert(coder->sampling_factor.vertical % coder->component[comp].sampling_factor.vertical == 0);
+        data->comp[comp].d_data = coder->component[comp].d_data;
+        data->comp[comp].sampling_factor.horizontal = coder->sampling_factor.horizontal / coder->component[comp].sampling_factor.horizontal;
+        data->comp[comp].sampling_factor.vertical = coder->sampling_factor.vertical / coder->component[comp].sampling_factor.vertical;
+        data->comp[comp].data_width = coder->component[comp].data_width;
+    }
 
     if (!gpujpeg_pixel_format_is_interleaved(coder->param_image.pixel_format) &&
             gpujpeg_preprocessor_decode_no_transform(coder)) {
@@ -421,24 +434,24 @@ gpujpeg_postprocessor_decoder_init(struct gpujpeg_coder* coder)
     // assert(coder->param.comp_count == 3 || coder->param.comp_count == 4);
 
     if (coder->param.color_space_internal == coder->param_image.color_space) {
-        coder->preprocessor = (void*)gpujpeg_preprocessor_select_decode_kernel<GPUJPEG_NONE>(coder);
+        coder->preprocessor.kernel = (void*)gpujpeg_preprocessor_select_decode_kernel<GPUJPEG_NONE>(coder);
     }
     else if (coder->param.color_space_internal == GPUJPEG_RGB) {
-        coder->preprocessor = (void*)gpujpeg_preprocessor_select_decode_kernel<GPUJPEG_RGB>(coder);
+        coder->preprocessor.kernel = (void*)gpujpeg_preprocessor_select_decode_kernel<GPUJPEG_RGB>(coder);
     }
     else if (coder->param.color_space_internal == GPUJPEG_YCBCR_BT601) {
-        coder->preprocessor = (void*)gpujpeg_preprocessor_select_decode_kernel<GPUJPEG_YCBCR_BT601>(coder);
+        coder->preprocessor.kernel = (void*)gpujpeg_preprocessor_select_decode_kernel<GPUJPEG_YCBCR_BT601>(coder);
     }
     else if (coder->param.color_space_internal == GPUJPEG_YCBCR_BT601_256LVLS) {
-        coder->preprocessor = (void*)gpujpeg_preprocessor_select_decode_kernel<GPUJPEG_YCBCR_BT601_256LVLS>(coder);
+        coder->preprocessor.kernel = (void*)gpujpeg_preprocessor_select_decode_kernel<GPUJPEG_YCBCR_BT601_256LVLS>(coder);
     }
     else if (coder->param.color_space_internal == GPUJPEG_YCBCR_BT709) {
-        coder->preprocessor = (void*)gpujpeg_preprocessor_select_decode_kernel<GPUJPEG_YCBCR_BT709>(coder);
+        coder->preprocessor.kernel = (void*)gpujpeg_preprocessor_select_decode_kernel<GPUJPEG_YCBCR_BT709>(coder);
     }
     else {
         assert(false);
     }
-    if (coder->preprocessor == NULL) {
+    if (coder->preprocessor.kernel == NULL) {
         return -1;
     }
     return 0;
@@ -487,12 +500,12 @@ gpujpeg_preprocessor_decoder_copy_planar_data(struct gpujpeg_coder * coder, cuda
 int
 gpujpeg_postprocessor_decode(struct gpujpeg_coder* coder, cudaStream_t stream)
 {
-    if (!coder->preprocessor) {
+    if (coder->preprocessor.kernel == nullptr) {
         return gpujpeg_preprocessor_decoder_copy_planar_data(coder, stream);
     }
 
     // Select kernel
-    gpujpeg_preprocessor_decode_kernel kernel = (gpujpeg_preprocessor_decode_kernel)coder->preprocessor;
+    gpujpeg_preprocessor_decode_kernel kernel = (gpujpeg_preprocessor_decode_kernel)coder->preprocessor.kernel;
     assert(kernel != NULL);
 
     int image_width = coder->param_image.width + coder->param_image.width_padding;
@@ -521,17 +534,8 @@ gpujpeg_postprocessor_decode(struct gpujpeg_coder* coder, cudaStream_t stream)
     }
 
     // Run kernel
-    struct gpujpeg_preprocessor_data data = {};
-    for ( int comp = 0; comp < coder->param.comp_count; comp++ ) {
-        assert(coder->sampling_factor.horizontal % coder->component[comp].sampling_factor.horizontal == 0);
-        assert(coder->sampling_factor.vertical % coder->component[comp].sampling_factor.vertical == 0);
-        data.comp[comp].d_data = coder->component[comp].d_data;
-        data.comp[comp].sampling_factor.horizontal = coder->sampling_factor.horizontal / coder->component[comp].sampling_factor.horizontal;
-        data.comp[comp].sampling_factor.vertical = coder->sampling_factor.vertical / coder->component[comp].sampling_factor.vertical;
-        data.comp[comp].data_width = coder->component[comp].data_width;
-    }
     kernel<<<grid, threads, 0, stream>>>(
-        data,
+        coder->preprocessor.data,
         coder->d_data_raw,
         image_width,
         image_height
diff --git a/src/gpujpeg_preprocessor.cu b/src/gpujpeg_preprocessor.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2011-2024, CESNET
+ * Copyright (c) 2011-2025, CESNET
  * Copyright (c) 2011, Silicon Genome, LLC.
  *
  * All rights reserved.
@@ -33,8 +33,10 @@
  * computational kernels. It also does color space transformations.
  */
 
-#include "gpujpeg_colorspace.h"
+#define PREPROCESSOR_INTERNAL_API
 #include "gpujpeg_preprocessor_common.cuh"
+
+#include "gpujpeg_colorspace.h"
 #include "gpujpeg_preprocessor.h"
 #include "gpujpeg_util.h"
 
@@ -325,30 +327,39 @@ gpujpeg_preprocessor_encode_no_transform(struct gpujpeg_coder *coder)
 int
 gpujpeg_preprocessor_encoder_init(struct gpujpeg_coder* coder)
 {
-    coder->preprocessor = NULL;
+    coder->preprocessor.kernel = nullptr;
+    struct gpujpeg_preprocessor_data *data = &coder->preprocessor.data;
+    for ( int comp = 0; comp < coder->param.comp_count; comp++ ) {
+        assert(coder->sampling_factor.horizontal % coder->component[comp].sampling_factor.horizontal == 0);
+        assert(coder->sampling_factor.vertical % coder->component[comp].sampling_factor.vertical == 0);
+        data->comp[comp].d_data = coder->component[comp].d_data;
+        data->comp[comp].sampling_factor.horizontal = coder->sampling_factor.horizontal / coder->component[comp].sampling_factor.horizontal;
+        data->comp[comp].sampling_factor.vertical = coder->sampling_factor.vertical / coder->component[comp].sampling_factor.vertical;
+        data->comp[comp].data_width = coder->component[comp].data_width;
+    }
 
     if ( gpujpeg_preprocessor_encode_no_transform(coder) ) {
         DEBUG_MSG(coder->param.verbose, "Matching format detected - not using preprocessor, using memcpy instead.\n");
         return 0;
     }
 
     if (coder->param.color_space_internal == GPUJPEG_NONE) {
-        coder->preprocessor = (void*)gpujpeg_preprocessor_select_encode_kernel<GPUJPEG_NONE>(coder);
+        coder->preprocessor.kernel = (void*)gpujpeg_preprocessor_select_encode_kernel<GPUJPEG_NONE>(coder);
     }
     else if (coder->param.color_space_internal == GPUJPEG_RGB) {
-        coder->preprocessor = (void*)gpujpeg_preprocessor_select_encode_kernel<GPUJPEG_RGB>(coder);
+        coder->preprocessor.kernel = (void*)gpujpeg_preprocessor_select_encode_kernel<GPUJPEG_RGB>(coder);
     }
     else if (coder->param.color_space_internal == GPUJPEG_YCBCR_BT601) {
-        coder->preprocessor = (void*)gpujpeg_preprocessor_select_encode_kernel<GPUJPEG_YCBCR_BT601>(coder);
+        coder->preprocessor.kernel = (void*)gpujpeg_preprocessor_select_encode_kernel<GPUJPEG_YCBCR_BT601>(coder);
     }
     else if (coder->param.color_space_internal == GPUJPEG_YCBCR_BT601_256LVLS) {
-        coder->preprocessor = (void*)gpujpeg_preprocessor_select_encode_kernel<GPUJPEG_YCBCR_BT601_256LVLS>(coder);
+        coder->preprocessor.kernel = (void*)gpujpeg_preprocessor_select_encode_kernel<GPUJPEG_YCBCR_BT601_256LVLS>(coder);
     }
     else if (coder->param.color_space_internal == GPUJPEG_YCBCR_BT709) {
-        coder->preprocessor = (void*)gpujpeg_preprocessor_select_encode_kernel<GPUJPEG_YCBCR_BT709>(coder);
+        coder->preprocessor.kernel = (void*)gpujpeg_preprocessor_select_encode_kernel<GPUJPEG_YCBCR_BT709>(coder);
     }
 
-    if ( coder->preprocessor == NULL ) {
+    if ( coder->preprocessor.kernel == NULL ) {
         return -1;
     }
 
@@ -361,7 +372,7 @@ gpujpeg_preprocessor_encode_interlaced(struct gpujpeg_encoder * encoder)
     struct gpujpeg_coder* coder = &encoder->coder;
 
     // Select kernel
-    gpujpeg_preprocessor_encode_kernel kernel = (gpujpeg_preprocessor_encode_kernel) coder->preprocessor;
+    gpujpeg_preprocessor_encode_kernel kernel = (gpujpeg_preprocessor_encode_kernel) coder->preprocessor.kernel;
     assert(kernel != NULL);
 
     int image_width = coder->param_image.width;
@@ -395,17 +406,8 @@ gpujpeg_preprocessor_encode_interlaced(struct gpujpeg_encoder * encoder)
     gpujpeg_const_div_prepare(image_width, width_div_mul, width_div_shift);
 
     // Run kernel
-    struct gpujpeg_preprocessor_data data;
-    for ( int comp = 0; comp < coder->param.comp_count; comp++ ) {
-        assert(coder->sampling_factor.horizontal % coder->component[comp].sampling_factor.horizontal == 0);
-        assert(coder->sampling_factor.vertical % coder->component[comp].sampling_factor.vertical == 0);
-        data.comp[comp].d_data = coder->component[comp].d_data;
-        data.comp[comp].sampling_factor.horizontal = coder->sampling_factor.horizontal / coder->component[comp].sampling_factor.horizontal;
-        data.comp[comp].sampling_factor.vertical = coder->sampling_factor.vertical / coder->component[comp].sampling_factor.vertical;
-        data.comp[comp].data_width = coder->component[comp].data_width;
-    }
     kernel<<<grid, threads, 0, encoder->stream>>>(
-        data,
+        coder->preprocessor.data,
         coder->d_data_raw,
         coder->param_image.width_padding,
         image_width,
@@ -465,8 +467,10 @@ gpujpeg_preprocessor_encode(struct gpujpeg_encoder * encoder)
 {
     struct gpujpeg_coder * coder = &encoder->coder;
     /// @todo support padding for other formats
-    assert(!coder->param_image.width_padding || (coder->param_image.pixel_format == GPUJPEG_444_U8_P012 && coder->preprocessor));
-    if (coder->preprocessor) {
+    assert(!coder->param_image.width_padding ||
+           (coder->param_image.pixel_format == GPUJPEG_444_U8_P012 && coder->preprocessor.kernel != nullptr));
+
+    if (coder->preprocessor.kernel != nullptr) {
             return gpujpeg_preprocessor_encode_interlaced(encoder);
     } else {
         return gpujpeg_preprocessor_encoder_copy_planar_data(encoder);
diff --git a/src/gpujpeg_preprocessor_common.cuh b/src/gpujpeg_preprocessor_common.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2011-2024, CESNET
+ * Copyright (c) 2011-2025, CESNET
  * Copyright (c) 2011, Silicon Genome, LLC.
  *
  * All rights reserved.
@@ -27,12 +27,11 @@
  * POSSIBILITY OF SUCH DAMAGE.
  */
 
+#ifndef GPUJPEG_PREPROCESSOR_COMMON_CUH_DCC657E3_2EDF_47E2_90F4_F7CA26829E81
+#define GPUJPEG_PREPROCESSOR_COMMON_CUH_DCC657E3_2EDF_47E2_90F4_F7CA26829E81
+
 #include "../libgpujpeg/gpujpeg_common.h"
 #include "../libgpujpeg/gpujpeg_type.h"
-#include "gpujpeg_common_internal.h"
-
-#include <cassert>
-#include <cstdint>
 
 #define RGB_8BIT_THREADS 256
 
@@ -54,6 +53,15 @@ struct gpujpeg_preprocessor_data
     struct gpujpeg_preprocessor_data_component comp[GPUJPEG_MAX_COMPONENT_COUNT];
 };
 
+struct gpujpeg_preprocessor {
+    void* kernel; // function poitner
+    struct gpujpeg_preprocessor_data data;
+};
+
+#ifdef PREPROCESSOR_INTERNAL_API
+#include "gpujpeg_common_internal.h"
+#include <cassert>
+#include <cstdint>
 /** Value that means that sampling factor has dynamic value */
 #define GPUJPEG_DYNAMIC 16
 
@@ -122,5 +130,7 @@ gpujpeg_preprocessor_make_sampling_factor_i(int comp_count, int numerator_h, int
            coder->component[1].sampling_factor.horizontal, coder->component[1].sampling_factor.vertical,               \
            coder->component[2].sampling_factor.horizontal, coder->component[2].sampling_factor.vertical,               \
            coder->component[3].sampling_factor.horizontal, coder->component[3].sampling_factor.vertical)
+#endif // defined PREPROCESSOR_INTERNAL_API
 
+#endif // defined GPUJPEG_PREPROCESSOR_COMMON_CUH_DCC657E3_2EDF_47E2_90F4_F7CA26829E81
 /* vi: set expandtab sw=4: */
