encoder: add option for channel remap

refer to GH-102

Author: MartinPulec

libgpujpeg/gpujpeg_encoder.h

@@ -234,6 +234,15 @@ gpujpeg_encoder_suggest_restart_interval(const struct gpujpeg_image_parameters*
 /// 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"
 
+/**
+ * remaps input channel order, format "XYZ" or "XYZW" where the letters stand for input channel
+ * indices (0-indexed) mapped to output position; so eg. remapping from ARGB to RGBA will be
+ * "1230". If number of channes doesn't equal the length of the string, the characters are invalid
+ * or out-of bound, the behavior is undefined. set to "" (empty string) to revert this setting.
+ * Placeholder'Z' or 'F' can be used instead of numbers to fill out channel with zeros or all-ones.
+ */
+#define GPUJPEG_ENC_OPT_CHANNEL_REMAP "enc_opt_channel_remap"
+
 /**
  * sets encoder option
  * @retval GPUJPEG_NOERR  option was sucessfully set

src/gpujpeg_encoder.c

@@ -652,6 +652,42 @@ 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)
+{
+    if ( strcmp(val, "help") == 0 ) {
+        printf("syntax for " GPUJPEG_ENC_OPT_CHANNEL_REMAP ":\n");
+        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");
+        printf("examples:\n");
+        printf("\t\"1230\" or \"123F\" to map ARGB to RGBA\n");
+        return GPUJPEG_ERROR;
+    }
+    if ( strlen(val) >= GPUJPEG_MAX_COMPONENT_COUNT ) {
+        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
+    while ( *val != '\0' ) {
+        encoder->coder.preprocessor.channel_remap >>= 4;
+        int src_chan = *val - '0';
+        if ( *val == 'F' ) {
+            src_chan = 4;
+        }
+        else if ( *val == 'Z' ) {
+            src_chan = 5;
+        }
+        else if ( src_chan >= GPUJPEG_MAX_COMPONENT_COUNT ) {
+            ERROR_MSG("Invalid channel %c for " GPUJPEG_ENC_OPT_CHANNEL_REMAP "!\n", *val);
+            return GPUJPEG_ERROR;
+        }
+        encoder->coder.preprocessor.channel_remap |= src_chan << 12;
+        val++;
+    }
+    return GPUJPEG_NOERR;
+}
+
 GPUJPEG_API int
 gpujpeg_encoder_set_option(struct gpujpeg_encoder* encoder, const char *opt, const char* val)
 {
@@ -699,7 +735,9 @@ gpujpeg_encoder_set_option(struct gpujpeg_encoder* encoder, const char *opt, con
             ERROR_MSG("Unknown option %s for " GPUJPEG_ENC_OPT_FLIPPED_BOOL "\n", val);
             return GPUJPEG_ERROR;
         }
-        return GPUJPEG_NOERR;
+    }
+    if ( strcmp(opt, GPUJPEG_ENC_OPT_CHANNEL_REMAP) == 0 ) {
+        return enc_opt_set_channel_remap(encoder, val);
     }
     ERROR_MSG("Invalid encoder option: %s!\n", opt);
     return GPUJPEG_ERROR;

src/gpujpeg_preprocessor.cu

@@ -494,6 +494,47 @@ flip_lines(struct gpujpeg_encoder* encoder)
     return 0;
 }
 
+template <enum gpujpeg_pixel_format>
+__global__ void
+channel_remap_kernel(uint8_t* data, int width, int height, unsigned int byte_map);
+
+template <>
+__global__ void
+channel_remap_kernel<GPUJPEG_4444_U8_P0123>(uint8_t* data, int width, int height, unsigned int byte_map)
+{
+    int x = blockIdx.x * blockDim.x + threadIdx.x; // column index
+    int y = blockIdx.y * blockDim.y + threadIdx.y; // row index
+
+    if ( x >= width || y >= height ) {
+        return;
+    }
+
+    uint32_t* data32 = (uint32_t*)data;
+    uint32_t val = data32[y * width + x];
+    data32[y * width + x] = __byte_perm(val, 0xFF, byte_map);
+}
+
+static int
+channel_remap(struct gpujpeg_encoder* encoder)
+{
+    struct gpujpeg_coder* coder = &encoder->coder;
+    if ( coder->param_image.pixel_format == GPUJPEG_4444_U8_P0123 ) {
+        dim3 block(16, 16);
+        int width = coder->param_image.width;
+        int height = coder->param_image.height;
+        dim3 grid((width + block.x - 1) / block.x, (height + block.y - 1) / block.y);
+        channel_remap_kernel<GPUJPEG_4444_U8_P0123><<<grid, block, 0, encoder->stream>>>(
+            encoder->coder.d_data_raw, width, height, coder->preprocessor.channel_remap);
+    }
+    else {
+        ERROR_MSG("Pixel format %s currently unsupported for channel remap!\n",
+                  gpujpeg_pixel_format_get_name(coder->param_image.pixel_format));
+        return -1;
+    }
+    gpujpeg_cuda_check_error("channel_remap_kernel failed", return -1);
+    return 0;
+}
+
 /* Documented at declaration */
 int
 gpujpeg_preprocessor_encode(struct gpujpeg_encoder * encoder)
@@ -503,6 +544,13 @@ gpujpeg_preprocessor_encode(struct gpujpeg_encoder * encoder)
     assert(coder->param_image.width_padding == 0 ||
            (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);
+        if ( ret != 0 ) {
+            return ret;
+        }
+    }
+
     int ret = coder->preprocessor.kernel != nullptr ? gpujpeg_preprocessor_encode_interlaced(encoder)
                                                     : gpujpeg_preprocessor_encoder_copy_planar_data(encoder);
     if (ret != 0) {

src/gpujpeg_preprocessor.h

@@ -69,7 +69,8 @@ struct gpujpeg_preprocessor_data
 struct gpujpeg_preprocessor
 {
     void* kernel;       // function poitner
-    bool input_flipped; // [preprocess only] input is flipped
+    bool input_flipped;         ///< [preprocess only] input buf is flipped
+    unsigned int channel_remap; ///< remap input channels if != 0; currently preproecss only
     struct gpujpeg_preprocessor_data data;
 };