Roi decode feature (#48)
* implemented ROI for NATIVE, YUV_PLANAR, Y, RGB and RGB_PLANAR * added the changes requested by Aryan in the PR * prelim check in of ROI * finished RGB and RGB_PLANAR ROI implementation and testing in rocjpeg_decoder.cpp, updated the versions to 0.6.0, updated jpegdecode.cpp and jpegdecodedbatched.cpp. Still need to modify jpegmultithreads.cpp. Need to run tests on JPEG 444 and 440. And need to add test to ctests and make test. Will update this PR when I've added everything mentioned here. * changed new_offset and new_uv_offset to roi_offset and roi_uv_offset in rocjpeg_decoder.cpp. Added ROI handling in jpegdecodemultithreads sample. Still need to run tests on jpegdecodemultithreads and jpegdecodebatched. * addressed all changes Aryan mentioned for PR 48 on August 12 * added tests to ctests and make tests and fixed conflict in jpegdecodemultithreads.cpp * addressed latest change requests * removed spaces after case VA_FOURCC_444P * updated ctests and make tests * fixed copy/paste error for ctests * fixed typo with extra $ * added print statement for cropped image dimensions * addressed latest change requests from Aryan. Ran make tests and ctests, all passed * added workaround for YUV440 to RGB conversion
Tento commit je obsažen v:
Pavel Tcherniaev
GitHub
+149
-63
@@ -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;
|
||||
}
|
||||
@@ -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
|
||||
|
||||
@@ -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<<<dim3(ceil(static_cast<float>(global_threads_x) / local_threads_x), ceil(static_cast<float>(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<<<dim3(ceil(static_cast<float>(global_threads_x) / local_threads_x), ceil(static_cast<float>(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);
|
||||
}
|
||||
|
||||
|
||||
@@ -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.
|
||||
|
||||
Odkázat v novém úkolu
Zablokovat Uživatele