add option for bottom-up images

refers to GH-100

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.6 LANGUAGES C CUDA)
+project(gpujpeg VERSION 0.27.7 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-07-10 - 0.27.7
+----------
+
+- add encoder option for vertically flipped images
+
 2025-07-08 - 0.27.6
 ----------
 

configure.ac

@@ -1,6 +1,6 @@
 AC_PREREQ([2.65])
 # change version also in CMakeLists.txt
-AC_INIT([libgpujpeg],[0.27.6],[https://github.com/CESNET/GPUJPEG/issues],[libgpujpeg],[https://github.com/CESNET/GPUJPEG])
+AC_INIT([libgpujpeg],[0.27.7],[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_encoder.h

@@ -231,6 +231,9 @@ gpujpeg_encoder_suggest_restart_interval(const struct gpujpeg_image_parameters*
 #define GPUJPEG_ENC_HDR_VAL_SPIFF    "SPIFF"
 /// @}
 
+/// input image is vertically flipped (bottom-up): values @ref GPUJPEG_VAL_TRUE or @ref GPUJPEG_VAL_FALSE
+#define GPUJPEG_ENC_OPT_FLIPPED_BOOL "enc_opt_flipped"
+
 /**
  * sets encoder option
  * @retval GPUJPEG_NOERR  option was sucessfully set

src/gpujpeg_encoder.c

@@ -688,6 +688,19 @@ 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;
+    }
     ERROR_MSG("Invalid encoder option: %s!\n", opt);
     return GPUJPEG_ERROR;
 }

src/gpujpeg_preprocessor.cu

@@ -461,20 +461,57 @@ gpujpeg_preprocessor_encoder_copy_planar_data(struct gpujpeg_encoder * encoder)
     return 0;
 }
 
+static __global__ void
+vertical_flip_kernel(uint32_t* data,
+                     int width, // image linesize/4
+                     int height // image height in pixels
+)
+{
+    int x = blockIdx.x * blockDim.x + threadIdx.x; // column index
+    int y = blockIdx.y * blockDim.y + threadIdx.y; // row index
+
+    if ( x < width ) {
+        // Flipped row index
+        int flipped_y = height - 1 - y;
+        uint32_t tmp = data[y * width + x];
+        data[y * width + x] = data[flipped_y * width + x];
+        data[flipped_y * width + x] = tmp;
+    }
+}
+
+static int
+flip_lines(struct gpujpeg_encoder* encoder)
+{
+    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);
+    }
+    gpujpeg_cuda_check_error("Preprocessor flip failed", return -1);
+    return 0;
+}
+
 /* Documented at declaration */
 int
 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 ||
+    assert(coder->param_image.width_padding == 0 ||
            (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);
+    int ret = coder->preprocessor.kernel != nullptr ? gpujpeg_preprocessor_encode_interlaced(encoder)
+                                                    : gpujpeg_preprocessor_encoder_copy_planar_data(encoder);
+    if (ret != 0) {
+        return ret;
+    }
+    if (coder->preprocessor.input_flipped) {
+        return flip_lines(encoder);
     }
+    return ret;
 }
 
 /* vi: set expandtab sw=4: */

src/gpujpeg_preprocessor_common.cuh

@@ -33,6 +33,10 @@
 #include "../libgpujpeg/gpujpeg_common.h"
 #include "../libgpujpeg/gpujpeg_type.h"
 
+#ifndef __cplusplus
+#include <stdbool.h>
+#endif
+
 #define RGB_8BIT_THREADS 256
 
 /**
@@ -55,6 +59,7 @@ struct gpujpeg_preprocessor_data
 
 struct gpujpeg_preprocessor {
     void* kernel; // function poitner
+    bool input_flipped; // [preprocess only] input is flipped
     struct gpujpeg_preprocessor_data data;
 };