From 8be6ab044351775b50b7ee2a6900b63f6f334e43 Mon Sep 17 00:00:00 2001 From: adam-ce <5292991+adam-ce@users.noreply.github.com> Date: Mon, 2 Nov 2026 22:11:52 +0100 Subject: [PATCH 1/5] encoder: average subsampled chroma The preprocessor stored the chroma of the top-left pixel of each subsampling window (point sampling), which aliases and shifts chroma by half a pixel from the centered siting that JFIF (and decoders) assume. Average the window instead if the input has full chroma resolution and Cb and Cr have the same sampling. On four 4096x4096 orthophotos (q95, 4:2:0, decoded by libjpeg-turbo with fancy upsampling) this improves PSNR by 0.17-0.49 dB, the JPEG is 0.4-0.6 % smaller and the preprocessor kernel takes ~0.01 ms longer. The previous behavior can be selected by GPUJPEG_ENC_OPT_CHROMA_FILTER_BOOL. --- libgpujpeg/gpujpeg_encoder.h | 3 +++ src/gpujpeg_common.c | 1 + src/gpujpeg_encoder.c | 6 +++++ src/gpujpeg_preprocessor.cu | 52 ++++++++++++++++++++++++++++++++++++ src/gpujpeg_preprocessor.h | 2 ++ 5 files changed, 64 insertions(+) diff --git a/libgpujpeg/gpujpeg_encoder.h b/libgpujpeg/gpujpeg_encoder.h index 581e0a2e..690d0b3e 100644 --- a/libgpujpeg/gpujpeg_encoder.h +++ b/libgpujpeg/gpujpeg_encoder.h @@ -225,6 +225,9 @@ 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" +/// average subsampled chroma over its sampling window (centered as JFIF assumes) instead of point sampling; values +/// @ref GPUJPEG_VAL_TRUE (default) or @ref GPUJPEG_VAL_FALSE +#define GPUJPEG_ENC_OPT_CHROMA_FILTER_BOOL "enc_opt_chroma_filter" /// custom exif tag in format :TYPE= #define GPUJPEG_ENC_OPT_EXIF_TAG "enc_exif_tag" /// set image orientation - syntax "=[-]" or "help"; only if header supports (Exif, SPIFF) diff --git a/src/gpujpeg_common.c b/src/gpujpeg_common.c index 2bf4b7a0..a42b5976 100644 --- a/src/gpujpeg_common.c +++ b/src/gpujpeg_common.c @@ -582,6 +582,7 @@ gpujpeg_coder_init(struct gpujpeg_coder * coder) coder->d_data_compressed = NULL; coder->d_temp_huffman = NULL; coder->data_compressed_allocated_size = 0; + coder->preprocessor.chroma_filter = true; GPUJPEG_CUSTOM_TIMER_CREATE(coder->duration_memory_to, return -1); GPUJPEG_CUSTOM_TIMER_CREATE(coder->duration_memory_from, return -1); diff --git a/src/gpujpeg_encoder.c b/src/gpujpeg_encoder.c index 067fbd5c..3cfd62bb 100644 --- a/src/gpujpeg_encoder.c +++ b/src/gpujpeg_encoder.c @@ -770,6 +770,10 @@ gpujpeg_encoder_set_option(struct gpujpeg_encoder* encoder, const char *opt, con if ( strcmp(opt, GPUJPEG_ENC_OPT_CHANNEL_REMAP) == 0 ) { return gpujpeg_opt_set_channel_remap(&encoder->coder, val, GPUJPEG_ENC_OPT_CHANNEL_REMAP); } + if ( strcmp(opt, GPUJPEG_ENC_OPT_CHROMA_FILTER_BOOL) == 0 ) { + return gpujpeg_parse_bool_opt(&encoder->coder.preprocessor.chroma_filter, val, + GPUJPEG_ENC_OPT_CHROMA_FILTER_BOOL); + } if ( strcmp(opt, GPUJPEG_ENC_OPT_EXIF_TAG) == 0 ) { encoder->header_type = GPUJPEG_HEADER_EXIF; return gpujpeg_exif_add_tag(&encoder->writer->exif_tags, val) ? GPUJPEG_NOERR : GPUJPEG_ERROR; @@ -790,6 +794,8 @@ gpujpeg_encoder_print_options() { "] - whether is the input image should be vertically flipped (prior encode)\n"); printf("\t" GPUJPEG_ENC_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"); + printf("\t" GPUJPEG_ENC_OPT_CHROMA_FILTER_BOOL "=[" GPUJPEG_VAL_FALSE "|" GPUJPEG_VAL_TRUE + "] - average subsampled chroma (default) instead of point sampling\n"); printf("\t" GPUJPEG_ENC_OPT_EXIF_TAG "==|help - custom EXIF tag (use help for syntax)\n"); printf("\t" GPUJPEG_ENC_OPT_METADATA "==|help - set image metadata\n"); } diff --git a/src/gpujpeg_preprocessor.cu b/src/gpujpeg_preprocessor.cu index 7289e747..f2ebda68 100644 --- a/src/gpujpeg_preprocessor.cu +++ b/src/gpujpeg_preprocessor.cu @@ -149,6 +149,38 @@ inline __device__ void raw_to_comp_load(const uint8_t* d_d } } +/** + * Set chroma of the top-left pixel of each subsampling window to the window average, so that the stored sample is + * centered as JFIF assumes (instead of point sampling). The window is clamped at the right and bottom image edges. + */ +template +static __device__ void +gpujpeg_preprocessor_chroma_box_filter(uchar4& r, const uint8_t* d_data_raw, int image_width_padding, int image_width, + int image_height, int x, int y, int fh, int fv) +{ + if ( x % fh != 0 || y % fv != 0 ) { + return; + } + int sum_cb = 0; + int sum_cr = 0; + for ( int dy = 0; dy < fv; ++dy ) { + for ( int dx = 0; dx < fh; ++dx ) { + int xx = min(x + dx, image_width - 1); + int yy = min(y + dy, image_height - 1); + int offset = (yy * image_width + xx) * unit_size() + image_width_padding * yy; + uchar4 q; + raw_to_comp_load(d_data_raw, image_width, image_height, offset, xx, yy, q); + gpujpeg_color_transform::perform(q); + sum_cb += q.y; + sum_cr += q.z; + } + } + const int n = fh * fv; + r.y = (sum_cb + n / 2) / n; + r.z = (sum_cr + n / 2) / n; +} + /** * Kernel - Copy raw image source data into three separated component buffers */ @@ -192,6 +224,25 @@ gpujpeg_preprocessor_raw_to_comp_kernel(struct gpujpeg_preprocessor_data data, c // Color transform gpujpeg_color_transform::perform(r); + // Average chroma if subsampled from full resolution (Cb and Cr with equal sampling) + constexpr bool full_chroma_input = pixel_format == GPUJPEG_444_U8_P012 || pixel_format == GPUJPEG_4444_U8_P0123 || + pixel_format == GPUJPEG_444_U8_P0P1P2; + if ( full_chroma_input && data.chroma_filter ) { + const int fh = + s_comp2_samp_factor_h == GPUJPEG_DYNAMIC ? data.comp[1].sampling_factor.horizontal : s_comp2_samp_factor_h; + const int fv = + s_comp2_samp_factor_v == GPUJPEG_DYNAMIC ? data.comp[1].sampling_factor.vertical : s_comp2_samp_factor_v; + const int fh_cr = + s_comp3_samp_factor_h == GPUJPEG_DYNAMIC ? data.comp[2].sampling_factor.horizontal : s_comp3_samp_factor_h; + const int fv_cr = + s_comp3_samp_factor_v == GPUJPEG_DYNAMIC ? data.comp[2].sampling_factor.vertical : s_comp3_samp_factor_v; + if ( fh * fv > 1 && fh == fh_cr && fv == fv_cr ) { + gpujpeg_preprocessor_chroma_box_filter( + r, d_data_raw, image_width_padding, image_width, image_height, image_position_x, image_position_y, fh, + fv); + } + } + // Store gpujpeg_preprocessor_raw_to_comp_storepreprocessor.data.chroma_filter = coder->preprocessor.chroma_filter; kernel<<stream>>>( coder->preprocessor.data, coder->d_data_raw, diff --git a/src/gpujpeg_preprocessor.h b/src/gpujpeg_preprocessor.h index d9c76025..f56a1ff7 100644 --- a/src/gpujpeg_preprocessor.h +++ b/src/gpujpeg_preprocessor.h @@ -61,6 +61,7 @@ struct gpujpeg_preprocessor_data_component struct gpujpeg_preprocessor_data { struct gpujpeg_preprocessor_data_component comp[GPUJPEG_MAX_COMPONENT_COUNT]; + bool chroma_filter; ///< copy of gpujpeg_preprocessor::chroma_filter set before kernel launch }; /** @@ -70,6 +71,7 @@ struct gpujpeg_preprocessor { void* kernel; // function poitner bool flipped; ///< flip image before encode or after decode + bool chroma_filter; ///< box-filter subsampled chroma (encode), interpolate it (decode); default true 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; From 7e4a5cedc977d53b06fd3271ce2b51555391d491 Mon Sep 17 00:00:00 2001 From: adam-ce <5292991+adam-ce@users.noreply.github.com> Date: Mon, 2 Nov 2026 22:11:52 +0100 Subject: [PATCH 2/5] decoder: interpolate subsampled chroma The postprocessor replicated 4:2:0, 4:2:2 and 4:4:0 chroma. For full-resolution output, interpolate it with the triangle filter that libjpeg uses by default ("fancy upsampling", same weights and rounding). Other layouts (dynamic kernels) and subsampled output formats are unchanged. Unlike libjpeg, no exception is made for chroma width <= 2. Decoding libjpeg-turbo q95 4:2:0 JPEGs of four 4096x4096 orthophotos, PSNR improves by 0.32-0.70 dB (now equal to libjpeg-turbo), the postprocessor kernel takes ~0.01 ms longer. The previous behavior can be selected by GPUJPEG_DEC_OPT_CHROMA_FILTER_BOOL. --- libgpujpeg/gpujpeg_decoder.h | 4 ++++ src/gpujpeg_decoder.c | 6 ++++++ src/gpujpeg_postprocessor.cu | 42 ++++++++++++++++++++++++++++++++++++ src/gpujpeg_preprocessor.h | 2 ++ 4 files changed, 54 insertions(+) diff --git a/libgpujpeg/gpujpeg_decoder.h b/libgpujpeg/gpujpeg_decoder.h index 015cab83..5ab56bdc 100644 --- a/libgpujpeg/gpujpeg_decoder.h +++ b/libgpujpeg/gpujpeg_decoder.h @@ -297,6 +297,10 @@ gpujpeg_decoder_get_image_info(uint8_t *image, size_t image_size, struct gpujpeg /// 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" +/// interpolate 4:2:0, 4:2:2 and 4:4:0 chroma for full-resolution output (triangle filter as libjpeg "fancy +/// upsampling") instead of replicating it; values @ref GPUJPEG_VAL_TRUE (default) or @ref GPUJPEG_VAL_FALSE +#define GPUJPEG_DEC_OPT_CHROMA_FILTER_BOOL "dec_opt_chroma_filter" + /// @copydoc GPUJPEG_ENC_OPT_CHANNEL_REMAP #define GPUJPEG_DEC_OPT_CHANNEL_REMAP "dec_opt_channel_remap" diff --git a/src/gpujpeg_decoder.c b/src/gpujpeg_decoder.c index 2dce5036..33d20151 100644 --- a/src/gpujpeg_decoder.c +++ b/src/gpujpeg_decoder.c @@ -504,6 +504,10 @@ gpujpeg_decoder_set_option(struct gpujpeg_decoder* decoder, const char *opt, con if ( strcmp(opt, GPUJPEG_DEC_OPT_CHANNEL_REMAP) == 0 ) { return gpujpeg_opt_set_channel_remap(&decoder->coder, val, GPUJPEG_DEC_OPT_CHANNEL_REMAP); } + if ( strcmp(opt, GPUJPEG_DEC_OPT_CHROMA_FILTER_BOOL) == 0 ) { + return gpujpeg_parse_bool_opt(&decoder->coder.preprocessor.chroma_filter, val, + GPUJPEG_DEC_OPT_CHROMA_FILTER_BOOL); + } if ( strcmp(opt, GPUJPEG_DEC_OPT_ALIGNMENT_BYTES_INT) == 0 ) { decoder->req_alignment = atoi(val); return GPUJPEG_NOERR; @@ -521,6 +525,8 @@ gpujpeg_decoder_print_options() "] - vertically flip the output image\n"); printf("\t" GPUJPEG_DEC_OPT_CHANNEL_REMAP "=XYZ[W] - ouput 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"); + printf("\t" GPUJPEG_DEC_OPT_CHROMA_FILTER_BOOL "=[" GPUJPEG_VAL_FALSE "|" GPUJPEG_VAL_TRUE + "] - interpolate subsampled chroma (default) instead of replicating it\n"); printf("\t" GPUJPEG_DEC_OPT_ALIGNMENT_BYTES_INT "= - output alignment in bytes,\n" "\t\tshould not be used for gpujpegtool (cannot process)\n"); } diff --git a/src/gpujpeg_postprocessor.cu b/src/gpujpeg_postprocessor.cu index 0a8dd846..b9650351 100644 --- a/src/gpujpeg_postprocessor.cu +++ b/src/gpujpeg_postprocessor.cu @@ -112,6 +112,31 @@ struct gpujpeg_preprocessor_comp_to_raw_load } }; +/** + * Interpolate a 2x subsampled component at full-resolution pixel (x, y) like libjpeg "fancy" upsampling (triangle + * filter, weights 3:1 per subsampled axis); neighbours are clamped to the component edges. + */ +template +static __device__ uint8_t +gpujpeg_postprocessor_chroma_interpolate(const struct gpujpeg_preprocessor_data_component& comp, int x, int y) +{ + const uint8_t* d = comp.d_data; + const int pitch = comp.data_width; + const int cx = x / fh; + const int cy = y / fv; + const int nx = fh == 2 ? ((x & 1) ? min(cx + 1, comp.width - 1) : max(cx - 1, 0)) : cx; + const int ny = fv == 2 ? ((y & 1) ? min(cy + 1, comp.height - 1) : max(cy - 1, 0)) : cy; + if ( fh == 2 && fv == 2 ) { + const int this_col = 3 * d[cy * pitch + cx] + d[ny * pitch + cx]; + const int next_col = 3 * d[cy * pitch + nx] + d[ny * pitch + nx]; + return (3 * this_col + next_col + ((x & 1) ? 7 : 8)) >> 4; + } + if ( fh == 2 ) { + return (3 * d[cy * pitch + cx] + d[cy * pitch + nx] + ((x & 1) ? 2 : 1)) >> 2; + } + return (3 * d[cy * pitch + cx] + d[ny * pitch + cx] + ((y & 1) ? 2 : 1)) >> 2; +} + /// set alpha that may not be included in input data but may be read by the CS conv template struct pre_load { @@ -207,6 +232,20 @@ gpujpeg_preprocessor_comp_to_raw_kernel(struct gpujpeg_preprocessor_data data, u gpujpeg_preprocessor_comp_to_raw_load::perform(r, image_position_x, image_position_y, data); post_load::perform(r, data); + // Interpolate 2x subsampled chroma (Cb and Cr with equal sampling) for full-resolution output, precompiled + // kernels only + constexpr bool full_chroma_output = pixel_format == GPUJPEG_444_U8_P012 || pixel_format == GPUJPEG_4444_U8_P0123 || + pixel_format == GPUJPEG_444_U8_P0P1P2; + constexpr bool interpolate = + s_comp2_samp_factor_h <= 2 && s_comp2_samp_factor_v <= 2 && s_comp2_samp_factor_h * s_comp2_samp_factor_v > 1 && + s_comp2_samp_factor_h == s_comp3_samp_factor_h && s_comp2_samp_factor_v == s_comp3_samp_factor_v; + if ( full_chroma_output && interpolate && data.chroma_filter ) { + constexpr int fh = interpolate ? s_comp2_samp_factor_h : 1; + constexpr int fv = interpolate ? s_comp2_samp_factor_v : 1; + r.y = gpujpeg_postprocessor_chroma_interpolate(data.comp[1], image_position_x, image_position_y); + r.z = gpujpeg_postprocessor_chroma_interpolate(data.comp[2], image_position_x, image_position_y); + } + // Color transform gpujpeg_color_transform::perform(r); @@ -359,6 +398,8 @@ gpujpeg_postprocessor_decoder_init(struct gpujpeg_coder* coder) data->comp[comp].sampling_factor.horizontal = coder->sampling_factor.horizontal / coder->component[comp].sampling_factor.horizontal; data->comp[comp].sampling_factor.vertical = coder->sampling_factor.vertical / coder->component[comp].sampling_factor.vertical; data->comp[comp].data_width = coder->component[comp].data_width; + data->comp[comp].width = coder->component[comp].width; + data->comp[comp].height = coder->component[comp].height; } if (!gpujpeg_pixel_format_is_interleaved(coder->param_image.pixel_format) && @@ -481,6 +522,7 @@ gpujpeg_postprocessor_decode(struct gpujpeg_coder* coder, cudaStream_t stream) } // Run kernel + coder->preprocessor.data.chroma_filter = coder->preprocessor.chroma_filter; kernel<<>>( coder->preprocessor.data, coder->d_data_raw, diff --git a/src/gpujpeg_preprocessor.h b/src/gpujpeg_preprocessor.h index f56a1ff7..5a0b6262 100644 --- a/src/gpujpeg_preprocessor.h +++ b/src/gpujpeg_preprocessor.h @@ -52,6 +52,8 @@ struct gpujpeg_preprocessor_data_component { uint8_t* d_data; int data_width; + int width; ///< component width without MCU padding (set for decoder) + int height; ///< component height without MCU padding (set for decoder) struct gpujpeg_component_sampling_factor sampling_factor; }; From f525110444d5ad59e2f434f4d7911f400003c2db Mon Sep 17 00:00:00 2001 From: adam-ce <5292991+adam-ce@users.noreply.github.com> Date: Mon, 2 Nov 2026 22:11:52 +0100 Subject: [PATCH 3/5] unit tests: chroma filters Compares the encoder's chroma planes and the decoder's full-resolution output with a CPU reference for 4:2:0, 4:2:2, 4:4:0 and 4:1:1 with the filters on and off, and the option change without reinitialization. --- test/unit/run_tests.c | 5 + test/unit/test_chroma_filter.c | 192 +++++++++++++++++++++++++++++++++ 2 files changed, 197 insertions(+) create mode 100644 test/unit/test_chroma_filter.c diff --git a/test/unit/run_tests.c b/test/unit/run_tests.c index a4ba0eac..c5e1d356 100644 --- a/test/unit/run_tests.c +++ b/test/unit/run_tests.c @@ -84,6 +84,10 @@ static void encode_gpu_mem_as_cpu() { void test_gh_95(); +// defined in test_chroma_filter.c +void +test_chroma_filter(); + static char* test_format_number_with_delim() { struct { @@ -127,6 +131,7 @@ int main() { encode_gpu_mem_as_cpu(); test_gh_95(); test_format_number_with_delim(); + test_chroma_filter(); printf("PASSED\n"); } diff --git a/test/unit/test_chroma_filter.c b/test/unit/test_chroma_filter.c new file mode 100644 index 00000000..744fa275 --- /dev/null +++ b/test/unit/test_chroma_filter.c @@ -0,0 +1,192 @@ +#include +#include +#include +#include +#include +#include +#include + +#include "../../src/gpujpeg_decoder_internal.h" +#include "../../src/gpujpeg_encoder_internal.h" + +/* + * Compares the chroma filters with a CPU reference: the encoder's component planes after preprocessing and the + * decoder's full-resolution output computed from its component planes. Input and output are in the JPEG color space, + * so no color transform is involved. + */ + +enum { + CHROMA_W = 37, + CHROMA_H = 21, + CHROMA_PAD = 3, ///< input line padding in bytes + CHROMA_PITCH = CHROMA_W * 3 + CHROMA_PAD, +}; + +#define CHROMA_CHECK(cond) \ + if ( !(cond) ) { \ + fprintf(stderr, "%s:%d: chroma filter check '%s' failed!\n", __FILE__, __LINE__, #cond); \ + abort(); \ + } + +static uint8_t* +chroma_plane(const struct gpujpeg_component* comp) +{ + size_t size = (size_t)comp->data_width * comp->height; + uint8_t* plane = malloc(size); + CHROMA_CHECK(cudaMemcpy(plane, comp->d_data, size, cudaMemcpyDeviceToHost) == cudaSuccess); + return plane; +} + +/// encoder reference: box average or point sample of channel ch for chroma sample (cx, cy), divisors h, v +static int +chroma_ref_down(const uint8_t* in, int ch, int cx, int cy, int h, int v, bool filter) +{ + if ( !filter ) { + return in[cy * v * CHROMA_PITCH + cx * h * 3 + ch]; + } + int sum = 0; + for ( int dy = 0; dy < v; ++dy ) { + for ( int dx = 0; dx < h; ++dx ) { + int x = MIN(cx * h + dx, CHROMA_W - 1); + int y = MIN(cy * v + dy, CHROMA_H - 1); + sum += in[y * CHROMA_PITCH + x * 3 + ch]; + } + } + return (sum + h * v / 2) / (h * v); +} + +/// decoder reference: libjpeg "fancy" upsampling (or replication) of plane p at pixel (x, y), divisors h, v +static int +chroma_ref_up(const uint8_t* p, const struct gpujpeg_component* comp, int x, int y, int h, int v, bool filter) +{ + const int pitch = comp->data_width; + const int cx = x / h; + const int cy = y / v; + if ( !filter ) { + return p[cy * pitch + cx]; + } + const int nx = h == 2 ? (x % 2 ? MIN(cx + 1, comp->width - 1) : (cx > 0 ? cx - 1 : 0)) : cx; + const int ny = v == 2 ? (y % 2 ? MIN(cy + 1, comp->height - 1) : (cy > 0 ? cy - 1 : 0)) : cy; + if ( h == 2 && v == 2 ) { + int this_col = 3 * p[cy * pitch + cx] + p[ny * pitch + cx]; + int next_col = 3 * p[cy * pitch + nx] + p[ny * pitch + nx]; + return (3 * this_col + next_col + (x % 2 ? 7 : 8)) >> 4; + } + if ( h == 2 ) { + return (3 * p[cy * pitch + cx] + p[cy * pitch + nx] + (x % 2 ? 2 : 1)) >> 2; + } + return (3 * p[cy * pitch + cx] + p[ny * pitch + cx] + (y % 2 ? 2 : 1)) >> 2; +} + +/// @param interpolated whether the decoder interpolates this subsampling (precompiled 4:2:0, 4:2:2, 4:4:0 kernels) +static void +chroma_case(struct gpujpeg_encoder* encoder, struct gpujpeg_decoder* decoder, const uint8_t* in, + gpujpeg_sampling_factor_t subsampling, bool interpolated, bool filter) +{ + const char* val = filter ? GPUJPEG_VAL_TRUE : GPUJPEG_VAL_FALSE; + CHROMA_CHECK(gpujpeg_encoder_set_option(encoder, GPUJPEG_ENC_OPT_CHROMA_FILTER_BOOL, val) == GPUJPEG_NOERR); + CHROMA_CHECK(gpujpeg_decoder_set_option(decoder, GPUJPEG_DEC_OPT_CHROMA_FILTER_BOOL, val) == GPUJPEG_NOERR); + + struct gpujpeg_parameters param = gpujpeg_default_parameters(); + param.verbose = GPUJPEG_LL_QUIET; + param.interleaved = 1; + gpujpeg_parameters_chroma_subsampling(¶m, subsampling); + struct gpujpeg_image_parameters param_image = gpujpeg_default_image_parameters(); + param_image.width = CHROMA_W; + param_image.height = CHROMA_H; + param_image.width_padding = CHROMA_PAD; + param_image.color_space = GPUJPEG_YCBCR_JPEG; + struct gpujpeg_encoder_input input = gpujpeg_encoder_input_image((uint8_t*)in); + uint8_t* jpeg = NULL; + size_t jpeg_len = 0; + CHROMA_CHECK(gpujpeg_encoder_encode(encoder, ¶m, ¶m_image, &input, &jpeg, &jpeg_len) == 0); + + struct gpujpeg_coder* coder = &encoder->coder; + for ( int c = 1; c < 3; ++c ) { + const struct gpujpeg_component* comp = &coder->component[c]; + const int h = coder->sampling_factor.horizontal / comp->sampling_factor.horizontal; + const int v = coder->sampling_factor.vertical / comp->sampling_factor.vertical; + uint8_t* plane = chroma_plane(comp); + for ( int cy = 0; cy < comp->height; ++cy ) { + for ( int cx = 0; cx < comp->width; ++cx ) { + CHROMA_CHECK(plane[cy * comp->data_width + cx] == chroma_ref_down(in, c, cx, cy, h, v, filter)); + } + } + free(plane); + } + + gpujpeg_decoder_set_output_format(decoder, GPUJPEG_YCBCR_JPEG, GPUJPEG_444_U8_P012); + struct gpujpeg_decoder_output output; + gpujpeg_decoder_output_set_default(&output); + CHROMA_CHECK(gpujpeg_decoder_decode(decoder, jpeg, jpeg_len, &output) == 0); + CHROMA_CHECK(output.data_size == (size_t)CHROMA_W * CHROMA_H * 3); + coder = &decoder->coder; + for ( int c = 1; c < 3; ++c ) { + const struct gpujpeg_component* comp = &coder->component[c]; + const int h = coder->sampling_factor.horizontal / comp->sampling_factor.horizontal; + const int v = coder->sampling_factor.vertical / comp->sampling_factor.vertical; + uint8_t* plane = chroma_plane(comp); + for ( int y = 0; y < CHROMA_H; ++y ) { + for ( int x = 0; x < CHROMA_W; ++x ) { + CHROMA_CHECK(output.data[(y * CHROMA_W + x) * 3 + c] == + chroma_ref_up(plane, comp, x, y, h, v, filter && interpolated)); + } + } + free(plane); + } +} + +static struct gpujpeg_decoder* +chroma_decoder_create(void) +{ + struct gpujpeg_decoder_init_parameters param = gpujpeg_decoder_default_init_parameters(); + param.verbose = GPUJPEG_LL_QUIET; + return gpujpeg_decoder_create_with_params(¶m); +} + +void +test_chroma_filter(); + +void +test_chroma_filter() +{ + printf("testing %s: ", __func__); + if ( getenv("NO_CUDA_DEV") ) { + printf("--\n"); + return; + } + static uint8_t in[CHROMA_H * CHROMA_PITCH]; + for ( unsigned i = 0; i < sizeof in; ++i ) { + in[i] = (uint8_t)((i * 2654435761U) >> 24); + } + const struct + { + gpujpeg_sampling_factor_t subsampling; + bool interpolated; + } cases[] = { + {GPUJPEG_SUBSAMPLING_420, true }, + {GPUJPEG_SUBSAMPLING_422, true }, + {GPUJPEG_SUBSAMPLING_440, true }, + {GPUJPEG_SUBSAMPLING_411, false}, // dynamic kernels: encoder averages, decoder replicates + }; + for ( unsigned i = 0; i < sizeof cases / sizeof cases[0]; ++i ) { + for ( int filter = 0; filter < 2; ++filter ) { + struct gpujpeg_encoder* encoder = gpujpeg_encoder_create(0); + struct gpujpeg_decoder* decoder = chroma_decoder_create(); + assert(encoder != NULL && decoder != NULL); + chroma_case(encoder, decoder, in, cases[i].subsampling, cases[i].interpolated, filter); + gpujpeg_decoder_destroy(decoder); + gpujpeg_encoder_destroy(encoder); + } + } + // the option applies to the next call without a reinitialization + struct gpujpeg_encoder* encoder = gpujpeg_encoder_create(0); + struct gpujpeg_decoder* decoder = chroma_decoder_create(); + assert(encoder != NULL && decoder != NULL); + chroma_case(encoder, decoder, in, GPUJPEG_SUBSAMPLING_420, true, true); + chroma_case(encoder, decoder, in, GPUJPEG_SUBSAMPLING_420, true, false); + chroma_case(encoder, decoder, in, GPUJPEG_SUBSAMPLING_420, true, true); + gpujpeg_decoder_destroy(decoder); + gpujpeg_encoder_destroy(encoder); + printf("Ok\n"); +} From 853814c1819513c9dd2862ebf31f932d1684047c Mon Sep 17 00:00:00 2001 From: adam-ce <5292991+adam-ce@users.noreply.github.com> Date: Tue, 3 Nov 2026 05:53:38 +0100 Subject: [PATCH 4/5] unit tests: chroma filter default, toggle per layout Check also the default (filter on) and the option change on the same coders for each layout (instead of only 4:2:0); express the 4:2:0 reference with the plain 9/3/3/1 weights. --- test/unit/test_chroma_filter.c | 41 +++++++++++++++------------------- 1 file changed, 18 insertions(+), 23 deletions(-) diff --git a/test/unit/test_chroma_filter.c b/test/unit/test_chroma_filter.c index 744fa275..53e4d2d8 100644 --- a/test/unit/test_chroma_filter.c +++ b/test/unit/test_chroma_filter.c @@ -67,10 +67,9 @@ chroma_ref_up(const uint8_t* p, const struct gpujpeg_component* comp, int x, int } const int nx = h == 2 ? (x % 2 ? MIN(cx + 1, comp->width - 1) : (cx > 0 ? cx - 1 : 0)) : cx; const int ny = v == 2 ? (y % 2 ? MIN(cy + 1, comp->height - 1) : (cy > 0 ? cy - 1 : 0)) : cy; - if ( h == 2 && v == 2 ) { - int this_col = 3 * p[cy * pitch + cx] + p[ny * pitch + cx]; - int next_col = 3 * p[cy * pitch + nx] + p[ny * pitch + nx]; - return (3 * this_col + next_col + (x % 2 ? 7 : 8)) >> 4; + if ( h == 2 && v == 2 ) { // weights 9 (nearest), 3, 3 and 1 (diagonal) + return (9 * p[cy * pitch + cx] + 3 * p[cy * pitch + nx] + 3 * p[ny * pitch + cx] + p[ny * pitch + nx] + + (x % 2 ? 7 : 8)) >> 4; } if ( h == 2 ) { return (3 * p[cy * pitch + cx] + p[cy * pitch + nx] + (x % 2 ? 2 : 1)) >> 2; @@ -79,13 +78,17 @@ chroma_ref_up(const uint8_t* p, const struct gpujpeg_component* comp, int x, int } /// @param interpolated whether the decoder interpolates this subsampling (precompiled 4:2:0, 4:2:2, 4:4:0 kernels) +/// @param filter option value to set, -1 to keep the default (on) static void chroma_case(struct gpujpeg_encoder* encoder, struct gpujpeg_decoder* decoder, const uint8_t* in, - gpujpeg_sampling_factor_t subsampling, bool interpolated, bool filter) + gpujpeg_sampling_factor_t subsampling, bool interpolated, int filter) { - const char* val = filter ? GPUJPEG_VAL_TRUE : GPUJPEG_VAL_FALSE; - CHROMA_CHECK(gpujpeg_encoder_set_option(encoder, GPUJPEG_ENC_OPT_CHROMA_FILTER_BOOL, val) == GPUJPEG_NOERR); - CHROMA_CHECK(gpujpeg_decoder_set_option(decoder, GPUJPEG_DEC_OPT_CHROMA_FILTER_BOOL, val) == GPUJPEG_NOERR); + if ( filter >= 0 ) { + const char* val = filter ? GPUJPEG_VAL_TRUE : GPUJPEG_VAL_FALSE; + CHROMA_CHECK(gpujpeg_encoder_set_option(encoder, GPUJPEG_ENC_OPT_CHROMA_FILTER_BOOL, val) == GPUJPEG_NOERR); + CHROMA_CHECK(gpujpeg_decoder_set_option(decoder, GPUJPEG_DEC_OPT_CHROMA_FILTER_BOOL, val) == GPUJPEG_NOERR); + } + filter = filter != 0; struct gpujpeg_parameters param = gpujpeg_default_parameters(); param.verbose = GPUJPEG_LL_QUIET; @@ -170,23 +173,15 @@ test_chroma_filter() {GPUJPEG_SUBSAMPLING_411, false}, // dynamic kernels: encoder averages, decoder replicates }; for ( unsigned i = 0; i < sizeof cases / sizeof cases[0]; ++i ) { - for ( int filter = 0; filter < 2; ++filter ) { - struct gpujpeg_encoder* encoder = gpujpeg_encoder_create(0); - struct gpujpeg_decoder* decoder = chroma_decoder_create(); - assert(encoder != NULL && decoder != NULL); + struct gpujpeg_encoder* encoder = gpujpeg_encoder_create(0); + struct gpujpeg_decoder* decoder = chroma_decoder_create(); + assert(encoder != NULL && decoder != NULL); + // default (on), then the option changed between calls without a reinitialization + for ( int filter = -1; filter < 2; ++filter ) { chroma_case(encoder, decoder, in, cases[i].subsampling, cases[i].interpolated, filter); - gpujpeg_decoder_destroy(decoder); - gpujpeg_encoder_destroy(encoder); } + gpujpeg_decoder_destroy(decoder); + gpujpeg_encoder_destroy(encoder); } - // the option applies to the next call without a reinitialization - struct gpujpeg_encoder* encoder = gpujpeg_encoder_create(0); - struct gpujpeg_decoder* decoder = chroma_decoder_create(); - assert(encoder != NULL && decoder != NULL); - chroma_case(encoder, decoder, in, GPUJPEG_SUBSAMPLING_420, true, true); - chroma_case(encoder, decoder, in, GPUJPEG_SUBSAMPLING_420, true, false); - chroma_case(encoder, decoder, in, GPUJPEG_SUBSAMPLING_420, true, true); - gpujpeg_decoder_destroy(decoder); - gpujpeg_encoder_destroy(encoder); printf("Ok\n"); } From 97f0a948216c8b9eaed8276215e38e7a4aa5cb2d Mon Sep 17 00:00:00 2001 From: adam-ce <5292991+adam-ce@users.noreply.github.com> Date: Tue, 3 Nov 2026 05:54:18 +0100 Subject: [PATCH 5/5] NEWS: chroma filters --- NEWS.md | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/NEWS.md b/NEWS.md index e4523783..1c1c9954 100644 --- a/NEWS.md +++ b/NEWS.md @@ -1,3 +1,12 @@ +unreleased +---------- + +- the encoder averages subsampled chroma (was point sampling) and the +decoder interpolates 4:2:0, 4:2:2 and 4:4:0 chroma like libjpeg for +full-resolution output (was replicated; also for chroma <= 2 samples wide, +unlike libjpeg) - output changes, the previous behavior can be selected by +GPUJPEG_ENC_OPT_CHROMA_FILTER_BOOL and GPUJPEG_DEC_OPT_CHROMA_FILTER_BOOL + 2026-03-19 - 0.27.13 ----------