From 1f53dedd4a7562a353e1c52c605e530e082133b8 Mon Sep 17 00:00:00 2001 From: jeffqjiangNew <142832361+jeffqjiangNew@users.noreply.github.com> Date: Tue, 18 Feb 2025 13:09:16 -0500 Subject: [PATCH] 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. --- CHANGELOG.md | 1 + samples/videoDecodeRGB/videodecrgb.cpp | 26 ++++++++++++++++---------- utils/rocvideodecode/roc_video_dec.cpp | 6 +++++- utils/rocvideodecode/roc_video_dec.h | 2 -- 4 files changed, 22 insertions(+), 13 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 08001da6fe..65e887b51f 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -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 diff --git a/samples/videoDecodeRGB/videodecrgb.cpp b/samples/videoDecodeRGB/videodecrgb.cpp index d6b2f89d11..0ccb357c80 100644 --- a/samples/videoDecodeRGB/videodecrgb.cpp +++ b/samples/videoDecodeRGB/videodecrgb.cpp @@ -66,7 +66,7 @@ std::condition_variable cv[frame_buffers_size]; void ColorSpaceConversionThread(std::atomic& 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& 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& 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& 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 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 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) { diff --git a/utils/rocvideodecode/roc_video_dec.cpp b/utils/rocvideodecode/roc_video_dec.cpp index d6bf15883e..a612cfd3cf 100644 --- a/utils/rocvideodecode/roc_video_dec.cpp +++ b/utils/rocvideodecode/roc_video_dec.cpp @@ -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; } diff --git a/utils/rocvideodecode/roc_video_dec.h b/utils/rocvideodecode/roc_video_dec.h index 25fafd0217..4bd7c1b4a7 100644 --- a/utils/rocvideodecode/roc_video_dec.h +++ b/utils/rocvideodecode/roc_video_dec.h @@ -219,8 +219,6 @@ class RocVideoDecoder { rocDecVideoCodec GetCodecId() { return codec_id_; } - hipStream_t GetStream() {return hip_stream_;} - /** * @brief Get the output frame width */