implement channel remap also for the decoder

+ make version 0.27.8

Author: MartinPulec

CMakeLists.txt

@@ -1,6 +1,6 @@
 cmake_minimum_required(VERSION 3.10.0 FATAL_ERROR)
 # change version also in configure.ac
-project(gpujpeg VERSION 0.27.7 LANGUAGES C CUDA)
+project(gpujpeg VERSION 0.27.8 LANGUAGES C CUDA)
 
 # options
 set(BUILD_OPENGL OFF CACHE STRING "Build with OpenGL support, options are: AUTO ON OFF")

NEWS.md

@@ -1,3 +1,8 @@
+2025-08-22 - 0.27.8
+----------
+
+- implement channel remapping via option in pre/postprocessor
+
 2025-07-10 - 0.27.7
 ----------
 

configure.ac

@@ -1,6 +1,6 @@
 AC_PREREQ([2.65])
 # change version also in CMakeLists.txt
-AC_INIT([libgpujpeg],[0.27.7],[https://github.com/CESNET/GPUJPEG/issues],[libgpujpeg],[https://github.com/CESNET/GPUJPEG])
+AC_INIT([libgpujpeg],[0.27.8],[https://github.com/CESNET/GPUJPEG/issues],[libgpujpeg],[https://github.com/CESNET/GPUJPEG])
 AC_CONFIG_MACRO_DIR([m4])
 AC_CONFIG_SRCDIR([src/main.c])
 AC_CONFIG_AUX_DIR([.])

libgpujpeg/gpujpeg_decoder.h

@@ -270,6 +270,10 @@ gpujpeg_decoder_get_image_info(uint8_t *image, size_t image_size, struct gpujpeg
 /// use RLE when writing TGA with gpujpeg_image_save_to_file (default is true), 0 or 1 please note that the option is
 /// global so it affects all decoder instances
 #define GPUJPEG_DEC_OPT_TGA_RLE_BOOL "dec_opt_tga_rle" ///< GPUJPEG_VAL_TRUE or GPUJPEG_VAL_FALSE
+
+/// @copydoc GPUJPEG_ENC_OPT_CHANNEL_REMAP
+#define GPUJPEG_DEC_OPT_CHANNEL_REMAP "dec_opt_channel_remap"
+
 /**
  * sets decoder option
  * @retval GPUJPEG_NOERR  option was sucessfully set

libgpujpeg/gpujpeg_encoder.h

@@ -235,7 +235,7 @@ gpujpeg_encoder_suggest_restart_interval(const struct gpujpeg_image_parameters*
 #define GPUJPEG_ENC_OPT_FLIPPED_BOOL "enc_opt_flipped"
 
 /**
- * remap input channel order
+ * remap channel order
  *
  * Format is "XYZ" or "XYZW" where the letters stand for input channel indices (0-indexed) mapped to output position;
  * so eg. for remapping from ARGB to RGBA the option value will be "1230".

src/gpujpeg_common_internal.h

@@ -539,6 +539,9 @@ gpujpeg_cuda_malloc_host(size_t size);
 void
 gpujpeg_init_term_colors();
 
+int
+gpujpeg_opt_set_channel_remap(struct gpujpeg_coder* coder, const char* val, const char* optname);
+
 #ifdef __cplusplus
 } // extern "C"
 #endif // __cplusplus

src/gpujpeg_decoder.c

@@ -493,6 +493,9 @@ gpujpeg_decoder_set_option(struct gpujpeg_decoder* decoder, const char *opt, con
         image_delegate_stbi_tga_set_rle(strcmp(val, GPUJPEG_VAL_TRUE) == 0);
         return GPUJPEG_NOERR;
     }
+    if ( strcmp(opt, GPUJPEG_DEC_OPT_CHANNEL_REMAP) == 0 ) {
+        return gpujpeg_opt_set_channel_remap(&decoder->coder, val, GPUJPEG_DEC_OPT_CHANNEL_REMAP);
+    }
     ERROR_MSG("Invalid decoder option: %s!\n", opt);
     return GPUJPEG_ERROR;
 }
@@ -502,6 +505,8 @@ gpujpeg_decoder_print_options()
 {
     printf("\t" GPUJPEG_DEC_OPT_TGA_RLE_BOOL "=[" GPUJPEG_VAL_FALSE "|" GPUJPEG_VAL_TRUE
            "] - set decoder option (not) to output RLE TGA\n");
+    printf("\t" GPUJPEG_DEC_OPT_CHANNEL_REMAP "=XYZ[W] - input channel mapping, eg. '210F' for GBRX,\n"
+        "\t\t'210' for GBR; special placeholders 'F' and 'Z' to set a channel to all-ones or all-zeros\n");
 }
 
 /* Documented at declaration */

src/gpujpeg_encoder.c

@@ -652,11 +652,11 @@ gpujpeg_encoder_set_jpeg_header(struct gpujpeg_encoder *encoder, enum gpujpeg_he
     encoder->header_type = header_type;
 }
 
-static int
-enc_opt_set_channel_remap(struct gpujpeg_encoder* encoder, const char* val)
+int
+gpujpeg_opt_set_channel_remap(struct gpujpeg_coder* coder, const char* val, const char *optname)
 {
     if ( strcmp(val, "help") == 0 ) {
-        printf("syntax for " GPUJPEG_ENC_OPT_CHANNEL_REMAP ":\n");
+        printf("syntax for %s:\n", optname);
         printf("\t\"XYZ\" or \"XYZW\" where the letters are input channel indices\n");
         printf("\tplaceholder 'Z' or 'F' can be used to set the channel to all-zeros or all-ones\n");
         printf("\n");
@@ -669,7 +669,7 @@ enc_opt_set_channel_remap(struct gpujpeg_encoder* encoder, const char* val)
         ERROR_MSG("Mapping for more than %d channels specified!\n", GPUJPEG_MAX_COMPONENT_COUNT);
         return GPUJPEG_ERROR;
     }
-    encoder->coder.preprocessor.channel_remap = 0; // clear old
+    coder->preprocessor.channel_remap = 0; // clear old
     const char *ptr = val + strlen(val) - 1;
     while ( ptr >= val ) {
         int src_chan = *ptr - '0';
@@ -680,15 +680,15 @@ enc_opt_set_channel_remap(struct gpujpeg_encoder* encoder, const char* val)
             src_chan = 5;
         }
         else if ( src_chan < 0 || src_chan >= mapped_count ) {
-            ERROR_MSG("Invalid channel index %c for " GPUJPEG_ENC_OPT_CHANNEL_REMAP " (mapping %d channels)!\n", *ptr,
-                      mapped_count);
+            ERROR_MSG("Invalid channel index %c for %s (mapping %d channels)!\n", *ptr,
+                      optname, mapped_count);
             return GPUJPEG_ERROR;
         }
-        encoder->coder.preprocessor.channel_remap <<= 4;
-        encoder->coder.preprocessor.channel_remap |= src_chan;
+        coder->preprocessor.channel_remap <<= 4;
+        coder->preprocessor.channel_remap |= src_chan;
         ptr--;
     }
-    encoder->coder.preprocessor.channel_remap |= mapped_count << 24;
+    coder->preprocessor.channel_remap |= mapped_count << 24;
     return GPUJPEG_NOERR;
 }
 
@@ -742,7 +742,7 @@ gpujpeg_encoder_set_option(struct gpujpeg_encoder* encoder, const char *opt, con
         return GPUJPEG_NOERR;
     }
     if ( strcmp(opt, GPUJPEG_ENC_OPT_CHANNEL_REMAP) == 0 ) {
-        return enc_opt_set_channel_remap(encoder, val);
+        return gpujpeg_opt_set_channel_remap(&encoder->coder, val, GPUJPEG_ENC_OPT_CHANNEL_REMAP);
     }
     ERROR_MSG("Invalid encoder option: %s!\n", opt);
     return GPUJPEG_ERROR;

src/gpujpeg_postprocessor.cu

@@ -479,6 +479,13 @@ gpujpeg_postprocessor_decode(struct gpujpeg_coder* coder, cudaStream_t stream)
     );
     gpujpeg_cuda_check_error("Preprocessor encoding failed", return -1);
 
+    if ( coder->preprocessor.channel_remap != 0 ) {
+        const int ret = gpujpeg_preprocessor_channel_remap(coder, stream);
+        if ( ret != 0 ) {
+            return ret;
+        }
+    }
+
     return 0;
 }
 

src/gpujpeg_preprocessor.cu

@@ -513,11 +513,9 @@ channel_remap_kernel(uint8_t* data, int width, int pitch, int height, unsigned i
     gpujpeg_comp_to_raw_store<pixel_format>(data, width, height, offset, x, y, r);
 }
 
-static int
-channel_remap(struct gpujpeg_encoder* encoder)
+int
+gpujpeg_preprocessor_channel_remap(struct gpujpeg_coder* coder, cudaStream_t stream)
 {
-    struct gpujpeg_coder* coder = &encoder->coder;
-
     const unsigned comp_count = gpujpeg_pixel_format_get_comp_count(coder->param_image.pixel_format);
     const unsigned mapped_count = coder->preprocessor.channel_remap >> 24;
     if (comp_count != mapped_count) {
@@ -551,7 +549,7 @@ channel_remap(struct gpujpeg_encoder* encoder)
         GPUJPEG_ASSERT(0 && "Preprocess from GPUJPEG_PIXFMT_NONE not allowed");
     }
 #undef SWITDH_KERNEL
-    kernel<<<grid, block, 0, encoder->stream>>>(encoder->coder.d_data_raw, width, pitch, height, mapping);
+    kernel<<<grid, block, 0, stream>>>(coder->d_data_raw, width, pitch, height, mapping);
     gpujpeg_cuda_check_error("channel_remap_kernel failed", return -1);
     return 0;
 }
@@ -566,7 +564,7 @@ gpujpeg_preprocessor_encode(struct gpujpeg_encoder * encoder)
     //        (coder->param_image.pixel_format == GPUJPEG_444_U8_P012 && coder->preprocessor.kernel != nullptr));
 
     if ( coder->preprocessor.channel_remap != 0 ) {
-        const int ret = channel_remap(encoder);
+        const int ret = gpujpeg_preprocessor_channel_remap(coder, encoder->stream);
         if ( ret != 0 ) {
             return ret;
         }