AI/rocm-systems
2
0
Bifurcation 0
Code Tickets Demandes d'ajout Actions Paquets Projets Publications Wiki Activité

Optimization: Reduced decode session start latency for certain output surface memory types. (#512)

Parcourir la source 
* *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.

[ROCm/rocdecode commit: 1f53dedd4a]
Cette révision appartient à :
jeffqjiangNew
2025-02-18 13:09:16 -05:00
révisé par GitHub
Parent 3b6a4d3293
révision 4ac1b8dd67
4 fichiers modifiés avec 22 ajouts et 13 suppressions
Télécharger le Fichier Patch Télécharger le Fichier des Différences Développer tous les fichiers Réduire tous les fichiers
projects/rocdecode/CHANGELOG.md
+1
Voir le fichier
@@ -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
projects/rocdecode/samples/videoDecodeRGB/videodecrgb.cpp
+16 -10
Voir le fichier
@@ -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) {
projects/rocdecode/utils/rocvideodecode/roc_video_dec.cpp
+5 -1
Voir le fichier
@@ -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;
}
projects/rocdecode/utils/rocvideodecode/roc_video_dec.h
-2
Voir le fichier
@@ -219,8 +219,6 @@ class RocVideoDecoder {
rocDecVideoCodec GetCodecId() { return codec_id_; }
hipStream_t GetStream() {return hip_stream_;}
/**
* @brief Get the output frame width
*/