diff --git a/projects/rocjpeg/CMakeLists.txt b/projects/rocjpeg/CMakeLists.txt index d0d89846f5..8bff63542f 100644 --- a/projects/rocjpeg/CMakeLists.txt +++ b/projects/rocjpeg/CMakeLists.txt @@ -23,7 +23,7 @@ cmake_minimum_required (VERSION 3.5) -set(VERSION "0.5.0") +set(VERSION "0.6.0") set(CMAKE_CXX_STANDARD 17) # Set Project Version and Language diff --git a/projects/rocjpeg/api/rocjpeg.h b/projects/rocjpeg/api/rocjpeg.h index 8142c15975..9bc012a683 100644 --- a/projects/rocjpeg/api/rocjpeg.h +++ b/projects/rocjpeg/api/rocjpeg.h @@ -157,7 +157,7 @@ typedef struct { int16_t top; /**< Top coordinate of the crop rectangle. */ int16_t right; /**< Right coordinate of the crop rectangle. */ int16_t bottom; /**< Bottom coordinate of the crop rectangle. */ - } crop_rectangle; /**< (future use) Defines the region of interest (ROI) to be copied into the RocJpegImage output buffers. */ + } crop_rectangle; /**< Defines the region of interest (ROI) to be copied into the RocJpegImage output buffers. */ struct { uint32_t width; /**< Target width of the picture to be resized. */ uint32_t height; /**< Target height of the picture to be resized. */ diff --git a/projects/rocjpeg/api/rocjpeg_version.h b/projects/rocjpeg/api/rocjpeg_version.h index 916998d4f8..8096cb1906 100644 --- a/projects/rocjpeg/api/rocjpeg_version.h +++ b/projects/rocjpeg/api/rocjpeg_version.h @@ -34,7 +34,7 @@ THE SOFTWARE. extern "C" { #endif #define ROCJPEG_MAJOR_VERSION 0 -#define ROCJPEG_MINOR_VERSION 5 +#define ROCJPEG_MINOR_VERSION 6 #define ROCJPEG_MICRO_VERSION 0 diff --git a/projects/rocjpeg/samples/CMakeLists.txt b/projects/rocjpeg/samples/CMakeLists.txt index 6556344989..700b761773 100644 --- a/projects/rocjpeg/samples/CMakeLists.txt +++ b/projects/rocjpeg/samples/CMakeLists.txt @@ -104,4 +104,76 @@ add_test( --build-generator "${CMAKE_GENERATOR}" --test-command "jpegdecodebatched" -i ${CMAKE_SOURCE_DIR}/data/images/ +) + +add_test( + NAME + jpeg-decode-crop-fmt-native + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${CMAKE_CURRENT_SOURCE_DIR}/jpegDecode" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecode" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecode" + -i ${CMAKE_SOURCE_DIR}/data/images/ -crop 960,540,2880,1620 +) + +add_test( + NAME + jpeg-decode-crop-fmt-yuv-planar + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${CMAKE_CURRENT_SOURCE_DIR}/jpegDecode" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecode" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecode" + -i ${CMAKE_SOURCE_DIR}/data/images/ -fmt yuv -crop 960,540,2880,1620 +) + +add_test( + NAME + jpeg-decode-crop-fmt-y + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${CMAKE_CURRENT_SOURCE_DIR}/jpegDecode" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecode" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecode" + -i ${CMAKE_SOURCE_DIR}/data/images/ -fmt y -crop 960,540,2880,1620 +) + +add_test( + NAME + jpeg-decode-crop-fmt-rgb + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${CMAKE_CURRENT_SOURCE_DIR}/jpegDecode" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecode" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecode" + -i ${CMAKE_SOURCE_DIR}/data/images/ -fmt rgb -crop 960,540,2880,1620 +) + +add_test( + NAME + jpeg-decode-crop-fmt-rgb-planar + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${CMAKE_CURRENT_SOURCE_DIR}/jpegDecode" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecode" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecode" + -i ${CMAKE_SOURCE_DIR}/data/images/ -fmt rgb_planar -crop 960,540,2880,1620 +) + +add_test( + NAME + jpeg-decode-crop-batch-fmt-native + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${CMAKE_CURRENT_SOURCE_DIR}/jpegDecodeBatched" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecodeBatched" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecodebatched" + -i ${CMAKE_SOURCE_DIR}/data/images/ -crop 960,540,2880,1620 ) \ No newline at end of file diff --git a/projects/rocjpeg/samples/jpegDecode/jpegdecode.cpp b/projects/rocjpeg/samples/jpegDecode/jpegdecode.cpp index 7382d3006c..e64ce4ec02 100644 --- a/projects/rocjpeg/samples/jpegDecode/jpegdecode.cpp +++ b/projects/rocjpeg/samples/jpegDecode/jpegdecode.cpp @@ -49,6 +49,13 @@ int main(int argc, char **argv) { RocJpegUtils rocjpeg_utils; RocJpegUtils::ParseCommandLine(input_path, output_file_path, save_images, device_id, rocjpeg_backend, decode_params, nullptr, nullptr, argc, argv); + + bool is_roi_valid = false; + uint32_t roi_width; + uint32_t roi_height; + roi_width = decode_params.crop_rectangle.right - decode_params.crop_rectangle.left; + roi_height = decode_params.crop_rectangle.bottom - decode_params.crop_rectangle.top; + if (!RocJpegUtils::GetFilePaths(input_path, file_paths, is_dir, is_file)) { std::cerr << "ERROR: Failed to get input file paths!" << std::endl; return EXIT_FAILURE; @@ -87,6 +94,10 @@ int main(int argc, char **argv) { CHECK_ROCJPEG(rocJpegStreamParse(reinterpret_cast(file_data.data()), file_size, rocjpeg_stream_handle)); CHECK_ROCJPEG(rocJpegGetImageInfo(rocjpeg_handle, rocjpeg_stream_handle, &num_components, &subsampling, widths, heights)); + if (roi_width > 0 && roi_height > 0 && roi_width <= widths[0] && roi_height <= heights[0]) { + is_roi_valid = true; + } + rocjpeg_utils.GetChromaSubsamplingStr(subsampling, chroma_sub_sampling); std::cout << "Input file name: " << base_file_name << std::endl; std::cout << "Input image resolution: " << widths[0] << "x" << heights[0] << std::endl; @@ -124,6 +135,9 @@ int main(int argc, char **argv) { } } + if (is_roi_valid) { + std::cout << "Cropped image resolution: " << roi_width << "x" << roi_height << std::endl; + } std::cout << "Decoding started, please wait! ... " << std::endl; auto start_time = std::chrono::high_resolution_clock::now(); CHECK_ROCJPEG(rocJpegDecode(rocjpeg_handle, rocjpeg_stream_handle, &decode_params, &output_image)); @@ -134,10 +148,13 @@ int main(int argc, char **argv) { if (save_images) { std::string image_save_path = output_file_path; + //if ROI is present, need to pass roi_width and roi_height + uint32_t width = is_roi_valid ? roi_width : widths[0]; + uint32_t height = is_roi_valid ? roi_height : heights[0]; if (is_dir) { - rocjpeg_utils.GetOutputFileExt(decode_params.output_format, base_file_name, widths[0], heights[0], subsampling, image_save_path); + rocjpeg_utils.GetOutputFileExt(decode_params.output_format, base_file_name, width, height, subsampling, image_save_path); } - rocjpeg_utils.SaveImage(image_save_path, &output_image, widths[0], heights[0], subsampling, decode_params.output_format); + rocjpeg_utils.SaveImage(image_save_path, &output_image, width, height, subsampling, decode_params.output_format); } std::cout << "Average processing time per image (ms): " << time_per_image_in_milli_sec << std::endl; diff --git a/projects/rocjpeg/samples/jpegDecodeBatched/jpegdecodebatched.cpp b/projects/rocjpeg/samples/jpegDecodeBatched/jpegdecodebatched.cpp index d2c192da6b..dae50032bb 100644 --- a/projects/rocjpeg/samples/jpegDecodeBatched/jpegdecodebatched.cpp +++ b/projects/rocjpeg/samples/jpegDecodeBatched/jpegdecodebatched.cpp @@ -52,6 +52,13 @@ int main(int argc, char **argv) { RocJpegUtils rocjpeg_utils; RocJpegUtils::ParseCommandLine(input_path, output_file_path, save_images, device_id, rocjpeg_backend, decode_params, nullptr, &batch_size, argc, argv); + + bool is_roi_valid = false; + uint32_t roi_width; + uint32_t roi_height; + roi_width = decode_params.crop_rectangle.right - decode_params.crop_rectangle.left; + roi_height = decode_params.crop_rectangle.bottom - decode_params.crop_rectangle.top; + if (!RocJpegUtils::GetFilePaths(input_path, file_paths, is_dir, is_file)) { std::cerr << "ERROR: Failed to get input file paths!" << std::endl; return EXIT_FAILURE; @@ -103,6 +110,10 @@ int main(int argc, char **argv) { CHECK_ROCJPEG(rocJpegStreamParse(reinterpret_cast(batch_images[index].data()), file_size, rocjpeg_stream_handles[index])); CHECK_ROCJPEG(rocJpegGetImageInfo(rocjpeg_handle, rocjpeg_stream_handles[index], &num_components, &subsamplings[index], widths[index].data(), heights[index].data())); + if (roi_width > 0 && roi_height > 0 && roi_width <= widths[index][0] && roi_height <= heights[index][0]) { + is_roi_valid = true; + } + rocjpeg_utils.GetChromaSubsamplingStr(subsamplings[index], chroma_sub_sampling); if (widths[index][0] < 64 || heights[index][0] < 64) { std::cerr << "The image resolution is not supported by VCN Hardware" << std::endl; @@ -156,10 +167,13 @@ int main(int argc, char **argv) { if (save_images) { for (int b = 0; b < current_batch_size; b++) { std::string image_save_path = output_file_path; + //if ROI is present, need to pass roi_width and roi_height + uint32_t width = is_roi_valid ? roi_width : widths[b][0]; + uint32_t height = is_roi_valid ? roi_height : heights[b][0]; if (is_dir) { - rocjpeg_utils.GetOutputFileExt(decode_params.output_format, base_file_names[b], widths[b][0], heights[b][0], subsamplings[b], image_save_path); + rocjpeg_utils.GetOutputFileExt(decode_params.output_format, base_file_names[b], width, height, subsamplings[b], image_save_path); } - rocjpeg_utils.SaveImage(image_save_path, &output_images[b], widths[b][0], heights[b][0], subsamplings[b], decode_params.output_format); + rocjpeg_utils.SaveImage(image_save_path, &output_images[b], width, height, subsamplings[b], decode_params.output_format); } } diff --git a/projects/rocjpeg/samples/jpegDecodeMultiThreads/jpegdecodemultithreads.cpp b/projects/rocjpeg/samples/jpegDecodeMultiThreads/jpegdecodemultithreads.cpp index 4f19aab7fe..77287cb32f 100644 --- a/projects/rocjpeg/samples/jpegDecodeMultiThreads/jpegdecodemultithreads.cpp +++ b/projects/rocjpeg/samples/jpegDecodeMultiThreads/jpegdecodemultithreads.cpp @@ -25,6 +25,12 @@ THE SOFTWARE. void ThreadFunction(std::vector& jpegFiles, RocJpegHandle rocjpeg_handle, RocJpegStreamHandle rocjpeg_stream, RocJpegUtils rocjpeg_util, RocJpegImage *output_image, std::mutex &mutex, RocJpegDecodeParams &decode_params, bool save_images, std::string &output_file_path, uint64_t *num_decoded_images, double *image_size_in_mpixels) { + bool is_roi_valid = false; + uint32_t roi_width; + uint32_t roi_height; + roi_width = decode_params.crop_rectangle.right - decode_params.crop_rectangle.left; + roi_height = decode_params.crop_rectangle.bottom - decode_params.crop_rectangle.top; + std::vector file_data; uint8_t num_components; uint32_t widths[ROCJPEG_MAX_COMPONENT] = {}; @@ -71,6 +77,10 @@ void ThreadFunction(std::vector& jpegFiles, RocJpegHandle rocjpeg_h CHECK_ROCJPEG(rocJpegStreamParse(reinterpret_cast(file_data.data()), file_size, rocjpeg_stream)); CHECK_ROCJPEG(rocJpegGetImageInfo(rocjpeg_handle, rocjpeg_stream, &num_components, &subsampling, widths, heights)); + if (roi_width > 0 && roi_height > 0 && roi_width <= widths[0] && roi_height <= heights[0]) { + is_roi_valid = true; + } + if (widths[0] < 64 || heights[0] < 64) { std::cerr << "The image resolution is not supported by VCN Hardware" << std::endl; std::cout << "Skipping decoding file " << base_file_name << std::endl; @@ -105,8 +115,11 @@ void ThreadFunction(std::vector& jpegFiles, RocJpegHandle rocjpeg_h if (save_images) { std::string image_save_path = output_file_path; - rocjpeg_util.GetOutputFileExt(decode_params.output_format, base_file_name, widths[0], heights[0], subsampling, image_save_path); - rocjpeg_util.SaveImage(image_save_path, output_image, widths[0], heights[0], subsampling, decode_params.output_format); + //if ROI is present, need to pass roi_width and roi_height + uint32_t width = is_roi_valid ? roi_width : widths[0]; + uint32_t height = is_roi_valid ? roi_height : heights[0]; + rocjpeg_util.GetOutputFileExt(decode_params.output_format, base_file_name, width, height, subsampling, image_save_path); + rocjpeg_util.SaveImage(image_save_path, output_image, width, height, subsampling, decode_params.output_format); } for (int i = 0; i < ROCJPEG_MAX_COMPONENT; i++) { diff --git a/projects/rocjpeg/src/rocjpeg_decoder.cpp b/projects/rocjpeg/src/rocjpeg_decoder.cpp index 57979e14e6..6a8bb9fbdf 100644 --- a/projects/rocjpeg/src/rocjpeg_decoder.cpp +++ b/projects/rocjpeg/src/rocjpeg_decoder.cpp @@ -118,46 +118,61 @@ RocJpegStatus RocJpegDecoder::Decode(RocJpegStreamHandle jpeg_stream_handle, con CHECK_ROCJPEG(jpeg_vaapi_decoder_.GetHipInteropMem(current_surface_id, hip_interop_dev_mem)); uint16_t chroma_height = 0; + uint16_t picture_width = 0; + uint16_t picture_height = 0; + bool is_roi_valid = false; + uint32_t roi_width; + uint32_t roi_height; + roi_width = decode_params->crop_rectangle.right - decode_params->crop_rectangle.left; + roi_height = decode_params->crop_rectangle.bottom - decode_params->crop_rectangle.top; + + if (roi_width > 0 && roi_height > 0 && roi_width <= jpeg_stream_params->picture_parameter_buffer.picture_width && roi_height <= jpeg_stream_params->picture_parameter_buffer.picture_height) { + is_roi_valid = true; + } + picture_width = is_roi_valid ? roi_width : jpeg_stream_params->picture_parameter_buffer.picture_width; + picture_height = is_roi_valid ? roi_height : jpeg_stream_params->picture_parameter_buffer.picture_height; + switch (decode_params->output_format) { case ROCJPEG_OUTPUT_NATIVE: // Copy the native decoded output buffers from interop memory directly to the destination buffers - CHECK_ROCJPEG(GetChromaHeight(hip_interop_dev_mem.surface_format, jpeg_stream_params->picture_parameter_buffer.picture_height, chroma_height)); + CHECK_ROCJPEG(GetChromaHeight(hip_interop_dev_mem.surface_format, picture_height, chroma_height)); + // Copy Luma (first channel) for any surface format - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, jpeg_stream_params->picture_parameter_buffer.picture_height, 0, destination)); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_height, 0, destination, decode_params, is_roi_valid)); + if (hip_interop_dev_mem.surface_format == VA_FOURCC_NV12) { // Copy the second channel (UV interleaved) for NV12 - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 1, destination)); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 1, destination, decode_params, is_roi_valid)); } else if (hip_interop_dev_mem.surface_format == VA_FOURCC_444P || hip_interop_dev_mem.surface_format == VA_FOURCC_422V) { // Copy the second and third channels for YUV444 and YUV440 (i.e., YUV422V) - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 1, destination)); - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 2, destination)); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 1, destination, decode_params, is_roi_valid)); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 2, destination, decode_params, is_roi_valid)); } break; case ROCJPEG_OUTPUT_YUV_PLANAR: - CHECK_ROCJPEG(GetChromaHeight(hip_interop_dev_mem.surface_format, jpeg_stream_params->picture_parameter_buffer.picture_height, chroma_height)); - CHECK_ROCJPEG(GetPlanarYUVOutputFormat(hip_interop_dev_mem, jpeg_stream_params->picture_parameter_buffer.picture_width, - jpeg_stream_params->picture_parameter_buffer.picture_height, chroma_height, destination)); + CHECK_ROCJPEG(GetChromaHeight(hip_interop_dev_mem.surface_format, picture_height, chroma_height)); + CHECK_ROCJPEG(GetPlanarYUVOutputFormat(hip_interop_dev_mem, picture_width, + picture_height, chroma_height, destination, decode_params, is_roi_valid)); break; case ROCJPEG_OUTPUT_Y: - CHECK_ROCJPEG(GetYOutputFormat(hip_interop_dev_mem, jpeg_stream_params->picture_parameter_buffer.picture_width, - jpeg_stream_params->picture_parameter_buffer.picture_height, destination)); + CHECK_ROCJPEG(GetYOutputFormat(hip_interop_dev_mem, picture_width, + picture_height, destination, decode_params, is_roi_valid)); break; case ROCJPEG_OUTPUT_RGB: - CHECK_ROCJPEG(ColorConvertToRGB(hip_interop_dev_mem, jpeg_stream_params->picture_parameter_buffer.picture_width, - jpeg_stream_params->picture_parameter_buffer.picture_height, destination)); + CHECK_ROCJPEG(ColorConvertToRGB(hip_interop_dev_mem, picture_width, + picture_height, destination, decode_params, is_roi_valid)); break; case ROCJPEG_OUTPUT_RGB_PLANAR: - CHECK_ROCJPEG(ColorConvertToRGBPlanar(hip_interop_dev_mem, jpeg_stream_params->picture_parameter_buffer.picture_width, - jpeg_stream_params->picture_parameter_buffer.picture_height, destination)); + CHECK_ROCJPEG(ColorConvertToRGBPlanar(hip_interop_dev_mem, picture_width, + picture_height, destination, decode_params, is_roi_valid)); break; default: break; } CHECK_HIP(hipStreamSynchronize(hip_stream_)); - return ROCJPEG_STATUS_SUCCESS; } @@ -202,39 +217,53 @@ RocJpegStatus RocJpegDecoder::DecodeBatched(RocJpegStreamHandle *jpeg_streams, i CHECK_ROCJPEG(jpeg_vaapi_decoder_.GetHipInteropMem(current_surface_id, hip_interop_dev_mem)); uint16_t chroma_height = 0; + uint16_t picture_width = 0; + uint16_t picture_height = 0; + bool is_roi_valid = false; + uint32_t roi_width; + uint32_t roi_height; + roi_width = decode_params->crop_rectangle.right - decode_params->crop_rectangle.left; + roi_height = decode_params->crop_rectangle.bottom - decode_params->crop_rectangle.top; + + if (roi_width > 0 && roi_height > 0 && roi_width <= jpeg_stream_params->picture_parameter_buffer.picture_width && roi_height <= jpeg_stream_params->picture_parameter_buffer.picture_height) { + is_roi_valid = true; + } + + picture_width = is_roi_valid ? roi_width : jpeg_stream_params->picture_parameter_buffer.picture_width; + picture_height = is_roi_valid ? roi_height : jpeg_stream_params->picture_parameter_buffer.picture_height; switch (decode_params->output_format) { case ROCJPEG_OUTPUT_NATIVE: // Copy the native decoded output buffers from interop memory directly to the destination buffers - CHECK_ROCJPEG(GetChromaHeight(hip_interop_dev_mem.surface_format, jpeg_stream_params->picture_parameter_buffer.picture_height, chroma_height)); + CHECK_ROCJPEG(GetChromaHeight(hip_interop_dev_mem.surface_format, picture_height, chroma_height)); // Copy Luma (first channel) for any surface format - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, jpeg_stream_params->picture_parameter_buffer.picture_height, 0, &destinations[k + i])); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_height, 0, &destinations[k + i], decode_params, is_roi_valid)); if (hip_interop_dev_mem.surface_format == VA_FOURCC_NV12) { // Copy the second channel (UV interleaved) for NV12 - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 1, &destinations[k + i])); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 1, &destinations[k + i], decode_params, is_roi_valid)); } else if (hip_interop_dev_mem.surface_format == VA_FOURCC_444P || hip_interop_dev_mem.surface_format == VA_FOURCC_422V) { // Copy the second and third channels for YUV444 and YUV440 (i.e., YUV422V) - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 1, &destinations[k + i])); - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 2, &destinations[k + i])); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 1, &destinations[k + i], decode_params, is_roi_valid)); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 2, &destinations[k + i], decode_params, is_roi_valid)); } break; case ROCJPEG_OUTPUT_YUV_PLANAR: - CHECK_ROCJPEG(GetChromaHeight(hip_interop_dev_mem.surface_format, jpeg_stream_params->picture_parameter_buffer.picture_height, chroma_height)); - CHECK_ROCJPEG(GetPlanarYUVOutputFormat(hip_interop_dev_mem, jpeg_stream_params->picture_parameter_buffer.picture_width, - jpeg_stream_params->picture_parameter_buffer.picture_height, chroma_height, &destinations[k + i])); + CHECK_ROCJPEG(GetChromaHeight(hip_interop_dev_mem.surface_format, picture_height, chroma_height)); + CHECK_ROCJPEG(GetPlanarYUVOutputFormat(hip_interop_dev_mem, picture_width, + picture_height, chroma_height, &destinations[k + i], decode_params, is_roi_valid)); break; case ROCJPEG_OUTPUT_Y: - CHECK_ROCJPEG(GetYOutputFormat(hip_interop_dev_mem, jpeg_stream_params->picture_parameter_buffer.picture_width, - jpeg_stream_params->picture_parameter_buffer.picture_height, &destinations[k + i])); + CHECK_ROCJPEG(GetYOutputFormat(hip_interop_dev_mem, picture_width, + picture_height, &destinations[k + i], decode_params, is_roi_valid)); break; case ROCJPEG_OUTPUT_RGB: - CHECK_ROCJPEG(ColorConvertToRGB(hip_interop_dev_mem, jpeg_stream_params->picture_parameter_buffer.picture_width, - jpeg_stream_params->picture_parameter_buffer.picture_height, &destinations[k + i])); + CHECK_ROCJPEG(ColorConvertToRGB(hip_interop_dev_mem, picture_width, + picture_height, &destinations[k + i], decode_params, is_roi_valid)); break; case ROCJPEG_OUTPUT_RGB_PLANAR: - CHECK_ROCJPEG(ColorConvertToRGBPlanar(hip_interop_dev_mem, jpeg_stream_params->picture_parameter_buffer.picture_width, - jpeg_stream_params->picture_parameter_buffer.picture_height, &destinations[k + i])); + CHECK_ROCJPEG(ColorConvertToRGBPlanar(hip_interop_dev_mem, picture_width, + picture_height, &destinations[k + i], decode_params, is_roi_valid)); break; default: break; @@ -325,13 +354,29 @@ RocJpegStatus RocJpegDecoder::GetImageInfo(RocJpegStreamHandle jpeg_stream_handl * @param destination The `RocJpegImage` object representing the destination image. * @return The status of the operation. Returns `ROCJPEG_STATUS_SUCCESS` if the channel was copied successfully. */ -RocJpegStatus RocJpegDecoder::CopyChannel(HipInteropDeviceMem& hip_interop_dev_mem, uint16_t channel_height, uint8_t channel_index, RocJpegImage *destination) { +RocJpegStatus RocJpegDecoder::CopyChannel(HipInteropDeviceMem& hip_interop_dev_mem, uint16_t channel_height, uint8_t channel_index, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { if (hip_interop_dev_mem.pitch[channel_index] != 0 && destination->pitch[channel_index] != 0 && destination->channel[channel_index] != nullptr) { + uint32_t roi_offset = 0; + if (is_roi_valid) { + int16_t top = decode_params->crop_rectangle.top; + int16_t left = decode_params->crop_rectangle.left; + // adjustments need to be made for these 3 pixel formats + switch (hip_interop_dev_mem.surface_format) { + case VA_FOURCC_NV12: + case VA_FOURCC_422V: + top = (channel_index == 1 || channel_index == 2) ? top >> 1 : top; + break; + case ROCJPEG_FOURCC_YUYV: + left *= 2; + break; + } + roi_offset = top * hip_interop_dev_mem.pitch[channel_index] + left; + } if (destination->pitch[channel_index] == hip_interop_dev_mem.pitch[channel_index]) { uint32_t channel_size = destination->pitch[channel_index] * channel_height; - CHECK_HIP(hipMemcpyDtoDAsync(destination->channel[channel_index], hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[channel_index], channel_size, hip_stream_)); + CHECK_HIP(hipMemcpyDtoDAsync(destination->channel[channel_index], hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[channel_index] + roi_offset, channel_size, hip_stream_)); } else { - CHECK_HIP(hipMemcpy2DAsync(destination->channel[channel_index], destination->pitch[channel_index], hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[channel_index], hip_interop_dev_mem.pitch[channel_index], + CHECK_HIP(hipMemcpy2DAsync(destination->channel[channel_index], destination->pitch[channel_index], hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[channel_index] + roi_offset, hip_interop_dev_mem.pitch[channel_index], destination->pitch[channel_index], channel_height, hipMemcpyDeviceToDevice, hip_stream_)); } } @@ -381,38 +426,50 @@ RocJpegStatus RocJpegDecoder::GetChromaHeight(uint32_t surface_format, uint16_t * parameter. * * @param hip_interop_dev_mem The HipInteropDeviceMem object containing the input image data. - * @param picture_width The width of the input image. - * @param picture_height The height of the input image. + * @param picture_width The width of the destination image. + * @param picture_height The height of the destination image. * @param destination Pointer to the RocJpegImage object where the converted image will be stored. * @return The status of the color conversion operation. Returns ROCJPEG_STATUS_SUCCESS if the conversion * is successful. Returns ROCJPEG_STATUS_JPEG_NOT_SUPPORTED if the surface format is not supported. */ -RocJpegStatus RocJpegDecoder::ColorConvertToRGB(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination) { +RocJpegStatus RocJpegDecoder::ColorConvertToRGB(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { + uint32_t roi_offset = 0; + uint32_t roi_uv_offset = 0; + int16_t top = decode_params->crop_rectangle.top; + int16_t left = decode_params->crop_rectangle.left; + if (is_roi_valid) { + if (hip_interop_dev_mem.surface_format == VA_FOURCC_422V || hip_interop_dev_mem.surface_format == VA_FOURCC_NV12){ + roi_uv_offset = (top >> 1) * hip_interop_dev_mem.pitch[1] + left; + } else if (hip_interop_dev_mem.surface_format == ROCJPEG_FOURCC_YUYV) { + left *= 2; + } + roi_offset = top * hip_interop_dev_mem.pitch[0] + left; + } switch (hip_interop_dev_mem.surface_format) { case VA_FOURCC_444P: ColorConvertYUV444ToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1] + roi_offset, hip_interop_dev_mem.offset[2] + roi_offset); break; case VA_FOURCC_422V: ColorConvertYUV440ToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1], hip_interop_dev_mem.offset[2]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1] /*+ roi_uv_offset*/, hip_interop_dev_mem.offset[2] /*+ roi_uv_offset*/); break; case ROCJPEG_FOURCC_YUYV: ColorConvertYUYVToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0]); break; case VA_FOURCC_NV12: ColorConvertNV12ToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[1], hip_interop_dev_mem.pitch[1]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], + hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[1] + roi_uv_offset, hip_interop_dev_mem.pitch[1]); break; case VA_FOURCC_Y800: ColorConvertYUV400ToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0]); break; case VA_FOURCC_RGBA: ColorConvertRGBAToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0]); break; default: ERR("ERROR! surface format is not supported!"); @@ -429,40 +486,52 @@ RocJpegStatus RocJpegDecoder::ColorConvertToRGB(HipInteropDeviceMem& hip_interop * The converted image is stored in the `destination` RocJpegImage object. * * @param hip_interop_dev_mem The HipInteropDeviceMem object containing the input image data. - * @param picture_width The width of the input image. - * @param picture_height The height of the input image. + * @param picture_width The width of the destination image. + * @param picture_height The height of the destination image. * @param destination Pointer to the RocJpegImage object where the converted image will be stored. * @return RocJpegStatus The status of the color conversion operation. * Returns ROCJPEG_STATUS_SUCCESS if the conversion is successful. * Returns ROCJPEG_STATUS_JPEG_NOT_SUPPORTED if the surface format is not supported. */ -RocJpegStatus RocJpegDecoder::ColorConvertToRGBPlanar(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination) { +RocJpegStatus RocJpegDecoder::ColorConvertToRGBPlanar(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { + uint32_t roi_offset = 0; + uint32_t roi_uv_offset = 0; + int16_t top = decode_params->crop_rectangle.top; + int16_t left = decode_params->crop_rectangle.left; + if (is_roi_valid) { + if (hip_interop_dev_mem.surface_format == VA_FOURCC_422V || hip_interop_dev_mem.surface_format == VA_FOURCC_NV12){ + roi_uv_offset = (top >> 1) * hip_interop_dev_mem.pitch[1] + left; + } else if (hip_interop_dev_mem.surface_format == ROCJPEG_FOURCC_YUYV) { + left *= 2; + } + roi_offset = top * hip_interop_dev_mem.pitch[0] + left; + } switch (hip_interop_dev_mem.surface_format) { case VA_FOURCC_444P: ColorConvertYUV444ToRGBPlanar(hip_stream_, picture_width, picture_height, destination->channel[0], destination->channel[1], destination->channel[2], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1] + roi_offset, hip_interop_dev_mem.offset[2] + roi_offset); break; case VA_FOURCC_422V: ColorConvertYUV440ToRGBPlanar(hip_stream_, picture_width, picture_height, destination->channel[0], destination->channel[1], destination->channel[2], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1], hip_interop_dev_mem.offset[2]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1] /*+ roi_uv_offset*/, hip_interop_dev_mem.offset[2] /*+ roi_uv_offset*/); break; case ROCJPEG_FOURCC_YUYV: ColorConvertYUYVToRGBPlanar(hip_stream_, picture_width, picture_height, destination->channel[0], destination->channel[1], destination->channel[2], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0]); break; case VA_FOURCC_NV12: ColorConvertNV12ToRGBPlanar(hip_stream_, picture_width, picture_height, destination->channel[0], destination->channel[1], destination->channel[2], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[1], hip_interop_dev_mem.pitch[1]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], + hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[1] + roi_uv_offset, hip_interop_dev_mem.pitch[1]); break; case VA_FOURCC_Y800: ColorConvertYUV400ToRGBPlanar(hip_stream_, picture_width, picture_height, destination->channel[0], destination->channel[1], destination->channel[2], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0]); break; case VA_FOURCC_RGBP: // Copy red, green, and blue channels from the interop memory into the destination for (uint8_t channel_index = 0; channel_index < 3; channel_index++) { - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_height, channel_index, destination)); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_height, channel_index, destination, decode_params, is_roi_valid)); } break; default: @@ -489,22 +558,32 @@ RocJpegStatus RocJpegDecoder::ColorConvertToRGBPlanar(HipInteropDeviceMem& hip_i * @param destination Pointer to the RocJpegImage object where the converted image data will be stored. * @return The status of the operation. Returns ROCJPEG_STATUS_SUCCESS if successful. */ -RocJpegStatus RocJpegDecoder::GetPlanarYUVOutputFormat(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, uint16_t chroma_height, RocJpegImage *destination) { +RocJpegStatus RocJpegDecoder::GetPlanarYUVOutputFormat(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, uint16_t chroma_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { + uint32_t roi_offset = 0; + if (is_roi_valid) { + int16_t top = decode_params->crop_rectangle.top; + int16_t left = decode_params->crop_rectangle.left; + if (hip_interop_dev_mem.surface_format == VA_FOURCC_NV12){ + roi_offset = (top >> 1) * hip_interop_dev_mem.pitch[1] + left; + } else if (hip_interop_dev_mem.surface_format == ROCJPEG_FOURCC_YUYV) { + roi_offset = top * hip_interop_dev_mem.pitch[0] + (left * 2); + } + } if (hip_interop_dev_mem.surface_format == ROCJPEG_FOURCC_YUYV) { - // Extract the packed YUYV and copy them into the first, second, and thrid channels of the destination. + // Extract the packed YUYV and copy them into the first, second, and third channels of the destination. ConvertPackedYUYVToPlanarYUV(hip_stream_, picture_width, picture_height, destination->channel[0], destination->channel[1], destination->channel[2], - destination->pitch[0], destination->pitch[1], hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0]); + destination->pitch[0], destination->pitch[1], hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0]); } else { // Copy Luma - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_height, 0, destination)); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_height, 0, destination, decode_params, is_roi_valid)); if (hip_interop_dev_mem.surface_format == VA_FOURCC_NV12) { - // Extract the interleaved UV channels and copy them into the second and thrid channels of the destination. + // Extract the interleaved UV channels and copy them into the second and third channels of the destination. ConvertInterleavedUVToPlanarUV(hip_stream_, picture_width >> 1, picture_height >> 1, destination->channel[1], destination->channel[2], - destination->pitch[1], hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[1] , hip_interop_dev_mem.pitch[1]); + destination->pitch[1], hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[1] + roi_offset, hip_interop_dev_mem.pitch[1]); } else if (hip_interop_dev_mem.surface_format == VA_FOURCC_444P || hip_interop_dev_mem.surface_format == VA_FOURCC_422V) { - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 1, destination)); - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 2, destination)); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 1, destination, decode_params, is_roi_valid)); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, chroma_height, 2, destination, decode_params, is_roi_valid)); } } return ROCJPEG_STATUS_SUCCESS; @@ -523,13 +602,20 @@ RocJpegStatus RocJpegDecoder::GetPlanarYUVOutputFormat(HipInteropDeviceMem& hip_ * @param destination Pointer to the RocJpegImage object where the extracted Y component will be stored. * @return The status of the operation. Returns ROCJPEG_STATUS_SUCCESS if successful. */ -RocJpegStatus RocJpegDecoder::GetYOutputFormat(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination) { +RocJpegStatus RocJpegDecoder::GetYOutputFormat(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { + uint32_t roi_offset = 0; if (hip_interop_dev_mem.surface_format == ROCJPEG_FOURCC_YUYV) { + // calculate offset and add to hip_mapped_device_mem + if (is_roi_valid) { + int16_t top = decode_params->crop_rectangle.top; + int16_t left = decode_params->crop_rectangle.left * 2; + roi_offset = top * hip_interop_dev_mem.pitch[0] + left; + } ExtractYFromPackedYUYV(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem, hip_interop_dev_mem.pitch[0]); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0]); } else { // Copy Luma - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_height, 0, destination)); + CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_height, 0, destination, decode_params, is_roi_valid)); } return ROCJPEG_STATUS_SUCCESS; } \ No newline at end of file diff --git a/projects/rocjpeg/src/rocjpeg_decoder.h b/projects/rocjpeg/src/rocjpeg_decoder.h index 05692a656e..443fae6729 100644 --- a/projects/rocjpeg/src/rocjpeg_decoder.h +++ b/projects/rocjpeg/src/rocjpeg_decoder.h @@ -124,7 +124,7 @@ private: * @param destination Pointer to the destination image. * @return The status of the operation. */ - RocJpegStatus CopyChannel(HipInteropDeviceMem& hip_interop, uint16_t channel_height, uint8_t channel_index, RocJpegImage *destination); + RocJpegStatus CopyChannel(HipInteropDeviceMem& hip_interop, uint16_t channel_height, uint8_t channel_index, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid); /** * @brief Converts the image to RGB color space. @@ -134,7 +134,7 @@ private: * @param destination Pointer to the destination image. * @return The status of the operation. */ - RocJpegStatus ColorConvertToRGB(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination); + RocJpegStatus ColorConvertToRGB(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid); /** * @brief Converts the image to RGB planar color space. @@ -144,7 +144,7 @@ private: * @param destination Pointer to the destination image. * @return The status of the operation. */ - RocJpegStatus ColorConvertToRGBPlanar(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination); + RocJpegStatus ColorConvertToRGBPlanar(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid); /** * @brief Retrieves the output format for planar YUV images. @@ -155,7 +155,7 @@ private: * @param destination Pointer to the destination image. * @return The status of the operation. */ - RocJpegStatus GetPlanarYUVOutputFormat(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, uint16_t chroma_height, RocJpegImage *destination); + RocJpegStatus GetPlanarYUVOutputFormat(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, uint16_t chroma_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid); /** * @brief Retrieves the output format for Y images. @@ -165,7 +165,7 @@ private: * @param destination Pointer to the destination image. * @return The status of the operation. */ - RocJpegStatus GetYOutputFormat(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination); + RocJpegStatus GetYOutputFormat(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid); int num_devices_; // Number of available devices int device_id_; // ID of the device to be used diff --git a/projects/rocjpeg/src/rocjpeg_hip_kernels.cpp b/projects/rocjpeg/src/rocjpeg_hip_kernels.cpp index 7c20a667a1..3cb89a3600 100644 --- a/projects/rocjpeg/src/rocjpeg_hip_kernels.cpp +++ b/projects/rocjpeg/src/rocjpeg_hip_kernels.cpp @@ -237,7 +237,7 @@ __global__ void ColorConvertYUV444ToRGBKernel(uint32_t dst_width, uint32_t dst_h */ void ColorConvertYUV444ToRGB(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset) { + uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset) { int32_t local_threads_x = 16; int32_t local_threads_y = 4; @@ -252,7 +252,7 @@ void ColorConvertYUV444ToRGB(hipStream_t stream, uint32_t dst_width, uint32_t ds ColorConvertYUV444ToRGBKernel<<(global_threads_x) / local_threads_x), ceil(static_cast(global_threads_y) / local_threads_y)), dim3(local_threads_x, local_threads_y), 0, stream>>>(dst_width, dst_height, dst_image, dst_image_stride_in_bytes, dst_image_stride_in_bytes_comp, src_yuv_image, src_yuv_image + src_u_image_offset, - src_yuv_image + (src_u_image_offset * 2), src_yuv_image_stride_in_bytes, + src_yuv_image + src_v_image_offset, src_yuv_image_stride_in_bytes, dst_width_comp, dst_height_comp, src_yuv_image_stride_in_bytes_comp); } @@ -467,7 +467,7 @@ __global__ void ColorConvertYUV444ToRGBPlanarKernel(uint32_t dst_width, uint32_t */ void ColorConvertYUV444ToRGBPlanar(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image_r, uint8_t *dst_image_g, uint8_t *dst_image_b, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset) { + uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset) { int32_t local_threads_x = 16; int32_t local_threads_y = 4; @@ -482,7 +482,7 @@ void ColorConvertYUV444ToRGBPlanar(hipStream_t stream, uint32_t dst_width, uint3 ColorConvertYUV444ToRGBPlanarKernel<<(global_threads_x) / local_threads_x), ceil(static_cast(global_threads_y) / local_threads_y)), dim3(local_threads_x, local_threads_y), 0, stream>>>(dst_width, dst_height, dst_image_r, dst_image_g, dst_image_b, dst_image_stride_in_bytes, dst_image_stride_in_bytes_comp, src_yuv_image, src_yuv_image + src_u_image_offset, - src_yuv_image + (src_u_image_offset * 2), src_yuv_image_stride_in_bytes, + src_yuv_image + src_v_image_offset, src_yuv_image_stride_in_bytes, dst_width_comp, dst_height_comp, src_yuv_image_stride_in_bytes_comp); } diff --git a/projects/rocjpeg/src/rocjpeg_hip_kernels.h b/projects/rocjpeg/src/rocjpeg_hip_kernels.h index 87e637ec39..d3864e9e63 100644 --- a/projects/rocjpeg/src/rocjpeg_hip_kernels.h +++ b/projects/rocjpeg/src/rocjpeg_hip_kernels.h @@ -43,7 +43,7 @@ THE SOFTWARE. */ void ColorConvertYUV444ToRGB(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset); + uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset); /** * @brief Converts YUV440 image to RGB image. @@ -156,7 +156,7 @@ void ColorConvertRGBAToRGB(hipStream_t stream, uint32_t dst_width, uint32_t dst_ */ void ColorConvertYUV444ToRGBPlanar(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image_r, uint8_t *dst_image_g, uint8_t *dst_image_b, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset); + uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset); /** * @brief Converts YUV440 image to RGB planar format. diff --git a/projects/rocjpeg/test/CMakeLists.txt b/projects/rocjpeg/test/CMakeLists.txt index 076af05992..cd24077ede 100644 --- a/projects/rocjpeg/test/CMakeLists.txt +++ b/projects/rocjpeg/test/CMakeLists.txt @@ -132,4 +132,76 @@ add_test( --build-generator "${CMAKE_GENERATOR}" --test-command "jpegdecodebatched" -i ${ROCM_PATH}/share/rocjpeg/images/ +) + +add_test( + NAME + jpeg-decode-crop-fmt-native + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${ROCM_PATH}/share/rocjpeg/samples/jpegDecode" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecode" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecode" + -i ${ROCM_PATH}/share/rocjpeg/images/ -crop 960,540,2880,1620 +) + +add_test( + NAME + jpeg-decode-crop-fmt-yuv-planar + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${ROCM_PATH}/share/rocjpeg/samples/jpegDecode" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecode" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecode" + -i ${ROCM_PATH}/share/rocjpeg/images/ -fmt yuv -crop 960,540,2880,1620 +) + +add_test( + NAME + jpeg-decode-crop-fmt-y + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${ROCM_PATH}/share/rocjpeg/samples/jpegDecode" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecode" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecode" + -i ${ROCM_PATH}/share/rocjpeg/images/ -fmt y -crop 960,540,2880,1620 +) + +add_test( + NAME + jpeg-decode-crop-fmt-rgb + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${ROCM_PATH}/share/rocjpeg/samples/jpegDecode" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecode" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecode" + -i ${ROCM_PATH}/share/rocjpeg/images/ -fmt rgb -crop 960,540,2880,1620 +) + +add_test( + NAME + jpeg-decode-crop-fmt-rgb-planar + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${ROCM_PATH}/share/rocjpeg/samples/jpegDecode" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecode" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecode" + -i ${ROCM_PATH}/share/rocjpeg/images/ -fmt rgb_planar -crop 960,540,2880,1620 +) + +add_test( + NAME + jpeg-decode-crop-batch-fmt-native + COMMAND + "${CMAKE_CTEST_COMMAND}" + --build-and-test "${ROCM_PATH}/share/rocjpeg/samples/jpegDecodeBatched" + "${CMAKE_CURRENT_BINARY_DIR}/jpegDecodeBatched" + --build-generator "${CMAKE_GENERATOR}" + --test-command "jpegdecodebatched" + -i ${ROCM_PATH}/share/rocjpeg/images/ -crop 960,540,2880,1620 ) \ No newline at end of file