decoder: add dec_opt_flipped

symmeetric to enc_opt_flipped

Author: MartinPulec

libgpujpeg/gpujpeg_decoder.h

@@ -271,6 +271,9 @@ gpujpeg_decoder_get_image_info(uint8_t *image, size_t image_size, struct gpujpeg
 /// 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
 
+/// output image should be vertically flipped (bottom-up): values @ref GPUJPEG_VAL_TRUE or @ref GPUJPEG_VAL_FALSE
+#define GPUJPEG_DEC_OPT_FLIPPED_BOOL "dec_opt_flipped"
+
 /// @copydoc GPUJPEG_ENC_OPT_CHANNEL_REMAP
 #define GPUJPEG_DEC_OPT_CHANNEL_REMAP "dec_opt_channel_remap"
 

src/gpujpeg_common.c

@@ -2318,4 +2318,19 @@ gpujpeg_init_term_colors()
     init = true;
 }
 
+int
+gpujpeg_parse_bool_opt(bool* out_var, const char* val, const char* opt)
+{
+    if ( strcasecmp(val, GPUJPEG_VAL_TRUE) == 0 ) {
+        *out_var = true;
+        return GPUJPEG_NOERR;
+    }
+    else if ( strcasecmp(val, GPUJPEG_VAL_FALSE) == 0 ) {
+        *out_var = false;
+        return GPUJPEG_NOERR;
+    }
+    ERROR_MSG("Unknown option %s for %s\n", val, opt);
+    return GPUJPEG_ERROR;
+}
+
 /* vi: set expandtab sw=4 : */

src/gpujpeg_common_internal.h

@@ -542,6 +542,9 @@ gpujpeg_init_term_colors();
 int
 gpujpeg_opt_set_channel_remap(struct gpujpeg_coder* coder, const char* val, const char* optname);
 
+int
+gpujpeg_parse_bool_opt(bool* out_var, const char* val, const char* opt);
+
 #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_FLIPPED_BOOL) == 0 ) {
+        return gpujpeg_parse_bool_opt(&decoder->coder.preprocessor.flipped, val, GPUJPEG_DEC_OPT_FLIPPED_BOOL);
+    }
     if ( strcmp(opt, GPUJPEG_DEC_OPT_CHANNEL_REMAP) == 0 ) {
         return gpujpeg_opt_set_channel_remap(&decoder->coder, val, GPUJPEG_DEC_OPT_CHANNEL_REMAP);
     }
@@ -505,6 +508,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_FLIPPED_BOOL "=[" GPUJPEG_VAL_FALSE "|" GPUJPEG_VAL_TRUE
+           "] - vertically flip the output image\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");
 }

src/gpujpeg_encoder.c

@@ -729,17 +729,7 @@ gpujpeg_encoder_set_option(struct gpujpeg_encoder* encoder, const char *opt, con
         return GPUJPEG_NOERR;
     }
     if ( strcmp(opt, GPUJPEG_ENC_OPT_FLIPPED_BOOL) == 0 ) {
-        if ( strcasecmp(val, GPUJPEG_VAL_TRUE) == 0 ) {
-            encoder->coder.preprocessor.input_flipped = true;
-        }
-        else if ( strcasecmp(val, GPUJPEG_VAL_FALSE) == 0 ) {
-            encoder->coder.preprocessor.input_flipped = false;
-        }
-        else {
-            ERROR_MSG("Unknown option %s for " GPUJPEG_ENC_OPT_FLIPPED_BOOL "\n", val);
-            return GPUJPEG_ERROR;
-        }
-        return GPUJPEG_NOERR;
+        return gpujpeg_parse_bool_opt(&encoder->coder.preprocessor.flipped, val, GPUJPEG_ENC_OPT_FLIPPED_BOOL);
     }
     if ( strcmp(opt, GPUJPEG_ENC_OPT_CHANNEL_REMAP) == 0 ) {
         return gpujpeg_opt_set_channel_remap(&encoder->coder, val, GPUJPEG_ENC_OPT_CHANNEL_REMAP);

src/gpujpeg_postprocessor.cu

@@ -436,6 +436,13 @@ 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.flipped ) {
+        int ret = gpujpeg_preprocessor_flip_lines(coder, stream);
+        if ( ret != 0 ) {
+            return ret;
+        }
+    }
+
     if ( coder->preprocessor.kernel == nullptr ) {
         return gpujpeg_preprocessor_decoder_copy_planar_data(coder, stream);
     }

src/gpujpeg_preprocessor.cu

@@ -470,16 +470,15 @@ vertical_flip_kernel(uint32_t* data,
     }
 }
 
-static int
-flip_lines(struct gpujpeg_encoder* encoder)
+int
+gpujpeg_preprocessor_flip_lines(struct gpujpeg_coder* coder, cudaStream_t stream)
 {
-    struct gpujpeg_coder* coder = &encoder->coder;
     for ( int i = 0; i < coder->param.comp_count; ++i ) {
         dim3 block(RGB_8BIT_THREADS, 1);
         int width = coder->component[i].data_width / 4;
         int height = coder->component[i].data_height;
         dim3 grid((width + block.x - 1) / block.x, height / 2); // only half of height
-        vertical_flip_kernel<<<grid, block, 0, encoder->stream>>>((uint32_t*)coder->component[i].d_data, width, height);
+        vertical_flip_kernel<<<grid, block, 0, stream>>>((uint32_t*)coder->component[i].d_data, width, height);
     }
     gpujpeg_cuda_check_error("Preprocessor flip failed", return -1);
     return 0;
@@ -579,8 +578,8 @@ gpujpeg_preprocessor_encode(struct gpujpeg_encoder * encoder)
     if (ret != 0) {
         return ret;
     }
-    if (coder->preprocessor.input_flipped) {
-        return flip_lines(encoder);
+    if ( coder->preprocessor.flipped ) {
+        return gpujpeg_preprocessor_flip_lines(coder, encoder->stream);
     }
     return ret;
 }

src/gpujpeg_preprocessor.h

@@ -69,7 +69,7 @@ struct gpujpeg_preprocessor_data
 struct gpujpeg_preprocessor
 {
     void* kernel;       // function poitner
-    bool input_flipped;         ///< [preprocess only] input buf is flipped
+    bool flipped;               ///< flip image before encode or after decode
     unsigned int channel_remap; ///< remap channels if != 0
                                 ///< format: count_8b | 00000000 | idx0_4b | idx1_4b | idx2_4b | idx3_4b
     struct gpujpeg_preprocessor_data data;
@@ -100,6 +100,9 @@ gpujpeg_preprocessor_encode(struct gpujpeg_encoder * encoder);
 // shared with postprocess
 int
 gpujpeg_preprocessor_channel_remap(struct gpujpeg_coder* coder, cudaStream_t stream);
+
+int
+gpujpeg_preprocessor_flip_lines(struct gpujpeg_coder* coder, cudaStream_t stream);
 
 #ifdef __cplusplus
 }