From 171d841b1d1154ee43656755745952e24a072fc1 Mon Sep 17 00:00:00 2001 From: adam-ce <5292991+adam-ce@users.noreply.github.com> Date: Mon, 2 Nov 2026 22:04:07 +0100 Subject: [PATCH 1/4] fix CPU Huffman buffer overflow after reconfiguration The host buffer for CPU Huffman coding was allocated with the current data_size but gpujpeg_coder_init_image() keeps it as long as the image fits data_allocated_size (the device buffer). After switching to a layout with larger data_size that still fits (eg. 4:2:0 -> 4:4:4 after a bigger image was coded), the CPU Huffman decoder wrote past the buffer and the host<->device copies exceeded it (encoder with restart interval 0, decoder with < 32 segments, eg. any JPEG without restart markers). --- src/gpujpeg_common.c | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/gpujpeg_common.c b/src/gpujpeg_common.c index eef236c2..c9c1ce82 100644 --- a/src/gpujpeg_common.c +++ b/src/gpujpeg_common.c @@ -641,7 +641,8 @@ reset_timers(struct gpujpeg_coder* coder) int gpujpeg_coder_allocate_cpu_huffman_buf(struct gpujpeg_coder * coder) { - cudaMallocHost((void**)&coder->data_quantized, coder->data_size * sizeof(int16_t)); + // sized like d_data_quantized, gpujpeg_coder_init_image() keeps both while the image fits + cudaMallocHost((void**)&coder->data_quantized, coder->data_allocated_size * sizeof(int16_t)); gpujpeg_cuda_check_error("Coder quantized data host allocation", return -1); int16_t* comp_data_quantized = coder->data_quantized; From c10450258f8afa951ca364ad8c0793fa92b6c2b0 Mon Sep 17 00:00:00 2001 From: adam-ce <5292991+adam-ce@users.noreply.github.com> Date: Mon, 2 Nov 2026 22:04:18 +0100 Subject: [PATCH 2/4] encoder: reinit preprocessor on JPEG layout change The preprocessor (kernel selection and component buffer pointers) was reinitialized only when the input image parameters changed. Changing eg. only the subsampling for the same input produced a wrong JPEG or even wrote to the freed device buffers (if they were reallocated). --- src/gpujpeg_common.c | 2 +- src/gpujpeg_common_internal.h | 6 ++++++ src/gpujpeg_encoder.c | 4 +++- 3 files changed, 10 insertions(+), 2 deletions(-) diff --git a/src/gpujpeg_common.c b/src/gpujpeg_common.c index c9c1ce82..ca275ae7 100644 --- a/src/gpujpeg_common.c +++ b/src/gpujpeg_common.c @@ -373,7 +373,7 @@ gpujpeg_parameters_chroma_subsampling(struct gpujpeg_parameters* param, * @ref gpujpeg_parameters members verbose, perf_stats and **quality** * ignored by the comparison */ -static bool +bool gpujpeg_parameters_equals(const struct gpujpeg_parameters *p1 , const struct gpujpeg_parameters *p2) { if ( p1->comp_count != p2->comp_count || diff --git a/src/gpujpeg_common_internal.h b/src/gpujpeg_common_internal.h index 2ecee77c..001c4919 100644 --- a/src/gpujpeg_common_internal.h +++ b/src/gpujpeg_common_internal.h @@ -515,6 +515,12 @@ gpujpeg_pixel_format_is_interleaved(enum gpujpeg_pixel_format pixel_format); bool gpujpeg_image_parameters_equals(const struct gpujpeg_image_parameters *p1 , const struct gpujpeg_image_parameters *p2); +/** + * @retval true parameters are the same (verbose, perf_stats and quality ignored) + */ +bool +gpujpeg_parameters_equals(const struct gpujpeg_parameters* p1, const struct gpujpeg_parameters* p2); + /** * returns difference between specified CUDA events * diff --git a/src/gpujpeg_encoder.c b/src/gpujpeg_encoder.c index 212eaa3e..1ffec13c 100644 --- a/src/gpujpeg_encoder.c +++ b/src/gpujpeg_encoder.c @@ -378,6 +378,8 @@ gpujpeg_encoder_encode(struct gpujpeg_encoder* encoder, const struct gpujpeg_par } gpujpeg_cuda_check_error("Quantization init", return -1); } + // preprocessor kernel and component buffers depend also on the JPEG layout (eg. subsampling) + const bool layout_changed = !gpujpeg_parameters_equals(&coder->param, ¶m_adjusted); if ( 0 == gpujpeg_coder_init_image(coder, ¶m_adjusted, param_image, coder->stream) ) { fprintf(stderr, "[GPUJPEG] [Error] Failed to init image encoding!\n"); return -1; @@ -389,7 +391,7 @@ gpujpeg_encoder_encode(struct gpujpeg_encoder* encoder, const struct gpujpeg_par return -1; } - if ( img_changed ) { + if ( img_changed || layout_changed ) { // (Re)initialize preprocessor if ( gpujpeg_preprocessor_encoder_init(&encoder->coder) != 0 ) { fprintf(stderr, "[GPUJPEG] [Error] Failed to init preprocessor!\n"); From 0db0e90c0c75380a5f0b4a602b5605aa39bd60bb Mon Sep 17 00:00:00 2001 From: adam-ce <5292991+adam-ce@users.noreply.github.com> Date: Mon, 2 Nov 2026 22:05:41 +0100 Subject: [PATCH 3/4] unit tests: coder reuse across layout changes Reuses one encoder and one decoder for a sequence of layouts (subsampling, restart interval -> GPU/CPU Huffman coder) and an output format change and compares the results with freshly created coders. --- test/unit/run_tests.c | 5 ++ test/unit/test_coder_reinit.c | 125 ++++++++++++++++++++++++++++++++++ 2 files changed, 130 insertions(+) create mode 100644 test/unit/test_coder_reinit.c diff --git a/test/unit/run_tests.c b/test/unit/run_tests.c index c5e1d356..1d00643f 100644 --- a/test/unit/run_tests.c +++ b/test/unit/run_tests.c @@ -88,6 +88,10 @@ test_gh_95(); void test_chroma_filter(); +// defined in test_coder_reinit.c +void +test_coder_reinit(); + static char* test_format_number_with_delim() { struct { @@ -132,6 +136,7 @@ int main() { test_gh_95(); test_format_number_with_delim(); test_chroma_filter(); + test_coder_reinit(); printf("PASSED\n"); } diff --git a/test/unit/test_coder_reinit.c b/test/unit/test_coder_reinit.c new file mode 100644 index 00000000..a6d4fb0a --- /dev/null +++ b/test/unit/test_coder_reinit.c @@ -0,0 +1,125 @@ +#include +#include +#include +#include +#include +#include + +/* + * Reuse one encoder and one decoder across JPEG layout and output format + * changes and compare every result with freshly created coders. + */ + +enum { + REINIT_W = 256, + REINIT_H = 256, +}; + +struct reinit_step +{ + gpujpeg_sampling_factor_t subsampling; + int restart_interval; + enum gpujpeg_color_space out_color_space; + enum gpujpeg_pixel_format out_pixel_format; + bool reuse_jpeg; ///< decode the previous JPEG again (only the output format changes) +}; + +static int +reinit_encode(struct gpujpeg_encoder* encoder, uint8_t* image, const struct reinit_step* step, uint8_t** out, + size_t* out_len) +{ + struct gpujpeg_parameters param = gpujpeg_default_parameters(); + param.verbose = GPUJPEG_LL_QUIET; + param.interleaved = 1; + param.restart_interval = step->restart_interval; + gpujpeg_parameters_chroma_subsampling(¶m, step->subsampling); + struct gpujpeg_image_parameters param_image = gpujpeg_default_image_parameters(); + param_image.width = REINIT_W; + param_image.height = REINIT_H; + struct gpujpeg_encoder_input input = gpujpeg_encoder_input_image(image); + return gpujpeg_encoder_encode(encoder, ¶m, ¶m_image, &input, out, out_len); +} + +static int +reinit_decode(struct gpujpeg_decoder* decoder, uint8_t* jpeg, size_t len, const struct reinit_step* step, + struct gpujpeg_decoder_output* output) +{ + gpujpeg_decoder_set_output_format(decoder, step->out_color_space, step->out_pixel_format); + gpujpeg_decoder_output_set_default(output); + return gpujpeg_decoder_decode(decoder, jpeg, len, output); +} + +static struct gpujpeg_decoder* +reinit_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); +} + +static void +reinit_check(bool ok, unsigned step, const char* what) +{ + if ( !ok ) { + fprintf(stderr, "step %u: reused %s differs from a fresh one!\n", step, what); + abort(); + } +} + +void +test_coder_reinit(); + +void +test_coder_reinit() +{ + printf("testing %s: ", __func__); + if ( getenv("NO_CUDA_DEV") ) { + printf("--\n"); + return; + } + static uint8_t image[REINIT_W * REINIT_H * 3]; + for ( unsigned i = 0; i < sizeof image; ++i ) { + image[i] = (uint8_t)(i / 3 % REINIT_W + i / (3 * REINIT_W) * 2 + ((i * 2654435761U) >> 28)); + } + const struct reinit_step steps[] = { + {GPUJPEG_SUBSAMPLING_444, 8, GPUJPEG_RGB, GPUJPEG_444_U8_P012, false}, // GPU Huffman coder + {GPUJPEG_SUBSAMPLING_420, 0, GPUJPEG_RGB, GPUJPEG_444_U8_P012, false}, // encoder subsampling change + {GPUJPEG_SUBSAMPLING_444, 0, GPUJPEG_RGB, GPUJPEG_444_U8_P012, false}, // CPU Huffman buffer growth + {GPUJPEG_SUBSAMPLING_420, 8, GPUJPEG_RGB, GPUJPEG_444_U8_P012, false}, + {GPUJPEG_SUBSAMPLING_420, 8, GPUJPEG_YCBCR_JPEG, GPUJPEG_420_U8_P0P1P2, true }, // output format change only + }; + struct gpujpeg_encoder* encoder = gpujpeg_encoder_create(0); + struct gpujpeg_decoder* decoder = reinit_decoder_create(); + assert(encoder != NULL && decoder != NULL); + uint8_t* jpeg = NULL; + size_t jpeg_len = 0; + for ( unsigned i = 0; i < sizeof steps / sizeof steps[0]; ++i ) { + const struct reinit_step* step = &steps[i]; + if ( !step->reuse_jpeg ) { + struct gpujpeg_encoder* fresh = gpujpeg_encoder_create(0); + uint8_t* out = NULL; + size_t out_len = 0; + reinit_check(reinit_encode(fresh, image, step, &out, &out_len) == 0, i, "encoder (fresh failed)"); + free(jpeg); + jpeg = malloc(out_len); + memcpy(jpeg, out, out_len); + jpeg_len = out_len; + gpujpeg_encoder_destroy(fresh); + reinit_check(reinit_encode(encoder, image, step, &out, &out_len) == 0 && out_len == jpeg_len && + memcmp(out, jpeg, jpeg_len) == 0, + i, "encoder"); + } + struct gpujpeg_decoder* fresh = reinit_decoder_create(); + struct gpujpeg_decoder_output fresh_out; + struct gpujpeg_decoder_output out; + reinit_check(reinit_decode(fresh, jpeg, jpeg_len, step, &fresh_out) == 0, i, "decoder (fresh failed)"); + reinit_check(reinit_decode(decoder, jpeg, jpeg_len, step, &out) == 0 && out.data_size == fresh_out.data_size && + memcmp(out.data, fresh_out.data, out.data_size) == 0, + i, "decoder"); + gpujpeg_decoder_destroy(fresh); + } + free(jpeg); + gpujpeg_decoder_destroy(decoder); + gpujpeg_encoder_destroy(encoder); + printf("Ok\n"); +} From ca24073927f0ed53f2e893accda5f7388b41c73c Mon Sep 17 00:00:00 2001 From: adam-ce <5292991+adam-ce@users.noreply.github.com> Date: Tue, 3 Nov 2026 05:55:14 +0100 Subject: [PATCH 4/4] unit tests: coder reinit - no JPEG copy Compare while the fresh encoder lives and decode the reused encoder's output, which stays valid until its next encode. --- test/unit/test_coder_reinit.c | 18 +++++++----------- 1 file changed, 7 insertions(+), 11 deletions(-) diff --git a/test/unit/test_coder_reinit.c b/test/unit/test_coder_reinit.c index a6d4fb0a..f120bff0 100644 --- a/test/unit/test_coder_reinit.c +++ b/test/unit/test_coder_reinit.c @@ -97,17 +97,14 @@ test_coder_reinit() const struct reinit_step* step = &steps[i]; if ( !step->reuse_jpeg ) { struct gpujpeg_encoder* fresh = gpujpeg_encoder_create(0); - uint8_t* out = NULL; - size_t out_len = 0; - reinit_check(reinit_encode(fresh, image, step, &out, &out_len) == 0, i, "encoder (fresh failed)"); - free(jpeg); - jpeg = malloc(out_len); - memcpy(jpeg, out, out_len); - jpeg_len = out_len; - gpujpeg_encoder_destroy(fresh); - reinit_check(reinit_encode(encoder, image, step, &out, &out_len) == 0 && out_len == jpeg_len && - memcmp(out, jpeg, jpeg_len) == 0, + uint8_t* fresh_jpeg = NULL; + size_t fresh_len = 0; + reinit_check(reinit_encode(fresh, image, step, &fresh_jpeg, &fresh_len) == 0, i, "encoder (fresh failed)"); + // jpeg (owned by the reused encoder) stays valid until its next encode + reinit_check(reinit_encode(encoder, image, step, &jpeg, &jpeg_len) == 0 && jpeg_len == fresh_len && + memcmp(jpeg, fresh_jpeg, jpeg_len) == 0, i, "encoder"); + gpujpeg_encoder_destroy(fresh); } struct gpujpeg_decoder* fresh = reinit_decoder_create(); struct gpujpeg_decoder_output fresh_out; @@ -118,7 +115,6 @@ test_coder_reinit() i, "decoder"); gpujpeg_decoder_destroy(fresh); } - free(jpeg); gpujpeg_decoder_destroy(decoder); gpujpeg_encoder_destroy(encoder); printf("Ok\n");