Initial commit of the rocJPEG library (#1)
* rocJPEG initial commit * Add gitignore file * clean up * code clean up * CMakeLists update * Update README * Update readme and setup script * add docs folder * update rocjepg headers * update the rocjepg header * update the copyright year * Add nullptr check for all APIs * update the hip kernels execution * follow google c++ style guide * code clean up * use google c++ style guide for hip kernels * add support for rocJpegGetErrorName API * Add ParseCommandLine function * update the rocjpeg sample * Update the readme, LIBVA instructions * udate the rocJPEG script * Update CMAKE and libdrm * Add support for ROCJPEG_OUTPUT_UNCHANGED and SaveImage function * code clean up * update the drm dependencies * remove the unused hipstream * code clean up * add -fmt option for selection the output format * update the GetImageInfo API * update the sample * update the save image function based onn the surface and putput format * add new sample images * Put hip interop mem into separate functions and some code clean up * Add support for CopyLuma and CopyChroma * Add support for the yuv output format for NV12 surface * Update HIP kernels * add support for rgbi format conversion for NV12 and YUv444 * Update the status * Add support for converting YUYV to RGB * Add support for unpacking YUYV format * Fix unpacking Y fron YUYV hip kernel * Add suppoort for extractign packed YUYV to YUV planar * hip code clean up * code clean up * Update the headers * code clean up * change hip kernel names * hip kernel clean up * hip clean up * code clean up * code clean up * code clean up * code clean up * code clean up * code clean up for jpege decoder class * code clean up * code clean up * code clean up * code clean up * update the dockers * code clean up * code clean up * hip kernels clean up * remove unused hip kernels * add additional test cases * update the APIs * add new hip yuv400torgbi kernel * update yuv400torgbi kernel * restructure files * code clean up * code clean up * add jenkins * code clean up * code clean up * update readme * update docker's README * make changes based on the reviewers comments * make changes based on the reviewers comments * return ROCJPEG_STATUS_JPEG_NOT_SUPPORTED if the resolution of the jpeg is not supported for HW decoding
This commit is contained in:
Aryan Salmanpour
GitHub
@@ -0,0 +1,171 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "rocjpeg_api_handle.h"
|
||||
#include "rocjpeg_commons.h"
|
||||
|
||||
/*****************************************************************************************************/
|
||||
//! \fn RocJpegStatus ROCJPEGAPI rocJpegCreate(RocJpegBackend backend, int device_id, RocJpegHandle *handle)
|
||||
//! Create the decoder object based on backend and device_id. A handle to the created decoder is returned
|
||||
/*****************************************************************************************************/
|
||||
RocJpegStatus ROCJPEGAPI rocJpegCreate(RocJpegBackend backend, int device_id, RocJpegHandle *handle) {
|
||||
if (handle == nullptr) {
|
||||
return ROCJPEG_STATUS_INVALID_PARAMETER;
|
||||
}
|
||||
RocJpegHandle rocjpeg_handle = nullptr;
|
||||
try {
|
||||
rocjpeg_handle = new RocJpegDecoderHandle(backend, device_id);
|
||||
} catch(const std::exception& e) {
|
||||
ERR(STR("Failed to init the rocJPEG handle, ") + STR(e.what()));
|
||||
return ROCJPEG_STATUS_NOT_INITIALIZED;
|
||||
}
|
||||
*handle = rocjpeg_handle;
|
||||
return static_cast<RocJpegDecoderHandle *>(rocjpeg_handle)->rocjpeg_decoder->InitializeDecoder();
|
||||
}
|
||||
|
||||
/*****************************************************************************************************/
|
||||
//! \fn RocJpegStatus ROCJPEGAPI rocJpegDestroy(RocJpegHandle handle)
|
||||
//! Release the decoder object and resources.
|
||||
/*****************************************************************************************************/
|
||||
RocJpegStatus ROCJPEGAPI rocJpegDestroy(RocJpegHandle handle) {
|
||||
if (handle == nullptr) {
|
||||
return ROCJPEG_STATUS_INVALID_PARAMETER;
|
||||
}
|
||||
auto rocjpeg_handle = static_cast<RocJpegDecoderHandle*>(handle);
|
||||
delete rocjpeg_handle;
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
/*****************************************************************************************************/
|
||||
//! \fn RocJpegStatus ROCJPEGAPI rocJpegGetImageInfo(RocJpegHandle handle, const uint8_t *data, size_t length,
|
||||
//! int *num_components, RocJpegChromaSubsampling *subsampling, int *widths, int *heights)
|
||||
//! Retrieve the image info, including channel, width and height of each component, and chroma subsampling.
|
||||
/*****************************************************************************************************/
|
||||
RocJpegStatus ROCJPEGAPI rocJpegGetImageInfo(RocJpegHandle handle, const uint8_t *data, size_t length, uint8_t *num_components,
|
||||
RocJpegChromaSubsampling *subsampling, uint32_t *widths, uint32_t *heights) {
|
||||
if (handle == nullptr || data == nullptr || num_components == nullptr ||
|
||||
subsampling == nullptr || widths == nullptr || heights == nullptr) {
|
||||
return ROCJPEG_STATUS_INVALID_PARAMETER;
|
||||
}
|
||||
RocJpegStatus rocjpeg_status = ROCJPEG_STATUS_SUCCESS;
|
||||
auto rocjpeg_handle = static_cast<RocJpegDecoderHandle*>(handle);
|
||||
try {
|
||||
rocjpeg_status = rocjpeg_handle->rocjpeg_decoder->GetImageInfo(data, length, num_components, subsampling, widths, heights);
|
||||
} catch (const std::exception& e) {
|
||||
rocjpeg_handle->CaptureError(e.what());
|
||||
ERR(e.what());
|
||||
return ROCJPEG_STATUS_RUNTIME_ERROR;
|
||||
}
|
||||
|
||||
return rocjpeg_status;
|
||||
}
|
||||
|
||||
/*****************************************************************************************************/
|
||||
//! \fn RocJpegStatus ROCJPEGAPI rocJpegDecode(RocJpegHandle handle, const uint8_t *data, size_t length, RocJpegOutputFormat output_format, RocJpegImage *destination, hipStream_t stream);
|
||||
//! \ingroup group_amd_rocjpeg
|
||||
//! Decodes single image based on the backend used to create the rocJpeg handle in rocJpegCreate API.
|
||||
//! Destination buffers should be large enough to be able to store output of specified format. These buffers should be pre-allocted by the user in the device memories.
|
||||
//! For each color plane (channel) sizes could be retrieved for image using rocJpegGetImageInfo API
|
||||
//! and minimum required memory buffer for each plane is plane_height * plane_pitch where plane_pitch >= plane_width for
|
||||
//! planar output formats and plane_pitch >= plane_width * num_components for interleaved output format.
|
||||
/*****************************************************************************************************/
|
||||
RocJpegStatus ROCJPEGAPI rocJpegDecode(RocJpegHandle handle, const uint8_t *data, size_t length, RocJpegOutputFormat output_format,
|
||||
RocJpegImage *destination) {
|
||||
|
||||
if (handle == nullptr || data == nullptr) {
|
||||
return ROCJPEG_STATUS_INVALID_PARAMETER;
|
||||
}
|
||||
RocJpegStatus rocjpeg_status = ROCJPEG_STATUS_SUCCESS;
|
||||
auto rocjpeg_handle = static_cast<RocJpegDecoderHandle*>(handle);
|
||||
try {
|
||||
rocjpeg_status = rocjpeg_handle->rocjpeg_decoder->Decode(data, length, output_format, destination);
|
||||
} catch (const std::exception& e) {
|
||||
rocjpeg_handle->CaptureError(e.what());
|
||||
ERR(e.what());
|
||||
return ROCJPEG_STATUS_RUNTIME_ERROR;
|
||||
}
|
||||
|
||||
return rocjpeg_status;
|
||||
}
|
||||
|
||||
/*****************************************************************************************************/
|
||||
//! \fn extern const char* ROCDECAPI rocJpegGetErrorName(RocJpegStatus rocjpeg_status);
|
||||
//! \ingroup group_amd_rocjpeg
|
||||
//! Return name of the specified error code in text form.
|
||||
/*****************************************************************************************************/
|
||||
extern const char* ROCJPEGAPI rocJpegGetErrorName(RocJpegStatus rocjpeg_status) {
|
||||
switch (rocjpeg_status) {
|
||||
case ROCJPEG_STATUS_SUCCESS:
|
||||
return "ROCJPEG_STATUS_SUCCESS";
|
||||
case ROCJPEG_STATUS_NOT_INITIALIZED:
|
||||
return "ROCJPEG_STATUS_NOT_INITIALIZED";
|
||||
case ROCJPEG_STATUS_INVALID_PARAMETER:
|
||||
return "ROCJPEG_STATUS_INVALID_PARAMETER";
|
||||
case ROCJPEG_STATUS_BAD_JPEG:
|
||||
return "ROCJPEG_STATUS_BAD_JPEG";
|
||||
case ROCJPEG_STATUS_JPEG_NOT_SUPPORTED:
|
||||
return "ROCJPEG_STATUS_JPEG_NOT_SUPPORTED";
|
||||
case ROCJPEG_STATUS_EXECUTION_FAILED:
|
||||
return "ROCJPEG_STATUS_EXECUTION_FAILED";
|
||||
case ROCJPEG_STATUS_ARCH_MISMATCH:
|
||||
return "ROCJPEG_STATUS_ARCH_MISMATCH";
|
||||
case ROCJPEG_STATUS_INTERNAL_ERROR:
|
||||
return "ROCJPEG_STATUS_INTERNAL_ERROR";
|
||||
case ROCJPEG_STATUS_IMPLEMENTATION_NOT_SUPPORTED:
|
||||
return "ROCJPEG_STATUS_IMPLEMENTATION_NOT_SUPPORTED";
|
||||
case ROCJPEG_STATUS_HW_JPEG_DECODER_NOT_SUPPORTED:
|
||||
return "ROCJPEG_STATUS_HW_JPEG_DECODER_NOT_SUPPORTED";
|
||||
case ROCJPEG_STATUS_RUNTIME_ERROR:
|
||||
return "ROCJPEG_STATUS_RUNTIME_ERROR";
|
||||
case ROCJPEG_STATUS_OUTOF_MEMORY:
|
||||
return "ROCJPEG_STATUS_OUTOF_MEMORY";
|
||||
case ROCJPEG_STATUS_NOT_IMPLEMENTED:
|
||||
return "ROCJPEG_STATUS_NOT_IMPLEMENTED";
|
||||
default:
|
||||
return "UNKNOWN_ERROR";
|
||||
}
|
||||
}
|
||||
|
||||
/*****************************************************************************************************/
|
||||
//! \fn RocJpegStatus ROCJPEGAPI rocJpegDecodeBatchedInitialize(RocJpegHandle handle, int batch_size, int max_cpu_threads, RocJpegOutputFormat output_format);
|
||||
//! \ingroup group_amd_rocjpeg
|
||||
//! Resets and initializes batch decoder for working on the batches of specified size
|
||||
//! Should be called once for decoding batches of this specific size, also use to reset failed batches
|
||||
//! \return ROCJPEG_STATUS_SUCCESS if successful
|
||||
/*****************************************************************************************************/
|
||||
RocJpegStatus ROCJPEGAPI rocJpegDecodeBatchedInitialize(RocJpegHandle handle, int batch_size, int max_cpu_threads, RocJpegOutputFormat output_format) {
|
||||
return ROCJPEG_STATUS_NOT_IMPLEMENTED;
|
||||
}
|
||||
|
||||
|
||||
/*****************************************************************************************************/
|
||||
//! \fn RocJpegStatus ROCJPEGAPI rocJpegDecodeBatched(RocJpegHandle handle, const uint8_t *data, const size_t *lengths, RocJpegImage *destinations, hipStream_t stream);
|
||||
//! \ingroup group_amd_rocjpeg
|
||||
//! Decodes batch of images. Output buffers should be large enough to be able to store
|
||||
//! outputs of specified format, see single image decoding description for details. Call to
|
||||
//! rocjpegDecodeBatchedInitialize() is required prior to this call, batch size is expected to be the same as
|
||||
//! parameter to this batch initialization function.
|
||||
//! \return ROCJPEG_STATUS_SUCCESS if successful
|
||||
/*****************************************************************************************************/
|
||||
RocJpegStatus ROCJPEGAPI rocJpegDecodeBatched(RocJpegHandle handle, const uint8_t *data, const size_t *lengths, RocJpegImage *destinations) {
|
||||
return ROCJPEG_STATUS_NOT_IMPLEMENTED;
|
||||
}
|
||||
@@ -0,0 +1,50 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#ifndef ROC_JPEG_HANDLE_H
|
||||
#define ROC_JPEG_HANDLE_H
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <memory>
|
||||
#include <string>
|
||||
|
||||
#include "rocjpeg_decoder.h"
|
||||
|
||||
/**
|
||||
* @brief RocJpegHandle class
|
||||
*
|
||||
*/
|
||||
class RocJpegDecoderHandle {
|
||||
public:
|
||||
explicit RocJpegDecoderHandle(RocJpegBackend backend, int device_id) : rocjpeg_decoder(std::make_shared<ROCJpegDecoder>(backend, device_id)) {};
|
||||
~RocJpegDecoderHandle() { ClearErrors(); }
|
||||
std::shared_ptr<ROCJpegDecoder> rocjpeg_decoder;
|
||||
bool NoError() { return error_.empty(); }
|
||||
const char* ErrorMsg() { return error_.c_str(); }
|
||||
void CaptureError(const std::string& err_msg) { error_ = err_msg; }
|
||||
private:
|
||||
void ClearErrors() { error_ = "";}
|
||||
std::string error_;
|
||||
};
|
||||
|
||||
#endif //ROC_JPEG_HANDLE_H
|
||||
@@ -0,0 +1,92 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#ifndef ROC_JPEG_COMMON_H_
|
||||
#define ROC_JPEG_COMMON_H_
|
||||
|
||||
#pragma once
|
||||
#include <stdexcept>
|
||||
#include <exception>
|
||||
#include <string>
|
||||
#include <iostream>
|
||||
#include <cstring>
|
||||
|
||||
#define TOSTR(X) std::to_string(static_cast<int>(X))
|
||||
#define STR(X) std::string(X)
|
||||
|
||||
#if DBGINFO
|
||||
#define INFO(X) std::clog << "[INF] " << " {" << __func__ <<"} " << " " << X << std::endl;
|
||||
#else
|
||||
#define INFO(X) ;
|
||||
#endif
|
||||
#define ERR(X) std::cerr << "[ERR] " << " {" << __func__ <<"} " << " " << X << std::endl;
|
||||
|
||||
#define CHECK_VAAPI(call) { \
|
||||
VAStatus va_status = (call); \
|
||||
if (va_status != VA_STATUS_SUCCESS) { \
|
||||
std::cerr << "VAAPI failure: " << #call << " failed with status: " << std::hex << "0x" << va_status << std::dec << " = '" << vaErrorStr(va_status) << "' at " << __FILE__ << ":" << __LINE__ << std::endl;\
|
||||
return ROCJPEG_STATUS_EXECUTION_FAILED; \
|
||||
} \
|
||||
}
|
||||
|
||||
#define CHECK_HIP(call) { \
|
||||
hipError_t hip_status = (call); \
|
||||
if (hip_status != hipSuccess) { \
|
||||
std::cerr << "HIP failure: 'status: " << hipGetErrorName(hip_status) << "' at " << __FILE__ << ":" << __LINE__ << std::endl;\
|
||||
return ROCJPEG_STATUS_EXECUTION_FAILED; \
|
||||
} \
|
||||
}
|
||||
|
||||
#define CHECK_ROCJPEG(call) { \
|
||||
RocJpegStatus rocjpeg_status = (call); \
|
||||
if (rocjpeg_status != ROCJPEG_STATUS_SUCCESS) { \
|
||||
std::cerr << #call << " returned " << rocJpegGetErrorName(rocjpeg_status) << " at " << __FILE__ << ":" << __LINE__ << std::endl;\
|
||||
return rocjpeg_status; \
|
||||
} \
|
||||
}
|
||||
|
||||
static bool GetEnv(const char *name, char *value, size_t valueSize) {
|
||||
const char *v = getenv(name);
|
||||
if (v) {
|
||||
strncpy(value, v, valueSize);
|
||||
value[valueSize - 1] = 0;
|
||||
}
|
||||
return v ? true : false;
|
||||
}
|
||||
|
||||
static inline int align(int value, int alignment) {
|
||||
return (value + alignment - 1) & ~(alignment - 1);
|
||||
}
|
||||
|
||||
class RocJpegException : public std::exception {
|
||||
public:
|
||||
explicit RocJpegException(const std::string& message):message_(message){}
|
||||
virtual const char* what() const throw() override {
|
||||
return message_.c_str();
|
||||
}
|
||||
private:
|
||||
std::string message_;
|
||||
};
|
||||
|
||||
#define THROW(X) throw RocJpegException(" { "+std::string(__func__)+" } " + X);
|
||||
|
||||
#endif //ROC_JPEG_COMMON_H_
|
||||
@@ -0,0 +1,331 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "rocjpeg_decoder.h"
|
||||
|
||||
ROCJpegDecoder::ROCJpegDecoder(RocJpegBackend backend, int device_id) :
|
||||
num_devices_{0}, device_id_ {device_id}, hip_stream_ {0}, backend_{backend}, hip_interop_{} {}
|
||||
|
||||
ROCJpegDecoder::~ROCJpegDecoder() {
|
||||
if (hip_stream_) {
|
||||
hipError_t hip_status = hipStreamDestroy(hip_stream_);
|
||||
}
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::InitHIP(int device_id) {
|
||||
hipError_t hip_status = hipSuccess;
|
||||
CHECK_HIP(hipGetDeviceCount(&num_devices_));
|
||||
if (num_devices_ < 1) {
|
||||
ERR("ERROR: Failed to find any GPU!");
|
||||
return ROCJPEG_STATUS_NOT_INITIALIZED;
|
||||
}
|
||||
if (device_id >= num_devices_) {
|
||||
ERR("ERROR: the requested device_id is not found!");
|
||||
return ROCJPEG_STATUS_INVALID_PARAMETER;
|
||||
}
|
||||
CHECK_HIP(hipSetDevice(device_id));
|
||||
CHECK_HIP(hipGetDeviceProperties(&hip_dev_prop_, device_id));
|
||||
CHECK_HIP(hipStreamCreate(&hip_stream_));
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::InitializeDecoder() {
|
||||
RocJpegStatus rocjpeg_status = ROCJPEG_STATUS_SUCCESS;
|
||||
rocjpeg_status = InitHIP(device_id_);
|
||||
if (rocjpeg_status != ROCJPEG_STATUS_SUCCESS) {
|
||||
ERR("ERROR: Failed to initilize the HIP!");
|
||||
return rocjpeg_status;
|
||||
}
|
||||
if (backend_ == ROCJPEG_BACKEND_HARDWARE) {
|
||||
rocjpeg_status = jpeg_vaapi_decoder_.InitializeDecoder(hip_dev_prop_.gcnArchName);
|
||||
if (rocjpeg_status != ROCJPEG_STATUS_SUCCESS) {
|
||||
ERR("ERROR: Failed to initialize the VA-API JPEG decoder!");
|
||||
return rocjpeg_status;
|
||||
}
|
||||
} else if (backend_ == ROCJPEG_BACKEND_HYBRID) {
|
||||
return ROCJPEG_STATUS_NOT_IMPLEMENTED;
|
||||
}
|
||||
return rocjpeg_status;
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::Decode(const uint8_t *data, size_t length, RocJpegOutputFormat output_format, RocJpegImage *destination) {
|
||||
std::lock_guard<std::mutex> lock(mutex_);
|
||||
RocJpegStatus rocjpeg_status = ROCJPEG_STATUS_SUCCESS;
|
||||
|
||||
if (!jpeg_parser_.ParseJpegStream(data, length)) {
|
||||
ERR("ERROR: Failed to parse the jpeg stream!");
|
||||
return ROCJPEG_STATUS_BAD_JPEG;
|
||||
}
|
||||
|
||||
const JpegStreamParameters *jpeg_stream_params = jpeg_parser_.GetJpegStreamParameters();
|
||||
VASurfaceID current_surface_id;
|
||||
CHECK_ROCJPEG(jpeg_vaapi_decoder_.SubmitDecode(jpeg_stream_params, current_surface_id));
|
||||
|
||||
if (destination != nullptr) {
|
||||
VADRMPRIMESurfaceDescriptor va_drm_prime_surface_desc = {};
|
||||
CHECK_ROCJPEG(jpeg_vaapi_decoder_.SyncSurface(current_surface_id));
|
||||
CHECK_ROCJPEG(jpeg_vaapi_decoder_.ExportSurface(current_surface_id, va_drm_prime_surface_desc));
|
||||
CHECK_ROCJPEG(GetHipInteropMem(va_drm_prime_surface_desc));
|
||||
|
||||
uint16_t chroma_height = 0;
|
||||
CHECK_ROCJPEG(GetChromaHeight(jpeg_stream_params->picture_parameter_buffer.picture_height, chroma_height));
|
||||
|
||||
switch (output_format) {
|
||||
case ROCJPEG_OUTPUT_NATIVE:
|
||||
// copy the native decoded output buffers from interop memory directly to the destination buffers
|
||||
CHECK_ROCJPEG(CopyLuma(destination, jpeg_stream_params->picture_parameter_buffer.picture_height));
|
||||
CHECK_ROCJPEG(CopyChroma(destination, chroma_height));
|
||||
break;
|
||||
case ROCJPEG_OUTPUT_YUV_PLANAR:
|
||||
CHECK_ROCJPEG(GetPlanarYUVOutputFormat(jpeg_stream_params->picture_parameter_buffer.picture_width,
|
||||
jpeg_stream_params->picture_parameter_buffer.picture_height, chroma_height, destination));
|
||||
break;
|
||||
case ROCJPEG_OUTPUT_Y:
|
||||
CHECK_ROCJPEG(GetYOutputFormat(jpeg_stream_params->picture_parameter_buffer.picture_width,
|
||||
jpeg_stream_params->picture_parameter_buffer.picture_height, destination));
|
||||
break;
|
||||
case ROCJPEG_OUTPUT_RGB:
|
||||
CHECK_ROCJPEG(ColorConvertToRGB(jpeg_stream_params->picture_parameter_buffer.picture_width,
|
||||
jpeg_stream_params->picture_parameter_buffer.picture_height, destination));
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
CHECK_HIP(hipStreamSynchronize(hip_stream_));
|
||||
|
||||
CHECK_ROCJPEG(ReleaseHipInteropMem(current_surface_id));
|
||||
}
|
||||
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::GetImageInfo(const uint8_t *data, size_t length, uint8_t *num_components, RocJpegChromaSubsampling *subsampling, uint32_t *widths, uint32_t *heights){
|
||||
std::lock_guard<std::mutex> lock(mutex_);
|
||||
if (widths == nullptr || heights == nullptr || num_components == nullptr) {
|
||||
return ROCJPEG_STATUS_INVALID_PARAMETER;
|
||||
}
|
||||
if (!jpeg_parser_.ParseJpegStream(data, length)) {
|
||||
ERR("ERROR: jpeg parser failed!");
|
||||
return ROCJPEG_STATUS_BAD_JPEG;
|
||||
}
|
||||
const JpegStreamParameters *jpeg_stream_params = jpeg_parser_.GetJpegStreamParameters();
|
||||
*num_components = jpeg_stream_params->picture_parameter_buffer.num_components;
|
||||
widths[0] = jpeg_stream_params->picture_parameter_buffer.picture_width;
|
||||
heights[0] = jpeg_stream_params->picture_parameter_buffer.picture_height;
|
||||
widths[3] = 0;
|
||||
heights[3] = 0;
|
||||
|
||||
switch (jpeg_stream_params->chroma_subsampling) {
|
||||
case CSS_444:
|
||||
*subsampling = ROCJPEG_CSS_444;
|
||||
widths[2] = widths[1] = widths[0];
|
||||
heights[2] = heights[1] = heights[0];
|
||||
break;
|
||||
case CSS_422:
|
||||
*subsampling = ROCJPEG_CSS_422;
|
||||
widths[2] = widths[1] = widths[0] >> 1;
|
||||
heights[2] = heights[1] = heights[0];
|
||||
break;
|
||||
case CSS_420:
|
||||
*subsampling = ROCJPEG_CSS_420;
|
||||
widths[2] = widths[1] = widths[0] >> 1;
|
||||
heights[2] = heights[1] = heights[0] >> 1;
|
||||
break;
|
||||
case CSS_400:
|
||||
*subsampling = ROCJPEG_CSS_400;
|
||||
widths[3] = widths[2] = widths[1] = 0;
|
||||
heights[3] = heights[2] = heights[1] = 0;
|
||||
break;
|
||||
case CSS_411:
|
||||
*subsampling = ROCJPEG_CSS_411;
|
||||
widths[2] = widths[1] = widths[0] >> 2;
|
||||
heights[2] = heights[1] = heights[0];
|
||||
break;
|
||||
case CSS_440:
|
||||
*subsampling = ROCJPEG_CSS_440;
|
||||
widths[2] = widths[1] = widths[0] >> 1;
|
||||
heights[2] = heights[1] = heights[0] >> 1;
|
||||
break;
|
||||
default:
|
||||
*subsampling = ROCJPEG_CSS_UNKNOWN;
|
||||
break;
|
||||
}
|
||||
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::GetHipInteropMem(VADRMPRIMESurfaceDescriptor &va_drm_prime_surface_desc) {
|
||||
hipExternalMemoryHandleDesc external_mem_handle_desc = {};
|
||||
hipExternalMemoryBufferDesc external_mem_buffer_desc = {};
|
||||
external_mem_handle_desc.type = hipExternalMemoryHandleTypeOpaqueFd;
|
||||
external_mem_handle_desc.handle.fd = va_drm_prime_surface_desc.objects[0].fd;
|
||||
external_mem_handle_desc.size = va_drm_prime_surface_desc.objects[0].size;
|
||||
|
||||
CHECK_HIP(hipImportExternalMemory(&hip_interop_.hip_ext_mem, &external_mem_handle_desc));
|
||||
external_mem_buffer_desc.size = va_drm_prime_surface_desc.objects[0].size;
|
||||
CHECK_HIP(hipExternalMemoryGetMappedBuffer((void**)&hip_interop_.hip_mapped_device_mem, hip_interop_.hip_ext_mem, &external_mem_buffer_desc));
|
||||
|
||||
hip_interop_.surface_format = va_drm_prime_surface_desc.fourcc;
|
||||
hip_interop_.width = va_drm_prime_surface_desc.width;
|
||||
hip_interop_.height = va_drm_prime_surface_desc.height;
|
||||
hip_interop_.offset[0] = va_drm_prime_surface_desc.layers[0].offset[0];
|
||||
hip_interop_.offset[1] = va_drm_prime_surface_desc.layers[1].offset[0];
|
||||
hip_interop_.offset[2] = va_drm_prime_surface_desc.layers[2].offset[0];
|
||||
hip_interop_.pitch[0] = va_drm_prime_surface_desc.layers[0].pitch[0];
|
||||
hip_interop_.pitch[1] = va_drm_prime_surface_desc.layers[1].pitch[0];
|
||||
hip_interop_.pitch[2] = va_drm_prime_surface_desc.layers[2].pitch[0];
|
||||
hip_interop_.num_layers = va_drm_prime_surface_desc.num_layers;
|
||||
|
||||
for (uint32_t i = 0; i < va_drm_prime_surface_desc.num_objects; ++i) {
|
||||
close(va_drm_prime_surface_desc.objects[i].fd);
|
||||
}
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::ReleaseHipInteropMem(VASurfaceID current_surface_id) {
|
||||
if (hip_interop_.hip_mapped_device_mem != nullptr) {
|
||||
CHECK_HIP(hipFree(hip_interop_.hip_mapped_device_mem));
|
||||
}
|
||||
if (hip_interop_.hip_ext_mem != nullptr) {
|
||||
CHECK_HIP(hipDestroyExternalMemory(hip_interop_.hip_ext_mem));
|
||||
}
|
||||
memset((void*)&hip_interop_, 0, sizeof(hip_interop_));
|
||||
|
||||
CHECK_ROCJPEG(jpeg_vaapi_decoder_.ReleaseSurface(current_surface_id));
|
||||
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::CopyLuma(RocJpegImage *destination, uint16_t picture_height) {
|
||||
if (hip_interop_.pitch[0] != 0 && destination->pitch[0] != 0 && destination->channel[0] != nullptr) {
|
||||
if (destination->pitch[0] == hip_interop_.pitch[0]) {
|
||||
uint32_t luma_size = destination->pitch[0] * picture_height;
|
||||
CHECK_HIP(hipMemcpyDtoDAsync(destination->channel[0], hip_interop_.hip_mapped_device_mem, luma_size, hip_stream_));
|
||||
} else {
|
||||
CHECK_HIP(hipMemcpy2DAsync(destination->channel[0], destination->pitch[0], hip_interop_.hip_mapped_device_mem, hip_interop_.pitch[0],
|
||||
destination->pitch[0], picture_height, hipMemcpyDeviceToDevice, hip_stream_));
|
||||
}
|
||||
}
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::CopyChroma(RocJpegImage *destination, uint16_t chroma_height) {
|
||||
// copy channel1
|
||||
if (hip_interop_.pitch[1] != 0 && destination->pitch[1] != 0 && destination->channel[1] != nullptr) {
|
||||
uint32_t chroma_size = destination->pitch[1] * chroma_height;
|
||||
uint8_t *layer1_mem = hip_interop_.hip_mapped_device_mem + hip_interop_.offset[1];
|
||||
if (destination->pitch[1] == hip_interop_.pitch[1]) {
|
||||
CHECK_HIP(hipMemcpyDtoDAsync(destination->channel[1], layer1_mem, chroma_size, hip_stream_));
|
||||
} else {
|
||||
CHECK_HIP(hipMemcpy2DAsync(destination->channel[1], destination->pitch[1], layer1_mem, hip_interop_.pitch[1],
|
||||
destination->pitch[1], chroma_height, hipMemcpyDeviceToDevice, hip_stream_));
|
||||
}
|
||||
}
|
||||
// copy channel2
|
||||
if (hip_interop_.pitch[2] != 0 && destination->pitch[2] != 0 && destination->channel[2] != nullptr) {
|
||||
uint32_t chroma_size = destination->pitch[2] * chroma_height;
|
||||
uint8_t *layer2_mem = hip_interop_.hip_mapped_device_mem + hip_interop_.offset[2];
|
||||
if (destination->pitch[2] == hip_interop_.pitch[2]) {
|
||||
CHECK_HIP(hipMemcpyDtoDAsync(destination->channel[2], layer2_mem, chroma_size, hip_stream_));
|
||||
} else {
|
||||
CHECK_HIP(hipMemcpy2DAsync(destination->channel[2], destination->pitch[2], layer2_mem, hip_interop_.pitch[2],
|
||||
destination->pitch[2], chroma_height, hipMemcpyDeviceToDevice, hip_stream_));
|
||||
}
|
||||
}
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::GetChromaHeight(uint16_t picture_height, uint16_t &chroma_height) {
|
||||
switch (hip_interop_.surface_format) {
|
||||
case VA_FOURCC_NV12: /*NV12: two-plane 8-bit YUV 4:2:0*/
|
||||
chroma_height = picture_height >> 1;
|
||||
break;
|
||||
case VA_FOURCC_444P: /*444P: three-plane 8-bit YUV 4:4:4*/
|
||||
chroma_height = picture_height;
|
||||
break;
|
||||
case VA_FOURCC_Y800: /*Y800: one-plane 8-bit greyscale YUV 4:0:0*/
|
||||
chroma_height = 0;
|
||||
break;
|
||||
case ROCJPEG_FOURCC_YUYV: /*YUYV: one-plane packed 8-bit YUV 4:2:2. Four bytes per pair of pixels: Y, U, Y, V*/
|
||||
chroma_height = picture_height;
|
||||
break;
|
||||
default:
|
||||
return ROCJPEG_STATUS_JPEG_NOT_SUPPORTED;
|
||||
}
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::ColorConvertToRGB(uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination) {
|
||||
switch (hip_interop_.surface_format) {
|
||||
case VA_FOURCC_444P:
|
||||
ColorConvertYUV444ToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0],
|
||||
hip_interop_.hip_mapped_device_mem, hip_interop_.pitch[0], hip_interop_.offset[1]);
|
||||
break;
|
||||
case ROCJPEG_FOURCC_YUYV:
|
||||
ColorConvertYUYVToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0],
|
||||
hip_interop_.hip_mapped_device_mem, hip_interop_.pitch[0]);
|
||||
break;
|
||||
case VA_FOURCC_NV12:
|
||||
ColorConvertNV12ToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0],
|
||||
hip_interop_.hip_mapped_device_mem, hip_interop_.pitch[0],
|
||||
hip_interop_.hip_mapped_device_mem + hip_interop_.offset[1], hip_interop_.pitch[1]);
|
||||
break;
|
||||
case VA_FOURCC_Y800:
|
||||
ColorConvertYUV400ToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0],
|
||||
hip_interop_.hip_mapped_device_mem, hip_interop_.pitch[0]);
|
||||
break;
|
||||
default:
|
||||
ERR("ERROR! surface format is not supported!");
|
||||
return ROCJPEG_STATUS_JPEG_NOT_SUPPORTED;
|
||||
}
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::GetPlanarYUVOutputFormat(uint32_t picture_width, uint32_t picture_height, uint16_t chroma_height, RocJpegImage *destination) {
|
||||
if (hip_interop_.surface_format == ROCJPEG_FOURCC_YUYV) {
|
||||
// Extract the packed YUYV and copy them into the first, second, and thrid 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_.hip_mapped_device_mem, hip_interop_.pitch[0]);
|
||||
} else {
|
||||
CHECK_ROCJPEG(CopyLuma(destination, picture_height));
|
||||
if (hip_interop_.surface_format == VA_FOURCC_NV12) {
|
||||
// Extract the interleaved UV channels and copy them into the second and thrid channels of the destination.
|
||||
ConvertInterleavedUVToPlanarUV(hip_stream_, picture_width >> 1, picture_height >> 1, destination->channel[1], destination->channel[2],
|
||||
destination->pitch[1], hip_interop_.hip_mapped_device_mem + hip_interop_.offset[1] , hip_interop_.pitch[1]);
|
||||
} else {
|
||||
CHECK_ROCJPEG(CopyChroma(destination, chroma_height));
|
||||
}
|
||||
}
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus ROCJpegDecoder::GetYOutputFormat(uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination) {
|
||||
if (hip_interop_.surface_format == ROCJPEG_FOURCC_YUYV) {
|
||||
ExtractYFromPackedYUYV(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0],
|
||||
hip_interop_.hip_mapped_device_mem, hip_interop_.pitch[0]);
|
||||
} else {
|
||||
CHECK_ROCJPEG(CopyLuma(destination, picture_height));
|
||||
}
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
@@ -0,0 +1,76 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#ifndef ROC_JPEG_DECODER_H_
|
||||
#define ROC_JPEG_DECODER_H_
|
||||
|
||||
#include <unistd.h>
|
||||
#include <vector>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <mutex>
|
||||
#include <queue>
|
||||
#include "../api/rocjpeg.h"
|
||||
#include "rocjpeg_parser.h"
|
||||
#include "rocjpeg_commons.h"
|
||||
#include "rocjpeg_vaapi_decoder.h"
|
||||
#include "rocjpeg_hip_kernels.h"
|
||||
|
||||
struct HipInteropDeviceMem {
|
||||
hipExternalMemory_t hip_ext_mem; // Interface to the vaapi-hip interop
|
||||
uint8_t* hip_mapped_device_mem; // Mapped device memory for the YUV plane
|
||||
uint32_t surface_format; // Pixel format fourcc of the whole surface
|
||||
uint32_t width; // Width of the surface in pixels.
|
||||
uint32_t height; // Height of the surface in pixels.
|
||||
uint32_t offset[3]; // Offset of each plane
|
||||
uint32_t pitch[3]; // Pitch of each plane
|
||||
uint32_t num_layers; // Number of layers making up the surface
|
||||
};
|
||||
|
||||
class ROCJpegDecoder {
|
||||
public:
|
||||
ROCJpegDecoder(RocJpegBackend backend = ROCJPEG_BACKEND_HARDWARE, int device_id = 0);
|
||||
~ROCJpegDecoder();
|
||||
RocJpegStatus InitializeDecoder();
|
||||
RocJpegStatus GetImageInfo(const uint8_t *data, size_t length, uint8_t *num_components, RocJpegChromaSubsampling *subsampling, uint32_t *widths, uint32_t *heights);
|
||||
RocJpegStatus Decode(const uint8_t *data, size_t length, RocJpegOutputFormat output_format, RocJpegImage *destination);
|
||||
private:
|
||||
RocJpegStatus InitHIP(int device_id);
|
||||
RocJpegStatus GetHipInteropMem(VADRMPRIMESurfaceDescriptor &va_drm_prime_surface_desc);
|
||||
RocJpegStatus ReleaseHipInteropMem(VASurfaceID current_surface_id);
|
||||
RocJpegStatus GetChromaHeight(uint16_t picture_height, uint16_t &chroma_height);
|
||||
RocJpegStatus CopyLuma(RocJpegImage *destination, uint16_t picture_height);
|
||||
RocJpegStatus CopyChroma(RocJpegImage *destination, uint16_t chroma_height);
|
||||
RocJpegStatus ColorConvertToRGB(uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination);
|
||||
RocJpegStatus GetPlanarYUVOutputFormat(uint32_t picture_width, uint32_t picture_height, uint16_t chroma_height, RocJpegImage *destination);
|
||||
RocJpegStatus GetYOutputFormat(uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination);
|
||||
int num_devices_;
|
||||
int device_id_;
|
||||
hipDeviceProp_t hip_dev_prop_;
|
||||
hipStream_t hip_stream_;
|
||||
std::mutex mutex_;
|
||||
JpegParser jpeg_parser_;
|
||||
RocJpegBackend backend_;
|
||||
RocJpegVappiDecoder jpeg_vaapi_decoder_;
|
||||
HipInteropDeviceMem hip_interop_;
|
||||
};
|
||||
|
||||
#endif //ROC_JPEG_DECODER_H_
|
||||
@@ -0,0 +1,961 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "rocjpeg_hip_kernels.h"
|
||||
|
||||
__device__ __forceinline__ uint32_t hipPack(float4 src) {
|
||||
return __builtin_amdgcn_cvt_pk_u8_f32(src.w, 3,
|
||||
__builtin_amdgcn_cvt_pk_u8_f32(src.z, 2,
|
||||
__builtin_amdgcn_cvt_pk_u8_f32(src.y, 1,
|
||||
__builtin_amdgcn_cvt_pk_u8_f32(src.x, 0, 0))));
|
||||
}
|
||||
|
||||
__device__ __forceinline__ float hipUnpack0(uint32_t src) {
|
||||
return (float)(src & 0xFF);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ float hipUnpack1(uint32_t src) {
|
||||
return (float)((src >> 8) & 0xFF);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ float hipUnpack2(uint32_t src) {
|
||||
return (float)((src >> 16) & 0xFF);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ float hipUnpack3(uint32_t src) {
|
||||
return (float)((src >> 24) & 0xFF);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ float4 hipUnpack(uint32_t src) {
|
||||
return make_float4(hipUnpack0(src), hipUnpack1(src), hipUnpack2(src), hipUnpack3(src));
|
||||
}
|
||||
|
||||
__global__ void ColorConvertYUV444ToRGBKernel(uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image, uint32_t dst_image_stride_in_bytes,
|
||||
uint32_t dst_image_stride_in_bytes_comp, const uint8_t *src_y_image, const uint8_t *src_u_image, const uint8_t *src_v_image,
|
||||
uint32_t src_yuv_image_stride_in_bytes, uint32_t dst_width_comp, uint32_t dst_height_comp, uint32_t src_yuv_image_stride_in_bytes_comp) {
|
||||
|
||||
int32_t x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
|
||||
int32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
|
||||
if ((x < dst_width_comp) && (y < dst_height_comp)) {
|
||||
uint32_t src_y0_idx = y * src_yuv_image_stride_in_bytes_comp + (x << 3);
|
||||
uint32_t src_y1_idx = src_y0_idx + src_yuv_image_stride_in_bytes;
|
||||
|
||||
|
||||
uint2 y0 = *((uint2 *)(&src_y_image[src_y0_idx]));
|
||||
uint2 y1 = *((uint2 *)(&src_y_image[src_y1_idx]));
|
||||
|
||||
uint2 u0 = *((uint2 *)(&src_u_image[src_y0_idx]));
|
||||
uint2 u1 = *((uint2 *)(&src_u_image[src_y1_idx]));
|
||||
|
||||
uint2 v0 = *((uint2 *)(&src_v_image[src_y0_idx]));
|
||||
uint2 v1 = *((uint2 *)(&src_v_image[src_y1_idx]));
|
||||
|
||||
uint32_t rgb0_idx = y * dst_image_stride_in_bytes_comp + (x * 24);
|
||||
uint32_t rgb1_idx = rgb0_idx + dst_image_stride_in_bytes;
|
||||
|
||||
float2 cr = make_float2( 0.0000f, 1.5748f);
|
||||
float2 cg = make_float2(-0.1873f, -0.4681f);
|
||||
float2 cb = make_float2( 1.8556f, 0.0000f);
|
||||
float3 yuv;
|
||||
DUINT6 rgb0, rgb1;
|
||||
float4 f;
|
||||
|
||||
yuv.x = hipUnpack0(y0.x);
|
||||
yuv.y = hipUnpack0(u0.x);
|
||||
yuv.z = hipUnpack0(v0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(y0.x);
|
||||
yuv.y = hipUnpack1(u0.x);
|
||||
yuv.z = hipUnpack1(v0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
rgb0.data[0] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(y0.x);
|
||||
yuv.y = hipUnpack2(u0.x);
|
||||
yuv.z = hipUnpack2(v0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
rgb0.data[1] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(y0.x);
|
||||
yuv.y = hipUnpack3(u0.x);
|
||||
yuv.z = hipUnpack3(v0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
rgb0.data[2] = hipPack(f);
|
||||
|
||||
yuv.x = hipUnpack0(y0.y);
|
||||
yuv.y = hipUnpack0(u0.y);
|
||||
yuv.z = hipUnpack0(v0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(y0.y);
|
||||
yuv.y = hipUnpack1(u0.y);
|
||||
yuv.z = hipUnpack1(v0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
rgb0.data[3] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(y0.y);
|
||||
yuv.y = hipUnpack2(u0.y);
|
||||
yuv.z = hipUnpack2(v0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
rgb0.data[4] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(y0.y);
|
||||
yuv.y = hipUnpack3(u0.y);
|
||||
yuv.z = hipUnpack3(v0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
rgb0.data[5] = hipPack(f);
|
||||
|
||||
yuv.x = hipUnpack0(y1.x);
|
||||
yuv.y = hipUnpack0(u1.x);
|
||||
yuv.z = hipUnpack0(v1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(y1.x);
|
||||
yuv.y = hipUnpack1(u1.x);
|
||||
yuv.z = hipUnpack1(v1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
rgb1.data[0] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(y1.x);
|
||||
yuv.y = hipUnpack2(u1.x);
|
||||
yuv.z = hipUnpack2(v1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
rgb1.data[1] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(y1.x);
|
||||
yuv.y = hipUnpack3(u1.x);
|
||||
yuv.z = hipUnpack3(v1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
rgb1.data[2] = hipPack(f);
|
||||
|
||||
yuv.x = hipUnpack0(y1.y);
|
||||
yuv.y = hipUnpack0(u1.y);
|
||||
yuv.z = hipUnpack0(v1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(y1.y);
|
||||
yuv.y = hipUnpack1(u1.y);
|
||||
yuv.z = hipUnpack1(v1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
rgb1.data[3] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(y1.y);
|
||||
yuv.y = hipUnpack2(u1.y);
|
||||
yuv.z = hipUnpack2(v1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
rgb1.data[4] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(y1.y);
|
||||
yuv.y = hipUnpack3(u1.y);
|
||||
yuv.z = hipUnpack3(v1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
rgb1.data[5] = hipPack(f);
|
||||
|
||||
*((DUINT6 *)(&dst_image[rgb0_idx])) = rgb0;
|
||||
*((DUINT6 *)(&dst_image[rgb1_idx])) = rgb1;
|
||||
}
|
||||
}
|
||||
|
||||
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) {
|
||||
|
||||
int32_t local_threads_x = 16;
|
||||
int32_t local_threads_y = 4;
|
||||
int32_t global_threads_x = (dst_width + 7) >> 3;
|
||||
int32_t global_threads_y = (dst_height + 1) >> 1;
|
||||
|
||||
uint32_t dst_width_comp = (dst_width + 7) / 8;
|
||||
uint32_t dst_height_comp = (dst_height + 1) / 2;
|
||||
uint32_t dst_image_stride_in_bytes_comp = dst_image_stride_in_bytes * 2;
|
||||
uint32_t src_yuv_image_stride_in_bytes_comp = src_yuv_image_stride_in_bytes * 2;
|
||||
|
||||
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, (uint8_t *)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,
|
||||
dst_width_comp, dst_height_comp, src_yuv_image_stride_in_bytes_comp);
|
||||
}
|
||||
|
||||
__global__ void ColorConvertYUYVToRGBKernel(uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *dst_image, uint32_t dst_image_stride_in_bytes, uint32_t dst_image_stride_in_bytes_comp,
|
||||
const uint8_t *src_image, uint32_t src_image_stride_in_bytes, uint32_t src_image_stride_in_bytes_comp,
|
||||
uint32_t dst_width_comp, uint32_t dst_height_comp) {
|
||||
|
||||
int32_t x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
|
||||
int32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
|
||||
if ((x < dst_width_comp) && (y < dst_height_comp)) {
|
||||
uint32_t l0_idx = y * src_image_stride_in_bytes_comp + (x << 4);
|
||||
uint32_t l1_idx = l0_idx + src_image_stride_in_bytes;
|
||||
uint4 l0 = *((uint4 *)(&src_image[l0_idx]));
|
||||
uint4 l1 = *((uint4 *)(&src_image[l1_idx]));
|
||||
|
||||
uint32_t rgb0_idx = y * dst_image_stride_in_bytes_comp + (x * 24);
|
||||
uint32_t rgb1_idx = rgb0_idx + dst_image_stride_in_bytes;
|
||||
|
||||
float4 f;
|
||||
|
||||
uint2 py0, py1;
|
||||
uint2 pu0, pu1;
|
||||
uint2 pv0, pv1;
|
||||
|
||||
py0.x = hipPack(make_float4(hipUnpack0(l0.x), hipUnpack2(l0.x), hipUnpack0(l0.y), hipUnpack2(l0.y)));
|
||||
py0.y = hipPack(make_float4(hipUnpack0(l0.z), hipUnpack2(l0.z), hipUnpack0(l0.w), hipUnpack2(l0.w)));
|
||||
py1.x = hipPack(make_float4(hipUnpack0(l1.x), hipUnpack2(l1.x), hipUnpack0(l1.y), hipUnpack2(l1.y)));
|
||||
py1.y = hipPack(make_float4(hipUnpack0(l1.z), hipUnpack2(l1.z), hipUnpack0(l1.w), hipUnpack2(l1.w)));
|
||||
pu0.x = hipPack(make_float4(hipUnpack1(l0.x), hipUnpack1(l0.x), hipUnpack1(l0.y), hipUnpack1(l0.y)));
|
||||
pu0.y = hipPack(make_float4(hipUnpack1(l0.z), hipUnpack1(l0.z), hipUnpack1(l0.w), hipUnpack1(l0.w)));
|
||||
pu1.x = hipPack(make_float4(hipUnpack1(l1.x), hipUnpack1(l1.x), hipUnpack1(l1.y), hipUnpack1(l1.y)));
|
||||
pu1.y = hipPack(make_float4(hipUnpack1(l1.z), hipUnpack1(l1.z), hipUnpack1(l1.w), hipUnpack1(l1.w)));
|
||||
pv0.x = hipPack(make_float4(hipUnpack3(l0.x), hipUnpack3(l0.x), hipUnpack3(l0.y), hipUnpack3(l0.y)));
|
||||
pv0.y = hipPack(make_float4(hipUnpack3(l0.z), hipUnpack3(l0.z), hipUnpack3(l0.w), hipUnpack3(l0.w)));
|
||||
pv1.x = hipPack(make_float4(hipUnpack3(l1.x), hipUnpack3(l1.x), hipUnpack3(l1.y), hipUnpack3(l1.y)));
|
||||
pv1.y = hipPack(make_float4(hipUnpack3(l1.z), hipUnpack3(l1.z), hipUnpack3(l1.w), hipUnpack3(l1.w)));
|
||||
|
||||
float2 cr = make_float2( 0.0000f, 1.5748f);
|
||||
float2 cg = make_float2(-0.1873f, -0.4681f);
|
||||
float2 cb = make_float2( 1.8556f, 0.0000f);
|
||||
float3 yuv;
|
||||
DUINT6 prgb0, prgb1;
|
||||
|
||||
yuv.x = hipUnpack0(py0.x);
|
||||
yuv.y = hipUnpack0(pu0.x);
|
||||
yuv.z = hipUnpack0(pv0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(py0.x);
|
||||
yuv.y = hipUnpack1(pu0.x);
|
||||
yuv.z = hipUnpack1(pv0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
prgb0.data[0] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(py0.x);
|
||||
yuv.y = hipUnpack2(pu0.x);
|
||||
yuv.z = hipUnpack2(pv0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
prgb0.data[1] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(py0.x);
|
||||
yuv.y = hipUnpack3(pu0.x);
|
||||
yuv.z = hipUnpack3(pv0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
prgb0.data[2] = hipPack(f);
|
||||
|
||||
yuv.x = hipUnpack0(py0.y);
|
||||
yuv.y = hipUnpack0(pu0.y);
|
||||
yuv.z = hipUnpack0(pv0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(py0.y);
|
||||
yuv.y = hipUnpack1(pu0.y);
|
||||
yuv.z = hipUnpack1(pv0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
prgb0.data[3] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(py0.y);
|
||||
yuv.y = hipUnpack2(pu0.y);
|
||||
yuv.z = hipUnpack2(pv0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
prgb0.data[4] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(py0.y);
|
||||
yuv.y = hipUnpack3(pu0.y);
|
||||
yuv.z = hipUnpack3(pv0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
prgb0.data[5] = hipPack(f);
|
||||
|
||||
yuv.x = hipUnpack0(py1.x);
|
||||
yuv.y = hipUnpack0(pu1.x);
|
||||
yuv.z = hipUnpack0(pv1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(py1.x);
|
||||
yuv.y = hipUnpack1(pu1.x);
|
||||
yuv.z = hipUnpack1(pv1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
prgb1.data[0] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(py1.x);
|
||||
yuv.y = hipUnpack2(pu1.x);
|
||||
yuv.z = hipUnpack2(pv1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
prgb1.data[1] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(py1.x);
|
||||
yuv.y = hipUnpack3(pu1.x);
|
||||
yuv.z = hipUnpack3(pv1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
prgb1.data[2] = hipPack(f);
|
||||
|
||||
yuv.x = hipUnpack0(py1.y);
|
||||
yuv.y = hipUnpack0(pu1.y);
|
||||
yuv.z = hipUnpack0(pv1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(py1.y);
|
||||
yuv.y = hipUnpack1(pu1.y);
|
||||
yuv.z = hipUnpack1(pv1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
prgb1.data[3] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(py1.y);
|
||||
yuv.y = hipUnpack2(pu1.y);
|
||||
yuv.z = hipUnpack2(pv1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
prgb1.data[4] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(py1.y);
|
||||
yuv.y = hipUnpack3(pu1.y);
|
||||
yuv.z = hipUnpack3(pv1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
prgb1.data[5] = hipPack(f);
|
||||
|
||||
*((DUINT6 *)(&dst_image[rgb0_idx])) = prgb0;
|
||||
*((DUINT6 *)(&dst_image[rgb1_idx])) = prgb1;
|
||||
}
|
||||
}
|
||||
|
||||
void ColorConvertYUYVToRGB(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_image, uint32_t src_image_stride_in_bytes) {
|
||||
int32_t local_threads_x = 16;
|
||||
int32_t local_threads_y = 4;
|
||||
int32_t global_threads_x = (dst_width + 7) >> 3;
|
||||
int32_t global_threads_y = (dst_height + 1) >> 1;
|
||||
|
||||
uint32_t dst_width_comp = (dst_width + 7) / 8;
|
||||
uint32_t dst_height_comp = (dst_height + 1) / 2;
|
||||
uint32_t dst_image_stride_in_bytes_comp = dst_image_stride_in_bytes * 2;
|
||||
uint32_t src_image_stride_in_bytes_comp = src_image_stride_in_bytes * 2;
|
||||
|
||||
ColorConvertYUYVToRGBKernel<<<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, (uint8_t *)dst_image,
|
||||
dst_image_stride_in_bytes, dst_image_stride_in_bytes_comp, src_image, src_image_stride_in_bytes,
|
||||
src_image_stride_in_bytes_comp, dst_width_comp, dst_height_comp);
|
||||
}
|
||||
|
||||
__global__ void ColorConvertNV12ToRGBKernel(uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *dst_image, uint32_t dst_image_stride_in_bytes, uint32_t dst_image_stride_in_bytes_comp,
|
||||
const uint8_t *src_luma_image, uint32_t src_luma_image_stride_in_bytes,
|
||||
const uint8_t *src_chroma_image, uint32_t src_chroma_image_stride_in_bytes,
|
||||
uint32_t dst_width_comp, uint32_t dst_height_comp, uint32_t src_luma_image_stride_in_bytes_comp) {
|
||||
|
||||
int32_t x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
|
||||
int32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
|
||||
if ((x < dst_width_comp) && (y < dst_height_comp)) {
|
||||
uint32_t src_y0_idx = y * src_luma_image_stride_in_bytes_comp + (x << 3);
|
||||
uint32_t src_y1_idx = src_y0_idx + src_luma_image_stride_in_bytes;
|
||||
uint32_t src_uv_idx = y * src_chroma_image_stride_in_bytes + (x << 3);
|
||||
uint2 y0 = *((uint2 *)(&src_luma_image[src_y0_idx]));
|
||||
uint2 y1 = *((uint2 *)(&src_luma_image[src_y1_idx]));
|
||||
uint2 uv = *((uint2 *)(&src_chroma_image[src_uv_idx]));
|
||||
|
||||
uint32_t rgb0_idx = y * dst_image_stride_in_bytes_comp + (x * 24);
|
||||
uint32_t rgb1_idx = rgb0_idx + dst_image_stride_in_bytes;
|
||||
|
||||
float4 f;
|
||||
uint2 u0, u1;
|
||||
uint2 v0, v1;
|
||||
|
||||
f.x = hipUnpack0(uv.x);
|
||||
f.y = f.x;
|
||||
f.z = hipUnpack2(uv.x);
|
||||
f.w = f.z;
|
||||
u0.x = hipPack(f);
|
||||
|
||||
f.x = hipUnpack0(uv.y);
|
||||
f.y = f.x;
|
||||
f.z = hipUnpack2(uv.y);
|
||||
f.w = f.z;
|
||||
u0.y = hipPack(f);
|
||||
|
||||
u1.x = u0.x;
|
||||
u1.y = u0.y;
|
||||
|
||||
f.x = hipUnpack1(uv.x);
|
||||
f.y = f.x;
|
||||
f.z = hipUnpack3(uv.x);
|
||||
f.w = f.z;
|
||||
v0.x = hipPack(f);
|
||||
|
||||
f.x = hipUnpack1(uv.y);
|
||||
f.y = f.x;
|
||||
f.z = hipUnpack3(uv.y);
|
||||
f.w = f.z;
|
||||
v0.y = hipPack(f);
|
||||
|
||||
v1.x = v0.x;
|
||||
v1.y = v0.y;
|
||||
|
||||
float2 cr = make_float2( 0.0000f, 1.5748f);
|
||||
float2 cg = make_float2(-0.1873f, -0.4681f);
|
||||
float2 cb = make_float2( 1.8556f, 0.0000f);
|
||||
float3 yuv;
|
||||
DUINT6 rgb0, rgb1;
|
||||
|
||||
yuv.x = hipUnpack0(y0.x);
|
||||
yuv.y = hipUnpack0(u0.x);
|
||||
yuv.z = hipUnpack0(v0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(y0.x);
|
||||
yuv.y = hipUnpack1(u0.x);
|
||||
yuv.z = hipUnpack1(v0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
rgb0.data[0] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(y0.x);
|
||||
yuv.y = hipUnpack2(u0.x);
|
||||
yuv.z = hipUnpack2(v0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
rgb0.data[1] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(y0.x);
|
||||
yuv.y = hipUnpack3(u0.x);
|
||||
yuv.z = hipUnpack3(v0.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
rgb0.data[2] = hipPack(f);
|
||||
|
||||
yuv.x = hipUnpack0(y0.y);
|
||||
yuv.y = hipUnpack0(u0.y);
|
||||
yuv.z = hipUnpack0(v0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(y0.y);
|
||||
yuv.y = hipUnpack1(u0.y);
|
||||
yuv.z = hipUnpack1(v0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
rgb0.data[3] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(y0.y);
|
||||
yuv.y = hipUnpack2(u0.y);
|
||||
yuv.z = hipUnpack2(v0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
rgb0.data[4] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(y0.y);
|
||||
yuv.y = hipUnpack3(u0.y);
|
||||
yuv.z = hipUnpack3(v0.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
rgb0.data[5] = hipPack(f);
|
||||
|
||||
yuv.x = hipUnpack0(y1.x);
|
||||
yuv.y = hipUnpack0(u1.x);
|
||||
yuv.z = hipUnpack0(v1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(y1.x);
|
||||
yuv.y = hipUnpack1(u1.x);
|
||||
yuv.z = hipUnpack1(v1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
rgb1.data[0] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(y1.x);
|
||||
yuv.y = hipUnpack2(u1.x);
|
||||
yuv.z = hipUnpack2(v1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
rgb1.data[1] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(y1.x);
|
||||
yuv.y = hipUnpack3(u1.x);
|
||||
yuv.z = hipUnpack3(v1.x);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
rgb1.data[2] = hipPack(f);
|
||||
|
||||
yuv.x = hipUnpack0(y1.y);
|
||||
yuv.y = hipUnpack0(u1.y);
|
||||
yuv.z = hipUnpack0(v1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.x = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.y = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.y = fmaf(cg.y, yuv.z, f.y);
|
||||
f.z = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack1(y1.y);
|
||||
yuv.y = hipUnpack1(u1.y);
|
||||
yuv.z = hipUnpack1(v1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.w = fmaf(cr.y, yuv.z, yuv.x);
|
||||
rgb1.data[3] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.x = fmaf(cg.y, yuv.z, f.x);
|
||||
f.y = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack2(y1.y);
|
||||
yuv.y = hipUnpack2(u1.y);
|
||||
yuv.z = hipUnpack2(v1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.z = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.w = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.w = fmaf(cg.y, yuv.z, f.w);
|
||||
rgb1.data[4] = hipPack(f);
|
||||
|
||||
f.x = fmaf(cb.x, yuv.y, yuv.x);
|
||||
yuv.x = hipUnpack3(y1.y);
|
||||
yuv.y = hipUnpack3(u1.y);
|
||||
yuv.z = hipUnpack3(v1.y);
|
||||
yuv.y -= 128.0f;
|
||||
yuv.z -= 128.0f;
|
||||
f.y = fmaf(cr.y, yuv.z, yuv.x);
|
||||
f.z = fmaf(cg.x, yuv.y, yuv.x);
|
||||
f.z = fmaf(cg.y, yuv.z, f.z);
|
||||
f.w = fmaf(cb.x, yuv.y, yuv.x);
|
||||
rgb1.data[5] = hipPack(f);
|
||||
|
||||
*((DUINT6 *)(&dst_image[rgb0_idx])) = rgb0;
|
||||
*((DUINT6 *)(&dst_image[rgb1_idx])) = rgb1;
|
||||
}
|
||||
}
|
||||
|
||||
void ColorConvertNV12ToRGB(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_luma_image, uint32_t src_luma_image_stride_in_bytes,
|
||||
const uint8_t *src_chroma_image, uint32_t src_chroma_image_stride_in_bytes) {
|
||||
int32_t local_threads_x = 16;
|
||||
int32_t local_threads_y = 4;
|
||||
int32_t global_threads_x = (dst_width + 7) >> 3;
|
||||
int32_t global_threads_y = (dst_height + 1) >> 1;
|
||||
|
||||
uint32_t dst_width_comp = (dst_width + 7) / 8;
|
||||
uint32_t dst_height_comp = (dst_height + 1) / 2;
|
||||
uint32_t dst_image_stride_in_bytes_comp = dst_image_stride_in_bytes * 2;
|
||||
uint32_t src_luma_image_stride_in_bytes_comp = src_luma_image_stride_in_bytes * 2;
|
||||
|
||||
ColorConvertNV12ToRGBKernel<<<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_luma_image, src_luma_image_stride_in_bytes, src_chroma_image,
|
||||
src_chroma_image_stride_in_bytes, dst_width_comp, dst_height_comp, src_luma_image_stride_in_bytes_comp);
|
||||
}
|
||||
|
||||
__global__ void ColorConvertYUV400ToRGBKernel(uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *dst_image, uint32_t dst_image_stride_in_bytes, uint32_t dst_image_stride_in_bytes_comp,
|
||||
const uint8_t *src_luma_image, uint32_t src_luma_image_stride_in_bytes,
|
||||
uint32_t dst_width_comp, uint32_t dst_height_comp, uint32_t src_luma_image_stride_in_bytes_comp) {
|
||||
|
||||
int32_t x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
|
||||
int32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
|
||||
if ((x < dst_width_comp) && (y < dst_height_comp)) {
|
||||
uint32_t src_y0_idx = y * src_luma_image_stride_in_bytes_comp + (x << 3);
|
||||
uint32_t src_y1_idx = src_y0_idx + src_luma_image_stride_in_bytes;
|
||||
|
||||
uint2 y0 = *((uint2 *)(&src_luma_image[src_y0_idx]));
|
||||
uint2 y1 = *((uint2 *)(&src_luma_image[src_y1_idx]));
|
||||
|
||||
uint32_t rgb0_idx = y * dst_image_stride_in_bytes_comp + (x * 24);
|
||||
uint32_t rgb1_idx = rgb0_idx + dst_image_stride_in_bytes;
|
||||
|
||||
DUINT6 rgb0, rgb1;
|
||||
|
||||
uint8_t y0_b0, y0_b1, y0_b2, y0_b3, y0_b4, y0_b5, y0_b6, y0_b7;
|
||||
uint8_t y1_b0, y1_b1, y1_b2, y1_b3, y1_b4, y1_b5, y1_b6, y1_b7;
|
||||
|
||||
y0_b0 = hipUnpack0(y0.x);
|
||||
y0_b1 = hipUnpack1(y0.x);
|
||||
y0_b2 = hipUnpack2(y0.x);
|
||||
y0_b3 = hipUnpack3(y0.x);
|
||||
y0_b4 = hipUnpack0(y0.y);
|
||||
y0_b5 = hipUnpack1(y0.y);
|
||||
y0_b6 = hipUnpack2(y0.y);
|
||||
y0_b7 = hipUnpack3(y0.y);
|
||||
|
||||
y1_b0 = hipUnpack0(y1.x);
|
||||
y1_b1 = hipUnpack1(y1.x);
|
||||
y1_b2 = hipUnpack2(y1.x);
|
||||
y1_b3 = hipUnpack3(y1.x);
|
||||
y1_b4 = hipUnpack0(y1.y);
|
||||
y1_b5 = hipUnpack1(y1.y);
|
||||
y1_b6 = hipUnpack2(y1.y);
|
||||
y1_b7 = hipUnpack3(y1.y);
|
||||
|
||||
rgb0.data[0] = hipPack(make_float4(y0_b0, y0_b0, y0_b0, y0_b1));
|
||||
rgb0.data[1] = hipPack(make_float4(y0_b1, y0_b1, y0_b2, y0_b2));
|
||||
rgb0.data[2] = hipPack(make_float4(y0_b2, y0_b3, y0_b3, y0_b3));
|
||||
rgb0.data[3] = hipPack(make_float4(y0_b4, y0_b4, y0_b4, y0_b5));
|
||||
rgb0.data[4] = hipPack(make_float4(y0_b5, y0_b5, y0_b6, y0_b6));
|
||||
rgb0.data[5] = hipPack(make_float4(y0_b6, y0_b7, y0_b7, y0_b7));
|
||||
|
||||
rgb1.data[0] = hipPack(make_float4(y1_b0, y1_b0, y1_b0, y1_b1));
|
||||
rgb1.data[1] = hipPack(make_float4(y1_b1, y1_b1, y1_b2, y1_b2));
|
||||
rgb1.data[2] = hipPack(make_float4(y1_b2, y1_b3, y1_b3, y1_b3));
|
||||
rgb1.data[3] = hipPack(make_float4(y1_b4, y1_b4, y1_b4, y1_b5));
|
||||
rgb1.data[4] = hipPack(make_float4(y1_b5, y1_b5, y1_b6, y1_b6));
|
||||
rgb1.data[5] = hipPack(make_float4(y1_b6, y1_b7, y1_b7, y1_b7));
|
||||
|
||||
*((DUINT6 *)(&dst_image[rgb0_idx])) = rgb0;
|
||||
*((DUINT6 *)(&dst_image[rgb1_idx])) = rgb1;
|
||||
}
|
||||
}
|
||||
|
||||
void ColorConvertYUV400ToRGB(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_luma_image, uint32_t src_luma_image_stride_in_bytes){
|
||||
|
||||
int32_t local_threads_x = 16;
|
||||
int32_t local_threads_y = 4;
|
||||
int32_t global_threads_x = (dst_width + 7) >> 3;
|
||||
int32_t global_threads_y = (dst_height + 1) >> 1;
|
||||
|
||||
uint32_t dst_width_comp = (dst_width + 7) / 8;
|
||||
uint32_t dst_height_comp = (dst_height + 1) / 2;
|
||||
uint32_t dst_image_stride_in_bytes_comp = dst_image_stride_in_bytes * 2;
|
||||
uint32_t src_luma_image_stride_in_bytes_comp = src_luma_image_stride_in_bytes * 2;
|
||||
|
||||
ColorConvertYUV400ToRGBKernel<<<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_luma_image, src_luma_image_stride_in_bytes, dst_width_comp, dst_height_comp,
|
||||
src_luma_image_stride_in_bytes_comp);
|
||||
|
||||
}
|
||||
|
||||
|
||||
__global__ void ConvertInterleavedUVToPlanarUVKernel(uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *dst_image1, uint8_t *dst_image2, uint32_t dst_image_stride_in_bytes,
|
||||
const uint8_t *src_image, uint32_t src_image_stride_in_bytes) {
|
||||
|
||||
int32_t x = (hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x) * 8;
|
||||
int32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
|
||||
if (x >= dst_width || y >= dst_height) {
|
||||
return;
|
||||
}
|
||||
|
||||
uint32_t src_idx = y * src_image_stride_in_bytes + x + x;
|
||||
uint32_t dst_idx = y * dst_image_stride_in_bytes + x;
|
||||
|
||||
uint4 src = *((uint4 *)(&src_image[src_idx]));
|
||||
uint2 dst1, dst2;
|
||||
|
||||
dst1.x = hipPack(make_float4(hipUnpack0(src.x), hipUnpack2(src.x), hipUnpack0(src.y), hipUnpack2(src.y)));
|
||||
dst1.y = hipPack(make_float4(hipUnpack0(src.z), hipUnpack2(src.z), hipUnpack0(src.w), hipUnpack2(src.w)));
|
||||
dst2.x = hipPack(make_float4(hipUnpack1(src.x), hipUnpack3(src.x), hipUnpack1(src.y), hipUnpack3(src.y)));
|
||||
dst2.y = hipPack(make_float4(hipUnpack1(src.z), hipUnpack3(src.z), hipUnpack1(src.w), hipUnpack3(src.w)));
|
||||
|
||||
*((uint2 *)(&dst_image1[dst_idx])) = dst1;
|
||||
*((uint2 *)(&dst_image2[dst_idx])) = dst2;
|
||||
|
||||
}
|
||||
void ConvertInterleavedUVToPlanarUV(hipStream_t stream, uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *dst_image1, uint8_t *dst_image2, uint32_t dst_image_stride_in_bytes,
|
||||
const uint8_t *src_image1, uint32_t src_image1_stride_in_bytes) {
|
||||
int32_t local_threads_x = 16, local_threads_y = 16;
|
||||
int32_t global_threads_x = (dst_width + 7) >> 3;
|
||||
int32_t global_threads_y = dst_height;
|
||||
|
||||
ConvertInterleavedUVToPlanarUVKernel<<<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_image1, dst_image2,
|
||||
dst_image_stride_in_bytes, src_image1, src_image1_stride_in_bytes);
|
||||
|
||||
}
|
||||
|
||||
__global__ void ExtractYFromPackedYUYVKernel(uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *destination_y, uint32_t dst_luma_stride_in_bytes,
|
||||
const uint8_t *src_image, uint32_t src_image_stride_in_bytes,
|
||||
uint32_t dst_width_comp) {
|
||||
|
||||
int32_t x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
|
||||
int32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
|
||||
if (x < dst_width_comp && y < dst_height) {
|
||||
uint32_t src_idx = y * src_image_stride_in_bytes + (x << 4);
|
||||
uint32_t dst_idx = y * dst_luma_stride_in_bytes + (x << 3);
|
||||
|
||||
uint4 src = *((uint4 *)(&src_image[src_idx]));
|
||||
uint2 dst_y;
|
||||
dst_y.x = hipPack(make_float4(hipUnpack0(src.x), hipUnpack2(src.x), hipUnpack0(src.y), hipUnpack2(src.y)));
|
||||
dst_y.y = hipPack(make_float4(hipUnpack0(src.z), hipUnpack2(src.z), hipUnpack0(src.w), hipUnpack2(src.w)));
|
||||
|
||||
*((uint2 *)(&destination_y[dst_idx])) = dst_y;
|
||||
}
|
||||
}
|
||||
|
||||
void ExtractYFromPackedYUYV(hipStream_t stream, uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *destination_y, uint32_t dst_luma_stride_in_bytes, const uint8_t *src_image, uint32_t src_image_stride_in_bytes) {
|
||||
int32_t local_threads_x = 16;
|
||||
int32_t local_threads_y = 4;
|
||||
int32_t global_threads_x = (dst_width + 7) >> 3;
|
||||
int32_t global_threads_y = dst_height;
|
||||
|
||||
uint32_t dst_width_comp = (dst_width + 7) / 8;
|
||||
|
||||
ExtractYFromPackedYUYVKernel<<<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, destination_y,
|
||||
dst_luma_stride_in_bytes, src_image, src_image_stride_in_bytes, dst_width_comp);
|
||||
}
|
||||
|
||||
__global__ void ConvertPackedYUYVToPlanarYUVKernel(uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *destination_y, uint8_t *destination_u, uint8_t *destination_v, uint32_t dst_luma_stride_in_bytes, uint32_t dst_chroma_stride_in_bytes,
|
||||
const uint8_t *src_image, uint32_t src_image_stride_in_bytes,
|
||||
uint32_t dst_width_comp) {
|
||||
|
||||
int32_t x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
|
||||
int32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
|
||||
if ((x < dst_width_comp && y < dst_height)) {
|
||||
uint32_t src_idx = y * src_image_stride_in_bytes + (x << 4);
|
||||
uint32_t dst_y_idx = y * dst_luma_stride_in_bytes + (x << 3);
|
||||
uint32_t dst_uv_idx = y * dst_chroma_stride_in_bytes + (x << 2);
|
||||
|
||||
uint4 src = *((uint4 *)(&src_image[src_idx]));
|
||||
uint2 dst_y;
|
||||
uint32_t dst_u, dst_v;
|
||||
|
||||
dst_y.x = hipPack(make_float4(hipUnpack0(src.x), hipUnpack2(src.x), hipUnpack0(src.y), hipUnpack2(src.y)));
|
||||
dst_y.y = hipPack(make_float4(hipUnpack0(src.z), hipUnpack2(src.z), hipUnpack0(src.w), hipUnpack2(src.w)));
|
||||
dst_u = hipPack(make_float4(hipUnpack1(src.x), hipUnpack1(src.y), hipUnpack1(src.z), hipUnpack1(src.w)));
|
||||
dst_v = hipPack(make_float4(hipUnpack3(src.x), hipUnpack3(src.y), hipUnpack3(src.z), hipUnpack3(src.w)));
|
||||
|
||||
*((uint2 *)(&destination_y[dst_y_idx])) = dst_y;
|
||||
*((uint32_t *)(&destination_u[dst_uv_idx])) = dst_u;
|
||||
*((uint32_t *)(&destination_v[dst_uv_idx])) = dst_v;
|
||||
}
|
||||
}
|
||||
|
||||
void ConvertPackedYUYVToPlanarYUV(hipStream_t stream, uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *destination_y, uint8_t *destination_u, uint8_t *destination_v, uint32_t dst_luma_stride_in_bytes, uint32_t dst_chroma_stride_in_bytes,
|
||||
const uint8_t *src_image, uint32_t src_image_stride_in_bytes) {
|
||||
|
||||
int32_t local_threads_x = 16;
|
||||
int32_t local_threads_y = 4;
|
||||
int32_t global_threads_x = (dst_width + 7) >> 3;
|
||||
int32_t global_threads_y = dst_height;
|
||||
uint32_t dst_width_comp = (dst_width + 7) / 8;
|
||||
|
||||
ConvertPackedYUYVToPlanarYUVKernel<<<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, destination_y, destination_u,
|
||||
destination_v, dst_luma_stride_in_bytes, dst_chroma_stride_in_bytes, src_image, src_image_stride_in_bytes, dst_width_comp);
|
||||
}
|
||||
@@ -0,0 +1,62 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#ifndef ROC_JPEG_HIP_KERNELS_H_
|
||||
#define ROC_JPEG_HIP_KERNELS_H_
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <hip/hip_runtime.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);
|
||||
|
||||
void ColorConvertYUYVToRGB(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_image, uint32_t src_image_stride_in_bytes);
|
||||
|
||||
void ColorConvertNV12ToRGB(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_luma_image, uint32_t src_luma_image_stride_in_bytes,
|
||||
const uint8_t *src_chroma_image, uint32_t src_chroma_image_stride_in_bytes);
|
||||
|
||||
void ColorConvertYUV400ToRGB(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_luma_image, uint32_t src_luma_image_stride_in_bytes);
|
||||
|
||||
void ConvertInterleavedUVToPlanarUV(hipStream_t stream, uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *dst_image1, uint8_t *dst_image2, uint32_t dst_image_stride_in_bytes,
|
||||
const uint8_t *src_image1, uint32_t src_image1_stride_in_bytes);
|
||||
|
||||
void ExtractYFromPackedYUYV(hipStream_t stream, uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *destination_y, uint32_t dst_luma_stride_in_bytes, const uint8_t *src_image, uint32_t src_image_stride_in_bytes);
|
||||
|
||||
void ConvertPackedYUYVToPlanarYUV(hipStream_t stream, uint32_t dst_width, uint32_t dst_height,
|
||||
uint8_t *destination_y, uint8_t *destination_u, uint8_t *destination_v, uint32_t dst_luma_stride_in_bytes,
|
||||
uint32_t dst_chroma_stride_in_bytes, const uint8_t *src_image, uint32_t src_image_stride_in_bytes);
|
||||
|
||||
typedef struct UINT6TYPE {
|
||||
uint data[6];
|
||||
} DUINT6;
|
||||
|
||||
#endif //ROC_JPEG_HIP_KERNELS_H_
|
||||
@@ -0,0 +1,387 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#include "rocjpeg_parser.h"
|
||||
|
||||
JpegParser::JpegParser() : stream_{nullptr}, stream_end_{nullptr}, stream_length_{0},
|
||||
jpeg_stream_parameters_{{}} {
|
||||
}
|
||||
|
||||
JpegParser::~JpegParser() {
|
||||
stream_ = nullptr;
|
||||
stream_end_ = nullptr;
|
||||
stream_length_ = 0;
|
||||
}
|
||||
|
||||
bool JpegParser::ParseJpegStream(const uint8_t *jpeg_stream, uint32_t jpeg_stream_size) {
|
||||
if (jpeg_stream == nullptr) {
|
||||
ERR("invalid argument!");
|
||||
return false;
|
||||
}
|
||||
|
||||
stream_ = jpeg_stream;
|
||||
stream_length_ = jpeg_stream_size;
|
||||
stream_end_ = stream_ + stream_length_;
|
||||
|
||||
jpeg_stream_parameters_ = {};
|
||||
bool soi_marker_found = false;
|
||||
bool sos_marker_found = false;
|
||||
bool dht_marker_found = false;
|
||||
bool dqt_marker_found = false;
|
||||
uint8_t marker;
|
||||
const uint8_t *next_chunck;
|
||||
int32_t chuck_len;
|
||||
|
||||
// The first two bytes of a JPEG must be 0XFFD8
|
||||
if (*stream_ != 0xFF || *(stream_ + 1) != SOI) {
|
||||
ERR("Invalid JPEG!");
|
||||
return false;
|
||||
}
|
||||
|
||||
soi_marker_found = ParseSOI();
|
||||
if (!soi_marker_found) {
|
||||
ERR("failed to find the SOI marker!");
|
||||
}
|
||||
|
||||
while (!sos_marker_found && stream_ <= stream_end_) {
|
||||
while ((*stream_ == 0xFF))
|
||||
stream_++;
|
||||
marker = *stream_++;
|
||||
chuck_len = swap_bytes(stream_);
|
||||
next_chunck = stream_ + chuck_len;
|
||||
|
||||
switch (marker) {
|
||||
case SOF:
|
||||
if (!ParseSOF())
|
||||
return false;
|
||||
break;
|
||||
case DHT:
|
||||
if (!ParseDHT())
|
||||
return false;
|
||||
dht_marker_found = true;
|
||||
break;
|
||||
case DQT:
|
||||
if (!ParseDQT())
|
||||
return false;
|
||||
dqt_marker_found = true;
|
||||
break;
|
||||
case DRI:
|
||||
if (!ParseDRI())
|
||||
return false;
|
||||
break;
|
||||
case SOS:
|
||||
if (!ParseSOS())
|
||||
return false;
|
||||
sos_marker_found = true;
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
stream_ = next_chunck;
|
||||
}
|
||||
|
||||
if (!dht_marker_found) {
|
||||
ERR("didn't find any Huffman table!");
|
||||
return false;
|
||||
}
|
||||
if (!dqt_marker_found) {
|
||||
ERR("didn't find any quantization table!");
|
||||
return false;
|
||||
}
|
||||
|
||||
if (!ParseEOI())
|
||||
return false;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool JpegParser::ParseSOI() {
|
||||
if (stream_ == nullptr) {
|
||||
return false;
|
||||
}
|
||||
while (!(*stream_ == 0xFF && *(stream_ + 1) == SOI)) {
|
||||
if (stream_ <= stream_end_) {
|
||||
stream_++;
|
||||
continue;
|
||||
} else
|
||||
return false;
|
||||
}
|
||||
stream_ += 2;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool JpegParser::ParseSOF() {
|
||||
uint32_t component_id, sampling_factor;
|
||||
uint8_t quantiser_table_selector;
|
||||
|
||||
if (stream_ == nullptr) {
|
||||
return false;
|
||||
}
|
||||
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.picture_height = swap_bytes(stream_ + 3);
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.picture_width = swap_bytes(stream_ + 5);
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.num_components = stream_[7];
|
||||
|
||||
if (jpeg_stream_parameters_.picture_parameter_buffer.num_components > NUM_COMPONENTS - 1) {
|
||||
ERR("invalid number of JPEG components!");
|
||||
return false;
|
||||
}
|
||||
|
||||
stream_ += 8;
|
||||
|
||||
for (int32_t i = 0; i < jpeg_stream_parameters_.picture_parameter_buffer.num_components; i++) {
|
||||
component_id = *stream_++;
|
||||
sampling_factor = *stream_++;
|
||||
quantiser_table_selector = *stream_++;
|
||||
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.components[i].component_id = component_id;
|
||||
if (quantiser_table_selector >= NUM_COMPONENTS) {
|
||||
ERR("invalid number of the quantization table!");
|
||||
return false;
|
||||
}
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.components[i].v_sampling_factor = sampling_factor & 0xF;
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.components[i].h_sampling_factor = sampling_factor >> 4;
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.components[i].quantiser_table_selector = quantiser_table_selector;
|
||||
}
|
||||
|
||||
uint8_t max_h_factor = jpeg_stream_parameters_.picture_parameter_buffer.components[0].h_sampling_factor;
|
||||
uint8_t max_v_factor = jpeg_stream_parameters_.picture_parameter_buffer.components[0].v_sampling_factor;
|
||||
|
||||
jpeg_stream_parameters_.slice_parameter_buffer.num_mcus = ((jpeg_stream_parameters_.picture_parameter_buffer.picture_width + max_h_factor * 8 - 1) / (max_h_factor * 8)) *
|
||||
((jpeg_stream_parameters_.picture_parameter_buffer.picture_height + max_v_factor * 8 - 1) / (max_v_factor * 8));
|
||||
|
||||
jpeg_stream_parameters_.chroma_subsampling = GetChromaSubsampling(jpeg_stream_parameters_.picture_parameter_buffer.components[0].h_sampling_factor,
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.components[1].h_sampling_factor,
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.components[2].h_sampling_factor,
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.components[0].v_sampling_factor,
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.components[1].v_sampling_factor,
|
||||
jpeg_stream_parameters_.picture_parameter_buffer.components[2].v_sampling_factor);
|
||||
return true;
|
||||
}
|
||||
|
||||
bool JpegParser::ParseDQT() {
|
||||
int32_t quantization_table_index = 0;
|
||||
const uint8_t *dqt_block_end;
|
||||
|
||||
if (stream_ == nullptr) {
|
||||
return false;
|
||||
}
|
||||
|
||||
dqt_block_end = stream_ + swap_bytes(stream_);
|
||||
stream_ += 2;
|
||||
|
||||
while (stream_ < dqt_block_end) {
|
||||
quantization_table_index = *stream_++;
|
||||
if (quantization_table_index >> 4) {
|
||||
ERR("16 bits quantization table is not supported!");
|
||||
return false;
|
||||
}
|
||||
if (quantization_table_index >= 4) {
|
||||
ERR("invalid number of quantization table!");
|
||||
return false;
|
||||
}
|
||||
|
||||
std::memcpy(jpeg_stream_parameters_.quantization_matrix_buffer.quantiser_table[quantization_table_index & 0x0F], stream_, 64);
|
||||
jpeg_stream_parameters_.quantization_matrix_buffer.load_quantiser_table[quantization_table_index & 0x0F] = 1;
|
||||
|
||||
stream_ += 64;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool JpegParser::ParseDHT() {
|
||||
uint32_t count, i;
|
||||
int32_t length, index;
|
||||
uint8_t ac_huffman_table, huffman_table_id;
|
||||
|
||||
if (stream_ == nullptr) {
|
||||
return false;
|
||||
}
|
||||
|
||||
length = swap_bytes(stream_) - 2;
|
||||
stream_ += 2;
|
||||
|
||||
while (length > 0) {
|
||||
index = *stream_++;
|
||||
|
||||
ac_huffman_table = index & 0xF0;
|
||||
huffman_table_id = index & 0x0F;
|
||||
|
||||
if (huffman_table_id >= HUFFMAN_TABLES) {
|
||||
ERR("invlaid number of Huffman table!");
|
||||
return false;
|
||||
}
|
||||
|
||||
if (ac_huffman_table) {
|
||||
std::memcpy(jpeg_stream_parameters_.huffman_table_buffer.huffman_table[huffman_table_id].num_ac_codes, stream_, 16);
|
||||
} else {
|
||||
std::memcpy(jpeg_stream_parameters_.huffman_table_buffer.huffman_table[huffman_table_id].num_dc_codes, stream_, 16);
|
||||
}
|
||||
|
||||
count = 0;
|
||||
for (i = 0; i < 16; i++) {
|
||||
count += *stream_++;
|
||||
}
|
||||
|
||||
if (ac_huffman_table) {
|
||||
if (count > AC_HUFFMAN_TABLE_VALUES_SIZE) {
|
||||
ERR("invalid AC Huffman table!");
|
||||
return false;
|
||||
}
|
||||
std::memcpy(jpeg_stream_parameters_.huffman_table_buffer.huffman_table[huffman_table_id].ac_values, stream_, count);
|
||||
jpeg_stream_parameters_.huffman_table_buffer.load_huffman_table[huffman_table_id] = 1;
|
||||
} else {
|
||||
if (count > DC_HUFFMAN_TABLE_VALUES_SIZE) {
|
||||
ERR("invlaid DC Huffman table!")
|
||||
return false;
|
||||
}
|
||||
std::memcpy(jpeg_stream_parameters_.huffman_table_buffer.huffman_table[huffman_table_id].dc_values, stream_, count);
|
||||
jpeg_stream_parameters_.huffman_table_buffer.load_huffman_table[huffman_table_id] = 1;
|
||||
}
|
||||
|
||||
length -= 1;
|
||||
length -= 16;
|
||||
length -= count;
|
||||
stream_ += count;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool JpegParser::ParseSOS() {
|
||||
uint32_t component_id, table;
|
||||
|
||||
if (stream_ == nullptr) {
|
||||
return false;
|
||||
}
|
||||
|
||||
uint32_t num_components = stream_[2];
|
||||
|
||||
if (num_components > NUM_COMPONENTS - 1) {
|
||||
ERR("invalid number of component!")
|
||||
return false;
|
||||
}
|
||||
jpeg_stream_parameters_.slice_parameter_buffer.num_components = num_components;
|
||||
|
||||
stream_ += 3;
|
||||
for (int32_t i = 0; i < num_components; i++) {
|
||||
component_id = *stream_++;
|
||||
table = *stream_++;
|
||||
jpeg_stream_parameters_.slice_parameter_buffer.components[i].component_selector = component_id;
|
||||
jpeg_stream_parameters_.slice_parameter_buffer.components[i].dc_table_selector = ((table >> 4) & 0x0F);
|
||||
jpeg_stream_parameters_.slice_parameter_buffer.components[i].ac_table_selector = (table & 0x0F);
|
||||
|
||||
if ((table & 0xF) >= 4) {
|
||||
ERR("invalid number of AC Huffman table!");
|
||||
return false;
|
||||
}
|
||||
if ((table >> 4) >= 4) {
|
||||
ERR("invalid number of DC Huffman table!");
|
||||
return false;
|
||||
}
|
||||
if (component_id != jpeg_stream_parameters_.picture_parameter_buffer.components[i].component_id) {
|
||||
ERR("component id mismatch between SOS and SOF marker!");
|
||||
return false;
|
||||
}
|
||||
}
|
||||
stream_ += 3;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
|
||||
bool JpegParser::ParseDRI() {
|
||||
uint32_t length;
|
||||
|
||||
if (stream_ == nullptr) {
|
||||
return false;
|
||||
}
|
||||
|
||||
length = swap_bytes(stream_);
|
||||
if (length != 4) {
|
||||
ERR("invalid size for DRI marker");
|
||||
return false;
|
||||
}
|
||||
|
||||
jpeg_stream_parameters_.slice_parameter_buffer.restart_interval = swap_bytes(stream_ + 2);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool JpegParser::ParseEOI() {
|
||||
|
||||
if (stream_ == nullptr) {
|
||||
return false;
|
||||
}
|
||||
|
||||
const uint8_t *stream_temp = stream_;
|
||||
while (stream_temp <= stream_end_ && !(*stream_temp == 0xFF && *(stream_temp + 1) == EOI)) {
|
||||
stream_temp++;
|
||||
continue;
|
||||
}
|
||||
|
||||
jpeg_stream_parameters_.slice_parameter_buffer.slice_data_size = stream_temp - stream_;
|
||||
jpeg_stream_parameters_.slice_data_buffer = stream_;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
ChromaSubsampling JpegParser::GetChromaSubsampling(uint8_t c1_h_sampling_factor, uint8_t c2_h_sampling_factor, uint8_t c3_h_sampling_factor,
|
||||
uint8_t c1_v_sampling_factor, uint8_t c2_v_sampling_factor, uint8_t c3_v_sampling_factor) {
|
||||
|
||||
ChromaSubsampling subsampling;
|
||||
|
||||
if ((c1_h_sampling_factor == 1 && c2_h_sampling_factor == 1 && c3_h_sampling_factor == 1 &&
|
||||
c1_v_sampling_factor == 1 && c2_v_sampling_factor == 1 && c3_v_sampling_factor == 1) ||
|
||||
(c1_h_sampling_factor == 2 && c2_h_sampling_factor == 2 && c3_h_sampling_factor == 2 &&
|
||||
c1_v_sampling_factor == 2 && c2_v_sampling_factor == 2 && c3_v_sampling_factor == 2) ||
|
||||
(c1_h_sampling_factor == 4 && c2_h_sampling_factor == 4 && c3_h_sampling_factor == 4 &&
|
||||
c1_v_sampling_factor == 4 && c2_v_sampling_factor == 4 && c3_v_sampling_factor == 4)) {
|
||||
subsampling = CSS_444;
|
||||
} else if (c1_h_sampling_factor == 1 && c2_h_sampling_factor == 1 && c3_h_sampling_factor == 1 &&
|
||||
c1_v_sampling_factor == 2 && c2_v_sampling_factor == 1 && c3_v_sampling_factor == 1) {
|
||||
subsampling = CSS_440;
|
||||
} else if ((c1_h_sampling_factor == 2 && c2_h_sampling_factor == 1 && c3_h_sampling_factor == 1 &&
|
||||
c1_v_sampling_factor == 1 && c2_v_sampling_factor == 1 && c3_v_sampling_factor == 1) ||
|
||||
(c1_h_sampling_factor == 2 && c2_h_sampling_factor == 1 && c3_h_sampling_factor == 1 &&
|
||||
c1_v_sampling_factor == 2 && c2_v_sampling_factor == 2 && c3_v_sampling_factor == 2) ||
|
||||
(c1_h_sampling_factor == 2 && c2_h_sampling_factor == 2 && c3_h_sampling_factor == 2 &&
|
||||
c1_v_sampling_factor == 2 && c2_v_sampling_factor == 1 && c3_v_sampling_factor == 1)) {
|
||||
subsampling = CSS_422;
|
||||
} else if (c1_h_sampling_factor == 2 && c2_h_sampling_factor == 1 && c3_h_sampling_factor == 1 &&
|
||||
c1_v_sampling_factor == 2 && c2_v_sampling_factor == 1 && c3_v_sampling_factor == 1) {
|
||||
subsampling = CSS_420;
|
||||
} else if (c1_h_sampling_factor == 4 && c2_h_sampling_factor == 1 && c3_h_sampling_factor == 1 &&
|
||||
c1_v_sampling_factor == 1 && c2_v_sampling_factor == 1 && c3_v_sampling_factor == 1) {
|
||||
subsampling = CSS_411;
|
||||
} else if ((c1_h_sampling_factor == 1 && c2_h_sampling_factor == 0 && c3_h_sampling_factor == 0 &&
|
||||
c1_v_sampling_factor == 1 && c2_v_sampling_factor == 0 && c3_v_sampling_factor == 0) ||
|
||||
(c1_h_sampling_factor == 4 && c2_h_sampling_factor == 0 && c3_h_sampling_factor == 0 &&
|
||||
c1_v_sampling_factor == 4 && c2_v_sampling_factor == 0 && c3_v_sampling_factor == 0)) {
|
||||
subsampling = CSS_400;
|
||||
} else {
|
||||
subsampling = CSS_UNKNOWN;
|
||||
}
|
||||
|
||||
return subsampling;
|
||||
}
|
||||
@@ -0,0 +1,177 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
|
||||
#ifndef ROC_JPEG_PARSER_H_
|
||||
#define ROC_JPEG_PARSER_H_
|
||||
|
||||
#include <stdint.h>
|
||||
#include <iostream>
|
||||
#include <cstring>
|
||||
#include "rocjpeg_commons.h"
|
||||
|
||||
#pragma once
|
||||
|
||||
#define NUM_COMPONENTS 4
|
||||
#define HUFFMAN_TABLES 2
|
||||
#define AC_HUFFMAN_TABLE_VALUES_SIZE 162
|
||||
#define DC_HUFFMAN_TABLE_VALUES_SIZE 12
|
||||
#define swap_bytes(x) (((x)[0] << 8) | (x)[1])
|
||||
|
||||
/***************************************************************/
|
||||
//! \enum enum JpegMarker
|
||||
//! common JPEG markers
|
||||
/***************************************************************/
|
||||
enum JpegMarkers {
|
||||
SOI = 0xD8, /* Start Of Image */
|
||||
SOF = 0xC0, /* Start Of Frame for a baseline DCT-based JPEG. */
|
||||
DHT = 0xC4, /* Define Huffman Table */
|
||||
DQT = 0xDB, /* Define Quantization Table */
|
||||
DRI = 0xDD, /* Define Restart Interval */
|
||||
SOS = 0xDA, /* Start of Scan */
|
||||
EOI = 0xD9, /* End Of Image */
|
||||
};
|
||||
|
||||
/***************************************************************/
|
||||
//! \struct Picture parameter for JPEG decoding.
|
||||
//! This structure holds information from the frame
|
||||
//! header and definitions from additional segments.
|
||||
/**************************************************************/
|
||||
typedef struct PictureParameterBufferType {
|
||||
uint16_t picture_width;
|
||||
uint16_t picture_height;
|
||||
struct {
|
||||
uint8_t component_id;
|
||||
uint8_t h_sampling_factor;
|
||||
uint8_t v_sampling_factor;
|
||||
uint8_t quantiser_table_selector;
|
||||
} components[255];
|
||||
uint8_t num_components;
|
||||
uint8_t color_space;
|
||||
uint32_t rotation;
|
||||
uint32_t reserved[7];
|
||||
} PictureParameterBuffer;
|
||||
|
||||
/***************************************************************/
|
||||
//! \struct Quantization table for JPEG decoding.
|
||||
//! This structure holds the quantization tables.
|
||||
//! The maximum number of quatization tables is four.
|
||||
//! The #load_quantization_table array can be used as a hint to notify
|
||||
//! which table(s) has valid values.
|
||||
//! The #quantiser_table values are specified in zig-zag scan order.
|
||||
/***************************************************************/
|
||||
typedef struct QuantizationMatrixBufferType {
|
||||
uint8_t load_quantiser_table[4];
|
||||
uint8_t quantiser_table[4][64];
|
||||
uint32_t reserved[4];
|
||||
} QuantizationMatrixBuffer;
|
||||
|
||||
/***************************************************************/
|
||||
//! \struct Huffman table for JPEG decoding.
|
||||
//! This structure holds the Huffman tables.
|
||||
//! The maximum number of Huffman tables is two.
|
||||
//! The #load_huffman_table array can be used as a hint to notify the
|
||||
//! which table(s) has valid values.
|
||||
/***************************************************************/
|
||||
typedef struct HuffmanTableBufferType {
|
||||
uint8_t load_huffman_table[2];
|
||||
struct {
|
||||
uint8_t num_dc_codes[16];
|
||||
uint8_t dc_values[12];
|
||||
uint8_t num_ac_codes[16];
|
||||
uint8_t ac_values[162];
|
||||
uint8_t pad[2];
|
||||
} huffman_table[2];
|
||||
uint32_t reserved[4];
|
||||
} HuffmanTableBuffer;
|
||||
|
||||
/***************************************************************/
|
||||
//! \struct Slice parameter for JPEG decoding.
|
||||
//! This structure holds information from the scan header, and
|
||||
//! definitions from additional segments.
|
||||
/***************************************************************/
|
||||
typedef struct SliceParameterBufferType {
|
||||
uint32_t slice_data_size;
|
||||
uint32_t slice_data_offset;
|
||||
uint32_t slice_data_flag;
|
||||
uint32_t slice_horizontal_position;
|
||||
uint32_t slice_vertical_position;
|
||||
struct {
|
||||
uint8_t component_selector;
|
||||
uint8_t dc_table_selector;
|
||||
uint8_t ac_table_selector;
|
||||
} components[4];
|
||||
uint8_t num_components;
|
||||
uint16_t restart_interval;
|
||||
uint32_t num_mcus;
|
||||
uint32_t reserved[4];
|
||||
} SliceParameterBuffer;
|
||||
|
||||
/***************************************************************/
|
||||
//! \struct Enum identifies image chroma subsampling values stored inside JPEG input stream
|
||||
/***************************************************************/
|
||||
typedef enum {
|
||||
CSS_444 = 0,
|
||||
CSS_440 = 1,
|
||||
CSS_422 = 2,
|
||||
CSS_420 = 3,
|
||||
CSS_411 = 4,
|
||||
CSS_400 = 5,
|
||||
CSS_UNKNOWN = -1
|
||||
} ChromaSubsampling;
|
||||
|
||||
/***************************************************************/
|
||||
//! \struct Jpeg stream parameters.
|
||||
//! This structure holds all information from a JPEG stream for decoding
|
||||
/***************************************************************/
|
||||
typedef struct JpegParameterBuffersType {
|
||||
PictureParameterBuffer picture_parameter_buffer;
|
||||
QuantizationMatrixBuffer quantization_matrix_buffer;
|
||||
HuffmanTableBuffer huffman_table_buffer;
|
||||
SliceParameterBuffer slice_parameter_buffer;
|
||||
ChromaSubsampling chroma_subsampling;
|
||||
const uint8_t* slice_data_buffer;
|
||||
} JpegStreamParameters;
|
||||
|
||||
class JpegParser {
|
||||
public:
|
||||
JpegParser();
|
||||
~JpegParser();
|
||||
bool ParseJpegStream(const uint8_t* jpeg_stream, uint32_t jpeg_stream_size);
|
||||
const JpegStreamParameters* GetJpegStreamParameters() const {return &jpeg_stream_parameters_;};
|
||||
private:
|
||||
bool ParseSOI();
|
||||
bool ParseSOF();
|
||||
bool ParseDQT();
|
||||
bool ParseSOS();
|
||||
bool ParseDHT();
|
||||
bool ParseDRI();
|
||||
bool ParseEOI();
|
||||
ChromaSubsampling GetChromaSubsampling(uint8_t c1_h_sampling_factor, uint8_t c2_h_sampling_factor, uint8_t c3_h_sampling_factor,
|
||||
uint8_t c1_v_sampling_factor, uint8_t c2_v_sampling_factor, uint8_t c3_v_sampling_factor);
|
||||
const uint8_t *stream_;
|
||||
const uint8_t *stream_end_;
|
||||
uint32_t stream_length_;
|
||||
JpegStreamParameters jpeg_stream_parameters_;
|
||||
};
|
||||
|
||||
#endif // ROC_JPEG_PARSER_H_
|
||||
@@ -0,0 +1,301 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "rocjpeg_vaapi_decoder.h"
|
||||
|
||||
RocJpegVappiDecoder::RocJpegVappiDecoder(int device_id) : device_id_{device_id}, drm_fd_{-1}, min_picture_width_{64}, min_picture_height_{64},
|
||||
max_picture_width_{4096}, max_picture_height_{4096}, va_display_{0}, va_config_attrib_{{}}, va_config_id_{0}, va_profile_{VAProfileJPEGBaseline},
|
||||
va_context_id_{0}, va_surface_ids_{}, va_picture_parameter_buf_id_{0}, va_quantization_matrix_buf_id_{0}, va_huffmantable_buf_id_{0},
|
||||
va_slice_param_buf_id_{0}, va_slice_data_buf_id_{0} {};
|
||||
|
||||
RocJpegVappiDecoder::~RocJpegVappiDecoder() {
|
||||
if (drm_fd_ != -1) {
|
||||
close(drm_fd_);
|
||||
}
|
||||
if (va_display_) {
|
||||
RocJpegStatus rocjpeg_status = DestroyDataBuffers();
|
||||
if (rocjpeg_status != ROCJPEG_STATUS_SUCCESS) {
|
||||
ERR("Error: Failed to destroy VAAPI buffer");
|
||||
}
|
||||
VAStatus va_status;
|
||||
if (va_surface_ids_.size() > 0) {
|
||||
va_status = vaDestroySurfaces(va_display_, va_surface_ids_.data(), va_surface_ids_.size());
|
||||
if (va_status != VA_STATUS_SUCCESS) {
|
||||
ERR("ERROR: vaDestroySurfaces failed!");
|
||||
}
|
||||
}
|
||||
if (va_context_id_) {
|
||||
va_status = vaDestroyContext(va_display_, va_context_id_);
|
||||
if (va_status != VA_STATUS_SUCCESS) {
|
||||
ERR("ERROR: vaDestroyContext failed!");
|
||||
}
|
||||
}
|
||||
if (va_config_id_) {
|
||||
va_status = vaDestroyConfig(va_display_, va_config_id_);
|
||||
if (va_status != VA_STATUS_SUCCESS) {
|
||||
ERR("ERROR: vaDestroyConfig failed!");
|
||||
}
|
||||
}
|
||||
va_status = vaTerminate(va_display_);
|
||||
if (va_status != VA_STATUS_SUCCESS) {
|
||||
ERR("ERROR: vaTerminate failed!");
|
||||
}
|
||||
|
||||
}
|
||||
}
|
||||
|
||||
RocJpegStatus RocJpegVappiDecoder::InitializeDecoder(std::string gcn_arch_name) {
|
||||
// There are 8 renderDXXX per physical device on gfx940, gfx941, and gfx942
|
||||
int num_render_cards_per_device = ((gcn_arch_name.compare("gfx940") == 0) ||
|
||||
(gcn_arch_name.compare("gfx941") == 0) ||
|
||||
(gcn_arch_name.compare("gfx942") == 0)) ? 8 : 1;
|
||||
std::string drm_node = "/dev/dri/renderD" + std::to_string(128 + device_id_ * num_render_cards_per_device);
|
||||
CHECK_ROCJPEG(InitVAAPI(drm_node));
|
||||
CHECK_ROCJPEG(CreateDecoderConfig());
|
||||
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus RocJpegVappiDecoder::InitVAAPI(std::string drm_node) {
|
||||
drm_fd_ = open(drm_node.c_str(), O_RDWR);
|
||||
if (drm_fd_ < 0) {
|
||||
ERR("ERROR: failed to open drm node " + drm_node);
|
||||
return ROCJPEG_STATUS_NOT_INITIALIZED;
|
||||
}
|
||||
va_display_ = vaGetDisplayDRM(drm_fd_);
|
||||
if (!va_display_) {
|
||||
ERR("ERROR: failed to create va_display!");
|
||||
return ROCJPEG_STATUS_NOT_INITIALIZED;
|
||||
}
|
||||
vaSetInfoCallback(va_display_, NULL, NULL);
|
||||
int major_version = 0, minor_version = 0;
|
||||
CHECK_VAAPI(vaInitialize(va_display_, &major_version, &minor_version))
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus RocJpegVappiDecoder::CreateDecoderConfig() {
|
||||
int max_num_entrypoints = vaMaxNumEntrypoints(va_display_);
|
||||
std::vector<VAEntrypoint> jpeg_entrypoint_list;
|
||||
jpeg_entrypoint_list.resize(max_num_entrypoints);
|
||||
int num_entrypoints = 0;
|
||||
CHECK_VAAPI(vaQueryConfigEntrypoints(va_display_, va_profile_, jpeg_entrypoint_list.data(), &num_entrypoints));
|
||||
bool hw_jpeg_decoder_supported = false;
|
||||
if (num_entrypoints > 0) {
|
||||
for (auto entry_point : jpeg_entrypoint_list) {
|
||||
if (entry_point == VAEntrypointVLD) {
|
||||
hw_jpeg_decoder_supported = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
return ROCJPEG_STATUS_HW_JPEG_DECODER_NOT_SUPPORTED;
|
||||
}
|
||||
|
||||
if (hw_jpeg_decoder_supported) {
|
||||
va_config_attrib_.resize(3);
|
||||
va_config_attrib_[0].type = VAConfigAttribRTFormat;
|
||||
va_config_attrib_[1].type = VAConfigAttribMaxPictureWidth;
|
||||
va_config_attrib_[2].type = VAConfigAttribMaxPictureHeight;
|
||||
CHECK_VAAPI(vaGetConfigAttributes(va_display_, va_profile_, VAEntrypointVLD, va_config_attrib_.data(), va_config_attrib_.size()));
|
||||
CHECK_VAAPI(vaCreateConfig(va_display_, va_profile_, VAEntrypointVLD, &va_config_attrib_[0], 1, &va_config_id_));
|
||||
if (va_config_attrib_[1].value != VA_ATTRIB_NOT_SUPPORTED) {
|
||||
max_picture_width_ = va_config_attrib_[1].value;
|
||||
}
|
||||
if (va_config_attrib_[2].value != VA_ATTRIB_NOT_SUPPORTED) {
|
||||
max_picture_height_ = va_config_attrib_[2].value;
|
||||
}
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
} else {
|
||||
return ROCJPEG_STATUS_HW_JPEG_DECODER_NOT_SUPPORTED;
|
||||
}
|
||||
}
|
||||
|
||||
RocJpegStatus RocJpegVappiDecoder::DestroyDataBuffers() {
|
||||
if (va_picture_parameter_buf_id_) {
|
||||
CHECK_VAAPI(vaDestroyBuffer(va_display_, va_picture_parameter_buf_id_));
|
||||
va_picture_parameter_buf_id_ = 0;
|
||||
}
|
||||
if (va_quantization_matrix_buf_id_) {
|
||||
CHECK_VAAPI(vaDestroyBuffer(va_display_, va_quantization_matrix_buf_id_));
|
||||
va_quantization_matrix_buf_id_ = 0;
|
||||
}
|
||||
if (va_huffmantable_buf_id_) {
|
||||
CHECK_VAAPI(vaDestroyBuffer(va_display_, va_huffmantable_buf_id_));
|
||||
va_huffmantable_buf_id_ = 0;
|
||||
}
|
||||
if (va_slice_param_buf_id_) {
|
||||
CHECK_VAAPI(vaDestroyBuffer(va_display_, va_slice_param_buf_id_));
|
||||
va_slice_param_buf_id_ = 0;
|
||||
}
|
||||
if (va_slice_data_buf_id_) {
|
||||
CHECK_VAAPI(vaDestroyBuffer(va_display_, va_slice_data_buf_id_));
|
||||
va_slice_data_buf_id_ = 0;
|
||||
}
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus RocJpegVappiDecoder::SubmitDecode(const JpegStreamParameters *jpeg_stream_params, uint32_t &surface_id) {
|
||||
if (jpeg_stream_params == nullptr) {
|
||||
return ROCJPEG_STATUS_INVALID_PARAMETER;
|
||||
}
|
||||
|
||||
if (sizeof(jpeg_stream_params->picture_parameter_buffer) != sizeof(VAPictureParameterBufferJPEGBaseline) ||
|
||||
sizeof(jpeg_stream_params->quantization_matrix_buffer) != sizeof(VAIQMatrixBufferJPEGBaseline) ||
|
||||
sizeof(jpeg_stream_params->huffman_table_buffer) != sizeof(VAHuffmanTableBufferJPEGBaseline) ||
|
||||
sizeof(jpeg_stream_params->slice_parameter_buffer) != sizeof(VASliceParameterBufferJPEGBaseline)) {
|
||||
return ROCJPEG_STATUS_INVALID_PARAMETER;
|
||||
}
|
||||
|
||||
if (jpeg_stream_params->picture_parameter_buffer.picture_width < min_picture_width_ ||
|
||||
jpeg_stream_params->picture_parameter_buffer.picture_height < min_picture_height_ ||
|
||||
jpeg_stream_params->picture_parameter_buffer.picture_width > max_picture_width_ ||
|
||||
jpeg_stream_params->picture_parameter_buffer.picture_height > max_picture_height_) {
|
||||
ERR("The JPEG image resolution is not supported!");
|
||||
return ROCJPEG_STATUS_JPEG_NOT_SUPPORTED;
|
||||
}
|
||||
|
||||
uint8_t surface_format;
|
||||
switch (jpeg_stream_params->chroma_subsampling) {
|
||||
case CSS_444:
|
||||
surface_format = VA_RT_FORMAT_YUV444;
|
||||
break;
|
||||
case CSS_422:
|
||||
surface_format = VA_RT_FORMAT_YUV422;
|
||||
break;
|
||||
case CSS_420:
|
||||
surface_format = VA_RT_FORMAT_YUV420;
|
||||
break;
|
||||
case CSS_400:
|
||||
surface_format = VA_RT_FORMAT_YUV400;
|
||||
break;
|
||||
default:
|
||||
ERR("ERROR: The chroma subsampling is not supported by the VCN hardware!");
|
||||
return ROCJPEG_STATUS_JPEG_NOT_SUPPORTED;
|
||||
break;
|
||||
}
|
||||
|
||||
VASurfaceID va_surface_id;
|
||||
CHECK_VAAPI(vaCreateSurfaces(va_display_, surface_format, jpeg_stream_params->picture_parameter_buffer.picture_width, jpeg_stream_params->picture_parameter_buffer.picture_height, &va_surface_id, 1, nullptr, 1));
|
||||
va_surface_ids_.push_back(va_surface_id);
|
||||
surface_id = va_surface_id;
|
||||
|
||||
if (va_context_id_) {
|
||||
vaDestroyContext(va_display_, va_context_id_);
|
||||
va_context_id_ = 0;
|
||||
}
|
||||
CHECK_VAAPI(vaCreateContext(va_display_, va_config_id_, jpeg_stream_params->picture_parameter_buffer.picture_width, jpeg_stream_params->picture_parameter_buffer.picture_height, VA_PROGRESSIVE, &va_surface_id, 1, &va_context_id_));
|
||||
|
||||
CHECK_ROCJPEG(DestroyDataBuffers());
|
||||
|
||||
CHECK_VAAPI(vaCreateBuffer(va_display_, va_context_id_, VAPictureParameterBufferType, sizeof(VAPictureParameterBufferJPEGBaseline), 1, (void *)&jpeg_stream_params->picture_parameter_buffer, &va_picture_parameter_buf_id_));
|
||||
CHECK_VAAPI(vaCreateBuffer(va_display_, va_context_id_, VAIQMatrixBufferType, sizeof(VAIQMatrixBufferJPEGBaseline), 1, (void *)&jpeg_stream_params->quantization_matrix_buffer, &va_quantization_matrix_buf_id_));
|
||||
CHECK_VAAPI(vaCreateBuffer(va_display_, va_context_id_, VAHuffmanTableBufferType, sizeof(VAHuffmanTableBufferJPEGBaseline), 1, (void *)&jpeg_stream_params->huffman_table_buffer, &va_huffmantable_buf_id_));
|
||||
CHECK_VAAPI(vaCreateBuffer(va_display_, va_context_id_, VASliceParameterBufferType, sizeof(VASliceParameterBufferJPEGBaseline), 1, (void *)&jpeg_stream_params->slice_parameter_buffer, &va_slice_param_buf_id_));
|
||||
CHECK_VAAPI(vaCreateBuffer(va_display_, va_context_id_, VASliceDataBufferType, jpeg_stream_params->slice_parameter_buffer.slice_data_size, 1, (void *)jpeg_stream_params->slice_data_buffer, &va_slice_data_buf_id_));
|
||||
|
||||
CHECK_VAAPI(vaBeginPicture(va_display_, va_context_id_, va_surface_id));
|
||||
CHECK_VAAPI(vaRenderPicture(va_display_, va_context_id_, &va_picture_parameter_buf_id_, 1));
|
||||
CHECK_VAAPI(vaRenderPicture(va_display_, va_context_id_, &va_quantization_matrix_buf_id_, 1));
|
||||
CHECK_VAAPI(vaRenderPicture(va_display_, va_context_id_, &va_huffmantable_buf_id_, 1));
|
||||
CHECK_VAAPI(vaRenderPicture(va_display_, va_context_id_, &va_slice_param_buf_id_, 1));
|
||||
CHECK_VAAPI(vaRenderPicture(va_display_, va_context_id_, &va_slice_data_buf_id_, 1));
|
||||
CHECK_VAAPI(vaEndPicture(va_display_, va_context_id_));
|
||||
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus RocJpegVappiDecoder::ExportSurface(VASurfaceID surface_id, VADRMPRIMESurfaceDescriptor &va_drm_prime_surface_desc) {
|
||||
|
||||
bool is_surface_id_found = false;
|
||||
int idx = 0;
|
||||
for (idx = 0; idx < va_surface_ids_.size(); idx++) {
|
||||
if (va_surface_ids_[idx] == surface_id) {
|
||||
is_surface_id_found = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!is_surface_id_found) {
|
||||
return ROCJPEG_STATUS_INVALID_PARAMETER;
|
||||
}
|
||||
CHECK_VAAPI(vaExportSurfaceHandle(va_display_, surface_id,
|
||||
VA_SURFACE_ATTRIB_MEM_TYPE_DRM_PRIME_2,
|
||||
VA_EXPORT_SURFACE_READ_ONLY |
|
||||
VA_EXPORT_SURFACE_SEPARATE_LAYERS,
|
||||
&va_drm_prime_surface_desc));
|
||||
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus RocJpegVappiDecoder::SyncSurface(VASurfaceID surface_id) {
|
||||
VASurfaceStatus surface_status;
|
||||
bool is_surface_id_found = false;
|
||||
int idx = 0;
|
||||
|
||||
for (idx = 0; idx < va_surface_ids_.size(); idx++) {
|
||||
if (va_surface_ids_[idx] == surface_id) {
|
||||
is_surface_id_found = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
if (!is_surface_id_found) {
|
||||
return ROCJPEG_STATUS_INVALID_PARAMETER;
|
||||
}
|
||||
|
||||
CHECK_VAAPI(vaQuerySurfaceStatus(va_display_, surface_id, &surface_status));
|
||||
while (surface_status != VASurfaceReady) {
|
||||
VAStatus va_status = vaSyncSurface(va_display_, surface_id);
|
||||
if (va_status != VA_STATUS_SUCCESS) {
|
||||
if (va_status == 0x26 /*VA_STATUS_ERROR_TIMEDOUT*/) {
|
||||
CHECK_VAAPI(vaQuerySurfaceStatus(va_display_, surface_id, &surface_status));
|
||||
} else {
|
||||
std::cout << "vaSyncSurface() failed with error code: 0x" << std::hex << va_status <<
|
||||
std::dec << "', status: " << vaErrorStr(va_status) << "' at " << __FILE__ << ":" << __LINE__ << std::endl;
|
||||
return ROCJPEG_STATUS_RUNTIME_ERROR;
|
||||
}
|
||||
} else {
|
||||
break;
|
||||
}
|
||||
}
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
RocJpegStatus RocJpegVappiDecoder::ReleaseSurface(VASurfaceID surface_id) {
|
||||
bool is_surface_id_found = false;
|
||||
int idx = 0;
|
||||
|
||||
for (idx = 0; idx < va_surface_ids_.size(); idx++) {
|
||||
if (va_surface_ids_[idx] == surface_id) {
|
||||
is_surface_id_found = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
if (!is_surface_id_found) {
|
||||
return ROCJPEG_STATUS_INVALID_PARAMETER;
|
||||
}
|
||||
|
||||
CHECK_VAAPI(vaDestroySurfaces(va_display_, &va_surface_ids_[idx], 1));
|
||||
va_surface_ids_.erase(va_surface_ids_.begin() + idx);
|
||||
|
||||
return ROCJPEG_STATUS_SUCCESS;
|
||||
}
|
||||
@@ -0,0 +1,75 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#ifndef ROC_JPEG_VAAPI_DECODER_H_
|
||||
#define ROC_JPEG_VAAPI_DECODER_H_
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <string>
|
||||
#include <fcntl.h>
|
||||
#include <unistd.h>
|
||||
#include <va/va.h>
|
||||
#include <va/va_drm.h>
|
||||
#include <va/va_drmcommon.h>
|
||||
#include "rocjpeg_commons.h"
|
||||
#include "rocjpeg_parser.h"
|
||||
#include "../api/rocjpeg.h"
|
||||
|
||||
/*Note: va.h doesn't have VA_FOURCC_YUYV defined but vaExportSurfaceHandle returns 0x56595559 for packed YUYV for YUV 4:2:2*/
|
||||
#define ROCJPEG_FOURCC_YUYV 0x56595559
|
||||
|
||||
class RocJpegVappiDecoder {
|
||||
public:
|
||||
RocJpegVappiDecoder(int device_id = 0);
|
||||
~RocJpegVappiDecoder();
|
||||
RocJpegStatus InitializeDecoder(std::string gcn_arch_name);
|
||||
RocJpegStatus SubmitDecode(const JpegStreamParameters *jpeg_stream_params, uint32_t &surface_id);
|
||||
RocJpegStatus ExportSurface(VASurfaceID surface_id, VADRMPRIMESurfaceDescriptor &va_drm_prime_surface_desc);
|
||||
RocJpegStatus SyncSurface(VASurfaceID surface_id);
|
||||
RocJpegStatus ReleaseSurface(VASurfaceID surface_id);
|
||||
private:
|
||||
int device_id_;
|
||||
int drm_fd_;
|
||||
uint32_t min_picture_width_;
|
||||
uint32_t min_picture_height_;
|
||||
uint32_t max_picture_width_;
|
||||
uint32_t max_picture_height_;
|
||||
VADisplay va_display_;
|
||||
std::vector<VAConfigAttrib> va_config_attrib_;
|
||||
VAConfigID va_config_id_;
|
||||
VAProfile va_profile_;
|
||||
VAContextID va_context_id_;
|
||||
std::vector<VASurfaceID> va_surface_ids_;
|
||||
VABufferID va_picture_parameter_buf_id_;
|
||||
VABufferID va_quantization_matrix_buf_id_;
|
||||
VABufferID va_huffmantable_buf_id_;
|
||||
VABufferID va_slice_param_buf_id_;
|
||||
VABufferID va_slice_data_buf_id_;
|
||||
RocJpegStatus InitVAAPI(std::string drm_node);
|
||||
RocJpegStatus CreateDecoderConfig();
|
||||
RocJpegStatus DestroyDataBuffers();
|
||||
};
|
||||
|
||||
#endif // ROC_JPEG_VAAPI_DECODER_H_
|
||||
Reference in New Issue
Block a user