diff --git a/src/gpujpeg_common.c b/src/gpujpeg_common.c index eef236c2..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 || @@ -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; 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"); 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..f120bff0 --- /dev/null +++ b/test/unit/test_coder_reinit.c @@ -0,0 +1,121 @@ +#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* 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; + 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); + } + gpujpeg_decoder_destroy(decoder); + gpujpeg_encoder_destroy(encoder); + printf("Ok\n"); +}