Optimization: Reduced decode session start latency for certain output surface memory types. (#512)
* *rocDecode/Optimization: Reduced decode session start latency for certain output surface memory types. * * rocDecode/Optimization: Video decode RGB sample creates itsown HIP stream. * * rocDecode/Optimization: Created separate HIP stream in each thread in video decode RGB sample. * * rocDecode/Optimization: Addressed the lower FPS report in RGB sample: move HIP stream creation out of FPS calculation scope. * * rocDecode/Optimize: Removed comments.
Этот коммит содержится в:
jeffqjiangNew
коммит произвёл
GitHub
GitHub
родитель
c4fe8f5258
Коммит
1f53dedd4a
@@ -19,6 +19,7 @@ Full documentation for rocDecode is available at [https://rocm.docs.amd.com/proj
|
||||
### Removed
|
||||
|
||||
* FFMPEG executable requirement for the package
|
||||
* GetStream() interface call from RocVideoDecoder utility class
|
||||
|
||||
## rocDecode 0.8.0 for ROCm 6.3
|
||||
|
||||
|
||||
@@ -66,7 +66,7 @@ std::condition_variable cv[frame_buffers_size];
|
||||
|
||||
void ColorSpaceConversionThread(std::atomic<bool>& continue_processing, bool convert_to_rgb, Dim *p_resize_dim, OutputSurfaceInfo **surf_info, OutputSurfaceInfo **res_surf_info,
|
||||
OutputFormatEnum e_output_format, uint8_t *p_rgb_dev_mem, uint8_t *p_resize_dev_mem, bool dump_output_frames,
|
||||
std::string &output_file_path, RocVideoDecoder &viddec, VideoPostProcess &post_proc, MD5Generator *md5_gen_handle, bool b_generate_md5, int device_id) {
|
||||
std::string &output_file_path, RocVideoDecoder &viddec, VideoPostProcess &post_proc, MD5Generator *md5_gen_handle, bool b_generate_md5, int device_id, hipStream_t hip_stream) {
|
||||
|
||||
size_t rgb_image_size, resize_image_size;
|
||||
hipError_t hip_status = hipSuccess;
|
||||
@@ -105,10 +105,10 @@ void ColorSpaceConversionThread(std::atomic<bool>& continue_processing, bool con
|
||||
// call resize kernel
|
||||
if ((*surf_info)->bytes_per_pixel == 2) {
|
||||
ResizeP016(p_resize_dev_mem, p_resize_dim->w * 2, p_resize_dim->w, p_resize_dim->h, frame, (*surf_info)->output_pitch, (*surf_info)->output_width,
|
||||
(*surf_info)->output_height, (frame + (*surf_info)->output_vstride * (*surf_info)->output_pitch), nullptr, viddec.GetStream());
|
||||
(*surf_info)->output_height, (frame + (*surf_info)->output_vstride * (*surf_info)->output_pitch), nullptr, hip_stream);
|
||||
} else {
|
||||
ResizeNv12(p_resize_dev_mem, p_resize_dim->w, p_resize_dim->w, p_resize_dim->h, frame, (*surf_info)->output_pitch, (*surf_info)->output_width,
|
||||
(*surf_info)->output_height, (frame + (*surf_info)->output_vstride * (*surf_info)->output_pitch), nullptr, viddec.GetStream());
|
||||
(*surf_info)->output_height, (frame + (*surf_info)->output_vstride * (*surf_info)->output_pitch), nullptr, hip_stream);
|
||||
}
|
||||
(*res_surf_info)->output_width = p_resize_dim->w;
|
||||
(*res_surf_info)->output_height = p_resize_dim->h;
|
||||
@@ -131,7 +131,7 @@ void ColorSpaceConversionThread(std::atomic<bool>& continue_processing, bool con
|
||||
return;
|
||||
}
|
||||
}
|
||||
post_proc.ColorConvertYUV2RGB(out_frame, p_surf_info, p_rgb_dev_mem, e_output_format, viddec.GetStream());
|
||||
post_proc.ColorConvertYUV2RGB(out_frame, p_surf_info, p_rgb_dev_mem, e_output_format, hip_stream);
|
||||
}
|
||||
if (dump_output_frames) {
|
||||
if (convert_to_rgb)
|
||||
@@ -143,14 +143,12 @@ void ColorSpaceConversionThread(std::atomic<bool>& continue_processing, bool con
|
||||
md5_gen_handle->UpdateMd5ForDataBuffer(p_rgb_dev_mem, rgb_image_size);
|
||||
}
|
||||
|
||||
|
||||
cv[current_frame_index].notify_one();
|
||||
current_frame_index = (current_frame_index + 1) % frame_buffers_size;
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char **argv) {
|
||||
|
||||
std::string input_file_path, output_file_path, md5_file_path;
|
||||
std::fstream ref_md5_file;
|
||||
bool b_generate_md5 = false;
|
||||
@@ -173,7 +171,8 @@ int main(int argc, char **argv) {
|
||||
int rgb_width;
|
||||
uint8_t* frame_buffers[frame_buffers_size] = {0};
|
||||
int current_frame_index = 0;
|
||||
|
||||
hipStream_t hip_stream_dec = 0;
|
||||
hipStream_t hip_stream_csc = 0;
|
||||
|
||||
// Parse command-line arguments
|
||||
if(argc <= 1) {
|
||||
@@ -276,13 +275,14 @@ int main(int argc, char **argv) {
|
||||
|
||||
std::string device_name, gcn_arch_name;
|
||||
int pci_bus_id, pci_domain_id, pci_device_id;
|
||||
hipStream_t stream = viddec.GetStream();
|
||||
|
||||
viddec.GetDeviceinfo(device_name, gcn_arch_name, pci_bus_id, pci_domain_id, pci_device_id);
|
||||
std::cout << "info: Using GPU device " << device_id << " " << device_name << "[" << gcn_arch_name << "] on PCI bus " <<
|
||||
std::setfill('0') << std::setw(2) << std::right << std::hex << pci_bus_id << ":" << std::setfill('0') << std::setw(2) <<
|
||||
std::right << std::hex << pci_domain_id << "." << pci_device_id << std::dec << std::endl;
|
||||
std::cout << "info: decoding started, please wait!" << std::endl;
|
||||
HIP_API_CALL(hipStreamCreate(&hip_stream_dec));
|
||||
HIP_API_CALL(hipStreamCreate(&hip_stream_csc));
|
||||
|
||||
if (b_generate_md5) {
|
||||
md5_generator = new MD5Generator();
|
||||
@@ -300,7 +300,7 @@ int main(int argc, char **argv) {
|
||||
convert_to_rgb = e_output_format != native;
|
||||
std::atomic<bool> continue_processing(true);
|
||||
std::thread color_space_conversion_thread(ColorSpaceConversionThread, std::ref(continue_processing), std::ref(convert_to_rgb), &resize_dim, &surf_info, &resize_surf_info, std::ref(e_output_format),
|
||||
std::ref(p_rgb_dev_mem), std::ref(p_resize_dev_mem), std::ref(dump_output_frames), std::ref(output_file_path), std::ref(viddec), std::ref(post_process), md5_generator, b_generate_md5, device_id);
|
||||
std::ref(p_rgb_dev_mem), std::ref(p_resize_dev_mem), std::ref(dump_output_frames), std::ref(output_file_path), std::ref(viddec), std::ref(post_process), md5_generator, b_generate_md5, device_id, hip_stream_csc);
|
||||
|
||||
auto startTime = std::chrono::high_resolution_clock::now();
|
||||
do {
|
||||
@@ -329,7 +329,7 @@ int main(int argc, char **argv) {
|
||||
std::unique_lock<std::mutex> lock(mutex[current_frame_index]);
|
||||
cv[current_frame_index].wait(lock, [&] {return frame_queue[current_frame_index].empty();});
|
||||
// copy the decoded frame into the frame_buffers at current_frame_index
|
||||
HIP_API_CALL(hipMemcpyDtoDAsync(frame_buffers[current_frame_index], p_frame, surf_info->output_surface_size_in_bytes, viddec.GetStream()));
|
||||
HIP_API_CALL(hipMemcpyDtoDAsync(frame_buffers[current_frame_index], p_frame, surf_info->output_surface_size_in_bytes, hip_stream_dec));
|
||||
frame_queue[current_frame_index].push(frame_buffers[current_frame_index]);
|
||||
}
|
||||
|
||||
@@ -368,6 +368,12 @@ int main(int argc, char **argv) {
|
||||
std::cout << "ERROR: hipFree failed! (" << hip_status << ")" << std::endl;
|
||||
}
|
||||
}
|
||||
if (hip_stream_dec) {
|
||||
HIP_API_CALL(hipStreamDestroy(hip_stream_dec));
|
||||
}
|
||||
if (hip_stream_csc) {
|
||||
HIP_API_CALL(hipStreamDestroy(hip_stream_csc));
|
||||
}
|
||||
|
||||
std::cout << "info: Total frame decoded: " << n_frame << std::endl;
|
||||
if (!dump_output_frames) {
|
||||
|
||||
@@ -1044,7 +1044,11 @@ bool RocVideoDecoder::InitHIP(int device_id) {
|
||||
}
|
||||
HIP_API_CALL(hipSetDevice(device_id));
|
||||
HIP_API_CALL(hipGetDeviceProperties(&hip_dev_prop_, device_id));
|
||||
HIP_API_CALL(hipStreamCreate(&hip_stream_));
|
||||
if (out_mem_type_ == OUT_SURFACE_MEM_DEV_INTERNAL || out_mem_type_ == OUT_SURFACE_MEM_NOT_MAPPED) {
|
||||
hip_stream_ = 0; // Null stream. For internal device or unmapped memory, we don't need to create a hip stream.
|
||||
} else {
|
||||
HIP_API_CALL(hipStreamCreate(&hip_stream_));
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
|
||||
@@ -219,8 +219,6 @@ class RocVideoDecoder {
|
||||
|
||||
rocDecVideoCodec GetCodecId() { return codec_id_; }
|
||||
|
||||
hipStream_t GetStream() {return hip_stream_;}
|
||||
|
||||
/**
|
||||
* @brief Get the output frame width
|
||||
*/
|
||||
|
||||
Ссылка в новой задаче
Block a user