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 ---------- 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/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_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_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_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.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..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; }; @@ -61,6 +63,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 +73,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; 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..53e4d2d8 --- /dev/null +++ b/test/unit/test_chroma_filter.c @@ -0,0 +1,187 @@ +#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 ) { // 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; + } + 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) +/// @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, int filter) +{ + 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; + 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 ) { + 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); + } + printf("Ok\n"); +}