From db115de7986bb1faa2cdb60b061a7c0bfee086c2 Mon Sep 17 00:00:00 2001 From: "iotg.worker" Date: Thu, 21 Mar 2024 14:33:50 +0100 Subject: [PATCH 1/7] Umat init --- mediapipe/examples/desktop/demo_run_graph_main.cc | 8 ++++---- mediapipe/framework/formats/image_frame_opencv.cc | 10 ++++++++++ mediapipe/framework/formats/image_frame_opencv.h | 5 +++++ 3 files changed, 19 insertions(+), 4 deletions(-) diff --git a/mediapipe/examples/desktop/demo_run_graph_main.cc b/mediapipe/examples/desktop/demo_run_graph_main.cc index 3b6465eec6..9ac4a88167 100644 --- a/mediapipe/examples/desktop/demo_run_graph_main.cc +++ b/mediapipe/examples/desktop/demo_run_graph_main.cc @@ -90,7 +90,7 @@ absl::Status RunMPPGraph() { while (grab_frames) { // Capture opencv camera or video frame. - cv::Mat camera_frame_raw; + cv::UMat camera_frame_raw; capture >> camera_frame_raw; if (camera_frame_raw.empty()) { if (!load_video) { @@ -101,7 +101,7 @@ absl::Status RunMPPGraph() { break; } count_frames+=1; - cv::Mat camera_frame; + cv::UMat camera_frame; cv::cvtColor(camera_frame_raw, camera_frame, cv::COLOR_BGR2RGB); if (!load_video) { cv::flip(camera_frame, camera_frame, /*flipcode=HORIZONTAL*/ 1); @@ -111,7 +111,7 @@ absl::Status RunMPPGraph() { auto input_frame = absl::make_unique( mediapipe::ImageFormat::SRGB, camera_frame.cols, camera_frame.rows, mediapipe::ImageFrame::kDefaultAlignmentBoundary); - cv::Mat input_frame_mat = mediapipe::formats::MatView(input_frame.get()); + cv::UMat input_frame_mat = mediapipe::formats::MatView(input_frame.get(), cv::USAGE_ALLOCATE_SHARED_MEMORY); camera_frame.copyTo(input_frame_mat); // Send image packet into the graph. @@ -127,7 +127,7 @@ absl::Status RunMPPGraph() { auto& output_frame = packet.Get(); // Convert back to opencv for display or saving. - cv::Mat output_frame_mat = mediapipe::formats::MatView(&output_frame); + cv::UMat output_frame_mat = mediapipe::formats::MatView(&output_frame, cv::USAGE_ALLOCATE_SHARED_MEMORY); cv::cvtColor(output_frame_mat, output_frame_mat, cv::COLOR_RGB2BGR); if (save_video) { if (!writer.isOpened()) { diff --git a/mediapipe/framework/formats/image_frame_opencv.cc b/mediapipe/framework/formats/image_frame_opencv.cc index 1ba8c719fa..d862c54079 100644 --- a/mediapipe/framework/formats/image_frame_opencv.cc +++ b/mediapipe/framework/formats/image_frame_opencv.cc @@ -11,6 +11,7 @@ // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // See the License for the specific language governing permissions and // limitations under the License. +#include #include "mediapipe/framework/formats/image_frame_opencv.h" @@ -92,5 +93,14 @@ cv::Mat MatView(const ImageFrame* image) { steps); } +// UMatUsageFlags +// USAGE_DEFAULT, USAGE_ALLOCATE_HOST_MEMORY, USAGE_ALLOCATE_DEVICE_MEMORY, USAGE_ALLOCATE_SHARED_MEMORY , __UMAT_USAGE_FLAGS_32BIT +cv::UMat MatView(const ImageFrame* image, cv::UMatUsageFlags usageFlags) { + cv::Mat mat = MatView(image); + cv::UMat umat = cv::UMat(usageFlags); + mat.copyTo(umat); + return umat; +} + } // namespace formats } // namespace mediapipe diff --git a/mediapipe/framework/formats/image_frame_opencv.h b/mediapipe/framework/formats/image_frame_opencv.h index d197d9a554..8af442bf46 100644 --- a/mediapipe/framework/formats/image_frame_opencv.h +++ b/mediapipe/framework/formats/image_frame_opencv.h @@ -18,6 +18,7 @@ #include "mediapipe/framework/formats/image_frame.h" #include "mediapipe/framework/port/opencv_core_inc.h" +#include namespace mediapipe { namespace formats { @@ -30,6 +31,10 @@ namespace formats { // even though the returned data is mutable. cv::Mat MatView(const ImageFrame* image); +// UMatUsageFlags +// USAGE_DEFAULT, USAGE_ALLOCATE_HOST_MEMORY, USAGE_ALLOCATE_DEVICE_MEMORY, USAGE_ALLOCATE_SHARED_MEMORY , __UMAT_USAGE_FLAGS_32BIT +cv::UMat MatView(const ImageFrame* image, cv::UMatUsageFlags usageFlags); + } // namespace formats } // namespace mediapipe From 20a1a48c84bf80b354c08e0e98422ef7ad56865f Mon Sep 17 00:00:00 2001 From: "iotg.worker" Date: Wed, 27 Mar 2024 19:12:04 +0100 Subject: [PATCH 2/7] ocl compile --- .bazelrc | 6 +- WORKSPACE | 6 + .../examples/desktop/demo_run_graph_main.cc | 48 +- mediapipe/framework/formats/BUILD | 6 +- mediapipe/framework/formats/helpers.cpp | 455 +++++++++++++++++ mediapipe/framework/formats/helpers.hpp | 471 ++++++++++++++++++ mediapipe/framework/formats/image_frame.cc | 12 +- .../framework/formats/image_frame_opencv.cc | 3 +- .../framework/formats/image_frame_opencv.h | 2 +- mediapipe/framework/port/BUILD | 8 + mediapipe/framework/port/opencv_opencl.h | 20 + 11 files changed, 1023 insertions(+), 14 deletions(-) create mode 100644 mediapipe/framework/formats/helpers.cpp create mode 100644 mediapipe/framework/formats/helpers.hpp create mode 100644 mediapipe/framework/port/opencv_opencl.h diff --git a/.bazelrc b/.bazelrc index af05ec14cf..1013f30f75 100644 --- a/.bazelrc +++ b/.bazelrc @@ -121,6 +121,6 @@ build --cxxopt=-DPYTHON_DISABLE=1 coverage --cxxopt=-DPYTHON_DISABLE=1 test --cxxopt=-DPYTHON_DISABLE=1 -build --cxxopt=-DOVMS_DUMP_TO_FILE=0 -coverage --cxxopt=-DOVMS_DUMP_TO_FILE=0 -test --cxxopt=-DOVMS_DUMP_TO_FILE=0 +build --cxxopt=-DOVMS_DUMP_TO_FILE=1 +coverage --cxxopt=-DOVMS_DUMP_TO_FILE=1 +test --cxxopt=-DOVMS_DUMP_TO_FILE=1 diff --git a/WORKSPACE b/WORKSPACE index 3bcaa674db..3316d30e3c 100644 --- a/WORKSPACE +++ b/WORKSPACE @@ -825,6 +825,12 @@ new_local_repository( path = "/opt/intel/openvino/runtime", ) +new_git_repository( + name = "openvino", + remote = "https:///github.com/openvinotoolkit/openvino/", + commit = "c3c409ee133ffb26bf8fd5570ef50a7c004839a4" +) + git_repository( name = "oneTBB", branch = "v2021.10.0", diff --git a/mediapipe/examples/desktop/demo_run_graph_main.cc b/mediapipe/examples/desktop/demo_run_graph_main.cc index 9ac4a88167..d33b19dd11 100644 --- a/mediapipe/examples/desktop/demo_run_graph_main.cc +++ b/mediapipe/examples/desktop/demo_run_graph_main.cc @@ -12,14 +12,17 @@ // See the License for the specific language governing permissions and // limitations under the License. // -// An example of sending OpenCV webcam frames into a MediaPipe graph. +// An example of sending OpenCV webcam frames INFO a MediaPipe graph. #include #include +#include + #include "absl/flags/flag.h" #include "absl/flags/parse.h" #include "mediapipe/framework/calculator_framework.h" #include "mediapipe/framework/formats/image_frame.h" +#include "mediapipe/framework/formats/helpers.hpp" #include "mediapipe/framework/formats/image_frame_opencv.h" #include "mediapipe/framework/port/file_helpers.h" #include "mediapipe/framework/port/opencv_highgui_inc.h" @@ -42,6 +45,16 @@ ABSL_FLAG(std::string, output_video_path, "", "Full path of where to save result (.mp4 only). " "If not provided, show result in a window."); +void dumpMatToFile(std::string& fileName, cv::UMat& umat){ + cv::FileStorage file(fileName, cv::FileStorage::WRITE); + cv::Mat input; + //file << "camera_frame_raw" << camera_frame_raw; + umat.copyTo(input); + file << "input_frame_mat" << input; + + file.release(); +} + absl::Status RunMPPGraph() { std::string calculator_graph_config_contents; MP_RETURN_IF_ERROR(mediapipe::file::GetContents( @@ -88,9 +101,29 @@ absl::Status RunMPPGraph() { int count_frames = 0; auto begin = std::chrono::high_resolution_clock::now(); + OpenClWrapper ocl; + + ocl.initOpenCL(); + + LOG(INFO) << "haveOpenCL " << cv::ocl::haveOpenCL() <> camera_frame_raw; if (camera_frame_raw.empty()) { if (!load_video) { @@ -101,20 +134,23 @@ absl::Status RunMPPGraph() { break; } count_frames+=1; - cv::UMat camera_frame; + cv::UMat camera_frame = cv::UMat(cv::USAGE_ALLOCATE_SHARED_MEMORY); cv::cvtColor(camera_frame_raw, camera_frame, cv::COLOR_BGR2RGB); if (!load_video) { cv::flip(camera_frame, camera_frame, /*flipcode=HORIZONTAL*/ 1); } - // Wrap Mat into an ImageFrame. + // Wrap Mat INFO an ImageFrame. auto input_frame = absl::make_unique( mediapipe::ImageFormat::SRGB, camera_frame.cols, camera_frame.rows, mediapipe::ImageFrame::kDefaultAlignmentBoundary); cv::UMat input_frame_mat = mediapipe::formats::MatView(input_frame.get(), cv::USAGE_ALLOCATE_SHARED_MEMORY); camera_frame.copyTo(input_frame_mat); + //std::string fileName = std::string("./input") + std::to_string(count_frames) + std::string("jpeg"); + + - // Send image packet into the graph. + // Send image packet INFO the graph. size_t frame_timestamp_us = (double)cv::getTickCount() / (double)cv::getTickFrequency() * 1e6; MP_RETURN_IF_ERROR(graph.AddPacketToInputStream( @@ -127,7 +163,7 @@ absl::Status RunMPPGraph() { auto& output_frame = packet.Get(); // Convert back to opencv for display or saving. - cv::UMat output_frame_mat = mediapipe::formats::MatView(&output_frame, cv::USAGE_ALLOCATE_SHARED_MEMORY); + cv::Mat output_frame_mat = mediapipe::formats::MatView(&output_frame); cv::cvtColor(output_frame_mat, output_frame_mat, cv::COLOR_RGB2BGR); if (save_video) { if (!writer.isOpened()) { diff --git a/mediapipe/framework/formats/BUILD b/mediapipe/framework/formats/BUILD index b23209f7d4..c87454a174 100644 --- a/mediapipe/framework/formats/BUILD +++ b/mediapipe/framework/formats/BUILD @@ -147,7 +147,9 @@ cc_library( cc_library( name = "image_frame", - srcs = ["image_frame.cc"], + srcs = ["image_frame.cc", + "helpers.hpp", + "helpers.cpp"], hdrs = ["image_frame.h"], deps = [ ":image_format_cc_proto", @@ -162,11 +164,13 @@ cc_library( "@com_google_absl//absl/base:core_headers", "@com_google_absl//absl/memory", "@com_google_absl//absl/strings", + "//mediapipe/framework/port:opencv_core", ] + select({ "//conditions:default": [ ], "//mediapipe/framework:disable_rtti_and_exceptions": [], }), + linkopts = ["-lOpenCL"], ) cc_library( diff --git a/mediapipe/framework/formats/helpers.cpp b/mediapipe/framework/formats/helpers.cpp new file mode 100644 index 0000000000..3b67913d62 --- /dev/null +++ b/mediapipe/framework/formats/helpers.cpp @@ -0,0 +1,455 @@ +/* +// The example of interoperability between OpenCL and OpenCV. +// This will loop through frames of video either from input media file +// or camera device and do processing of these data in OpenCL and then +// in OpenCV. In OpenCL it does inversion of pixels in left half of frame and +// in OpenCV it does blurring in the right half of frame. +*/ +#include +#include +#include +#include +#include +#include +#include +#include + +#define CL_USE_DEPRECATED_OPENCL_1_1_APIS +#define CL_USE_DEPRECATED_OPENCL_1_2_APIS +#define CL_USE_DEPRECATED_OPENCL_2_0_APIS // eliminate build warning +#define CL_TARGET_OPENCL_VERSION 200 // 2.0 + +#ifdef __APPLE__ +#define CL_SILENCE_DEPRECATION +#include +#else +#include +#endif +#include "mediapipe/framework/formats/helpers.hpp" + +#include +#include +#include +#include +#include + + +using namespace std; +using namespace cv; + +OpenClWrapper::OpenClWrapper() +{ + cout << "\nPress ESC to exit\n" << endl; + cout << "\n 'p' to toggle ON/OFF processing\n" << endl; + cout << "\n SPACE to switch between OpenCL buffer/image\n" << endl; + + m_camera_id = 0; //cmd.get("camera"); + m_file_name = "0"; //cmd.get("video"); + + m_running = false; + m_process = false; + m_use_buffer = false; + + m_t0 = 0; + m_t1 = 0; + m_time = 0.0; + m_frequency = (float)cv::getTickFrequency(); + + m_context = 0; + m_device_id = 0; + m_queue = 0; + m_program = 0; + m_kernelBuf = 0; + m_kernelImg = 0; + m_img_src = 0; + m_mem_obj = 0; +} // ctor + + +OpenClWrapper::~OpenClWrapper() +{ + if (m_queue) + { + clFinish(m_queue); + clReleaseCommandQueue(m_queue); + m_queue = 0; + } + + if (m_program) + { + clReleaseProgram(m_program); + m_program = 0; + } + + if (m_img_src) + { + clReleaseMemObject(m_img_src); + m_img_src = 0; + } + + if (m_mem_obj) + { + clReleaseMemObject(m_mem_obj); + m_mem_obj = 0; + } + + if (m_kernelBuf) + { + clReleaseKernel(m_kernelBuf); + m_kernelBuf = 0; + } + + if (m_kernelImg) + { + clReleaseKernel(m_kernelImg); + m_kernelImg = 0; + } + + if (m_device_id) + { + clReleaseDevice(m_device_id); + m_device_id = 0; + } + + if (m_context) + { + clReleaseContext(m_context); + m_context = 0; + } +} // dtor + + +int OpenClWrapper::initOpenCL() +{ + cl_int res = CL_SUCCESS; + cl_uint num_entries = 0; + + res = clGetPlatformIDs(0, 0, &num_entries); + if (CL_SUCCESS != res) + return -1; + + m_platform_ids.resize(num_entries); + + res = clGetPlatformIDs(num_entries, &m_platform_ids[0], 0); + if (CL_SUCCESS != res) + return -1; + + unsigned int i; + + // create context from first platform with GPU device + for (i = 0; i < m_platform_ids.size(); i++) + { + cl_context_properties props[] = + { + CL_CONTEXT_PLATFORM, + (cl_context_properties)(m_platform_ids[i]), + 0 + }; + + m_context = clCreateContextFromType(props, CL_DEVICE_TYPE_GPU, 0, 0, &res); + if (0 == m_context || CL_SUCCESS != res) + continue; + + res = clGetContextInfo(m_context, CL_CONTEXT_DEVICES, sizeof(cl_device_id), &m_device_id, 0); + if (CL_SUCCESS != res) + return -1; + + m_queue = clCreateCommandQueue(m_context, m_device_id, 0, &res); + if (0 == m_queue || CL_SUCCESS != res) + return -1; + + const char* kernelSrc = + "__kernel " + "void bitwise_inv_buf_8uC1(" + " __global unsigned char* pSrcDst," + " int srcDstStep," + " int rows," + " int cols)" + "{" + " int x = get_global_id(0);" + " int y = get_global_id(1);" + " int idx = mad24(y, srcDstStep, x);" + " pSrcDst[idx] = ~pSrcDst[idx];" + "}" + "__kernel " + "void bitwise_inv_img_8uC1(" + " read_only image2d_t srcImg," + " write_only image2d_t dstImg)" + "{" + " int x = get_global_id(0);" + " int y = get_global_id(1);" + " int2 coord = (int2)(x, y);" + " uint4 val = read_imageui(srcImg, coord);" + " val.x = (~val.x) & 0x000000FF;" + " write_imageui(dstImg, coord, val);" + "}"; + size_t len = strlen(kernelSrc); + m_program = clCreateProgramWithSource(m_context, 1, &kernelSrc, &len, &res); + if (0 == m_program || CL_SUCCESS != res) + return -1; + + res = clBuildProgram(m_program, 1, &m_device_id, 0, 0, 0); + if (CL_SUCCESS != res) + return -1; + + m_kernelBuf = clCreateKernel(m_program, "bitwise_inv_buf_8uC1", &res); + if (0 == m_kernelBuf || CL_SUCCESS != res) + return -1; + + m_kernelImg = clCreateKernel(m_program, "bitwise_inv_img_8uC1", &res); + if (0 == m_kernelImg || CL_SUCCESS != res) + return -1; + + m_platformInfo.QueryInfo(m_platform_ids[i]); + m_deviceInfo.QueryInfo(m_device_id); + + // attach OpenCL context to OpenCV + cv::ocl::attachContext(m_platformInfo.Name(), m_platform_ids[i], m_context, m_device_id); + + break; + } + + return m_context != 0 ? CL_SUCCESS : -1; +} // initOpenCL() + + +// this function is an example of "typical" OpenCL processing pipeline +// It creates OpenCL buffer or image, depending on use_buffer flag, +// from input media frame and process these data +// (inverts each pixel value in half of frame) with OpenCL kernel +int OpenClWrapper::process_frame_with_open_cl(cv::Mat& frame, bool use_buffer, cl_mem* mem_obj) +{ + cl_int res = CL_SUCCESS; + + CV_Assert(mem_obj); + + cl_kernel kernel = 0; + cl_mem mem = mem_obj[0]; + + if (0 == mem || 0 == m_img_src) + { + // allocate/delete cl memory objects every frame for the simplicity. + // in real application more efficient pipeline can be built. + + if (use_buffer) + { + cl_mem_flags flags = CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR; + + mem = clCreateBuffer(m_context, flags, frame.total(), frame.ptr(), &res); + if (0 == mem || CL_SUCCESS != res) + return -1; + + res = clSetKernelArg(m_kernelBuf, 0, sizeof(cl_mem), &mem); + if (CL_SUCCESS != res) + return -1; + + res = clSetKernelArg(m_kernelBuf, 1, sizeof(int), &frame.step[0]); + if (CL_SUCCESS != res) + return -1; + + res = clSetKernelArg(m_kernelBuf, 2, sizeof(int), &frame.rows); + if (CL_SUCCESS != res) + return -1; + + int cols2 = frame.cols / 2; + res = clSetKernelArg(m_kernelBuf, 3, sizeof(int), &cols2); + if (CL_SUCCESS != res) + return -1; + + kernel = m_kernelBuf; + } + else + { + cl_mem_flags flags_src = CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR; + + cl_image_format fmt; + fmt.image_channel_order = CL_R; + fmt.image_channel_data_type = CL_UNSIGNED_INT8; + + cl_image_desc desc_src; + desc_src.image_type = CL_MEM_OBJECT_IMAGE2D; + desc_src.image_width = frame.cols; + desc_src.image_height = frame.rows; + desc_src.image_depth = 0; + desc_src.image_array_size = 0; + desc_src.image_row_pitch = frame.step[0]; + desc_src.image_slice_pitch = 0; + desc_src.num_mip_levels = 0; + desc_src.num_samples = 0; + desc_src.buffer = 0; + m_img_src = clCreateImage(m_context, flags_src, &fmt, &desc_src, frame.ptr(), &res); + if (0 == m_img_src || CL_SUCCESS != res) + return -1; + + cl_mem_flags flags_dst = CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR; + + cl_image_desc desc_dst; + desc_dst.image_type = CL_MEM_OBJECT_IMAGE2D; + desc_dst.image_width = frame.cols; + desc_dst.image_height = frame.rows; + desc_dst.image_depth = 0; + desc_dst.image_array_size = 0; + desc_dst.image_row_pitch = 0; + desc_dst.image_slice_pitch = 0; + desc_dst.num_mip_levels = 0; + desc_dst.num_samples = 0; + desc_dst.buffer = 0; + mem = clCreateImage(m_context, flags_dst, &fmt, &desc_dst, 0, &res); + if (0 == mem || CL_SUCCESS != res) + return -1; + + size_t origin[] = { 0, 0, 0 }; + size_t region[] = { (size_t)frame.cols, (size_t)frame.rows, 1 }; + cl_event asyncEvent = 0; + res = clEnqueueCopyImage(m_queue, m_img_src, mem, origin, origin, region, 0, 0, &asyncEvent); + if (CL_SUCCESS != res) + return -1; + + res = clWaitForEvents(1, &asyncEvent); + clReleaseEvent(asyncEvent); + if (CL_SUCCESS != res) + return -1; + + res = clSetKernelArg(m_kernelImg, 0, sizeof(cl_mem), &m_img_src); + if (CL_SUCCESS != res) + return -1; + + res = clSetKernelArg(m_kernelImg, 1, sizeof(cl_mem), &mem); + if (CL_SUCCESS != res) + return -1; + + kernel = m_kernelImg; + } + } + + // process left half of frame in OpenCL + size_t size[] = { (size_t)frame.cols / 2, (size_t)frame.rows }; + cl_event asyncEvent = 0; + res = clEnqueueNDRangeKernel(m_queue, kernel, 2, 0, size, 0, 0, 0, &asyncEvent); + if (CL_SUCCESS != res) + return -1; + + res = clWaitForEvents(1, &asyncEvent); + clReleaseEvent(asyncEvent); + if (CL_SUCCESS != res) + return -1; + + mem_obj[0] = mem; + + return 0; +} + + +// this function is an example of interoperability between OpenCL buffer +// and OpenCV UMat objects. It converts (without copying data) OpenCL buffer +// to OpenCV UMat and then do blur on these data +int OpenClWrapper::process_cl_buffer_with_opencv(cl_mem buffer, size_t step, int rows, int cols, int type, cv::UMat& u) +{ + cv::ocl::convertFromBuffer(buffer, step, rows, cols, type, u); + + // process right half of frame in OpenCV + cv::Point pt(u.cols / 2, 0); + cv::Size sz(u.cols / 2, u.rows); + cv::Rect roi(pt, sz); + cv::UMat uroi(u, roi); + cv::blur(uroi, uroi, cv::Size(7, 7), cv::Point(-3, -3)); + + if (buffer) + clReleaseMemObject(buffer); + m_mem_obj = 0; + + return 0; +} + + +// this function is an example of interoperability between OpenCL image +// and OpenCV UMat objects. It converts OpenCL image +// to OpenCV UMat and then do blur on these data +int OpenClWrapper::process_cl_image_with_opencv(cl_mem image, cv::UMat& u) +{ + cv::ocl::convertFromImage(image, u); + + // process right half of frame in OpenCV + cv::Point pt(u.cols / 2, 0); + cv::Size sz(u.cols / 2, u.rows); + cv::Rect roi(pt, sz); + cv::UMat uroi(u, roi); + cv::blur(uroi, uroi, cv::Size(7, 7), cv::Point(-3, -3)); + + if (image) + clReleaseMemObject(image); + m_mem_obj = 0; + + if (m_img_src) + clReleaseMemObject(m_img_src); + m_img_src = 0; + + return 0; +} + + +int OpenClWrapper::run() +{ + if (0 != initOpenCL()) + return -1; + + if (0 != initVideoSource()) + return -1; + + Mat img_to_show; + + // set running state until ESC pressed + setRunning(true); + // set process flag to show some data processing + // can be toggled on/off by 'p' button + setDoProcess(true); + // set use buffer flag, + // when it is set to true, will demo interop opencl buffer and cv::Umat, + // otherwise demo interop opencl image and cv::UMat + // can be switched on/of by SPACE button + setUseBuffer(true); + + // Iterate over all frames + while (isRunning() && nextFrame(m_frame)) + { + cv::cvtColor(m_frame, m_frameGray, COLOR_BGR2GRAY); + + UMat uframe; + + // work + timerStart(); + + if (doProcess()) + { + process_frame_with_open_cl(m_frameGray, useBuffer(), &m_mem_obj); + + if (useBuffer()) + process_cl_buffer_with_opencv( + m_mem_obj, m_frameGray.step[0], m_frameGray.rows, m_frameGray.cols, m_frameGray.type(), uframe); + else + process_cl_image_with_opencv(m_mem_obj, uframe); + } + else + { + m_frameGray.copyTo(uframe); + } + + timerEnd(); + + uframe.copyTo(img_to_show); + + putText(img_to_show, "Version : " + m_platformInfo.Version(), Point(5, 30), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); + putText(img_to_show, "Name : " + m_platformInfo.Name(), Point(5, 60), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); + putText(img_to_show, "Device : " + m_deviceInfo.Name(), Point(5, 90), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); + cv::String memtype = useBuffer() ? "buffer" : "image"; + putText(img_to_show, "interop with OpenCL " + memtype, Point(5, 120), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); + putText(img_to_show, "Time : " + timeStr() + " msec", Point(5, 150), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); + + imshow("opencl_interop", img_to_show); + + handleKey((char)waitKey(3)); + } + + return 0; +} + diff --git a/mediapipe/framework/formats/helpers.hpp b/mediapipe/framework/formats/helpers.hpp new file mode 100644 index 0000000000..951300e113 --- /dev/null +++ b/mediapipe/framework/formats/helpers.hpp @@ -0,0 +1,471 @@ +/* +// The example of interoperability between OpenCL and OpenCV. +// This will loop through frames of video either from input media file +// or camera device and do processing of these data in OpenCL and then +// in OpenCV. In OpenCL it does inversion of pixels in left half of frame and +// in OpenCV it does blurring in the right half of frame. +*/ +#include +#include +#include +#include +#include +#include +#include +#include + +#define CL_USE_DEPRECATED_OPENCL_1_1_APIS +#define CL_USE_DEPRECATED_OPENCL_1_2_APIS +#define CL_USE_DEPRECATED_OPENCL_2_0_APIS // eliminate build warning +#define CL_TARGET_OPENCL_VERSION 200 // 2.0 + +#ifdef __APPLE__ +#define CL_SILENCE_DEPRECATION +#include +#else +#include +#endif + +#include +#include +#include +#include +#include + + +using namespace std; +using namespace cv; + +namespace opencl { + +class PlatformInfo +{ +public: + PlatformInfo() + {} + + ~PlatformInfo() + {} + + cl_int QueryInfo(cl_platform_id id) + { + query_param(id, CL_PLATFORM_PROFILE, m_profile); + query_param(id, CL_PLATFORM_VERSION, m_version); + query_param(id, CL_PLATFORM_NAME, m_name); + query_param(id, CL_PLATFORM_VENDOR, m_vendor); + query_param(id, CL_PLATFORM_EXTENSIONS, m_extensions); + return CL_SUCCESS; + } + + std::string Profile() { return m_profile; } + std::string Version() { return m_version; } + std::string Name() { return m_name; } + std::string Vendor() { return m_vendor; } + std::string Extensions() { return m_extensions; } + +private: + cl_int query_param(cl_platform_id id, cl_platform_info param, std::string& paramStr) + { + cl_int res; + + size_t psize; + cv::AutoBuffer buf; + + res = clGetPlatformInfo(id, param, 0, 0, &psize); + if (CL_SUCCESS != res) + throw std::runtime_error(std::string("clGetPlatformInfo failed")); + + buf.resize(psize); + res = clGetPlatformInfo(id, param, psize, buf, 0); + if (CL_SUCCESS != res) + throw std::runtime_error(std::string("clGetPlatformInfo failed")); + + // just in case, ensure trailing zero for ASCIIZ string + buf[psize] = 0; + + paramStr = buf; + + return CL_SUCCESS; + } + +private: + std::string m_profile; + std::string m_version; + std::string m_name; + std::string m_vendor; + std::string m_extensions; +}; + + +class DeviceInfo +{ +public: + DeviceInfo() + {} + + ~DeviceInfo() + {} + + cl_int QueryInfo(cl_device_id id) + { + query_param(id, CL_DEVICE_TYPE, m_type); + query_param(id, CL_DEVICE_VENDOR_ID, m_vendor_id); + query_param(id, CL_DEVICE_MAX_COMPUTE_UNITS, m_max_compute_units); + query_param(id, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, m_max_work_item_dimensions); + query_param(id, CL_DEVICE_MAX_WORK_ITEM_SIZES, m_max_work_item_sizes); + query_param(id, CL_DEVICE_MAX_WORK_GROUP_SIZE, m_max_work_group_size); + query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR, m_preferred_vector_width_char); + query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT, m_preferred_vector_width_short); + query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, m_preferred_vector_width_int); + query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG, m_preferred_vector_width_long); + query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT, m_preferred_vector_width_float); + query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE, m_preferred_vector_width_double); +#if defined(CL_VERSION_1_1) + query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF, m_preferred_vector_width_half); + query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR, m_native_vector_width_char); + query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT, m_native_vector_width_short); + query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, m_native_vector_width_int); + query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG, m_native_vector_width_long); + query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT, m_native_vector_width_float); + query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE, m_native_vector_width_double); + query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF, m_native_vector_width_half); +#endif + query_param(id, CL_DEVICE_MAX_CLOCK_FREQUENCY, m_max_clock_frequency); + query_param(id, CL_DEVICE_ADDRESS_BITS, m_address_bits); + query_param(id, CL_DEVICE_MAX_MEM_ALLOC_SIZE, m_max_mem_alloc_size); + query_param(id, CL_DEVICE_IMAGE_SUPPORT, m_image_support); + query_param(id, CL_DEVICE_MAX_READ_IMAGE_ARGS, m_max_read_image_args); + query_param(id, CL_DEVICE_MAX_WRITE_IMAGE_ARGS, m_max_write_image_args); +#if defined(CL_VERSION_2_0) + query_param(id, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS, m_max_read_write_image_args); +#endif + query_param(id, CL_DEVICE_IMAGE2D_MAX_WIDTH, m_image2d_max_width); + query_param(id, CL_DEVICE_IMAGE2D_MAX_HEIGHT, m_image2d_max_height); + query_param(id, CL_DEVICE_IMAGE3D_MAX_WIDTH, m_image3d_max_width); + query_param(id, CL_DEVICE_IMAGE3D_MAX_HEIGHT, m_image3d_max_height); + query_param(id, CL_DEVICE_IMAGE3D_MAX_DEPTH, m_image3d_max_depth); +#if defined(CL_VERSION_1_2) + query_param(id, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE, m_image_max_buffer_size); + query_param(id, CL_DEVICE_IMAGE_MAX_ARRAY_SIZE, m_image_max_array_size); +#endif + query_param(id, CL_DEVICE_MAX_SAMPLERS, m_max_samplers); +#if defined(CL_VERSION_1_2) + query_param(id, CL_DEVICE_IMAGE_PITCH_ALIGNMENT, m_image_pitch_alignment); + query_param(id, CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT, m_image_base_address_alignment); +#endif +#if defined(CL_VERSION_2_0) + query_param(id, CL_DEVICE_MAX_PIPE_ARGS, m_max_pipe_args); + query_param(id, CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS, m_pipe_max_active_reservations); + query_param(id, CL_DEVICE_PIPE_MAX_PACKET_SIZE, m_pipe_max_packet_size); +#endif + query_param(id, CL_DEVICE_MAX_PARAMETER_SIZE, m_max_parameter_size); + query_param(id, CL_DEVICE_MEM_BASE_ADDR_ALIGN, m_mem_base_addr_align); + query_param(id, CL_DEVICE_SINGLE_FP_CONFIG, m_single_fp_config); +#if defined(CL_VERSION_1_2) + query_param(id, CL_DEVICE_DOUBLE_FP_CONFIG, m_double_fp_config); +#endif + query_param(id, CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, m_global_mem_cache_type); + query_param(id, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, m_global_mem_cacheline_size); + query_param(id, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, m_global_mem_cache_size); + query_param(id, CL_DEVICE_GLOBAL_MEM_SIZE, m_global_mem_size); + query_param(id, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, m_max_constant_buffer_size); + query_param(id, CL_DEVICE_MAX_CONSTANT_ARGS, m_max_constant_args); +#if defined(CL_VERSION_2_0) + query_param(id, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE, m_max_global_variable_size); + query_param(id, CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE, m_global_variable_preferred_total_size); +#endif + query_param(id, CL_DEVICE_LOCAL_MEM_TYPE, m_local_mem_type); + query_param(id, CL_DEVICE_LOCAL_MEM_SIZE, m_local_mem_size); + query_param(id, CL_DEVICE_ERROR_CORRECTION_SUPPORT, m_error_correction_support); +#if defined(CL_VERSION_1_1) + query_param(id, CL_DEVICE_HOST_UNIFIED_MEMORY, m_host_unified_memory); +#endif + query_param(id, CL_DEVICE_PROFILING_TIMER_RESOLUTION, m_profiling_timer_resolution); + query_param(id, CL_DEVICE_ENDIAN_LITTLE, m_endian_little); + query_param(id, CL_DEVICE_AVAILABLE, m_available); + query_param(id, CL_DEVICE_COMPILER_AVAILABLE, m_compiler_available); +#if defined(CL_VERSION_1_2) + query_param(id, CL_DEVICE_LINKER_AVAILABLE, m_linker_available); +#endif + query_param(id, CL_DEVICE_EXECUTION_CAPABILITIES, m_execution_capabilities); + query_param(id, CL_DEVICE_QUEUE_PROPERTIES, m_queue_properties); +#if defined(CL_VERSION_2_0) + query_param(id, CL_DEVICE_QUEUE_ON_HOST_PROPERTIES, m_queue_on_host_properties); + query_param(id, CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES, m_queue_on_device_properties); + query_param(id, CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE, m_queue_on_device_preferred_size); + query_param(id, CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE, m_queue_on_device_max_size); + query_param(id, CL_DEVICE_MAX_ON_DEVICE_QUEUES, m_max_on_device_queues); + query_param(id, CL_DEVICE_MAX_ON_DEVICE_EVENTS, m_max_on_device_events); +#endif +#if defined(CL_VERSION_1_2) + query_param(id, CL_DEVICE_BUILT_IN_KERNELS, m_built_in_kernels); +#endif + query_param(id, CL_DEVICE_PLATFORM, m_platform); + query_param(id, CL_DEVICE_NAME, m_name); + query_param(id, CL_DEVICE_VENDOR, m_vendor); + query_param(id, CL_DRIVER_VERSION, m_driver_version); + query_param(id, CL_DEVICE_PROFILE, m_profile); + query_param(id, CL_DEVICE_VERSION, m_version); +#if defined(CL_VERSION_1_1) + query_param(id, CL_DEVICE_OPENCL_C_VERSION, m_opencl_c_version); +#endif + query_param(id, CL_DEVICE_EXTENSIONS, m_extensions); +#if defined(CL_VERSION_1_2) + query_param(id, CL_DEVICE_PRINTF_BUFFER_SIZE, m_printf_buffer_size); + query_param(id, CL_DEVICE_PREFERRED_INTEROP_USER_SYNC, m_preferred_interop_user_sync); + query_param(id, CL_DEVICE_PARENT_DEVICE, m_parent_device); + query_param(id, CL_DEVICE_PARTITION_MAX_SUB_DEVICES, m_partition_max_sub_devices); + query_param(id, CL_DEVICE_PARTITION_PROPERTIES, m_partition_properties); + query_param(id, CL_DEVICE_PARTITION_AFFINITY_DOMAIN, m_partition_affinity_domain); + query_param(id, CL_DEVICE_PARTITION_TYPE, m_partition_type); + query_param(id, CL_DEVICE_REFERENCE_COUNT, m_reference_count); +#endif + return CL_SUCCESS; + } + + std::string Name() { return m_name; } + +private: + template + cl_int query_param(cl_device_id id, cl_device_info param, T& value) + { + cl_int res; + size_t size = 0; + + res = clGetDeviceInfo(id, param, 0, 0, &size); + if (CL_SUCCESS != res && size != 0) + throw std::runtime_error(std::string("clGetDeviceInfo failed")); + + if (0 == size) + return CL_SUCCESS; + + if (sizeof(T) != size) + throw std::runtime_error(std::string("clGetDeviceInfo: param size mismatch")); + + res = clGetDeviceInfo(id, param, size, &value, 0); + if (CL_SUCCESS != res) + throw std::runtime_error(std::string("clGetDeviceInfo failed")); + + return CL_SUCCESS; + } + + template + cl_int query_param(cl_device_id id, cl_device_info param, std::vector& value) + { + cl_int res; + size_t size; + + res = clGetDeviceInfo(id, param, 0, 0, &size); + if (CL_SUCCESS != res) + throw std::runtime_error(std::string("clGetDeviceInfo failed")); + + if (0 == size) + return CL_SUCCESS; + + value.resize(size / sizeof(T)); + + res = clGetDeviceInfo(id, param, size, &value[0], 0); + if (CL_SUCCESS != res) + throw std::runtime_error(std::string("clGetDeviceInfo failed")); + + return CL_SUCCESS; + } + + cl_int query_param(cl_device_id id, cl_device_info param, std::string& value) + { + cl_int res; + size_t size; + + res = clGetDeviceInfo(id, param, 0, 0, &size); + if (CL_SUCCESS != res) + throw std::runtime_error(std::string("clGetDeviceInfo failed")); + + value.resize(size + 1); + + res = clGetDeviceInfo(id, param, size, &value[0], 0); + if (CL_SUCCESS != res) + throw std::runtime_error(std::string("clGetDeviceInfo failed")); + + // just in case, ensure trailing zero for ASCIIZ string + value[size] = 0; + + return CL_SUCCESS; + } + +private: + cl_device_type m_type; + cl_uint m_vendor_id; + cl_uint m_max_compute_units; + cl_uint m_max_work_item_dimensions; + std::vector m_max_work_item_sizes; + size_t m_max_work_group_size; + cl_uint m_preferred_vector_width_char; + cl_uint m_preferred_vector_width_short; + cl_uint m_preferred_vector_width_int; + cl_uint m_preferred_vector_width_long; + cl_uint m_preferred_vector_width_float; + cl_uint m_preferred_vector_width_double; +#if defined(CL_VERSION_1_1) + cl_uint m_preferred_vector_width_half; + cl_uint m_native_vector_width_char; + cl_uint m_native_vector_width_short; + cl_uint m_native_vector_width_int; + cl_uint m_native_vector_width_long; + cl_uint m_native_vector_width_float; + cl_uint m_native_vector_width_double; + cl_uint m_native_vector_width_half; +#endif + cl_uint m_max_clock_frequency; + cl_uint m_address_bits; + cl_ulong m_max_mem_alloc_size; + cl_bool m_image_support; + cl_uint m_max_read_image_args; + cl_uint m_max_write_image_args; +#if defined(CL_VERSION_2_0) + cl_uint m_max_read_write_image_args; +#endif + size_t m_image2d_max_width; + size_t m_image2d_max_height; + size_t m_image3d_max_width; + size_t m_image3d_max_height; + size_t m_image3d_max_depth; +#if defined(CL_VERSION_1_2) + size_t m_image_max_buffer_size; + size_t m_image_max_array_size; +#endif + cl_uint m_max_samplers; +#if defined(CL_VERSION_1_2) + cl_uint m_image_pitch_alignment; + cl_uint m_image_base_address_alignment; +#endif +#if defined(CL_VERSION_2_0) + cl_uint m_max_pipe_args; + cl_uint m_pipe_max_active_reservations; + cl_uint m_pipe_max_packet_size; +#endif + size_t m_max_parameter_size; + cl_uint m_mem_base_addr_align; + cl_device_fp_config m_single_fp_config; +#if defined(CL_VERSION_1_2) + cl_device_fp_config m_double_fp_config; +#endif + cl_device_mem_cache_type m_global_mem_cache_type; + cl_uint m_global_mem_cacheline_size; + cl_ulong m_global_mem_cache_size; + cl_ulong m_global_mem_size; + cl_ulong m_max_constant_buffer_size; + cl_uint m_max_constant_args; +#if defined(CL_VERSION_2_0) + size_t m_max_global_variable_size; + size_t m_global_variable_preferred_total_size; +#endif + cl_device_local_mem_type m_local_mem_type; + cl_ulong m_local_mem_size; + cl_bool m_error_correction_support; +#if defined(CL_VERSION_1_1) + cl_bool m_host_unified_memory; +#endif + size_t m_profiling_timer_resolution; + cl_bool m_endian_little; + cl_bool m_available; + cl_bool m_compiler_available; +#if defined(CL_VERSION_1_2) + cl_bool m_linker_available; +#endif + cl_device_exec_capabilities m_execution_capabilities; + cl_command_queue_properties m_queue_properties; +#if defined(CL_VERSION_2_0) + cl_command_queue_properties m_queue_on_host_properties; + cl_command_queue_properties m_queue_on_device_properties; + cl_uint m_queue_on_device_preferred_size; + cl_uint m_queue_on_device_max_size; + cl_uint m_max_on_device_queues; + cl_uint m_max_on_device_events; +#endif +#if defined(CL_VERSION_1_2) + std::string m_built_in_kernels; +#endif + cl_platform_id m_platform; + std::string m_name; + std::string m_vendor; + std::string m_driver_version; + std::string m_profile; + std::string m_version; +#if defined(CL_VERSION_1_1) + std::string m_opencl_c_version; +#endif + std::string m_extensions; +#if defined(CL_VERSION_1_2) + size_t m_printf_buffer_size; + cl_bool m_preferred_interop_user_sync; + cl_device_id m_parent_device; + cl_uint m_partition_max_sub_devices; + std::vector m_partition_properties; + cl_device_affinity_domain m_partition_affinity_domain; + std::vector m_partition_type; + cl_uint m_reference_count; +#endif +}; + +} // namespace opencl + + +class OpenClWrapper +{ +public: + OpenClWrapper(); + ~OpenClWrapper(); + + int initOpenCL(); + int initVideoSource(); + + int process_frame_with_open_cl(cv::Mat& frame, bool use_buffer, cl_mem* cl_buffer); + int process_cl_buffer_with_opencv(cl_mem buffer, size_t step, int rows, int cols, int type, cv::UMat& u); + int process_cl_image_with_opencv(cl_mem image, cv::UMat& u); + + int run(); + + bool isRunning() { return m_running; } + bool doProcess() { return m_process; } + bool useBuffer() { return m_use_buffer; } + + void setRunning(bool running) { m_running = running; } + void setDoProcess(bool process) { m_process = process; } + void setUseBuffer(bool use_buffer) { m_use_buffer = use_buffer; } + +protected: + bool nextFrame(cv::Mat& frame) { return m_cap.read(frame); } + void handleKey(char key); + void timerStart(); + void timerEnd(); + std::string timeStr() const; + std::string message() const; + +private: + bool m_running; + bool m_process; + bool m_use_buffer; + + int64 m_t0; + int64 m_t1; + float m_time; + float m_frequency; + + string m_file_name; + int m_camera_id; + cv::VideoCapture m_cap; + cv::Mat m_frame; + cv::Mat m_frameGray; + + opencl::PlatformInfo m_platformInfo; + opencl::DeviceInfo m_deviceInfo; + std::vector m_platform_ids; + cl_context m_context; + cl_device_id m_device_id; + cl_command_queue m_queue; + cl_program m_program; + cl_kernel m_kernelBuf; + cl_kernel m_kernelImg; + cl_mem m_img_src; // used as src in case processing of cl image + cl_mem m_mem_obj; +}; \ No newline at end of file diff --git a/mediapipe/framework/formats/image_frame.cc b/mediapipe/framework/formats/image_frame.cc index 2de819a35b..a3a982c4ea 100644 --- a/mediapipe/framework/formats/image_frame.cc +++ b/mediapipe/framework/formats/image_frame.cc @@ -20,6 +20,8 @@ #include #include +#include +#include #include #include @@ -28,6 +30,7 @@ #include "mediapipe/framework/port/aligned_malloc_and_free.h" #include "mediapipe/framework/port/logging.h" #include "mediapipe/framework/port/proto_ns.h" +#include "mediapipe/framework/formats/helpers.hpp" namespace mediapipe { @@ -46,6 +49,10 @@ int CountOnes(uint32_t n) { } // namespace +const char* DEVICE_GPU = ":GPU:0"; + +static cv::ocl::Context context = cv::ocl::Context::create(DEVICE_GPU); + const ImageFrame::Deleter ImageFrame::PixelDataDeleter::kArrayDelete = std::default_delete(); const ImageFrame::Deleter ImageFrame::PixelDataDeleter::kFree = free; @@ -93,6 +100,7 @@ ImageFrame& ImageFrame::operator=(ImageFrame&& move_from) { return *this; } + void ImageFrame::Reset(ImageFormat::Format format, int width, int height, uint32_t alignment_boundary) { format_ = format; @@ -101,7 +109,7 @@ void ImageFrame::Reset(ImageFormat::Format format, int width, int height, CHECK_NE(ImageFormat::UNKNOWN, format_); CHECK(IsValidAlignmentNumber(alignment_boundary)); width_step_ = width * NumberOfChannels() * ByteDepth(); - if (alignment_boundary == 1) { + /*if (alignment_boundary == 1) { pixel_data_ = {new uint8_t[height * width_step_], PixelDataDeleter::kArrayDelete}; } else { @@ -113,7 +121,7 @@ void ImageFrame::Reset(ImageFormat::Format format, int width, int height, pixel_data_ = {reinterpret_cast(aligned_malloc( height * width_step_, alignment_boundary)), PixelDataDeleter::kAlignedFree}; - } + }*/ } void ImageFrame::AdoptPixelData(ImageFormat::Format format, int width, diff --git a/mediapipe/framework/formats/image_frame_opencv.cc b/mediapipe/framework/formats/image_frame_opencv.cc index d862c54079..3831e9742e 100644 --- a/mediapipe/framework/formats/image_frame_opencv.cc +++ b/mediapipe/framework/formats/image_frame_opencv.cc @@ -95,10 +95,11 @@ cv::Mat MatView(const ImageFrame* image) { // UMatUsageFlags // USAGE_DEFAULT, USAGE_ALLOCATE_HOST_MEMORY, USAGE_ALLOCATE_DEVICE_MEMORY, USAGE_ALLOCATE_SHARED_MEMORY , __UMAT_USAGE_FLAGS_32BIT -cv::UMat MatView(const ImageFrame* image, cv::UMatUsageFlags usageFlags) { +cv::UMat MatView(ImageFrame* image, cv::UMatUsageFlags usageFlags) { cv::Mat mat = MatView(image); cv::UMat umat = cv::UMat(usageFlags); mat.copyTo(umat); + //*image->MutablePixelData() = static_cast(umat.u->data); return umat; } diff --git a/mediapipe/framework/formats/image_frame_opencv.h b/mediapipe/framework/formats/image_frame_opencv.h index 8af442bf46..c472108242 100644 --- a/mediapipe/framework/formats/image_frame_opencv.h +++ b/mediapipe/framework/formats/image_frame_opencv.h @@ -33,7 +33,7 @@ cv::Mat MatView(const ImageFrame* image); // UMatUsageFlags // USAGE_DEFAULT, USAGE_ALLOCATE_HOST_MEMORY, USAGE_ALLOCATE_DEVICE_MEMORY, USAGE_ALLOCATE_SHARED_MEMORY , __UMAT_USAGE_FLAGS_32BIT -cv::UMat MatView(const ImageFrame* image, cv::UMatUsageFlags usageFlags); +cv::UMat MatView(ImageFrame* image, cv::UMatUsageFlags usageFlags); } // namespace formats } // namespace mediapipe diff --git a/mediapipe/framework/port/BUILD b/mediapipe/framework/port/BUILD index 5894e4715c..ad6f78871c 100644 --- a/mediapipe/framework/port/BUILD +++ b/mediapipe/framework/port/BUILD @@ -242,6 +242,14 @@ cc_library( ], ) +cc_library( + name = "opencv_opencl", + hdrs = ["opencv_opencl.h"], + deps = [ + "//third_party:opencv", + ], +) + cc_library( name = "opencv_imgproc", hdrs = ["opencv_imgproc_inc.h"], diff --git a/mediapipe/framework/port/opencv_opencl.h b/mediapipe/framework/port/opencv_opencl.h new file mode 100644 index 0000000000..c5455b20e7 --- /dev/null +++ b/mediapipe/framework/port/opencv_opencl.h @@ -0,0 +1,20 @@ +// Copyright 2019 The MediaPipe Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef MEDIAPIPE_PORT_OPENCV_CORE_INC_H_ +#define MEDIAPIPE_PORT_OPENCV_CORE_INC_H_ + +#include + +#endif // MEDIAPIPE_PORT_OPENCV_CORE_INC_H_ From 6f6131ae8e2a1c512d231442ab3a6f65b2a1a081 Mon Sep 17 00:00:00 2001 From: "iotg.worker" Date: Tue, 2 Apr 2024 16:08:35 +0200 Subject: [PATCH 3/7] Full pass segfault --- Dockerfile.openvino | 20 +++ Makefile | 4 + .../image/image_transformation_calculator.cc | 124 +++++++++++++- .../util/annotation_overlay_calculator.cc | 24 ++- .../examples/desktop/demo_run_graph_main.cc | 32 ++-- mediapipe/framework/formats/BUILD | 4 +- mediapipe/framework/formats/helpers.cpp | 160 +++++++++++------- mediapipe/framework/formats/helpers.hpp | 23 ++- mediapipe/framework/formats/image_frame.cc | 59 ++++++- mediapipe/framework/formats/image_frame.h | 10 ++ .../framework/formats/image_frame_opencv.cc | 12 +- 11 files changed, 364 insertions(+), 108 deletions(-) diff --git a/Dockerfile.openvino b/Dockerfile.openvino index 64dfebdc71..b4f4a06e6e 100644 --- a/Dockerfile.openvino +++ b/Dockerfile.openvino @@ -170,6 +170,26 @@ RUN cd flatbuffers && mkdir build && cd build && cmake -G "Unix Makefiles" -DCMA ## End of base image +## GPU OCL support +ARG INSTALL_DRIVER_VERSION="23.13.26032" +ARG GPU=0 +if [ "$GPU" == "1" ] ; then \ + apt-get update && apt-get install -y ocl-icd-opencl-dev libnuma1 ocl-icd-libopencl1 --no-install-recommends && rm -rf /var/lib/apt/lists/* && \ + case $INSTALL_DRIVER_VERSION in \ + "23.22.26516") \ + mkdir /tmp/gpu_deps && cd /tmp/gpu_deps ; \ + curl -L -O https://github.com/intel/compute-runtime/releases/download/23.22.26516.18/intel-level-zero-gpu_1.3.26516.18_amd64.deb ; \ + curl -L -O https://github.com/intel/intel-graphics-compiler/releases/download/igc-1.0.14062.11/intel-igc-core_1.0.14062.11_amd64.deb ; \ + curl -L -O https://github.com/intel/intel-graphics-compiler/releases/download/igc-1.0.14062.11/intel-igc-opencl_1.0.14062.11_amd64.deb ; \ + curl -L -O https://github.com/intel/compute-runtime/releases/download/23.22.26516.18/intel-opencl-icd_23.22.26516.18_amd64.deb ; \ + curl -L -O https://github.com/intel/compute-runtime/releases/download/23.22.26516.18/libigdgmm12_22.3.0_amd64.deb ; \ + dpkg -i *.deb && rm -Rf /tmp/gpu_deps ; \ + ;; \ + esac ; \ +fi ; \ +apt-get clean ; \ +rm -rf /var/lib/apt/lists/* && rm -rf /tmp/* ; + FROM base as build RUN wget -O video.mp4 "https://www.pexels.com/download/video/3044127/?fps=24.0&h=1080&w=1920" diff --git a/Makefile b/Makefile index 1fc4abf41b..8b8084cbba 100644 --- a/Makefile +++ b/Makefile @@ -23,6 +23,8 @@ INPUT_VIDEO_LINK ?= "https://www.pexels.com/download/video/3044127/?fps=24.0&h=1 OVMS_COMMIT ?="7bf5d7fc217a81ce4b0a089c7679ecb77e6c544c" JOBS ?= $(shell python3 -c 'import multiprocessing as mp; print(mp.cpu_count())') DLDT_PACKAGE_URL ?= https://storage.openvinotoolkit.org/repositories/openvino/packages/2024.0/linux/l_openvino_toolkit_ubuntu20_2024.0.0.14509.34caeefd078_x86_64.tgz +INSTALL_DRIVER_VERSION ?= "23.22.26516" +GPU ?= 0 # Targets to use outside running mediapipe_ovms container docker_build: @@ -32,6 +34,8 @@ docker_build: --build-arg JOBS=$(JOBS) . \ --build-arg OVMS_COMMIT=$(OVMS_COMMIT) \ -t $(OVMS_MEDIA_DOCKER_IMAGE):$(OVMS_MEDIA_IMAGE_TAG) + --build-arg INSTALL_DRIVER_VERSION=$(INSTALL_DRIVER_VERSION)\ + --build-arg GPU=$(GPU)\ tests: run_unit_tests run_hello_world run_hello_ovms run_hello_ovms: diff --git a/mediapipe/calculators/image/image_transformation_calculator.cc b/mediapipe/calculators/image/image_transformation_calculator.cc index dbf8f73374..51b97c57d9 100644 --- a/mediapipe/calculators/image/image_transformation_calculator.cc +++ b/mediapipe/calculators/image/image_transformation_calculator.cc @@ -190,6 +190,7 @@ class ImageTransformationCalculator : public CalculatorBase { private: absl::Status RenderCpu(CalculatorContext* cc); absl::Status RenderGpu(CalculatorContext* cc); + absl::Status RenderOpenCl(CalculatorContext* cc); absl::Status GlSetup(); void ComputeOutputDimensions(int input_width, int input_height, @@ -411,7 +412,8 @@ absl::Status ImageTransformationCalculator::Process(CalculatorContext* cc) { if (cc->Inputs().Tag(kImageFrameTag).IsEmpty()) { return absl::OkStatus(); } - return RenderCpu(cc); + //return RenderCpu(cc); + return RenderOpenCl(cc); } return absl::OkStatus(); } @@ -442,6 +444,126 @@ absl::Status ImageTransformationCalculator::Close(CalculatorContext* cc) { return absl::OkStatus(); } +absl::Status ImageTransformationCalculator::RenderOpenCl(CalculatorContext* cc) { + // TODO UMAT + cv::Mat input_mat; + mediapipe::ImageFormat::Format format; + + const auto& input = cc->Inputs().Tag(kImageFrameTag).Get(); + std::cout << " cc->Inputs().Tag(kImageFrameTag).Get(); " << std::endl; + // input_mat = formats::MatView(const_cast(&input), cv::USAGE_ALLOCATE_SHARED_MEMORY); + input_mat = formats::MatView(&input); + std::string fileName = std::string("./imageTrans/input") + std::to_string(cc->InputTimestamp().Value()); + //dumpMatToFile(fileName, input_mat); + + format = input.Format(); + + const int input_width = input_mat.cols; + const int input_height = input_mat.rows; + int output_width; + int output_height; + ComputeOutputDimensions(input_width, input_height, &output_width, + &output_height); + + if (output_width_ > 0 && output_height_ > 0) { + // TODO umat + cv::Mat scaled_mat; + if (scale_mode_ == mediapipe::ScaleMode_Mode_STRETCH) { + int scale_flag = + input_mat.cols > output_width_ && input_mat.rows > output_height_ + ? cv::INTER_AREA + : cv::INTER_LINEAR; + cv::resize(input_mat, scaled_mat, cv::Size(output_width_, output_height_), + 0, 0, scale_flag); + } else { + const float scale = + std::min(static_cast(output_width_) / input_width, + static_cast(output_height_) / input_height); + const int target_width = std::round(input_width * scale); + const int target_height = std::round(input_height * scale); + int scale_flag = scale < 1.0f ? cv::INTER_AREA : cv::INTER_LINEAR; + if (scale_mode_ == mediapipe::ScaleMode_Mode_FIT) { + cv::UMat intermediate_mat; + cv::resize(input_mat, intermediate_mat, + cv::Size(target_width, target_height), 0, 0, scale_flag); + const int top = (output_height_ - target_height) / 2; + const int bottom = output_height_ - target_height - top; + const int left = (output_width_ - target_width) / 2; + const int right = output_width_ - target_width - left; + cv::copyMakeBorder(intermediate_mat, scaled_mat, top, bottom, left, + right, + options_.constant_padding() ? cv::BORDER_CONSTANT + : cv::BORDER_REPLICATE, + padding_color_); + } else { + cv::resize(input_mat, scaled_mat, cv::Size(target_width, target_height), + 0, 0, scale_flag); + output_width = target_width; + output_height = target_height; + } + } + input_mat = scaled_mat; + } + + if (cc->Outputs().HasTag("LETTERBOX_PADDING")) { + auto padding = absl::make_unique>(); + ComputeOutputLetterboxPadding(input_width, input_height, output_width, + output_height, padding.get()); + cc->Outputs() + .Tag("LETTERBOX_PADDING") + .Add(padding.release(), cc->InputTimestamp()); + } + + //TODO umat + cv::Mat rotated_mat; + cv::Size rotated_size(output_width, output_height); + if (input_mat.size() == rotated_size) { + const int angle = RotationModeToDegrees(rotation_); + cv::Point2f src_center(input_mat.cols / 2.0, input_mat.rows / 2.0); + // TODO UMAT + cv::Mat rotation_mat = cv::getRotationMatrix2D(src_center, angle, 1.0); + cv::warpAffine(input_mat, rotated_mat, rotation_mat, rotated_size); + } else { + switch (rotation_) { + case mediapipe::RotationMode_Mode_UNKNOWN: + case mediapipe::RotationMode_Mode_ROTATION_0: + rotated_mat = input_mat; + break; + case mediapipe::RotationMode_Mode_ROTATION_90: + cv::rotate(input_mat, rotated_mat, cv::ROTATE_90_COUNTERCLOCKWISE); + break; + case mediapipe::RotationMode_Mode_ROTATION_180: + cv::rotate(input_mat, rotated_mat, cv::ROTATE_180); + break; + case mediapipe::RotationMode_Mode_ROTATION_270: + cv::rotate(input_mat, rotated_mat, cv::ROTATE_90_CLOCKWISE); + break; + } + } + + // TODO UMAT + cv::Mat flipped_mat; + if (flip_horizontally_ || flip_vertically_) { + const int flip_code = + flip_horizontally_ && flip_vertically_ ? -1 : flip_horizontally_; + cv::flip(rotated_mat, flipped_mat, flip_code); + } else { + flipped_mat = rotated_mat; + } + + std::unique_ptr output_frame( + new ImageFrame(flipped_mat, format, output_width, output_height)); + + fileName = std::string("./imageTrans/output") + std::to_string(cc->InputTimestamp().Value()); + //dumpMatToFile(fileName, flipped_mat); + + cc->Outputs() + .Tag(kImageFrameTag) + .Add(output_frame.release(), cc->InputTimestamp()); + + return absl::OkStatus(); +} + absl::Status ImageTransformationCalculator::RenderCpu(CalculatorContext* cc) { cv::Mat input_mat; mediapipe::ImageFormat::Format format; diff --git a/mediapipe/calculators/util/annotation_overlay_calculator.cc b/mediapipe/calculators/util/annotation_overlay_calculator.cc index 5afede99d2..4033823a0e 100644 --- a/mediapipe/calculators/util/annotation_overlay_calculator.cc +++ b/mediapipe/calculators/util/annotation_overlay_calculator.cc @@ -143,6 +143,9 @@ class AnnotationOverlayCalculator : public CalculatorBase { absl::Status RenderToCpu(CalculatorContext* cc, const ImageFormat::Format& target_format, uchar* data_image); + absl::Status RenderToOpencl( + CalculatorContext* cc, const ImageFormat::Format& target_format, + cv::Mat& input_mat); absl::Status GlRender(CalculatorContext* cc); template @@ -345,8 +348,10 @@ absl::Status AnnotationOverlayCalculator::Process(CalculatorContext* cc) { #endif // !MEDIAPIPE_DISABLE_GPU } else { // Copy the rendered image to output. - uchar* image_mat_ptr = image_mat->data; - MP_RETURN_IF_ERROR(RenderToCpu(cc, target_format, image_mat_ptr)); + //uchar* image_mat_ptr = image_mat->data; + //MP_RETURN_IF_ERROR(RenderToCpu(cc, target_format, image_mat_ptr)); + + MP_RETURN_IF_ERROR(RenderToOpencl(cc, target_format, *image_mat.get())); } return absl::OkStatus(); @@ -365,6 +370,21 @@ absl::Status AnnotationOverlayCalculator::Close(CalculatorContext* cc) { return absl::OkStatus(); } +absl::Status AnnotationOverlayCalculator::RenderToOpencl( + CalculatorContext* cc, const ImageFormat::Format& target_format, + cv::Mat& input_mat) { + std::unique_ptr output_frame( + new ImageFrame(input_mat, target_format, renderer_->GetImageWidth(), renderer_->GetImageHeight())); + + if (cc->Outputs().HasTag(kImageFrameTag)) { + cc->Outputs() + .Tag(kImageFrameTag) + .Add(output_frame.release(), cc->InputTimestamp()); + } + + return absl::OkStatus(); +} + absl::Status AnnotationOverlayCalculator::RenderToCpu( CalculatorContext* cc, const ImageFormat::Format& target_format, uchar* data_image) { diff --git a/mediapipe/examples/desktop/demo_run_graph_main.cc b/mediapipe/examples/desktop/demo_run_graph_main.cc index d33b19dd11..1cbacf638e 100644 --- a/mediapipe/examples/desktop/demo_run_graph_main.cc +++ b/mediapipe/examples/desktop/demo_run_graph_main.cc @@ -22,7 +22,6 @@ #include "absl/flags/parse.h" #include "mediapipe/framework/calculator_framework.h" #include "mediapipe/framework/formats/image_frame.h" -#include "mediapipe/framework/formats/helpers.hpp" #include "mediapipe/framework/formats/image_frame_opencv.h" #include "mediapipe/framework/port/file_helpers.h" #include "mediapipe/framework/port/opencv_highgui_inc.h" @@ -45,16 +44,6 @@ ABSL_FLAG(std::string, output_video_path, "", "Full path of where to save result (.mp4 only). " "If not provided, show result in a window."); -void dumpMatToFile(std::string& fileName, cv::UMat& umat){ - cv::FileStorage file(fileName, cv::FileStorage::WRITE); - cv::Mat input; - //file << "camera_frame_raw" << camera_frame_raw; - umat.copyTo(input); - file << "input_frame_mat" << input; - - file.release(); -} - absl::Status RunMPPGraph() { std::string calculator_graph_config_contents; MP_RETURN_IF_ERROR(mediapipe::file::GetContents( @@ -101,9 +90,9 @@ absl::Status RunMPPGraph() { int count_frames = 0; auto begin = std::chrono::high_resolution_clock::now(); - OpenClWrapper ocl; - - ocl.initOpenCL(); + //OpenClWrapper ocl; + //ocl.initOpenCL(); + //ocl.printInfo(); LOG(INFO) << "haveOpenCL " << cv::ocl::haveOpenCL() <> camera_frame_raw; if (camera_frame_raw.empty()) { if (!load_video) { @@ -134,7 +123,7 @@ absl::Status RunMPPGraph() { break; } count_frames+=1; - cv::UMat camera_frame = cv::UMat(cv::USAGE_ALLOCATE_SHARED_MEMORY); + cv::Mat camera_frame; cv::cvtColor(camera_frame_raw, camera_frame, cv::COLOR_BGR2RGB); if (!load_video) { cv::flip(camera_frame, camera_frame, /*flipcode=HORIZONTAL*/ 1); @@ -142,12 +131,13 @@ absl::Status RunMPPGraph() { // Wrap Mat INFO an ImageFrame. auto input_frame = absl::make_unique( + camera_frame, mediapipe::ImageFormat::SRGB, camera_frame.cols, camera_frame.rows, mediapipe::ImageFrame::kDefaultAlignmentBoundary); - cv::UMat input_frame_mat = mediapipe::formats::MatView(input_frame.get(), cv::USAGE_ALLOCATE_SHARED_MEMORY); - camera_frame.copyTo(input_frame_mat); - //std::string fileName = std::string("./input") + std::to_string(count_frames) + std::string("jpeg"); - + //cv::UMat input_frame_mat = mediapipe::formats::MatView(input_frame.get(), cv::USAGE_ALLOCATE_SHARED_MEMORY); + //camera_frame.copyTo(input_frame_mat); + std::string fileName = std::string("./imageDemo/input") + std::to_string(count_frames); + dumpMatToFile(fileName, camera_frame); // Send image packet INFO the graph. diff --git a/mediapipe/framework/formats/BUILD b/mediapipe/framework/formats/BUILD index c87454a174..a4590f9d9c 100644 --- a/mediapipe/framework/formats/BUILD +++ b/mediapipe/framework/formats/BUILD @@ -148,9 +148,9 @@ cc_library( cc_library( name = "image_frame", srcs = ["image_frame.cc", - "helpers.hpp", "helpers.cpp"], - hdrs = ["image_frame.h"], + hdrs = ["image_frame.h", + "helpers.hpp"], deps = [ ":image_format_cc_proto", "//mediapipe/framework:port", diff --git a/mediapipe/framework/formats/helpers.cpp b/mediapipe/framework/formats/helpers.cpp index 3b67913d62..32d14cf45f 100644 --- a/mediapipe/framework/formats/helpers.cpp +++ b/mediapipe/framework/formats/helpers.cpp @@ -9,6 +9,7 @@ #include #include #include +#include #include #include #include @@ -36,16 +37,34 @@ using namespace std; using namespace cv; +static std::mutex g_dump_mutex; +static std::mutex g_dump_mutex2; + +void dumpMatToFile(std::string& fileName, cv::UMat& umat){ + std::lock_guard guard(g_dump_mutex); + cv::FileStorage file(fileName, cv::FileStorage::WRITE); + cv::Mat input; + //file << "camera_frame_raw" << camera_frame_raw; + umat.copyTo(input); + file << "input_frame_mat" << input; + + file.release(); +} + +void dumpMatToFile(std::string& fileName, cv::Mat& umat){ + std::lock_guard guard(g_dump_mutex2); + cv::FileStorage file(fileName, cv::FileStorage::WRITE); + file << "input_frame_mat" << umat; + + file.release(); +} OpenClWrapper::OpenClWrapper() { - cout << "\nPress ESC to exit\n" << endl; - cout << "\n 'p' to toggle ON/OFF processing\n" << endl; - cout << "\n SPACE to switch between OpenCL buffer/image\n" << endl; - m_camera_id = 0; //cmd.get("camera"); m_file_name = "0"; //cmd.get("video"); + m_is_initialized = false; m_running = false; m_process = false; m_use_buffer = false; @@ -56,7 +75,7 @@ OpenClWrapper::OpenClWrapper() m_frequency = (float)cv::getTickFrequency(); m_context = 0; - m_device_id = 0; + //m_device_id = 0; m_queue = 0; m_program = 0; m_kernelBuf = 0; @@ -105,11 +124,11 @@ OpenClWrapper::~OpenClWrapper() m_kernelImg = 0; } - if (m_device_id) + /*if (m_device_id) { clReleaseDevice(m_device_id); m_device_id = 0; - } + }*/ if (m_context) { @@ -119,8 +138,18 @@ OpenClWrapper::~OpenClWrapper() } // dtor +bool OpenClWrapper::m_is_initialized = false; +cl_context OpenClWrapper::m_context = nullptr; +cl_command_queue OpenClWrapper::m_queue = nullptr; + + int OpenClWrapper::initOpenCL() { + if (m_is_initialized){ + std::cout << "OpenCL already initialized." << std::endl; + return CL_SUCCESS; + } + cl_int res = CL_SUCCESS; cl_uint num_entries = 0; @@ -128,6 +157,12 @@ int OpenClWrapper::initOpenCL() if (CL_SUCCESS != res) return -1; + opencl::PlatformInfo m_platformInfo; + opencl::DeviceInfo m_deviceInfo; + cl_device_id m_device_id; + + std::vector m_platform_ids; + m_platform_ids.resize(num_entries); res = clGetPlatformIDs(num_entries, &m_platform_ids[0], 0); @@ -158,48 +193,6 @@ int OpenClWrapper::initOpenCL() if (0 == m_queue || CL_SUCCESS != res) return -1; - const char* kernelSrc = - "__kernel " - "void bitwise_inv_buf_8uC1(" - " __global unsigned char* pSrcDst," - " int srcDstStep," - " int rows," - " int cols)" - "{" - " int x = get_global_id(0);" - " int y = get_global_id(1);" - " int idx = mad24(y, srcDstStep, x);" - " pSrcDst[idx] = ~pSrcDst[idx];" - "}" - "__kernel " - "void bitwise_inv_img_8uC1(" - " read_only image2d_t srcImg," - " write_only image2d_t dstImg)" - "{" - " int x = get_global_id(0);" - " int y = get_global_id(1);" - " int2 coord = (int2)(x, y);" - " uint4 val = read_imageui(srcImg, coord);" - " val.x = (~val.x) & 0x000000FF;" - " write_imageui(dstImg, coord, val);" - "}"; - size_t len = strlen(kernelSrc); - m_program = clCreateProgramWithSource(m_context, 1, &kernelSrc, &len, &res); - if (0 == m_program || CL_SUCCESS != res) - return -1; - - res = clBuildProgram(m_program, 1, &m_device_id, 0, 0, 0); - if (CL_SUCCESS != res) - return -1; - - m_kernelBuf = clCreateKernel(m_program, "bitwise_inv_buf_8uC1", &res); - if (0 == m_kernelBuf || CL_SUCCESS != res) - return -1; - - m_kernelImg = clCreateKernel(m_program, "bitwise_inv_img_8uC1", &res); - if (0 == m_kernelImg || CL_SUCCESS != res) - return -1; - m_platformInfo.QueryInfo(m_platform_ids[i]); m_deviceInfo.QueryInfo(m_device_id); @@ -209,9 +202,57 @@ int OpenClWrapper::initOpenCL() break; } + m_is_initialized = true; + std::cout << "OpenCL initialized for the first time." << std::endl; + cout << "Version : " << m_platformInfo.Version() << std::endl; + cout << "Name : " << m_platformInfo.Name()<< std::endl; + cout << "Device : " << m_deviceInfo.Name()<< std::endl; + + if (m_device_id) + { + clReleaseDevice(m_device_id); + m_device_id = 0; + } + + //printInfo(); return m_context != 0 ? CL_SUCCESS : -1; } // initOpenCL() +int OpenClWrapper::createMemObject(cl_mem* mem_obj, cv::Mat& inputData){ + cl_int res = CL_SUCCESS; + cl_mem mem = mem_obj[0]; + + if (inputData.ptr() == nullptr) + { + std::cout << "Error:createMemObject nupptr as input ptr" << std::endl; + return -1; + } + + cl_image_format fmt; + fmt.image_channel_order = CL_R; + fmt.image_channel_data_type = CL_UNSIGNED_INT8; + + cl_mem_flags flags_dst = CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR; + + cl_image_desc desc_dst; + desc_dst.image_type = CL_MEM_OBJECT_IMAGE2D; + desc_dst.image_width = inputData.cols; + desc_dst.image_height = inputData.rows; + desc_dst.image_depth = 0; + desc_dst.image_array_size = 0; + desc_dst.image_row_pitch = inputData.step[0]; + desc_dst.image_slice_pitch = 0; + desc_dst.num_mip_levels = 0; + desc_dst.num_samples = 0; + desc_dst.buffer = 0; + mem = clCreateImage(m_context, flags_dst, &fmt, &desc_dst, inputData.ptr(), &res); + if (0 == mem || CL_SUCCESS != res) + return -1; + + mem_obj[0] = mem; + + return 0; +} // this function is an example of "typical" OpenCL processing pipeline // It creates OpenCL buffer or image, depending on use_buffer flag, @@ -393,8 +434,8 @@ int OpenClWrapper::run() if (0 != initOpenCL()) return -1; - if (0 != initVideoSource()) - return -1; + //if (0 != initVideoSource()) + // return -1; Mat img_to_show; @@ -417,7 +458,7 @@ int OpenClWrapper::run() UMat uframe; // work - timerStart(); + //timerStart(); if (doProcess()) { @@ -434,22 +475,19 @@ int OpenClWrapper::run() m_frameGray.copyTo(uframe); } - timerEnd(); + //timerEnd(); uframe.copyTo(img_to_show); - putText(img_to_show, "Version : " + m_platformInfo.Version(), Point(5, 30), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); - putText(img_to_show, "Name : " + m_platformInfo.Name(), Point(5, 60), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); - putText(img_to_show, "Device : " + m_deviceInfo.Name(), Point(5, 90), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); - cv::String memtype = useBuffer() ? "buffer" : "image"; - putText(img_to_show, "interop with OpenCL " + memtype, Point(5, 120), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); - putText(img_to_show, "Time : " + timeStr() + " msec", Point(5, 150), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2); - imshow("opencl_interop", img_to_show); - handleKey((char)waitKey(3)); + //handleKey((char)waitKey(3)); } return 0; } +void OpenClWrapper::printInfo() +{ + // TODO +} \ No newline at end of file diff --git a/mediapipe/framework/formats/helpers.hpp b/mediapipe/framework/formats/helpers.hpp index 951300e113..4243a9aa44 100644 --- a/mediapipe/framework/formats/helpers.hpp +++ b/mediapipe/framework/formats/helpers.hpp @@ -409,6 +409,8 @@ class DeviceInfo } // namespace opencl +void dumpMatToFile(std::string& fileName, cv::UMat& umat); +void dumpMatToFile(std::string& fileName, cv::Mat& umat); class OpenClWrapper { @@ -416,9 +418,9 @@ class OpenClWrapper OpenClWrapper(); ~OpenClWrapper(); - int initOpenCL(); - int initVideoSource(); + static int initOpenCL(); + int createMemObject(cl_mem* mem_obj, cv::Mat& inputData); int process_frame_with_open_cl(cv::Mat& frame, bool use_buffer, cl_mem* cl_buffer); int process_cl_buffer_with_opencv(cl_mem buffer, size_t step, int rows, int cols, int type, cv::UMat& u); int process_cl_image_with_opencv(cl_mem image, cv::UMat& u); @@ -432,13 +434,11 @@ class OpenClWrapper void setRunning(bool running) { m_running = running; } void setDoProcess(bool process) { m_process = process; } void setUseBuffer(bool use_buffer) { m_use_buffer = use_buffer; } + void printInfo(); + cl_mem m_mem_obj; protected: bool nextFrame(cv::Mat& frame) { return m_cap.read(frame); } - void handleKey(char key); - void timerStart(); - void timerEnd(); - std::string timeStr() const; std::string message() const; private: @@ -456,16 +456,13 @@ class OpenClWrapper cv::VideoCapture m_cap; cv::Mat m_frame; cv::Mat m_frameGray; + + static bool m_is_initialized; + static cl_context m_context; + static cl_command_queue m_queue; - opencl::PlatformInfo m_platformInfo; - opencl::DeviceInfo m_deviceInfo; - std::vector m_platform_ids; - cl_context m_context; - cl_device_id m_device_id; - cl_command_queue m_queue; cl_program m_program; cl_kernel m_kernelBuf; cl_kernel m_kernelImg; cl_mem m_img_src; // used as src in case processing of cl image - cl_mem m_mem_obj; }; \ No newline at end of file diff --git a/mediapipe/framework/formats/image_frame.cc b/mediapipe/framework/formats/image_frame.cc index a3a982c4ea..0f75a5ca00 100644 --- a/mediapipe/framework/formats/image_frame.cc +++ b/mediapipe/framework/formats/image_frame.cc @@ -30,7 +30,6 @@ #include "mediapipe/framework/port/aligned_malloc_and_free.h" #include "mediapipe/framework/port/logging.h" #include "mediapipe/framework/port/proto_ns.h" -#include "mediapipe/framework/formats/helpers.hpp" namespace mediapipe { @@ -49,9 +48,9 @@ int CountOnes(uint32_t n) { } // namespace -const char* DEVICE_GPU = ":GPU:0"; +//const char* DEVICE_GPU = ":GPU:0"; -static cv::ocl::Context context = cv::ocl::Context::create(DEVICE_GPU); +//static cv::ocl::Context context = cv::ocl::Context::create(DEVICE_GPU); const ImageFrame::Deleter ImageFrame::PixelDataDeleter::kArrayDelete = std::default_delete(); @@ -65,22 +64,40 @@ const uint32_t ImageFrame::kDefaultAlignmentBoundary; const uint32_t ImageFrame::kGlDefaultAlignmentBoundary; ImageFrame::ImageFrame() - : format_(ImageFormat::UNKNOWN), width_(0), height_(0), width_step_(0) {} + : format_(ImageFormat::UNKNOWN), width_(0), height_(0), width_step_(0) { + CHECK_NE(ocl.initOpenCL(), -1); + } + +ImageFrame::ImageFrame(cv::Mat& inputData, ImageFormat::Format format, int width, int height, + uint32_t alignment_boundary) + : format_(format), width_(width), height_(height) { + CHECK_NE(ocl.initOpenCL(), -1); + Reset(inputData, format, width, height, alignment_boundary); +} ImageFrame::ImageFrame(ImageFormat::Format format, int width, int height, uint32_t alignment_boundary) : format_(format), width_(width), height_(height) { + CHECK_NE(ocl.initOpenCL(), -1); Reset(format, width, height, alignment_boundary); } ImageFrame::ImageFrame(ImageFormat::Format format, int width, int height) : format_(format), width_(width), height_(height) { + CHECK_NE(ocl.initOpenCL(), -1); Reset(format, width, height, kDefaultAlignmentBoundary); } +ImageFrame::ImageFrame(cv::Mat& inputData, ImageFormat::Format format, int width, int height) + : format_(format), width_(width), height_(height) { + CHECK_NE(ocl.initOpenCL(), -1); + Reset(inputData, format, width, height, kDefaultAlignmentBoundary); +} + ImageFrame::ImageFrame(ImageFormat::Format format, int width, int height, int width_step, uint8_t* pixel_data, ImageFrame::Deleter deleter) { + CHECK_NE(ocl.initOpenCL(), -1); AdoptPixelData(format, width, height, width_step, pixel_data, deleter); } @@ -92,6 +109,7 @@ ImageFrame& ImageFrame::operator=(ImageFrame&& move_from) { width_ = move_from.width_; height_ = move_from.height_; width_step_ = move_from.width_step_; + ocl = move_from.ocl; move_from.format_ = ImageFormat::UNKNOWN; move_from.width_ = 0; @@ -100,8 +118,7 @@ ImageFrame& ImageFrame::operator=(ImageFrame&& move_from) { return *this; } - -void ImageFrame::Reset(ImageFormat::Format format, int width, int height, +void ImageFrame::Reset(cv::Mat& inputData, ImageFormat::Format format, int width, int height, uint32_t alignment_boundary) { format_ = format; width_ = width; @@ -122,6 +139,36 @@ void ImageFrame::Reset(ImageFormat::Format format, int width, int height, height * width_step_, alignment_boundary)), PixelDataDeleter::kAlignedFree}; }*/ + + width_step_ = ((width_step_ - 1) | (alignment_boundary - 1)) + 1; + CHECK_NE(ocl.createMemObject(&ocl.m_mem_obj, inputData), -1); + + pixel_data_ = { reinterpret_cast(ocl.m_mem_obj) , PixelDataDeleter::kNone}; + } + +void ImageFrame::Reset(ImageFormat::Format format, int width, int height, + uint32_t alignment_boundary) { + format_ = format; + width_ = width; + height_ = height; + // TODO graph execution check + CHECK_NE(0,0); + CHECK_NE(ImageFormat::UNKNOWN, format_); + CHECK(IsValidAlignmentNumber(alignment_boundary)); + width_step_ = width * NumberOfChannels() * ByteDepth(); + if (alignment_boundary == 1) { + pixel_data_ = {new uint8_t[height * width_step_], + PixelDataDeleter::kArrayDelete}; + } else { + // Increase width_step_ to the smallest multiple of alignment_boundary + // which is large enough to hold all the data. This is done by + // twiddling bits. alignment_boundary - 1 is a mask which sets all + // the low order bits. + width_step_ = ((width_step_ - 1) | (alignment_boundary - 1)) + 1; + pixel_data_ = {reinterpret_cast(aligned_malloc( + height * width_step_, alignment_boundary)), + PixelDataDeleter::kAlignedFree}; + } } void ImageFrame::AdoptPixelData(ImageFormat::Format format, int width, diff --git a/mediapipe/framework/formats/image_frame.h b/mediapipe/framework/formats/image_frame.h index 6fcefbd38f..acc966891e 100644 --- a/mediapipe/framework/formats/image_frame.h +++ b/mediapipe/framework/formats/image_frame.h @@ -43,6 +43,7 @@ #include "mediapipe/framework/port.h" #include "mediapipe/framework/port/integral_types.h" #include "mediapipe/framework/tool/type_util.h" +#include "mediapipe/framework/formats/helpers.hpp" #define IMAGE_FRAME_RAW_IMAGE MEDIAPIPE_HAS_RTTI @@ -61,6 +62,7 @@ namespace mediapipe { // // Do not assume that the pixel data is stored contiguously. It may be // stored with row padding for alignment purposes. + class ImageFrame { public: typedef std::function Deleter; @@ -105,8 +107,11 @@ class ImageFrame { // be stored contiguously). ImageFrame(ImageFormat::Format format, int width, int height, uint32 alignment_boundary); + ImageFrame(cv::Mat& inputData, ImageFormat::Format format, int width, int height, + uint32_t alignment_boundary); // Same as above, but use kDefaultAlignmentBoundary for alignment_boundary. ImageFrame(ImageFormat::Format format, int width, int height); + ImageFrame(cv::Mat& inputData, ImageFormat::Format format, int width, int height); // Acquires ownership of pixel_data. Sets the deletion method // to use on pixel_data with deletion_method (which defaults @@ -209,6 +214,9 @@ class ImageFrame { void Reset(ImageFormat::Format format, int width, int height, uint32 alignment_boundary); + void Reset(cv::Mat& inputData, ImageFormat::Format format, int width, int height, + uint32_t alignment_boundary); + // Relinquishes ownership of the pixel data. Notice that the unique_ptr // uses a non-standard deleter. std::unique_ptr Release(); @@ -250,6 +258,8 @@ class ImageFrame { int width_step_; std::unique_ptr pixel_data_; + + OpenClWrapper ocl; }; } // namespace mediapipe diff --git a/mediapipe/framework/formats/image_frame_opencv.cc b/mediapipe/framework/formats/image_frame_opencv.cc index 3831e9742e..08d7990506 100644 --- a/mediapipe/framework/formats/image_frame_opencv.cc +++ b/mediapipe/framework/formats/image_frame_opencv.cc @@ -88,6 +88,7 @@ cv::Mat MatView(const ImageFrame* image) { CV_MAKETYPE(GetMatType(image->Format()), image->NumberOfChannels()); const size_t steps[] = {static_cast(image->WidthStep()), static_cast(image->ByteDepth())}; + // Use ImageFrame to initialize in-place. ImageFrame still owns memory. return cv::Mat(dims, sizes, type, const_cast(image->PixelData()), steps); @@ -96,10 +97,17 @@ cv::Mat MatView(const ImageFrame* image) { // UMatUsageFlags // USAGE_DEFAULT, USAGE_ALLOCATE_HOST_MEMORY, USAGE_ALLOCATE_DEVICE_MEMORY, USAGE_ALLOCATE_SHARED_MEMORY , __UMAT_USAGE_FLAGS_32BIT cv::UMat MatView(ImageFrame* image, cv::UMatUsageFlags usageFlags) { - cv::Mat mat = MatView(image); + //cv::Mat mat = MatView(image); cv::UMat umat = cv::UMat(usageFlags); - mat.copyTo(umat); + // mat.copyTo(umat); //*image->MutablePixelData() = static_cast(umat.u->data); + //const int type = CV_MAKETYPE(GetMatType(image->Format()), image->NumberOfChannels()); + + //TODO - works only for buffers:CL_MEM_OBJECT_BUFFER == mem_type + //cv::ocl::convertFromBuffer(reinterpret_cast(image->MutablePixelData()), static_cast(image->WidthStep()), image->Height(), image->Width(), type, umat); + + // CL_MEM_OBJECT_IMAGE2D == mem_type + cv::ocl::convertFromImage(reinterpret_cast(image->MutablePixelData()), umat); return umat; } From 512f81ac19f1c179446eb41f9840cf0e4fc2486f Mon Sep 17 00:00:00 2001 From: "iotg.worker" Date: Wed, 3 Apr 2024 10:28:08 +0200 Subject: [PATCH 4/7] OCL helpers --- mediapipe/framework/formats/helpers.cpp | 195 +++++++++++++++++++++++- setup_opencv.sh | 2 +- 2 files changed, 191 insertions(+), 6 deletions(-) diff --git a/mediapipe/framework/formats/helpers.cpp b/mediapipe/framework/formats/helpers.cpp index 32d14cf45f..ccbdde1b0d 100644 --- a/mediapipe/framework/formats/helpers.cpp +++ b/mediapipe/framework/formats/helpers.cpp @@ -59,6 +59,150 @@ void dumpMatToFile(std::string& fileName, cv::Mat& umat){ file.release(); } +cl_channel_order GetChannelOrderFromMatType(int cvMatType) { + switch (cvMatType) { + case CV_8UC1: + case CV_8SC1: + std::cout<< "Using CL_R " << endl; + return CL_R; + case CV_8UC2: + case CV_8SC2: + std::cout<< "Using CL_RG " << endl; + return CL_RG; + case CV_8UC3: + case CV_8SC3: + std::cout<< "Using CL_RGB " << endl; + return CL_RGB; + case CV_8UC4: + case CV_8SC4: + std::cout<< "Using CL_RGBA " << endl; + return CL_RGBA; + default: + { + std::cout<< "Using default CL_R for CV type: " << cvMatType << endl; + return CL_R; // Default to single channel + } + } +} + + +cl_channel_type GetChannelDataTypeFromOrder(cl_channel_order cl_order) { + switch (cl_order) { + case CL_R: + case CL_A: + return CL_UNSIGNED_INT8; // Or your desired data type + std::cout<< "Using CL_UNSIGNED_INT8 " << endl; + case CL_RG: + return CL_UNSIGNED_INT16; // Or your desired data type + case CL_RGB: + std::cout<< "Using CL_HALF_FLOAT " << endl; + return CL_HALF_FLOAT; // Or your desired data type + case CL_RGBA: + std::cout<< "Using CL_FLOAT " << endl; + return CL_FLOAT; // Or your desired data type + // Add more cases for other channel orders as needed + default: + // Handle unsupported channel order + std::cout<< "Using default CL_UNSIGNED_INT8 for channel order: " << cl_order << endl; + break; + } +} + +const char* clGetErrorString(int errorCode) { + switch (errorCode) { + case 0: return "CL_SUCCESS"; + case -1: return "CL_DEVICE_NOT_FOUND"; + case -2: return "CL_DEVICE_NOT_AVAILABLE"; + case -3: return "CL_COMPILER_NOT_AVAILABLE"; + case -4: return "CL_MEM_OBJECT_ALLOCATION_FAILURE"; + case -5: return "CL_OUT_OF_RESOURCES"; + case -6: return "CL_OUT_OF_HOST_MEMORY"; + case -7: return "CL_PROFILING_INFO_NOT_AVAILABLE"; + case -8: return "CL_MEM_COPY_OVERLAP"; + case -9: return "CL_IMAGE_FORMAT_MISMATCH"; + case -10: return "CL_IMAGE_FORMAT_NOT_SUPPORTED"; + case -12: return "CL_MAP_FAILURE"; + case -13: return "CL_MISALIGNED_SUB_BUFFER_OFFSET"; + case -14: return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"; + case -15: return "CL_COMPILE_PROGRAM_FAILURE"; + case -16: return "CL_LINKER_NOT_AVAILABLE"; + case -17: return "CL_LINK_PROGRAM_FAILURE"; + case -18: return "CL_DEVICE_PARTITION_FAILED"; + case -19: return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE"; + case -30: return "CL_INVALID_VALUE"; + case -31: return "CL_INVALID_DEVICE_TYPE"; + case -32: return "CL_INVALID_PLATFORM"; + case -33: return "CL_INVALID_DEVICE"; + case -34: return "CL_INVALID_CONTEXT"; + case -35: return "CL_INVALID_QUEUE_PROPERTIES"; + case -36: return "CL_INVALID_COMMAND_QUEUE"; + case -37: return "CL_INVALID_HOST_PTR"; + case -38: return "CL_INVALID_MEM_OBJECT"; + case -39: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"; + case -40: return "CL_INVALID_IMAGE_SIZE"; + case -41: return "CL_INVALID_SAMPLER"; + case -42: return "CL_INVALID_BINARY"; + case -43: return "CL_INVALID_BUILD_OPTIONS"; + case -44: return "CL_INVALID_PROGRAM"; + case -45: return "CL_INVALID_PROGRAM_EXECUTABLE"; + case -46: return "CL_INVALID_KERNEL_NAME"; + case -47: return "CL_INVALID_KERNEL_DEFINITION"; + case -48: return "CL_INVALID_KERNEL"; + case -49: return "CL_INVALID_ARG_INDEX"; + case -50: return "CL_INVALID_ARG_VALUE"; + case -51: return "CL_INVALID_ARG_SIZE"; + case -52: return "CL_INVALID_KERNEL_ARGS"; + case -53: return "CL_INVALID_WORK_DIMENSION"; + case -54: return "CL_INVALID_WORK_GROUP_SIZE"; + case -55: return "CL_INVALID_WORK_ITEM_SIZE"; + case -56: return "CL_INVALID_GLOBAL_OFFSET"; + case -57: return "CL_INVALID_EVENT_WAIT_LIST"; + case -58: return "CL_INVALID_EVENT"; + case -59: return "CL_INVALID_OPERATION"; + case -60: return "CL_INVALID_GL_OBJECT"; + case -61: return "CL_INVALID_BUFFER_SIZE"; + case -62: return "CL_INVALID_MIP_LEVEL"; + case -63: return "CL_INVALID_GLOBAL_WORK_SIZE"; + case -64: return "CL_INVALID_PROPERTY"; + case -65: return "CL_INVALID_IMAGE_DESCRIPTOR"; + case -66: return "CL_INVALID_COMPILER_OPTIONS"; + case -67: return "CL_INVALID_LINKER_OPTIONS"; + case -68: return "CL_INVALID_DEVICE_PARTITION_COUNT"; + case -69: return "CL_INVALID_PIPE_SIZE"; + case -70: return "CL_INVALID_DEVICE_QUEUE"; + case -71: return "CL_INVALID_SPEC_ID"; + case -72: return "CL_MAX_SIZE_RESTRICTION_EXCEEDED"; + case -1002: return "CL_INVALID_D3D10_DEVICE_KHR"; + case -1003: return "CL_INVALID_D3D10_RESOURCE_KHR"; + case -1004: return "CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR"; + case -1005: return "CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR"; + case -1006: return "CL_INVALID_D3D11_DEVICE_KHR"; + case -1007: return "CL_INVALID_D3D11_RESOURCE_KHR"; + case -1008: return "CL_D3D11_RESOURCE_ALREADY_ACQUIRED_KHR"; + case -1009: return "CL_D3D11_RESOURCE_NOT_ACQUIRED_KHR"; + case -1010: return "CL_INVALID_DX9_MEDIA_ADAPTER_KHR"; + case -1011: return "CL_INVALID_DX9_MEDIA_SURFACE_KHR"; + case -1012: return "CL_DX9_MEDIA_SURFACE_ALREADY_ACQUIRED_KHR"; + case -1013: return "CL_DX9_MEDIA_SURFACE_NOT_ACQUIRED_KHR"; + case -1093: return "CL_INVALID_EGL_OBJECT_KHR"; + case -1092: return "CL_EGL_RESOURCE_NOT_ACQUIRED_KHR"; + case -1001: return "CL_PLATFORM_NOT_FOUND_KHR"; + case -1057: return "CL_DEVICE_PARTITION_FAILED_EXT"; + case -1058: return "CL_INVALID_PARTITION_COUNT_EXT"; + case -1059: return "CL_INVALID_PARTITION_NAME_EXT"; + case -1094: return "CL_INVALID_ACCELERATOR_INTEL"; + case -1095: return "CL_INVALID_ACCELERATOR_TYPE_INTEL"; + case -1096: return "CL_INVALID_ACCELERATOR_DESCRIPTOR_INTEL"; + case -1097: return "CL_ACCELERATOR_TYPE_NOT_SUPPORTED_INTEL"; + case -1000: return "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR"; + case -1098: return "CL_INVALID_VA_API_MEDIA_ADAPTER_INTEL"; + case -1099: return "CL_INVALID_VA_API_MEDIA_SURFACE_INTEL"; + case -1100: return "CL_VA_API_MEDIA_SURFACE_ALREADY_ACQUIRED_INTEL"; + case -1101: return "CL_VA_API_MEDIA_SURFACE_NOT_ACQUIRED_INTEL"; + default: return "CL_UNKNOWN_ERROR"; + } +} + OpenClWrapper::OpenClWrapper() { m_camera_id = 0; //cmd.get("camera"); @@ -137,12 +281,10 @@ OpenClWrapper::~OpenClWrapper() } } // dtor - bool OpenClWrapper::m_is_initialized = false; cl_context OpenClWrapper::m_context = nullptr; cl_command_queue OpenClWrapper::m_queue = nullptr; - int OpenClWrapper::initOpenCL() { if (m_is_initialized){ @@ -181,6 +323,18 @@ int OpenClWrapper::initOpenCL() 0 }; + cl_device_id device; + // Get the first available device on the platform + res = clGetDeviceIDs(m_platform_ids[i], CL_DEVICE_TYPE_GPU, 1, &device, NULL); + if (res != CL_SUCCESS) { + printf("Error getting device IDs\n"); + return 1; + } + + // Query the device for supported image formats + cl_image_format supportedFormats[128]; // Assuming maximum 128 supported formats + cl_uint numSupportedFormats; + m_context = clCreateContextFromType(props, CL_DEVICE_TYPE_GPU, 0, 0, &res); if (0 == m_context || CL_SUCCESS != res) continue; @@ -189,6 +343,28 @@ int OpenClWrapper::initOpenCL() if (CL_SUCCESS != res) return -1; + res = clGetSupportedImageFormats(m_context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, 128, supportedFormats, &numSupportedFormats); + if (res != CL_SUCCESS) { + printf("Error getting supported image formats\n"); + return 1; + } + + // Check if CL_HALF_FLOAT is supported + cl_bool supportsHalfFloat = CL_FALSE; + for (cl_uint i = 0; i < numSupportedFormats; ++i) { + if (supportedFormats[i].image_channel_data_type == CL_HALF_FLOAT) { + supportsHalfFloat = CL_TRUE; + break; + } + } + + // Print the result + if (supportsHalfFloat == CL_TRUE) { + printf("Device supports CL_HALF_FLOAT format\n"); + } else { + printf("Device does not support CL_HALF_FLOAT format\n"); + } + m_queue = clCreateCommandQueue(m_context, m_device_id, 0, &res); if (0 == m_queue || CL_SUCCESS != res) return -1; @@ -229,25 +405,34 @@ int OpenClWrapper::createMemObject(cl_mem* mem_obj, cv::Mat& inputData){ } cl_image_format fmt; - fmt.image_channel_order = CL_R; - fmt.image_channel_data_type = CL_UNSIGNED_INT8; + cl_channel_order channelOrder = GetChannelOrderFromMatType(inputData.type()); + fmt.image_channel_order = channelOrder; + cl_channel_type channelType = GetChannelDataTypeFromOrder(channelOrder); + fmt.image_channel_data_type = channelType; cl_mem_flags flags_dst = CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR; + // TODO: fix CL_INVALID_IMAGE_FORMAT_DESCRIPTOR now - Hardcoded values + fmt.image_channel_order = CL_R; + fmt.image_channel_data_type = CL_UNSIGNED_INT8; cl_image_desc desc_dst; + memset(&desc_dst, 0, sizeof(cl_image_desc)); desc_dst.image_type = CL_MEM_OBJECT_IMAGE2D; desc_dst.image_width = inputData.cols; desc_dst.image_height = inputData.rows; desc_dst.image_depth = 0; desc_dst.image_array_size = 0; + // desc_dst.image_row_pitch = inputData.step[0]; desc_dst.image_row_pitch = inputData.step[0]; desc_dst.image_slice_pitch = 0; desc_dst.num_mip_levels = 0; desc_dst.num_samples = 0; desc_dst.buffer = 0; mem = clCreateImage(m_context, flags_dst, &fmt, &desc_dst, inputData.ptr(), &res); - if (0 == mem || CL_SUCCESS != res) + if (0 == mem || CL_SUCCESS != res){ + std::cout <<"Error: " << clGetErrorString(res) << std::endl; return -1; + } mem_obj[0] = mem; diff --git a/setup_opencv.sh b/setup_opencv.sh index e83a2295ba..3c26a44ad9 100755 --- a/setup_opencv.sh +++ b/setup_opencv.sh @@ -60,7 +60,7 @@ if [ -z "$1" ] cd ../opencv git checkout 4.7.0 cd release - cmake .. -DCMAKE_BUILD_TYPE=RELEASE -DCMAKE_INSTALL_PREFIX=/usr/local \ + cmake .. -DCMAKE_BUILD_TYPE=DEBUG -DCMAKE_INSTALL_PREFIX=/usr/local \ -DBUILD_LIST=core,improc,imgcodecs,calib3d,features2d,highgui,imgproc,video,videoio,optflow \ -DBUILD_TESTS=OFF -DBUILD_PERF_TESTS=OFF -DBUILD_opencv_ts=OFF \ -DOPENCV_EXTRA_MODULES_PATH=/tmp/build_opencv/opencv_contrib/modules \ From ee0e075da6892b25158800f959e236a4d79c9a7a Mon Sep 17 00:00:00 2001 From: "iotg.worker" Date: Tue, 9 Apr 2024 16:42:09 +0200 Subject: [PATCH 5/7] Working tranformation --- .../image/image_transformation_calculator.cc | 14 ++++----- .../util/annotation_overlay_calculator.cc | 6 +++- .../examples/desktop/demo_run_graph_main.cc | 22 +++++++++----- mediapipe/framework/formats/helpers.cpp | 12 ++++++-- mediapipe/framework/formats/helpers.hpp | 1 + mediapipe/framework/formats/image_frame.cc | 24 +++++++-------- mediapipe/framework/formats/image_frame.h | 8 ++--- .../framework/formats/image_frame_opencv.cc | 6 ++-- .../object_detection/transformation.pbtxt | 29 +++++++++++++++++++ 9 files changed, 85 insertions(+), 37 deletions(-) create mode 100644 mediapipe/graphs/object_detection/transformation.pbtxt diff --git a/mediapipe/calculators/image/image_transformation_calculator.cc b/mediapipe/calculators/image/image_transformation_calculator.cc index 51b97c57d9..b81d4a48c6 100644 --- a/mediapipe/calculators/image/image_transformation_calculator.cc +++ b/mediapipe/calculators/image/image_transformation_calculator.cc @@ -446,13 +446,13 @@ absl::Status ImageTransformationCalculator::Close(CalculatorContext* cc) { absl::Status ImageTransformationCalculator::RenderOpenCl(CalculatorContext* cc) { // TODO UMAT - cv::Mat input_mat; + cv::UMat input_mat; mediapipe::ImageFormat::Format format; const auto& input = cc->Inputs().Tag(kImageFrameTag).Get(); - std::cout << " cc->Inputs().Tag(kImageFrameTag).Get(); " << std::endl; - // input_mat = formats::MatView(const_cast(&input), cv::USAGE_ALLOCATE_SHARED_MEMORY); - input_mat = formats::MatView(&input); + //std::cout << " cc->Inputs().Tag(kImageFrameTag).Get(); " << std::endl; + input_mat = formats::MatView(const_cast(&input), cv::USAGE_ALLOCATE_SHARED_MEMORY); + //input_mat = formats::MatView(&input, cv::USAGE_ALLOCATE_SHARED_MEMORY); std::string fileName = std::string("./imageTrans/input") + std::to_string(cc->InputTimestamp().Value()); //dumpMatToFile(fileName, input_mat); @@ -467,7 +467,7 @@ absl::Status ImageTransformationCalculator::RenderOpenCl(CalculatorContext* cc) if (output_width_ > 0 && output_height_ > 0) { // TODO umat - cv::Mat scaled_mat; + cv::UMat scaled_mat; if (scale_mode_ == mediapipe::ScaleMode_Mode_STRETCH) { int scale_flag = input_mat.cols > output_width_ && input_mat.rows > output_height_ @@ -515,7 +515,7 @@ absl::Status ImageTransformationCalculator::RenderOpenCl(CalculatorContext* cc) } //TODO umat - cv::Mat rotated_mat; + cv::UMat rotated_mat; cv::Size rotated_size(output_width, output_height); if (input_mat.size() == rotated_size) { const int angle = RotationModeToDegrees(rotation_); @@ -542,7 +542,7 @@ absl::Status ImageTransformationCalculator::RenderOpenCl(CalculatorContext* cc) } // TODO UMAT - cv::Mat flipped_mat; + cv::UMat flipped_mat; if (flip_horizontally_ || flip_vertically_) { const int flip_code = flip_horizontally_ && flip_vertically_ ? -1 : flip_horizontally_; diff --git a/mediapipe/calculators/util/annotation_overlay_calculator.cc b/mediapipe/calculators/util/annotation_overlay_calculator.cc index 4033823a0e..5f6571d02a 100644 --- a/mediapipe/calculators/util/annotation_overlay_calculator.cc +++ b/mediapipe/calculators/util/annotation_overlay_calculator.cc @@ -373,8 +373,12 @@ absl::Status AnnotationOverlayCalculator::Close(CalculatorContext* cc) { absl::Status AnnotationOverlayCalculator::RenderToOpencl( CalculatorContext* cc, const ImageFormat::Format& target_format, cv::Mat& input_mat) { + // TODO switch to umat + cv::UMat copy; + input_mat.copyTo(copy); + std::unique_ptr output_frame( - new ImageFrame(input_mat, target_format, renderer_->GetImageWidth(), renderer_->GetImageHeight())); + new ImageFrame(copy, target_format, renderer_->GetImageWidth(), renderer_->GetImageHeight())); if (cc->Outputs().HasTag(kImageFrameTag)) { cc->Outputs() diff --git a/mediapipe/examples/desktop/demo_run_graph_main.cc b/mediapipe/examples/desktop/demo_run_graph_main.cc index 1cbacf638e..30a2ec99b7 100644 --- a/mediapipe/examples/desktop/demo_run_graph_main.cc +++ b/mediapipe/examples/desktop/demo_run_graph_main.cc @@ -90,10 +90,11 @@ absl::Status RunMPPGraph() { int count_frames = 0; auto begin = std::chrono::high_resolution_clock::now(); - //OpenClWrapper ocl; - //ocl.initOpenCL(); + OpenClWrapper ocl; + ocl.initOpenCL(); //ocl.printInfo(); + /* LOG(INFO) << "haveOpenCL " << cv::ocl::haveOpenCL() <(); // Convert back to opencv for display or saving. - cv::Mat output_frame_mat = mediapipe::formats::MatView(&output_frame); - cv::cvtColor(output_frame_mat, output_frame_mat, cv::COLOR_RGB2BGR); + cv::UMat output_frame_mat = mediapipe::formats::MatView(const_cast(&output_frame), cv::USAGE_ALLOCATE_SHARED_MEMORY); + // SEGFAULT icv_k0_ownsCopy_8u_repE9 cv::cvtColor(output_frame_mat, output_frame_mat, cv::COLOR_RGB2BGR); if (save_video) { if (!writer.isOpened()) { LOG(INFO) << "Prepare video writer."; diff --git a/mediapipe/framework/formats/helpers.cpp b/mediapipe/framework/formats/helpers.cpp index ccbdde1b0d..f6acecd2f1 100644 --- a/mediapipe/framework/formats/helpers.cpp +++ b/mediapipe/framework/formats/helpers.cpp @@ -104,6 +104,7 @@ cl_channel_type GetChannelDataTypeFromOrder(cl_channel_order cl_order) { default: // Handle unsupported channel order std::cout<< "Using default CL_UNSIGNED_INT8 for channel order: " << cl_order << endl; + return CL_UNSIGNED_INT8; break; } } @@ -394,7 +395,14 @@ int OpenClWrapper::initOpenCL() return m_context != 0 ? CL_SUCCESS : -1; } // initOpenCL() +int OpenClWrapper::createMemObject(cl_mem* mem_obj, cv::UMat& inputData){ + //cv::ocl::Image2D image = cv::ocl::Image2D(inputData); + //mem_obj[0] = image.handle; + + return 0; +} int OpenClWrapper::createMemObject(cl_mem* mem_obj, cv::Mat& inputData){ + // OLD IMPLEMENTATION cl_int res = CL_SUCCESS; cl_mem mem = mem_obj[0]; @@ -412,8 +420,8 @@ int OpenClWrapper::createMemObject(cl_mem* mem_obj, cv::Mat& inputData){ cl_mem_flags flags_dst = CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR; // TODO: fix CL_INVALID_IMAGE_FORMAT_DESCRIPTOR now - Hardcoded values - fmt.image_channel_order = CL_R; - fmt.image_channel_data_type = CL_UNSIGNED_INT8; + // fmt.image_channel_order = CL_R; + // fmt.image_channel_data_type = CL_UNSIGNED_INT8; cl_image_desc desc_dst; memset(&desc_dst, 0, sizeof(cl_image_desc)); diff --git a/mediapipe/framework/formats/helpers.hpp b/mediapipe/framework/formats/helpers.hpp index 4243a9aa44..3ab16cfd43 100644 --- a/mediapipe/framework/formats/helpers.hpp +++ b/mediapipe/framework/formats/helpers.hpp @@ -420,6 +420,7 @@ class OpenClWrapper static int initOpenCL(); + int createMemObject(cl_mem* mem_obj, cv::UMat& inputData); int createMemObject(cl_mem* mem_obj, cv::Mat& inputData); int process_frame_with_open_cl(cv::Mat& frame, bool use_buffer, cl_mem* cl_buffer); int process_cl_buffer_with_opencv(cl_mem buffer, size_t step, int rows, int cols, int type, cv::UMat& u); diff --git a/mediapipe/framework/formats/image_frame.cc b/mediapipe/framework/formats/image_frame.cc index 0f75a5ca00..bd6ee8f67a 100644 --- a/mediapipe/framework/formats/image_frame.cc +++ b/mediapipe/framework/formats/image_frame.cc @@ -65,39 +65,39 @@ const uint32_t ImageFrame::kGlDefaultAlignmentBoundary; ImageFrame::ImageFrame() : format_(ImageFormat::UNKNOWN), width_(0), height_(0), width_step_(0) { - CHECK_NE(ocl.initOpenCL(), -1); + //CHECK_NE(ocl.initOpenCL(), -1); } -ImageFrame::ImageFrame(cv::Mat& inputData, ImageFormat::Format format, int width, int height, +ImageFrame::ImageFrame(cv::UMat& inputData, ImageFormat::Format format, int width, int height, uint32_t alignment_boundary) : format_(format), width_(width), height_(height) { - CHECK_NE(ocl.initOpenCL(), -1); + //CHECK_NE(ocl.initOpenCL(), -1); Reset(inputData, format, width, height, alignment_boundary); } ImageFrame::ImageFrame(ImageFormat::Format format, int width, int height, uint32_t alignment_boundary) : format_(format), width_(width), height_(height) { - CHECK_NE(ocl.initOpenCL(), -1); + //CHECK_NE(ocl.initOpenCL(), -1); Reset(format, width, height, alignment_boundary); } ImageFrame::ImageFrame(ImageFormat::Format format, int width, int height) : format_(format), width_(width), height_(height) { - CHECK_NE(ocl.initOpenCL(), -1); + //CHECK_NE(ocl.initOpenCL(), -1); Reset(format, width, height, kDefaultAlignmentBoundary); } -ImageFrame::ImageFrame(cv::Mat& inputData, ImageFormat::Format format, int width, int height) +ImageFrame::ImageFrame(cv::UMat& inputData, ImageFormat::Format format, int width, int height) : format_(format), width_(width), height_(height) { - CHECK_NE(ocl.initOpenCL(), -1); + //CHECK_NE(ocl.initOpenCL(), -1); Reset(inputData, format, width, height, kDefaultAlignmentBoundary); } ImageFrame::ImageFrame(ImageFormat::Format format, int width, int height, int width_step, uint8_t* pixel_data, ImageFrame::Deleter deleter) { - CHECK_NE(ocl.initOpenCL(), -1); + //CHECK_NE(ocl.initOpenCL(), -1); AdoptPixelData(format, width, height, width_step, pixel_data, deleter); } @@ -109,7 +109,7 @@ ImageFrame& ImageFrame::operator=(ImageFrame&& move_from) { width_ = move_from.width_; height_ = move_from.height_; width_step_ = move_from.width_step_; - ocl = move_from.ocl; + //ocl = move_from.ocl; move_from.format_ = ImageFormat::UNKNOWN; move_from.width_ = 0; @@ -118,7 +118,7 @@ ImageFrame& ImageFrame::operator=(ImageFrame&& move_from) { return *this; } -void ImageFrame::Reset(cv::Mat& inputData, ImageFormat::Format format, int width, int height, +void ImageFrame::Reset(cv::UMat& inputData, ImageFormat::Format format, int width, int height, uint32_t alignment_boundary) { format_ = format; width_ = width; @@ -141,9 +141,9 @@ void ImageFrame::Reset(cv::Mat& inputData, ImageFormat::Format format, int width }*/ width_step_ = ((width_step_ - 1) | (alignment_boundary - 1)) + 1; - CHECK_NE(ocl.createMemObject(&ocl.m_mem_obj, inputData), -1); + //CHECK_NE(ocl.createMemObject(&ocl.m_mem_obj, inputData), -1); - pixel_data_ = { reinterpret_cast(ocl.m_mem_obj) , PixelDataDeleter::kNone}; + pixel_data_ = { reinterpret_cast(inputData.handle(ACCESS_RW)) , PixelDataDeleter::kNone}; } void ImageFrame::Reset(ImageFormat::Format format, int width, int height, diff --git a/mediapipe/framework/formats/image_frame.h b/mediapipe/framework/formats/image_frame.h index acc966891e..ce7ea974f5 100644 --- a/mediapipe/framework/formats/image_frame.h +++ b/mediapipe/framework/formats/image_frame.h @@ -107,11 +107,11 @@ class ImageFrame { // be stored contiguously). ImageFrame(ImageFormat::Format format, int width, int height, uint32 alignment_boundary); - ImageFrame(cv::Mat& inputData, ImageFormat::Format format, int width, int height, + ImageFrame(cv::UMat& inputData, ImageFormat::Format format, int width, int height, uint32_t alignment_boundary); // Same as above, but use kDefaultAlignmentBoundary for alignment_boundary. ImageFrame(ImageFormat::Format format, int width, int height); - ImageFrame(cv::Mat& inputData, ImageFormat::Format format, int width, int height); + ImageFrame(cv::UMat& inputData, ImageFormat::Format format, int width, int height); // Acquires ownership of pixel_data. Sets the deletion method // to use on pixel_data with deletion_method (which defaults @@ -214,7 +214,7 @@ class ImageFrame { void Reset(ImageFormat::Format format, int width, int height, uint32 alignment_boundary); - void Reset(cv::Mat& inputData, ImageFormat::Format format, int width, int height, + void Reset(cv::UMat& inputData, ImageFormat::Format format, int width, int height, uint32_t alignment_boundary); // Relinquishes ownership of the pixel data. Notice that the unique_ptr @@ -259,7 +259,7 @@ class ImageFrame { std::unique_ptr pixel_data_; - OpenClWrapper ocl; + //OpenClWrapper ocl; }; } // namespace mediapipe diff --git a/mediapipe/framework/formats/image_frame_opencv.cc b/mediapipe/framework/formats/image_frame_opencv.cc index 08d7990506..d79c8372ba 100644 --- a/mediapipe/framework/formats/image_frame_opencv.cc +++ b/mediapipe/framework/formats/image_frame_opencv.cc @@ -101,13 +101,13 @@ cv::UMat MatView(ImageFrame* image, cv::UMatUsageFlags usageFlags) { cv::UMat umat = cv::UMat(usageFlags); // mat.copyTo(umat); //*image->MutablePixelData() = static_cast(umat.u->data); - //const int type = CV_MAKETYPE(GetMatType(image->Format()), image->NumberOfChannels()); + const int type = CV_MAKETYPE(GetMatType(image->Format()), image->NumberOfChannels()); //TODO - works only for buffers:CL_MEM_OBJECT_BUFFER == mem_type - //cv::ocl::convertFromBuffer(reinterpret_cast(image->MutablePixelData()), static_cast(image->WidthStep()), image->Height(), image->Width(), type, umat); + cv::ocl::convertFromBuffer(reinterpret_cast(image->MutablePixelData()), static_cast(image->WidthStep()), image->Height(), image->Width(), type, umat); // CL_MEM_OBJECT_IMAGE2D == mem_type - cv::ocl::convertFromImage(reinterpret_cast(image->MutablePixelData()), umat); + //cv::ocl::convertFromImage(reinterpret_cast(image->MutablePixelData()), umat); return umat; } diff --git a/mediapipe/graphs/object_detection/transformation.pbtxt b/mediapipe/graphs/object_detection/transformation.pbtxt new file mode 100644 index 0000000000..60dd9ae5c4 --- /dev/null +++ b/mediapipe/graphs/object_detection/transformation.pbtxt @@ -0,0 +1,29 @@ +# MediaPipe graph that performs object detection on desktop with TensorFlow Lite +# on CPU. +# Used in the example in +# mediapipe/examples/desktop/object_detection:object_detection_openvino. + +# max_queue_size limits the number of packets enqueued on any input stream +# by throttling inputs to the graph. This makes the graph only process one +# frame per time. +max_queue_size: 1 + +input_stream: "input_video" +output_stream: "output_video" + +# Transforms the input image on CPU to a 320x320 image. To scale the image, by +# default it uses the STRETCH scale mode that maps the entire input image to the +# entire transformed image. As a result, image aspect ratio may be changed and +# objects in the image may be deformed (stretched or squeezed), but the object +# detection model used in this graph is agnostic to that deformation. +node: { + calculator: "ImageTransformationCalculator" + input_stream: "IMAGE:input_video" + output_stream: "IMAGE:output_video" + node_options: { + [type.googleapis.com/mediapipe.ImageTransformationCalculatorOptions] { + output_width: 320 + output_height: 320 + } + } +} From bf0cc26cd92fbd8de2998e2cc232066899e52cd6 Mon Sep 17 00:00:00 2001 From: "iotg.worker" Date: Tue, 16 Apr 2024 12:43:23 +0200 Subject: [PATCH 6/7] Performance measurments --- .../image/image_transformation_calculator.cc | 23 ++++++++++---- .../examples/desktop/demo_run_graph_main.cc | 14 ++++++--- .../desktop/demo_run_graph_main_gpu.cc | 9 ++++++ .../examples/desktop/face_detection/BUILD | 1 + .../face_detection/transformation_cpu.pbtxt | 31 +++++++++++++++++++ .../face_detection/transformation_gpu.pbtxt | 31 +++++++++++++++++++ .../transformation_umat.pbtxt | 30 ++++++++++++++++++ 7 files changed, 129 insertions(+), 10 deletions(-) create mode 100644 mediapipe/graphs/face_detection/transformation_cpu.pbtxt create mode 100644 mediapipe/graphs/face_detection/transformation_gpu.pbtxt create mode 100644 mediapipe/graphs/object_detection/transformation_umat.pbtxt diff --git a/mediapipe/calculators/image/image_transformation_calculator.cc b/mediapipe/calculators/image/image_transformation_calculator.cc index b81d4a48c6..6c9be60978 100644 --- a/mediapipe/calculators/image/image_transformation_calculator.cc +++ b/mediapipe/calculators/image/image_transformation_calculator.cc @@ -446,12 +446,13 @@ absl::Status ImageTransformationCalculator::Close(CalculatorContext* cc) { absl::Status ImageTransformationCalculator::RenderOpenCl(CalculatorContext* cc) { // TODO UMAT - cv::UMat input_mat; + cv::UMatUsageFlags usageFlags = cv::USAGE_ALLOCATE_HOST_MEMORY; + cv::UMat input_mat = cv::UMat(usageFlags); mediapipe::ImageFormat::Format format; const auto& input = cc->Inputs().Tag(kImageFrameTag).Get(); //std::cout << " cc->Inputs().Tag(kImageFrameTag).Get(); " << std::endl; - input_mat = formats::MatView(const_cast(&input), cv::USAGE_ALLOCATE_SHARED_MEMORY); + input_mat = formats::MatView(const_cast(&input), usageFlags); //input_mat = formats::MatView(&input, cv::USAGE_ALLOCATE_SHARED_MEMORY); std::string fileName = std::string("./imageTrans/input") + std::to_string(cc->InputTimestamp().Value()); //dumpMatToFile(fileName, input_mat); @@ -467,7 +468,7 @@ absl::Status ImageTransformationCalculator::RenderOpenCl(CalculatorContext* cc) if (output_width_ > 0 && output_height_ > 0) { // TODO umat - cv::UMat scaled_mat; + cv::UMat scaled_mat = cv::UMat(usageFlags); if (scale_mode_ == mediapipe::ScaleMode_Mode_STRETCH) { int scale_flag = input_mat.cols > output_width_ && input_mat.rows > output_height_ @@ -475,7 +476,10 @@ absl::Status ImageTransformationCalculator::RenderOpenCl(CalculatorContext* cc) : cv::INTER_LINEAR; cv::resize(input_mat, scaled_mat, cv::Size(output_width_, output_height_), 0, 0, scale_flag); + // DEFAULT + // std::cout << "SCALE - STRETCH" <(output_width_) / input_width, static_cast(output_height_) / input_height); @@ -483,7 +487,7 @@ absl::Status ImageTransformationCalculator::RenderOpenCl(CalculatorContext* cc) const int target_height = std::round(input_height * scale); int scale_flag = scale < 1.0f ? cv::INTER_AREA : cv::INTER_LINEAR; if (scale_mode_ == mediapipe::ScaleMode_Mode_FIT) { - cv::UMat intermediate_mat; + cv::UMat intermediate_mat = cv::UMat(usageFlags); cv::resize(input_mat, intermediate_mat, cv::Size(target_width, target_height), 0, 0, scale_flag); const int top = (output_height_ - target_height) / 2; @@ -515,7 +519,7 @@ absl::Status ImageTransformationCalculator::RenderOpenCl(CalculatorContext* cc) } //TODO umat - cv::UMat rotated_mat; + cv::UMat rotated_mat = cv::UMat(usageFlags); cv::Size rotated_size(output_width, output_height); if (input_mat.size() == rotated_size) { const int angle = RotationModeToDegrees(rotation_); @@ -523,10 +527,13 @@ absl::Status ImageTransformationCalculator::RenderOpenCl(CalculatorContext* cc) // TODO UMAT cv::Mat rotation_mat = cv::getRotationMatrix2D(src_center, angle, 1.0); cv::warpAffine(input_mat, rotated_mat, rotation_mat, rotated_size); + // DEFAULT + // std::cout << "ROTATE - MATRIX" < output_frame( new ImageFrame(flipped_mat, format, output_width, output_height)); diff --git a/mediapipe/examples/desktop/demo_run_graph_main.cc b/mediapipe/examples/desktop/demo_run_graph_main.cc index 30a2ec99b7..8b796c5198 100644 --- a/mediapipe/examples/desktop/demo_run_graph_main.cc +++ b/mediapipe/examples/desktop/demo_run_graph_main.cc @@ -88,7 +88,7 @@ absl::Status RunMPPGraph() { LOG(INFO) << "Start grabbing and processing frames."; bool grab_frames = true; int count_frames = 0; - auto begin = std::chrono::high_resolution_clock::now(); + OpenClWrapper ocl; ocl.initOpenCL(); @@ -110,7 +110,9 @@ absl::Status RunMPPGraph() { LOG(INFO) <<"CL device doubleFPConfig() " << device.doubleFPConfig() <(); // Convert back to opencv for display or saving. - cv::UMat output_frame_mat = mediapipe::formats::MatView(const_cast(&output_frame), cv::USAGE_ALLOCATE_SHARED_MEMORY); + cv::UMat output_frame_mat = mediapipe::formats::MatView(const_cast(&output_frame), usageFlags); // SEGFAULT icv_k0_ownsCopy_8u_repE9 cv::cvtColor(output_frame_mat, output_frame_mat, cv::COLOR_RGB2BGR); if (save_video) { if (!writer.isOpened()) { @@ -176,6 +178,10 @@ absl::Status RunMPPGraph() { const int pressed_key = cv::waitKey(5); if (pressed_key >= 0 && pressed_key != 255) grab_frames = false; } + + if (count_frames == max_frame) { + break; + } } auto duration = std::chrono::duration_cast(std::chrono::high_resolution_clock::now() - begin); auto totalTime = duration.count(); diff --git a/mediapipe/examples/desktop/demo_run_graph_main_gpu.cc b/mediapipe/examples/desktop/demo_run_graph_main_gpu.cc index 8336e56705..ecfb9ba5fd 100644 --- a/mediapipe/examples/desktop/demo_run_graph_main_gpu.cc +++ b/mediapipe/examples/desktop/demo_run_graph_main_gpu.cc @@ -94,8 +94,11 @@ absl::Status RunMPPGraph() { LOG(INFO) << "Start grabbing and processing frames."; bool grab_frames = true; + auto begin = std::chrono::high_resolution_clock::now(); + int count_frames = 0; while (grab_frames) { // Capture opencv camera or video frame. + count_frames+=1; cv::Mat camera_frame_raw; capture >> camera_frame_raw; if (camera_frame_raw.empty()) { @@ -183,7 +186,13 @@ absl::Status RunMPPGraph() { if (pressed_key >= 0 && pressed_key != 255) grab_frames = false; } } + auto duration = std::chrono::duration_cast(std::chrono::high_resolution_clock::now() - begin); + auto totalTime = duration.count(); + float avgFps = (1000000 * (float)(count_frames) / (float)totalTime); + float avgLatencyms = 1000 / avgFps; + LOG(INFO) << "Frames:" << count_frames << ", Duration [ms]:" << totalTime / 1000 << ", FPS:" << avgFps << ", Avg latency [ms]:" << avgLatencyms; + LOG(INFO) << "Shutting down."; if (writer.isOpened()) writer.release(); MP_RETURN_IF_ERROR(graph.CloseInputStream(kInputStream)); diff --git a/mediapipe/examples/desktop/face_detection/BUILD b/mediapipe/examples/desktop/face_detection/BUILD index cf647e0de6..1b21398014 100644 --- a/mediapipe/examples/desktop/face_detection/BUILD +++ b/mediapipe/examples/desktop/face_detection/BUILD @@ -40,6 +40,7 @@ cc_binary( name = "face_detection_gpu", data = ["//mediapipe/modules/face_detection:face_detection_short_range.tflite"], deps = [ + "@ovms//src:ovms_lib", "//mediapipe/examples/desktop:demo_run_graph_main_gpu", "//mediapipe/graphs/face_detection:desktop_live_gpu_calculators", ], diff --git a/mediapipe/graphs/face_detection/transformation_cpu.pbtxt b/mediapipe/graphs/face_detection/transformation_cpu.pbtxt new file mode 100644 index 0000000000..66d50c7ddf --- /dev/null +++ b/mediapipe/graphs/face_detection/transformation_cpu.pbtxt @@ -0,0 +1,31 @@ +profiler_config { + trace_enabled: true + enable_profiler: true + trace_log_interval_count: 200 + trace_log_path: "/mediapipe/tracing/cpu/" +} +input_stream: "input_video" +output_stream: "output_video" +node: { + calculator: "ImageTransformationCalculator" + input_stream: "IMAGE:input_video" + output_stream: "IMAGE:tmp_video" + node_options: { + [type.googleapis.com/mediapipe.ImageTransformationCalculatorOptions] { + output_width: 320 + output_height: 320 + } + } +} +node: { + calculator: "ImageTransformationCalculator" + input_stream: "IMAGE:tmp_video" + output_stream: "IMAGE:output_video" + node_options: { + [type.googleapis.com/mediapipe.ImageTransformationCalculatorOptions] { + output_width: 640 + output_height: 640 + } + } +} + diff --git a/mediapipe/graphs/face_detection/transformation_gpu.pbtxt b/mediapipe/graphs/face_detection/transformation_gpu.pbtxt new file mode 100644 index 0000000000..c85487480c --- /dev/null +++ b/mediapipe/graphs/face_detection/transformation_gpu.pbtxt @@ -0,0 +1,31 @@ +profiler_config { + trace_enabled: true + enable_profiler: true + trace_log_interval_count: 200 + trace_log_path: "/mediapipe/tracing/gpu/" +} +input_stream: "input_video" +output_stream: "output_video" +node: { + calculator: "ImageTransformationCalculator" + input_stream: "IMAGE_GPU:input_video" + output_stream: "IMAGE_GPU:tmp_video" + node_options: { + [type.googleapis.com/mediapipe.ImageTransformationCalculatorOptions] { + output_width: 320 + output_height: 320 + } + } +} + +node: { + calculator: "ImageTransformationCalculator" + input_stream: "IMAGE_GPU:tmp_video" + output_stream: "IMAGE_GPU:output_video" + node_options: { + [type.googleapis.com/mediapipe.ImageTransformationCalculatorOptions] { + output_width: 640 + output_height: 640 + } + } +} diff --git a/mediapipe/graphs/object_detection/transformation_umat.pbtxt b/mediapipe/graphs/object_detection/transformation_umat.pbtxt new file mode 100644 index 0000000000..17394754c8 --- /dev/null +++ b/mediapipe/graphs/object_detection/transformation_umat.pbtxt @@ -0,0 +1,30 @@ +profiler_config { + trace_enabled: true + enable_profiler: true + trace_log_interval_count: 200 + trace_log_path: "/mediapipe/tracing/" +} +input_stream: "input_video" +output_stream: "output_video" +node: { + calculator: "ImageTransformationCalculator" + input_stream: "IMAGE:input_video" + output_stream: "IMAGE:tmp_video" + node_options: { + [type.googleapis.com/mediapipe.ImageTransformationCalculatorOptions] { + output_width: 320 + output_height: 320 + } + } +} +node: { + calculator: "ImageTransformationCalculator" + input_stream: "IMAGE:tmp_video" + output_stream: "IMAGE:output_video" + node_options: { + [type.googleapis.com/mediapipe.ImageTransformationCalculatorOptions] { + output_width: 640 + output_height: 640 + } + } +} \ No newline at end of file From 00f97d0ab47c07b160983f80bf23d3f610571be7 Mon Sep 17 00:00:00 2001 From: "iotg.worker" Date: Tue, 16 Apr 2024 12:57:00 +0200 Subject: [PATCH 7/7] Cleanup --- .../examples/desktop/demo_run_graph_main.cc | 23 +++---------------- mediapipe/framework/formats/image_frame.cc | 17 -------------- setup_opencv.sh | 2 +- 3 files changed, 4 insertions(+), 38 deletions(-) diff --git a/mediapipe/examples/desktop/demo_run_graph_main.cc b/mediapipe/examples/desktop/demo_run_graph_main.cc index 8b796c5198..fd261f6f38 100644 --- a/mediapipe/examples/desktop/demo_run_graph_main.cc +++ b/mediapipe/examples/desktop/demo_run_graph_main.cc @@ -12,7 +12,7 @@ // See the License for the specific language governing permissions and // limitations under the License. // -// An example of sending OpenCV webcam frames INFO a MediaPipe graph. +// An example of sending OpenCV webcam frames into a MediaPipe graph. #include #include @@ -92,24 +92,7 @@ absl::Status RunMPPGraph() { OpenClWrapper ocl; ocl.initOpenCL(); - //ocl.printInfo(); - - /* - LOG(INFO) << "haveOpenCL " << cv::ocl::haveOpenCL() <( camera_frame, mediapipe::ImageFormat::SRGB, camera_frame.cols, camera_frame.rows, diff --git a/mediapipe/framework/formats/image_frame.cc b/mediapipe/framework/formats/image_frame.cc index bd6ee8f67a..61ba25a86e 100644 --- a/mediapipe/framework/formats/image_frame.cc +++ b/mediapipe/framework/formats/image_frame.cc @@ -48,10 +48,6 @@ int CountOnes(uint32_t n) { } // namespace -//const char* DEVICE_GPU = ":GPU:0"; - -//static cv::ocl::Context context = cv::ocl::Context::create(DEVICE_GPU); - const ImageFrame::Deleter ImageFrame::PixelDataDeleter::kArrayDelete = std::default_delete(); const ImageFrame::Deleter ImageFrame::PixelDataDeleter::kFree = free; @@ -126,19 +122,6 @@ void ImageFrame::Reset(cv::UMat& inputData, ImageFormat::Format format, int widt CHECK_NE(ImageFormat::UNKNOWN, format_); CHECK(IsValidAlignmentNumber(alignment_boundary)); width_step_ = width * NumberOfChannels() * ByteDepth(); - /*if (alignment_boundary == 1) { - pixel_data_ = {new uint8_t[height * width_step_], - PixelDataDeleter::kArrayDelete}; - } else { - // Increase width_step_ to the smallest multiple of alignment_boundary - // which is large enough to hold all the data. This is done by - // twiddling bits. alignment_boundary - 1 is a mask which sets all - // the low order bits. - width_step_ = ((width_step_ - 1) | (alignment_boundary - 1)) + 1; - pixel_data_ = {reinterpret_cast(aligned_malloc( - height * width_step_, alignment_boundary)), - PixelDataDeleter::kAlignedFree}; - }*/ width_step_ = ((width_step_ - 1) | (alignment_boundary - 1)) + 1; //CHECK_NE(ocl.createMemObject(&ocl.m_mem_obj, inputData), -1); diff --git a/setup_opencv.sh b/setup_opencv.sh index 3c26a44ad9..e83a2295ba 100755 --- a/setup_opencv.sh +++ b/setup_opencv.sh @@ -60,7 +60,7 @@ if [ -z "$1" ] cd ../opencv git checkout 4.7.0 cd release - cmake .. -DCMAKE_BUILD_TYPE=DEBUG -DCMAKE_INSTALL_PREFIX=/usr/local \ + cmake .. -DCMAKE_BUILD_TYPE=RELEASE -DCMAKE_INSTALL_PREFIX=/usr/local \ -DBUILD_LIST=core,improc,imgcodecs,calib3d,features2d,highgui,imgproc,video,videoio,optflow \ -DBUILD_TESTS=OFF -DBUILD_PERF_TESTS=OFF -DBUILD_opencv_ts=OFF \ -DOPENCV_EXTRA_MODULES_PATH=/tmp/build_opencv/opencv_contrib/modules \