tmp

Author: MartinPulec

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,6 +461,37 @@ 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(256, 1);
+        size_t 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>>>((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)
@@ -470,11 +501,15 @@ gpujpeg_preprocessor_encode(struct gpujpeg_encoder * encoder)
     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);
+    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;
 };