channel remap: support for 444_U8_P012

Author: MartinPulec

libgpujpeg/gpujpeg_encoder.h

@@ -243,7 +243,7 @@ gpujpeg_encoder_suggest_restart_interval(const struct gpujpeg_image_parameters*
  * The number of image channes must equal the length of the string. Set to "" (empty string) to revert this setting.
  * Letters 'Z' or 'F' can be used instead of indices to fill given output channel with zeros or all-ones.
  *
- * Currently, the remapping is implemented for GPUJPEG_4444_U8_P0123 only.
+ * Currently, the remapping is implemented for GPUJPEG_4444_U8_P0123 and GPUJPEG_444_U8_P012.
  */
 #define GPUJPEG_ENC_OPT_CHANNEL_REMAP "enc_opt_channel_remap"
 

src/gpujpeg_encoder.c

@@ -670,22 +670,23 @@ enc_opt_set_channel_remap(struct gpujpeg_encoder* encoder, const char* val)
         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' ) {
+    const char *ptr = val + strlen(val) - 1;
+    while ( ptr >= val ) {
+        int src_chan = *ptr - '0';
+        if ( *ptr == 'F' ) {
             src_chan = 4;
         }
-        else if ( *val == 'Z' ) {
+        else if ( *ptr == 'Z' ) {
             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", *val,
+            ERROR_MSG("Invalid channel index %c for " GPUJPEG_ENC_OPT_CHANNEL_REMAP " (mapping %d channels)!\n", *ptr,
                       mapped_count);
             return GPUJPEG_ERROR;
         }
-        encoder->coder.preprocessor.channel_remap |= src_chan << 12;
-        val++;
+        encoder->coder.preprocessor.channel_remap <<= 4;
+        encoder->coder.preprocessor.channel_remap |= src_chan;
+        ptr--;
     }
     encoder->coder.preprocessor.channel_remap |= mapped_count << 24;
     return GPUJPEG_NOERR;

src/gpujpeg_preprocessor.cu

@@ -523,6 +523,24 @@ channel_remap_kernel_4444_u8_p0123(uint8_t* data, int width, int pitch, int heig
     }
 }
 
+__global__ void
+channel_remap_kernel_444_u8_p012(uint8_t* data, int width, int pitch, 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;
+    }
+
+    data += y * pitch + x * 3;
+    uint32_t val = data[2] << 16 | data[1] << 8 | data[0];
+    val = __byte_perm(val, 0xFF, byte_map);
+    data[0] = val & 0xFF;
+    data[1] = (val >> 8) & 0xFF;
+    data[2] = val >> 16;
+}
+
 static int
 channel_remap(struct gpujpeg_encoder* encoder)
 {
@@ -537,12 +555,13 @@ channel_remap(struct gpujpeg_encoder* encoder)
     }
     const unsigned mapping = coder->preprocessor.channel_remap & 0xFFFF;
 
+    dim3 block(16, 16);
+    int width = coder->param_image.width;
+    int height = coder->param_image.height;
+    int pitch = (width * gpujpeg_pixel_format_get_comp_count(coder->param_image.pixel_format)) +
+                coder->param_image.width_padding;
+    dim3 grid((width + block.x - 1) / block.x, (height + block.y - 1) / block.y);
     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);
-        int pitch = width * sizeof(uint32_t) + coder->param_image.width_padding;
         bool aligned = coder->param_image.width_padding % sizeof(uint32_t) == 0;
         if (aligned) {
             channel_remap_kernel_4444_u8_p0123<true><<<grid, block, 0, encoder->stream>>>(encoder->coder.d_data_raw, width,
@@ -552,8 +571,13 @@ channel_remap(struct gpujpeg_encoder* encoder)
                                                                                     pitch, height, mapping);
         }
     }
+    else if ( coder->param_image.pixel_format == GPUJPEG_444_U8_P012 ) {
+        channel_remap_kernel_444_u8_p012<<<grid, block, 0, encoder->stream>>>(encoder->coder.d_data_raw, width, pitch,
+                                                                              height, mapping);
+    }
     else {
-        ERROR_MSG("Pixel format %s currently unsupported for channel remap!\n",
+        ERROR_MSG("Pixel format %s currently not supported for channel remap!\n"
+                  "444-u8-p012 and 4444-u8-p0123 are currently supported.\n",
                   gpujpeg_pixel_format_get_name(coder->param_image.pixel_format));
         return -1;
     }