diff --git a/CMakeLists.txt b/CMakeLists.txt index fc5dc047..bbb74803 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -131,7 +131,7 @@ set(GPU_MODE CUDA CACHE STRING "Select the acceleration GPU library or CPU other # else () # set_property(CACHE GPU_MODE PROPERTY STRINGS CPU_ONLY) # endif (CUDA_FOUND AND OpenCL_FOUND) -set_property(CACHE GPU_MODE PROPERTY STRINGS CUDA CPU_ONLY) +set_property(CACHE GPU_MODE PROPERTY STRINGS CUDA OPENCL CPU_ONLY) # Look for CUDA set(CUDA_FOUND FALSE) @@ -234,6 +234,7 @@ option(USE_ASIO "Include Asio header-only library." OFF) # More options option(BUILD_EXAMPLES "Build OpenPose examples." ON) option(BUILD_DOCS "Build OpenPose documentation." OFF) +option(BUILD_PYTHON "Build OpenPose python." OFF) # Build as shared library option(BUILD_SHARED_LIBS "Build as shared lib" ON) @@ -794,6 +795,12 @@ download_model("hand" ${DOWNLOAD_HAND_MODEL} hand/pose_iter_102000.caffemodel message(STATUS "Models Downloaded.") +### PYTHON +if (Caffe_FOUND) + if(BUILD_PYTHON) + add_subdirectory(python) + endif () +endif() ### GENERATE DOCUMENTATION if (UNIX OR APPLE) diff --git a/doc/installation.md b/doc/installation.md index a98f8fde..e73d1423 100644 --- a/doc/installation.md +++ b/doc/installation.md @@ -14,13 +14,14 @@ OpenPose - Installation 1. [Profiling Speed](#profiling-speed) 2. [COCO and MPI Model](#coco-and-mpi-model) 3. [CPU Version](#cpu-version) - 4. [3D Reconstruction Module](#3d-reconstruction-module) - 5. [Calibration Module](#calibration-module) - 6. [Compiling without cuDNN](#compiling-without-cudnn) - 7. [Custom Caffe (Ubuntu Only)](#custom-caffe-ubuntu-only) - 8. [Custom OpenCV (Ubuntu Only)](#custom-opencv-ubuntu-only) - 9. [Doxygen Documentation Autogeneration (Ubuntu Only)](#doxygen-documentation-autogeneration-ubuntu-only) - 10. [CMake Command Line Configuration (Ubuntu Only)](#cmake-command-line-configuration-ubuntu-only) + 4. [OpenCL Version](#opencl-version) + 5. [3D Reconstruction Module](#3d-reconstruction-module) + 6. [Calibration Module](#calibration-module) + 7. [Compiling without cuDNN](#compiling-without-cudnn) + 8. [Custom Caffe (Ubuntu Only)](#custom-caffe-ubuntu-only) + 9. [Custom OpenCV (Ubuntu Only)](#custom-opencv-ubuntu-only) + 10. [Doxygen Documentation Autogeneration (Ubuntu Only)](#doxygen-documentation-autogeneration-ubuntu-only) + 11. [CMake Command Line Configuration (Ubuntu Only)](#cmake-command-line-configuration-ubuntu-only) @@ -49,6 +50,9 @@ Requirements for the default configuration (you might need more resources with a - NVIDIA graphics card with at least 1.6 GB available (the `nvidia-smi` command checks the available GPU memory in Ubuntu). - At least 2.5 GB of free RAM memory for BODY_25 model or 2 GB for COCO model (assuming cuDNN installed). - Highly recommended: cuDNN. + - AMD GPU version: + - Vega series graphics card + - At least 2 GB of free RAM memory. - CPU version: - Around 8GB of free RAM memory. - Highly recommended: a CPU with at least 8 cores. @@ -105,13 +109,19 @@ The instructions in this section describe the steps to build OpenPose using CMak 2. [**cuDNN 5.1**](https://developer.nvidia.com/cudnn): - Ubuntu: Run `sudo ubuntu/install_cudnn.sh` or alternatively download and install it from their website. - Windows (and Ubuntu if manual installation): In order to manually install it, just unzip it and copy (merge) the contents on the CUDA folder, usually `/usr/local/cuda/` in Ubuntu and `C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0` in Windows. -3. Ubuntu - Other prerequisites: +3. AMD GPU version prerequisites: + 1. [**AMD - Windows**](https://support.amd.com/en-us/download): + - Download official AMD drivers for Windows + 2. [**AMD - OpenCL**](https://rocm.github.io/ROCmInstall.html): + - Download 3rd party ROCM driver for Ubuntu + 3. AMD Drivers have not been tested on OSX. Please email us if you wish to test it. This has only been tested on Vega series cards +4. Ubuntu - Other prerequisites: - Caffe prerequisites: By default, OpenPose uses Caffe under the hood. If you have not used Caffe previously, install its dependencies by running `sudo bash ./ubuntu/install_cmake.sh`. - OpenCV must be already installed on your machine. It can be installed with `apt-get install libopencv-dev`. You can also use your own compiled OpenCV version. -4. Windows - **Microsoft Visual Studio (VS) 2015 Enterprise Update 3**: +5. Windows - **Microsoft Visual Studio (VS) 2015 Enterprise Update 3**: - If **Visual Studio 2017 Community** is desired, we do not officially support it, but it might be compiled by firstly [enabling CUDA 8.0 in VS2017](https://stackoverflow.com/questions/43745099/using-cuda-with-visual-studio-2017?answertab=active#tab-top) or use **VS2017 with CUDA 9** by checking the `.vcxproj` file and changing the necessary paths from CUDA 8 to 9. - VS 2015 Enterprise Update 1 will give some compiler errors and VS 2015 Community has not been tested. -5. Windows - **Caffe, OpenCV, and Caffe prerequisites**: +6. Windows - **Caffe, OpenCV, and Caffe prerequisites**: - CMake automatically downloads all the Windows DLLs. Alternatively, you might prefer to download them manually: - Models: - [COCO model](http://posefs1.perception.cs.cmu.edu/OpenPose/models/pose/coco/pose_iter_440000.caffemodel): download in `models/pose/coco/`. @@ -123,10 +133,10 @@ The instructions in this section describe the steps to build OpenPose using CMak - [Caffe](http://posefs1.perception.cs.cmu.edu/OpenPose/3rdparty/windows/caffe_2018_01_18.zip): Unzip as `3rdparty/windows/caffe/`. - [Caffe dependencies](http://posefs1.perception.cs.cmu.edu/OpenPose/3rdparty/windows/caffe3rdparty_2017_07_14.zip): Unzip as `3rdparty/windows/caffe3rdparty/`. - [OpenCV 3.1](http://posefs1.perception.cs.cmu.edu/OpenPose/3rdparty/windows/opencv_310.zip): Unzip as `3rdparty/windows/opencv/`. -6. Mac - **Caffe, OpenCV, and Caffe prerequisites**: +7. Mac - **Caffe, OpenCV, and Caffe prerequisites**: - Install deps by running `bash 3rdparty/osx/install_deps.sh` on your terminal. - Current OSX has only been tested with the CPU Version, and hence must be compiled with the `-DGPU_MODE=CPU_ONLY` flag. -7. **Eigen prerequisite**: +8. **Eigen prerequisite**: - Note: This step is optional, only required for some specific extra functionality, such as extrinsic camera calibration. - If you enable the `WITH_EIGEN` flag when running CMake. You can either: 1. Do not do anything if you set the `WITH_EIGEN` flag to `BUILD`, CMake will automatically download Eigen. Alternatively, you might prefer to download it manually: @@ -262,6 +272,15 @@ Do note that increasing the number of threads results in more memory use. You ca +#### OpenCL Version +If you have an AMD graphics card, you can compile OpenPose with the OpenCL option. To manually select the OpenCL Version, open CMake GUI mentioned above, and set the `GPU_MODE` flag to `OPENCL`. + +The OpenCL version has been tested on Ubuntu and Windows. The OpenCL version requires `libviennacl-dev` which can be installed via apt-get on Ubuntu. This comes packaged with Windows. This has been tested only on AMD Vega series and NVIDIA 10 series graphics cards. Please email us if you have issues with other operating systems or graphics cards. + +Lastly, OpenCL version does not support unfixed resolution. So a folder of images of different resolutions with openpose, requires the `--net_resolution 656x368` flag for example. This should be fixed by the Caffe author in a future patch. + + + #### 3D Reconstruction Module You can include the 3D reconstruction module by: diff --git a/doc/python_module.md b/doc/python_module.md new file mode 100644 index 00000000..d88bd7ff --- /dev/null +++ b/doc/python_module.md @@ -0,0 +1,52 @@ +OpenPose Python Module +============================================= + +## Contents +1. [Introduction](#introduction) +2. [Testing and Installation](#testing-and-installation) + + +## Introduction +This experimental module exposes a Python API for OpenPose. This allows you to construct an OpenPose object, pass in a numpy array for an image, and get a numpy array of the pose positions. This API also exposes an API that allows you to directly pass in heatmaps from a network and extract poses out of it. + +## Testing and Installation +To install the API so that it can be used globally, ensure that the `BUILD_PYTHON` flag is turned on, and run `make install` after compilation. This will install the python library at your desired installation path. (default is `/usr/local/python`) Ensure that this is in your python path in order to use it. + +Two examples can be found in `build/examples/tutorial_python` in your build folder. Navigate directly to this path to run examples. `1_extract_pose` demonstrates a simple use of the API. `2_pose_from_heatmaps` demonstrates constructing pose from heatmaps from the caffe network. + +```python +import sys +import cv2 +import os +# Remember to add your installation path here +sys.path.append('/usr/local/python') +from openpose import * + +# Parameters for OpenPose. Take a look at C++ OpenPose example for meaning of components. Ensure all below are filled +params = dict() +params["logging_level"] = 3 +params["output_resolution"] = "-1x-1" +params["net_resolution"] = "-1x368" +params["model_pose"] = "COCO" +params["alpha_pose"] = 0.6 +params["scale_gap"] = 0.3 +params["scale_number"] = 1 +params["render_threshold"] = 0.05 +params["num_gpu_start"] = 0 +# If GPU version is built, and multiple GPUs are available, set the ID here +params["disable_blending"] = False +params["default_model_folder"] = "/home/user/openpose/models" +# Construct OpenPose object allocates GPU memory +openpose = OpenPose(params) +img = cv2.imread("image.png") +arr, output_image = openpose.forward(img, True) +print arr + +while 1: + cv2.imshow("output", output_image) + cv2.waitKey(15) + +``` + + + diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 083738a7..716e5f4d 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -2,6 +2,7 @@ add_subdirectory(calibration) add_subdirectory(openpose) add_subdirectory(tutorial_add_module) add_subdirectory(tutorial_pose) +add_subdirectory(tutorial_python) add_subdirectory(tutorial_thread) add_subdirectory(tutorial_wrapper) if (UNIX OR APPLE) diff --git a/examples/tests/clTest.cpp b/examples/tests/clTest.cpp new file mode 100644 index 00000000..23db4b79 --- /dev/null +++ b/examples/tests/clTest.cpp @@ -0,0 +1,233 @@ +// ------------------------- OpenPose Resize Layer Testing ------------------------- + +#include +#ifdef USE_OPENCL +#include +#include +#include // `std::chrono::` functions and classes, e.g. std::chrono::milliseconds +// GFlags: DEFINE_bool, _int32, _int64, _uint64, _double, _string +#include +// Allow Google Flags in Ubuntu 14 +#ifndef GFLAGS_GFLAGS_H_ +namespace gflags = google; +#endif +#ifdef USE_CAFFE +#include +#endif + +DEFINE_string(image_path, "examples/media/COCO_val2014_000000000192.jpg", "Process the desired image."); + +// cv::Mat gpuResize(cv::Mat& img, const cv::Size& newSize) +// { +// #ifdef USE_CUDA +// // Upload to Source to GPU +// float* cpuPtr = &img.at(0); +// float* gpuPtr; +// cudaMallocHost((void **)&gpuPtr, img.size().width * img.size().height * sizeof(float)); +// cudaMemcpy(gpuPtr, cpuPtr, img.size().width * img.size().height * sizeof(float), +// cudaMemcpyHostToDevice); + +// // Upload to Dest to GPU +// cv::Mat newImg = cv::Mat(newSize,CV_32FC1,cv::Scalar(0)); +// float* newCpuPtr = &newImg.at(0); +// float* newGpuPtr; +// cudaMallocHost((void **)&newGpuPtr, newSize.width * newSize.height * sizeof(float)); +// cudaMemcpy(newGpuPtr, newCpuPtr, newSize.width * newSize.height * sizeof(float), +// cudaMemcpyHostToDevice); + +// std::vector sourcePtrs; +// sourcePtrs.emplace_back(gpuPtr); +// std::array targetSize = {1,1,newImg.size().height,newImg.size().width}; +// std::array sourceSize = {1,1,img.size().height,img.size().width}; +// std::vector> sourceSizes; +// sourceSizes.emplace_back(sourceSize); +// op::resizeAndMergeGpu(newGpuPtr, sourcePtrs, targetSize, sourceSizes); +// cudaMemcpy(newCpuPtr, newGpuPtr, newImg.size().width * newImg.size().height * sizeof(float), +// cudaMemcpyDeviceToHost); + +// cudaFree(gpuPtr); +// cudaFree(newGpuPtr); +// return newImg; +// #else +// UNUSED(img); +// UNUSED(newSize); +// op::error("OpenPose must be compiled with the `USE_CAFFE` & `USE_CUDA` macro definitions in order to run" +// " this functionality.", __LINE__, __FUNCTION__, __FILE__); +// #endif +// } + +// cv::Mat cpuResize(cv::Mat& img, cv::Size newSize) +// { +// // Upload to Source to GPU +// float* cpuPtr = &img.at(0); + +// // Upload to Dest to GPU +// cv::Mat newImg = cv::Mat(newSize,CV_32FC1,cv::Scalar(0)); + +// std::vector sourcePtrs; +// sourcePtrs.emplace_back(cpuPtr); +// std::array targetSize = {1,1,newImg.size().height,newImg.size().width}; +// std::array sourceSize = {1,1,img.size().height,img.size().width}; +// std::vector> sourceSizes; +// sourceSizes.emplace_back(sourceSize); +// op::resizeAndMergeCpu(&newImg.at(0), sourcePtrs, targetSize, sourceSizes); + +// return newImg; +// } + +typedef cl::KernelFunctor ScaleFunctor; +const std::string scaleKernelString = MULTI_LINE_STRING( + __kernel void scaleKernel(__global float* targetPtr, const int targetWidth, const int targetHeight, const float scale) +{ + int x = get_global_id(0); + int y = get_global_id(1); + int c = get_global_id(2); + + __global float* targetPtrC = &targetPtr[c*targetWidth*targetHeight]; + targetPtrC[y*targetWidth+x] *= scale; + } + ); + +int clTest() +{ + try + { + // logging_level + cv::Mat img = cv::imread(FLAGS_image_path); + if(img.empty()) + op::error("Could not open or find the image: " + FLAGS_image_path, __LINE__, __FUNCTION__, __FILE__); + cv::Mat imgResize; cv::resize(img, imgResize, cv::Size(368,368)); + cv::Mat imgFloat; imgResize.convertTo(imgFloat, CV_32FC3); + imgFloat /= 255.; + int imageVolume = imgFloat.size().width * imgFloat.size().height * imgFloat.channels(); + std::cout << imgFloat.channels() << std::endl; + + // Setup caffe + caffe::Caffe::set_mode(caffe::Caffe::GPU); + std::vector devices; + const int maxNumberGpu = op::OpenCL::getTotalGPU(); + for (auto i = 0; i < maxNumberGpu; i++){ + devices.emplace_back(i); + std::cout << i << std::endl; + } + caffe::Caffe::SetDevices(devices); + + // Load model + std::unique_ptr> upCaffeNet; + caffe::Caffe::set_mode(caffe::Caffe::GPU); + caffe::Caffe::SelectDevice(0, true); + upCaffeNet.reset(new caffe::Net{"models/pose/coco/pose_deploy_linevec.prototxt", caffe::TEST, caffe::Caffe::GetDefaultDevice()}); + upCaffeNet->CopyTrainedLayersFrom("models/pose/coco/pose_iter_440000.caffemodel"); + op::OpenCL::getInstance(0, CL_DEVICE_TYPE_GPU, true); + + // Reshape net to image size + upCaffeNet->blobs()[0]->Reshape({1,imgFloat.channels(),imgResize.size().width,imgResize.size().height}); + upCaffeNet->Reshape(); + + // Convert to caffe image + caffe::BlobProto blob_proto; + blob_proto.set_channels(3); + blob_proto.set_height(imgResize.size().width); + blob_proto.set_width(imgResize.size().height); + blob_proto.clear_data(); + for (int c = 0; c < 3; ++c) { + for (int h = 0; h < imgResize.size().height; ++h) { + for (int w = 0; w < imgResize.size().width; ++w) { + blob_proto.add_data(imgResize.at(h, w)[c]); + } + } + } + blob_proto.set_num(1); + caffe::Blob* input_layer = upCaffeNet->input_blobs()[0]; + input_layer->FromProto(blob_proto); + upCaffeNet->Forward(0); + + boost::shared_ptr > output_blob = upCaffeNet->blob_by_name("net_output"); + + // Test + cl::Device& device = op::OpenCL::getInstance(0)->getDevice(); + cl_uint mem_align; + clGetDeviceInfo(device.get(), CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof(mem_align), &mem_align, nullptr); + std::cout << "Alignment in bits of the base address : " << mem_align << std::endl; + + // GPU Test + cv::Mat finalImage = imgFloat; + try{ + + // Get + float* gpuPtr = output_blob->mutable_gpu_data(); + cl::Buffer outputBuffer((cl_mem)gpuPtr, true); + + // Read it + // Read back image to GPU + float* heatmaps = new float[output_blob->shape()[1] * output_blob->shape()[2] * output_blob->shape()[3]]; + op::OpenCL::getInstance(0)->getQueue().enqueueReadBuffer(outputBuffer, CL_TRUE, 0, + output_blob->shape()[1] * output_blob->shape()[2] * output_blob->shape()[3] * sizeof(float), heatmaps); + + int heatmapChannels = output_blob->shape()[1]; + int shape = output_blob->shape()[2] * output_blob->shape()[3]; + for(int i=0; ishape()[2], output_blob->shape()[3]), CV_32FC1); + // Read subbuffer + cl_buffer_region sourceRegion; + op::OpenCL::getBufferRegion(sourceRegion, i * shape, shape); + cl::Buffer regionBuffer = outputBuffer.createSubBuffer(CL_MEM_READ_WRITE, + CL_BUFFER_CREATE_TYPE_REGION, + &sourceRegion); + } + } + #if defined(USE_OPENCL) && defined(CL_HPP_ENABLE_EXCEPTIONS) + catch (const cl::Error& e) + { + op::error(std::string(e.what()) + " : " + op::OpenCL::clErrorToString(e.err()) + " ID: " + + std::to_string(0), __LINE__, __FUNCTION__, __FILE__); + } + #endif + catch (const std::exception& e) + { + op::error(e.what(), __LINE__, __FUNCTION__, __FILE__); + } + + cv::imshow("win", finalImage); + cv::waitKey(0); + + // Load model + + // img.convertTo(img, CV_32FC1); + // img = cpuResize(img, cv::Size(img.size().width/4,img.size().height/4)); + // img*=0.005; + + // cv::Mat gpuImg = gpuResize(img, cv::Size(img.size().width*8,img.size().height*8)); + // cv::Mat cpuImg = cpuResize(img, cv::Size(img.size().width*8,img.size().height*8)); + // cv::imshow("gpuImg", gpuImg); + // cv::imshow("cpuImg", cpuImg); + + // op::log("Done"); + // cv::waitKey(0); + + return 0; + } + catch (const std::exception& e) + { + op::error(e.what(), __LINE__, __FUNCTION__, __FILE__); + return -1; + } +} +#endif + +int main(int argc, char *argv[]) +{ +#ifdef USE_OPENCL + // Parsing command line flags + gflags::ParseCommandLineFlags(&argc, &argv, true); + + // Running handFromJsonTest + std::thread t(&clTest); + t.join(); + return 0; +#else + op::error("OpenPose must be compiled with the `USE_CAFFE` & `USE_OPENCL` macro definitions in order to run" + " this functionality.", __LINE__, __FUNCTION__, __FILE__); + return 0; +#endif +} diff --git a/examples/tutorial_python/1_extract_pose.py b/examples/tutorial_python/1_extract_pose.py new file mode 100644 index 00000000..ccf9a7fe --- /dev/null +++ b/examples/tutorial_python/1_extract_pose.py @@ -0,0 +1,27 @@ +import sys +import cv2 +import os +dir_path = os.path.dirname(os.path.realpath(__file__)) +sys.path.append('../../python') +from openpose import * + +params = dict() +params["logging_level"] = 3 +params["output_resolution"] = "-1x-1" +params["net_resolution"] = "-1x368" +params["model_pose"] = "COCO" +params["alpha_pose"] = 0.6 +params["scale_gap"] = 0.3 +params["scale_number"] = 1 +params["render_threshold"] = 0.05 +params["num_gpu_start"] = 0 +params["disable_blending"] = False +params["default_model_folder"] = dir_path + "/../../../models/" +openpose = OpenPose(params) +img = cv2.imread(dir_path + "/../../../examples/media/COCO_val2014_000000000192.jpg") +arr, output_image = openpose.forward(img, True) +print arr + +while 1: + cv2.imshow("output", output_image) + cv2.waitKey(15) diff --git a/examples/tutorial_python/2_pose_from_heatmaps.py b/examples/tutorial_python/2_pose_from_heatmaps.py new file mode 100644 index 00000000..a0783707 --- /dev/null +++ b/examples/tutorial_python/2_pose_from_heatmaps.py @@ -0,0 +1,85 @@ +import os +os.environ["GLOG_minloglevel"] = "1" +import caffe +import cv2 +import numpy as np +import sys +import time +dir_path = os.path.dirname(os.path.realpath(__file__)) +sys.path.append('../../python') +dir_path + "/../../models/" +from openpose import OpenPose + +# Params for change +defRes = 736 +scales = [1,0.5] +class Param: + caffemodel = dir_path + "/../../../models/pose/coco/pose_iter_440000.caffemodel" + prototxt = dir_path + "/../../../models/pose/coco/pose_deploy_linevec.prototxt" + +# Load OpenPose object and Caffe Nets +params = dict() +params["logging_level"] = 3 +params["output_resolution"] = "-1x-1" +params["net_resolution"] = "-1x"+str(defRes) +params["model_pose"] = "COCO" +params["alpha_pose"] = 0.6 +params["scale_gap"] = 0.5 +params["scale_number"] = len(scales) +params["render_threshold"] = 0.05 +params["num_gpu_start"] = 0 +params["disable_blending"] = False +params["default_model_folder"] = dir_path + "/../../../models/" +openpose = OpenPose(params) +caffe.set_mode_gpu() +caffe.set_device(0) +nets = [] +for scale in scales: + nets.append(caffe.Net(Param.prototxt, Param.caffemodel, caffe.TEST)) +print "Net loaded" + +# Test Function +first_run = True +def func(frame): + + # Get image processed for network, and scaled image + imagesForNet, imagesOrig = OpenPose.process_frames(frame, defRes, scales) + + # Reshape + global first_run + if first_run: + for i in range(0, len(scales)): + net = nets[i] + imageForNet = imagesForNet[i] + in_shape = net.blobs['image'].data.shape + in_shape = (1, 3, imageForNet.shape[1], imageForNet.shape[2]) + net.blobs['image'].reshape(*in_shape) + net.reshape() + + first_run = False + print "Reshaped" + + # Forward pass to get heatmaps + heatmaps = [] + for i in range(0, len(scales)): + net = nets[i] + imageForNet = imagesForNet[i] + net.blobs['image'].data[0,:,:,:] = imageForNet + net.forward() + heatmaps.append(net.blobs['net_output'].data[:,:,:,:]) + + # Pose from HM Test + array, frame = openpose.poseFromHM(frame, heatmaps, scales) + + # Draw Heatmaps instead + #hm = heatmaps[0][:,0:18,:,:]; frame = OpenPose.draw_all(imagesOrig[0], hm, -1, 1, True) + #paf = heatmaps[0][:,20:,:,:]; frame = OpenPose.draw_all(imagesOrig[0], paf, -1, 4, False) + + return frame + + +img = cv2.imread(dir_path + "/../../../examples/media/COCO_val2014_000000000192.jpg") +frame = func(img) +while 1: + cv2.imshow("output", frame) + cv2.waitKey(15) diff --git a/examples/tutorial_python/CMakeLists.txt b/examples/tutorial_python/CMakeLists.txt new file mode 100644 index 00000000..834e2bb1 --- /dev/null +++ b/examples/tutorial_python/CMakeLists.txt @@ -0,0 +1,3 @@ +### Add Python Test +configure_file(1_extract_pose.py 1_extract_pose.py) +configure_file(2_pose_from_heatmaps.py 2_pose_from_heatmaps.py) diff --git a/include/openpose/gpu/opencl.hcl b/include/openpose/gpu/opencl.hcl index 47b31577..bdc44b1d 100644 --- a/include/openpose/gpu/opencl.hcl +++ b/include/openpose/gpu/opencl.hcl @@ -23,6 +23,8 @@ namespace cl { class CommandQueue; class Kernel; + class Device; + class Context; } // Singleton structure @@ -39,6 +41,10 @@ namespace op cl::CommandQueue& getQueue(); + cl::Device& getDevice(); + + cl::Context& getContext(); + template bool buildKernelIntoManager(const std::string& kernelName, const std::string& src = "", bool isFile = false); @@ -53,6 +59,8 @@ namespace op template static void getBufferRegion(cl_buffer_region& region, const int origin, const int size); + int getAlignment(); + static std::string clErrorToString(int err); static int getTotalGPU(); diff --git a/include/openpose/net/resizeAndMergeBase.hpp b/include/openpose/net/resizeAndMergeBase.hpp index f72baa4e..c90b5698 100644 --- a/include/openpose/net/resizeAndMergeBase.hpp +++ b/include/openpose/net/resizeAndMergeBase.hpp @@ -19,6 +19,7 @@ namespace op template OP_API void resizeAndMergeOcl(T* targetPtr, const std::vector& sourcePtrs, + std::vector& sourceTempPtrs, const std::array& targetSize, const std::vector>& sourceSizes, const std::vector& scaleInputToNetInputs = {1.f}, diff --git a/include/openpose/net/resizeAndMergeCaffe.hpp b/include/openpose/net/resizeAndMergeCaffe.hpp index 66c768db..e25b4025 100644 --- a/include/openpose/net/resizeAndMergeCaffe.hpp +++ b/include/openpose/net/resizeAndMergeCaffe.hpp @@ -44,6 +44,7 @@ namespace op const std::vector*>& bottom); private: + std::vector mTempGPUData; std::vector mScaleRatios; std::vector> mBottomSizes; std::array mTopSize; diff --git a/include/openpose/pose/poseParameters.hpp b/include/openpose/pose/poseParameters.hpp index 4a8a4f02..9dc1e69e 100644 --- a/include/openpose/pose/poseParameters.hpp +++ b/include/openpose/pose/poseParameters.hpp @@ -8,7 +8,8 @@ namespace op { // Constant Global Parameters - const auto POSE_MAX_PEOPLE = 96u; + // For OpenCL-NMS, (POSE_MAX_PEOPLE+1)*3(x,y,score) must be divisible by 32. Easy fix: POSE_MAX_PEOPLE = 32n - 1 + const auto POSE_MAX_PEOPLE = 127u; // Model functions OP_API const std::map& getPoseBodyPartMapping(const PoseModel poseModel); diff --git a/python/CMakeLists.txt b/python/CMakeLists.txt new file mode 100644 index 00000000..c7f3cc71 --- /dev/null +++ b/python/CMakeLists.txt @@ -0,0 +1,2 @@ +add_subdirectory(openpose) + diff --git a/python/openpose/CMakeLists.txt b/python/openpose/CMakeLists.txt new file mode 100644 index 00000000..c266faa3 --- /dev/null +++ b/python/openpose/CMakeLists.txt @@ -0,0 +1,14 @@ +set(PYTHON_FILES + openpose.py + __init__.py + _openpose.cpp) + +add_library(_openpose SHARED ${PYTHON_FILES}) +target_link_libraries(_openpose openpose ${GLOG_LIBRARY} ${GFLAGS_LIBRARY} ${Caffe_LIBS} ${MKL_LIBS} ${GLUT_LIBRARY} ${SPINNAKER_LIB} ${OpenCL_LIBRARIES} ${CMAKE_THREAD_LIBS_INIT}) +SET_TARGET_PROPERTIES(_openpose PROPERTIES PREFIX "") +configure_file(openpose.py openpose.py) +configure_file(__init__.py __init__.py) + +#install(TARGETS _openpose DESTINATION python) +install(DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/ DESTINATION python/openpose FILES_MATCHING PATTERN "*.so") +install(DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/ DESTINATION python/openpose FILES_MATCHING PATTERN "*.py") diff --git a/python/openpose/__init__.py b/python/openpose/__init__.py new file mode 100644 index 00000000..53bfb853 --- /dev/null +++ b/python/openpose/__init__.py @@ -0,0 +1 @@ +from openpose import * diff --git a/python/openpose/_openpose.cpp b/python/openpose/_openpose.cpp new file mode 100644 index 00000000..0b95fcb1 --- /dev/null +++ b/python/openpose/_openpose.cpp @@ -0,0 +1,343 @@ +#ifndef OPENPOSE_PYTHON_HPP +#define OPENPOSE_PYTHON_HPP + +// OpenPose dependencies +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#define default_logging_level 3 +#define default_output_resolution "-1x-1" +#define default_net_resolution "-1x368" +#define default_model_pose "COCO" +#define default_alpha_pose 0.6 +#define default_scale_gap 0.3 +#define default_scale_number 1 +#define default_render_threshold 0.05 +#define default_num_gpu_start 0 +#define default_disable_blending false +#define default_model_folder "models/" + +// Todo, have GPU Number, handle, OpenCL/CPU Cases + +class OpenPose{ +public: + std::unique_ptr poseExtractorCaffe; + std::unique_ptr poseRenderer; + std::unique_ptr frameDisplayer; + std::unique_ptr scaleAndSizeExtractor; + + std::unique_ptr> resizeAndMergeCaffe; + std::unique_ptr> nmsCaffe; + std::unique_ptr> bodyPartConnectorCaffe; + std::shared_ptr> heatMapsBlob; + std::shared_ptr> peaksBlob; + op::Array mPoseKeypoints; + op::Array mPoseScores; + op::PoseModel poseModel; + int mGpuID; + + OpenPose(int FLAGS_logging_level = default_logging_level, + std::string FLAGS_output_resolution = default_output_resolution, + std::string FLAGS_net_resolution = default_net_resolution, + std::string FLAGS_model_pose = default_model_pose, + float FLAGS_alpha_pose = default_alpha_pose, + float FLAGS_scale_gap = default_scale_gap, + int FLAGS_scale_number = default_scale_number, + float FLAGS_render_threshold = default_render_threshold, + int FLAGS_num_gpu_start = default_num_gpu_start, + int FLAGS_disable_blending = default_disable_blending, + std::string FLAGS_model_folder = default_model_folder + ){ + mGpuID = FLAGS_num_gpu_start; + #ifdef USE_CUDA + caffe::Caffe::set_mode(caffe::Caffe::GPU); + caffe::Caffe::SetDevice(mGpuID); + #elif USE_OPENCL + caffe::Caffe::set_mode(caffe::Caffe::GPU); + std::vector devices; + const int maxNumberGpu = op::OpenCL::getTotalGPU(); + for (auto i = 0; i < maxNumberGpu; i++) + devices.emplace_back(i); + caffe::Caffe::SetDevices(devices); + caffe::Caffe::SelectDevice(mGpuID, true); + op::OpenCL::getInstance(mGpuID, CL_DEVICE_TYPE_GPU, true); + #else + caffe::Caffe::set_mode(caffe::Caffe::CPU); + #endif + op::log("OpenPose Library Python Wrapper", op::Priority::High); + // ------------------------- INITIALIZATION ------------------------- + // Step 1 - Set logging level + // - 0 will output all the logging messages + // - 255 will output nothing + op::ConfigureLog::setPriorityThreshold((op::Priority)FLAGS_logging_level); + op::log("", op::Priority::Low, __LINE__, __FUNCTION__, __FILE__); + // Step 2 - Read Google flags (user defined configuration) + // outputSize + const auto outputSize = op::flagsToPoint(FLAGS_output_resolution, "-1x-1"); + // netInputSize + const auto netInputSize = op::flagsToPoint(FLAGS_net_resolution, "-1x368"); + // poseModel + poseModel = op::flagsToPoseModel(FLAGS_model_pose); + // Check no contradictory flags enabled + if (FLAGS_alpha_pose < 0. || FLAGS_alpha_pose > 1.) + op::error("Alpha value for blending must be in the range [0,1].", __LINE__, __FUNCTION__, __FILE__); + if (FLAGS_scale_gap <= 0. && FLAGS_scale_number > 1) + op::error("Incompatible flag configuration: scale_gap must be greater than 0 or scale_number = 1.", + __LINE__, __FUNCTION__, __FILE__); + // Logging + op::log("", op::Priority::Low, __LINE__, __FUNCTION__, __FILE__); + // Step 3 - Initialize all required classes + scaleAndSizeExtractor = std::unique_ptr(new op::ScaleAndSizeExtractor(netInputSize, outputSize, FLAGS_scale_number, FLAGS_scale_gap)); + + poseExtractorCaffe = std::unique_ptr(new op::PoseExtractorCaffe{poseModel, FLAGS_model_folder, FLAGS_num_gpu_start}); + + poseRenderer = std::unique_ptr(new op::PoseCpuRenderer{poseModel, (float)FLAGS_render_threshold, !FLAGS_disable_blending, + (float)FLAGS_alpha_pose}); + frameDisplayer = std::unique_ptr(new op::FrameDisplayer{"OpenPose Tutorial - Example 1", outputSize}); + + // Custom + resizeAndMergeCaffe = std::unique_ptr>(new op::ResizeAndMergeCaffe{}); + nmsCaffe = std::unique_ptr>(new op::NmsCaffe{}); + bodyPartConnectorCaffe = std::unique_ptr>(new op::BodyPartConnectorCaffe{}); + heatMapsBlob = {std::make_shared>(1,1,1,1)}; + peaksBlob = {std::make_shared>(1,1,1,1)}; + bodyPartConnectorCaffe->setPoseModel(poseModel); + + // Step 4 - Initialize resources on desired thread (in this case single thread, i.e. we init resources here) + poseExtractorCaffe->initializationOnThread(); + poseRenderer->initializationOnThread(); + } + + std::vector*> caffeNetSharedToPtr( + std::vector>>& caffeNetOutputBlob) + { + try + { + // Prepare spCaffeNetOutputBlobss + std::vector*> caffeNetOutputBlobs(caffeNetOutputBlob.size()); + for (auto i = 0u ; i < caffeNetOutputBlobs.size() ; i++) + caffeNetOutputBlobs[i] = caffeNetOutputBlob[i].get(); + return caffeNetOutputBlobs; + } + catch (const std::exception& e) + { + op::error(e.what(), __LINE__, __FUNCTION__, __FILE__); + return {}; + } + } + + void forward(const cv::Mat& inputImage, op::Array& poseKeypoints, cv::Mat& displayImage, bool display = false){ + op::OpOutputToCvMat opOutputToCvMat; + op::CvMatToOpInput cvMatToOpInput; + op::CvMatToOpOutput cvMatToOpOutput; + if(inputImage.empty()) + op::error("Could not open or find the image: ", __LINE__, __FUNCTION__, __FILE__); + const op::Point imageSize{inputImage.cols, inputImage.rows}; + // Step 2 - Get desired scale sizes + std::vector scaleInputToNetInputs; + std::vector> netInputSizes; + double scaleInputToOutput; + op::Point outputResolution; + std::tie(scaleInputToNetInputs, netInputSizes, scaleInputToOutput, outputResolution) + = scaleAndSizeExtractor->extract(imageSize); + // Step 3 - Format input image to OpenPose input and output formats + const auto netInputArray = cvMatToOpInput.createArray(inputImage, scaleInputToNetInputs, netInputSizes); + + // Step 4 - Estimate poseKeypoints + poseExtractorCaffe->forwardPass(netInputArray, imageSize, scaleInputToNetInputs); + poseKeypoints = poseExtractorCaffe->getPoseKeypoints(); + + if(display){ + auto outputArray = cvMatToOpOutput.createArray(inputImage, scaleInputToOutput, outputResolution); + // Step 5 - Render poseKeypoints + poseRenderer->renderPose(outputArray, poseKeypoints, scaleInputToOutput); + // Step 6 - OpenPose output format to cv::Mat + displayImage = opOutputToCvMat.formatToCvMat(outputArray); + } + } + + void poseFromHeatmap(const cv::Mat& inputImage, std::vector>>& caffeNetOutputBlob, op::Array& poseKeypoints, cv::Mat& displayImage, std::vector>& imageSizes){ + // Get Scale + const op::Point inputDataSize{inputImage.cols, inputImage.rows}; + + // Convert to Ptr + //std::vector>> a; + //caffeNetOutputBlob.emplace_back(caffeHmPtr); + const auto caffeNetOutputBlobs = caffeNetSharedToPtr(caffeNetOutputBlob); + + // To be called once only + resizeAndMergeCaffe->Reshape(caffeNetOutputBlobs, {heatMapsBlob.get()}, + op::getPoseNetDecreaseFactor(poseModel), 1.f/1.f, true, + 0); + nmsCaffe->Reshape({heatMapsBlob.get()}, {peaksBlob.get()}, op::getPoseMaxPeaks(poseModel), + op::getPoseNumberBodyParts(poseModel), 0); + bodyPartConnectorCaffe->Reshape({heatMapsBlob.get(), peaksBlob.get()}); + + // Normal + op::OpOutputToCvMat opOutputToCvMat; + op::CvMatToOpInput cvMatToOpInput; + op::CvMatToOpOutput cvMatToOpOutput; + if(inputImage.empty()) + op::error("Could not open or find the image: ", __LINE__, __FUNCTION__, __FILE__); + const op::Point imageSize{inputImage.cols, inputImage.rows}; + // Step 2 - Get desired scale sizes + std::vector scaleInputToNetInputs; + std::vector> netInputSizes; + double scaleInputToOutput; + op::Point outputResolution; + + std::tie(scaleInputToNetInputs, netInputSizes, scaleInputToOutput, outputResolution) + = scaleAndSizeExtractor->extract(imageSize); + + const auto netInputArray = cvMatToOpInput.createArray(inputImage, scaleInputToNetInputs, netInputSizes); + + // Run the modes + const std::vector floatScaleRatios(scaleInputToNetInputs.begin(), scaleInputToNetInputs.end()); + resizeAndMergeCaffe->setScaleRatios(floatScaleRatios); + std::vector*> heatMapsBlobs{heatMapsBlob.get()}; + std::vector*> peaksBlobs{peaksBlob.get()}; + #ifdef USE_CUDA + resizeAndMergeCaffe->Forward_gpu(caffeNetOutputBlobs, heatMapsBlobs); // ~5ms + #elif USE_OPENCL + resizeAndMergeCaffe->Forward_ocl(caffeNetOutputBlobs, heatMapsBlobs); // ~5ms + #else + resizeAndMergeCaffe->Forward_cpu(caffeNetOutputBlobs, heatMapsBlobs); // ~5ms + #endif + + nmsCaffe->setThreshold((float)poseExtractorCaffe->get(op::PoseProperty::NMSThreshold)); + #ifdef USE_CUDA + nmsCaffe->Forward_gpu(heatMapsBlobs, peaksBlobs);// ~2ms + #elif USE_OPENCL + nmsCaffe->Forward_ocl(heatMapsBlobs, peaksBlobs);// ~2ms + #else + nmsCaffe->Forward_cpu(heatMapsBlobs, peaksBlobs);// ~2ms + #endif + op::cudaCheck(__LINE__, __FUNCTION__, __FILE__); + + float mScaleNetToOutput = 1./scaleInputToNetInputs[0]; + bodyPartConnectorCaffe->setScaleNetToOutput(mScaleNetToOutput); + bodyPartConnectorCaffe->setInterMinAboveThreshold( + (float)poseExtractorCaffe->get(op::PoseProperty::ConnectInterMinAboveThreshold) + ); + bodyPartConnectorCaffe->setInterThreshold((float)poseExtractorCaffe->get(op::PoseProperty::ConnectInterThreshold)); + bodyPartConnectorCaffe->setMinSubsetCnt((int)poseExtractorCaffe->get(op::PoseProperty::ConnectMinSubsetCnt)); + bodyPartConnectorCaffe->setMinSubsetScore((float)poseExtractorCaffe->get(op::PoseProperty::ConnectMinSubsetScore)); + + bodyPartConnectorCaffe->Forward_cpu({heatMapsBlob.get(), + peaksBlob.get()}, + mPoseKeypoints, mPoseScores); + poseKeypoints = mPoseKeypoints; + + auto outputArray = cvMatToOpOutput.createArray(inputImage, scaleInputToOutput, outputResolution); + // Step 5 - Render poseKeypoints + poseRenderer->renderPose(outputArray, mPoseKeypoints, scaleInputToOutput); + // Step 6 - OpenPose output format to cv::Mat + displayImage = opOutputToCvMat.formatToCvMat(outputArray); + } +}; + +#ifdef __cplusplus +extern "C" { +#endif + +typedef void* c_OP; +op::Array output; + +c_OP newOP(int logging_level, + char* output_resolution, + char* net_resolution, + char* model_pose, + float alpha_pose, + float scale_gap, + int scale_number, + float render_threshold, + int num_gpu_start, + bool disable_blending, + char* model_folder + ){ + return new OpenPose(logging_level, output_resolution, net_resolution, model_pose, alpha_pose, + scale_gap, scale_number, render_threshold, num_gpu_start, disable_blending, model_folder); +} +void delOP(c_OP op){ + delete (OpenPose *)op; +} +void forward(c_OP op, unsigned char* img, size_t rows, size_t cols, int* size, unsigned char* displayImg, bool display){ + OpenPose* openPose = (OpenPose*)op; + cv::Mat image(rows, cols, CV_8UC3, img); + cv::Mat displayImage(rows, cols, CV_8UC3, displayImg); + openPose->forward(image, output, displayImage, display); + if(output.getSize().size()){ + size[0] = output.getSize()[0]; + size[1] = output.getSize()[1]; + size[2] = output.getSize()[2]; + }else{ + size[0] = 0; size[1] = 0; size[2] = 0; + } + if(display) memcpy(displayImg, displayImage.ptr(), sizeof(unsigned char)*rows*cols*3); +} +void getOutputs(c_OP op, float* array){ + if(output.getSize().size()) + memcpy(array, output.getPtr(), output.getSize()[0]*output.getSize()[1]*output.getSize()[2]*sizeof(float)); +} + +void poseFromHeatmap(c_OP op, unsigned char* img, size_t rows, size_t cols, unsigned char* displayImg, float* hm, int* size, float* ratios){ + OpenPose* openPose = (OpenPose*)op; + cv::Mat image(rows, cols, CV_8UC3, img); + cv::Mat displayImage(rows, cols, CV_8UC3, displayImg); + + std::vector>> caffeNetOutputBlob; + + for(int i=0; i> caffeHmPtr(new caffe::Blob()); + caffeHmPtr->Reshape(1,size[1],size[2]*((float)ratios[i]/(float)ratios[0]),size[3]*((float)ratios[i]/(float)ratios[0])); + float* startIndex = &hm[i*size[1]*size[2]*size[3]]; + for(int d=0; dshape()[1]; d++){ + for(int r=0; rshape()[2]; r++){ + for(int c=0; cshape()[3]; c++){ + int toI = d*caffeHmPtr->shape()[2]*caffeHmPtr->shape()[3] + r*caffeHmPtr->shape()[3] + c; + int fromI = d*size[2]*size[3] + r*size[3] + c; + caffeHmPtr->mutable_cpu_data()[toI] = startIndex[fromI]; + } + } + } + caffeNetOutputBlob.emplace_back(caffeHmPtr); + } + + std::vector> imageSizes; + for(int i=0; i point(cols*ratios[i], rows*ratios[i]); + imageSizes.emplace_back(point); + } + + openPose->poseFromHeatmap(image, caffeNetOutputBlob, output, displayImage, imageSizes); + memcpy(displayImg, displayImage.ptr(), sizeof(unsigned char)*rows*cols*3); + // Copy back kp size + if(output.getSize().size()){ + size[0] = output.getSize()[0]; + size[1] = output.getSize()[1]; + size[2] = output.getSize()[2]; + }else{ + size[0] = 0; size[1] = 0; size[2] = 0; + } +} + +#ifdef __cplusplus +} +#endif + +#endif diff --git a/python/openpose/openpose.py b/python/openpose/openpose.py new file mode 100644 index 00000000..59139933 --- /dev/null +++ b/python/openpose/openpose.py @@ -0,0 +1,227 @@ +""" +Wrap the OpenPose library with Python. +To install run `make install` and library will be stored in /usr/local/python +""" + +import numpy as np +import ctypes as ct +import cv2 +import os +dir_path = os.path.dirname(os.path.realpath(__file__)) + +class OpenPose(object): + """ + Ctypes linkage + """ + _libop= np.ctypeslib.load_library('_openpose', dir_path+'/_openpose.so') + _libop.newOP.argtypes = [ + ct.c_int, ct.c_char_p, ct.c_char_p, ct.c_char_p, ct.c_float, ct.c_float, ct.c_int, ct.c_float, ct.c_int, ct.c_bool, ct.c_char_p] + _libop.newOP.restype = ct.c_void_p + _libop.delOP.argtypes = [ct.c_void_p] + _libop.delOP.restype = None + + _libop.forward.argtypes = [ + ct.c_void_p, np.ctypeslib.ndpointer(dtype=np.uint8), + ct.c_size_t, ct.c_size_t, + np.ctypeslib.ndpointer(dtype=np.int32), np.ctypeslib.ndpointer(dtype=np.uint8), ct.c_bool] + _libop.forward.restype = None + + _libop.getOutputs.argtypes = [ + ct.c_void_p, np.ctypeslib.ndpointer(dtype=np.float32)] + _libop.getOutputs.restype = None + + _libop.poseFromHeatmap.argtypes = [ + ct.c_void_p, np.ctypeslib.ndpointer(dtype=np.uint8), + ct.c_size_t, ct.c_size_t, + np.ctypeslib.ndpointer(dtype=np.uint8), + np.ctypeslib.ndpointer(dtype=np.float32), np.ctypeslib.ndpointer(dtype=np.int32), np.ctypeslib.ndpointer(dtype=np.float32)] + _libop.poseFromHeatmap.restype = None + + def __init__(self, params): + """ + OpenPose Constructor: Prepares OpenPose object + + Parameters + ---------- + params : dict of required parameters. refer to openpose example for more details + + Returns + ------- + outs: OpenPose object + """ + self.op = self._libop.newOP(params["logging_level"], + params["output_resolution"], + params["net_resolution"], + params["model_pose"], + params["alpha_pose"], + params["scale_gap"], + params["scale_number"], + params["render_threshold"], + params["num_gpu_start"], + params["disable_blending"], + params["default_model_folder"]) + + def __del__(self): + """ + OpenPose Destructor: Destroys OpenPose object + """ + self._libop.delOP(self.op) + + def forward(self, image, display = False): + """ + Forward: Takes in an image and returns the human 2D poses, along with drawn image if required + + Parameters + ---------- + image : color image of type ndarray + display : If set to true, we return both the pose and an annotated image for visualization + + Returns + ------- + array: ndarray of human 2D poses [People * BodyPart * XYConfidence] + displayImage : image for visualization + """ + shape = image.shape + displayImage = np.zeros(shape=(image.shape),dtype=np.uint8) + size = np.zeros(shape=(3),dtype=np.int32) + self._libop.forward(self.op, image, shape[0], shape[1], size, displayImage, display) + array = np.zeros(shape=(size),dtype=np.float32) + self._libop.getOutputs(self.op, array) + if display: + return array, displayImage + return array + + def poseFromHM(self, image, hm, ratios=[1]): + """ + Pose From Heatmap: Takes in an image, computed heatmaps, and require scales and computes pose + + Parameters + ---------- + image : color image of type ndarray + hm : heatmap of type ndarray with heatmaps and part affinity fields + ratios : scaling ration if needed to fuse multiple scales + + Returns + ------- + array: ndarray of human 2D poses [People * BodyPart * XYConfidence] + displayImage : image for visualization + """ + if len(ratios) != len(hm): + raise Exception("Ratio shape mismatch") + + # Find largest + hm_combine = np.zeros(shape=(len(hm), hm[0].shape[1], hm[0].shape[2], hm[0].shape[3]),dtype=np.float32) + i=0 + for h in hm: + hm_combine[i,:,0:h.shape[2],0:h.shape[3]] = h + i+=1 + hm = hm_combine + + ratios = np.array(ratios,dtype=np.float32) + + shape = image.shape + displayImage = np.zeros(shape=(image.shape),dtype=np.uint8) + size = np.zeros(shape=(4),dtype=np.int32) + size[0] = hm.shape[0] + size[1] = hm.shape[1] + size[2] = hm.shape[2] + size[3] = hm.shape[3] + + self._libop.poseFromHeatmap(self.op, image, shape[0], shape[1], displayImage, hm, size, ratios) + array = np.zeros(shape=(size[0],size[1],size[2]),dtype=np.float32) + self._libop.getOutputs(self.op, array) + return array, displayImage + + @staticmethod + def process_frames(frame, boxsize = 368, scales = [1]): + base_net_res = None + imagesForNet = [] + imagesOrig = [] + for idx, scale in enumerate(scales): + # Calculate net resolution (width, height) + if idx == 0: + net_res = (16 * int((boxsize * frame.shape[1] / float(frame.shape[0]) / 16) + 0.5), boxsize) + base_net_res = net_res + else: + net_res = ((min(base_net_res[0], max(1, int((base_net_res[0] * scale)+0.5)/16*16))), + (min(base_net_res[1], max(1, int((base_net_res[1] * scale)+0.5)/16*16)))) + input_res = [frame.shape[1], frame.shape[0]] + scale_factor = min((net_res[0] - 1) / float(input_res[0] - 1), (net_res[1] - 1) / float(input_res[1] - 1)) + warp_matrix = np.array([[scale_factor,0,0], + [0,scale_factor,0]]) + if scale_factor != 1: + imageForNet = cv2.warpAffine(frame, warp_matrix, net_res, flags=(cv2.INTER_AREA if scale_factor < 1. else cv2.INTER_CUBIC), borderMode=cv2.BORDER_CONSTANT, borderValue=(0,0,0)) + else: + imageForNet = frame.copy() + + imageOrig = imageForNet.copy() + imageForNet = imageForNet.astype(float) + imageForNet = imageForNet/256. - 0.5 + imageForNet = np.transpose(imageForNet, (2,0,1)) + + imagesForNet.append(imageForNet) + imagesOrig.append(imageOrig) + + return imagesForNet, imagesOrig + + @staticmethod + def draw_all(imageForNet, heatmaps, currIndex, div=4., norm=False): + netDecreaseFactor = float(imageForNet.shape[0]) / float(heatmaps.shape[2]) # 8 + resized_heatmaps = np.zeros(shape=(heatmaps.shape[0], heatmaps.shape[1], imageForNet.shape[0], imageForNet.shape[1])) + num_maps = heatmaps.shape[1] + combined = None + for i in range(0, num_maps): + heatmap = heatmaps[0,i,:,:] + resizedHeatmap = cv2.resize(heatmap, (0,0), fx=netDecreaseFactor, fy=netDecreaseFactor) + + minVal, maxVal, minLoc, maxLoc = cv2.minMaxLoc(resizedHeatmap) + + if i==currIndex and currIndex >=0: + resizedHeatmap = np.abs(resizedHeatmap) + resizedHeatmap = (resizedHeatmap*255.).astype(dtype='uint8') + im_color = cv2.applyColorMap(resizedHeatmap, cv2.COLORMAP_JET) + resizedHeatmap = cv2.addWeighted(imageForNet, 1, im_color, 0.3, 0) + cv2.circle(resizedHeatmap, (int(maxLoc[0]),int(maxLoc[1])), 5, (255,0,0), -1) + return resizedHeatmap + else: + resizedHeatmap = np.abs(resizedHeatmap) + if combined is None: + combined = np.copy(resizedHeatmap); + else: + if i <= num_maps-2: + combined += resizedHeatmap; + if norm: + combined = np.maximum(0, np.minimum(1, combined)); + + if currIndex < 0: + combined /= div + combined = (combined*255.).astype(dtype='uint8') + im_color = cv2.applyColorMap(combined, cv2.COLORMAP_JET) + combined = cv2.addWeighted(imageForNet, 0.5, im_color, 0.5, 0) + cv2.circle(combined, (int(maxLoc[0]),int(maxLoc[1])), 5, (255,0,0), -1) + return combined + + +if __name__ == "__main__": + params = dict() + params["logging_level"] = 3 + params["output_resolution"] = "-1x-1" + params["net_resolution"] = "-1x736" + params["model_pose"] = "COCO" + params["alpha_pose"] = 0.6 + params["scale_gap"] = 0.3 + params["scale_number"] = 2 + params["render_threshold"] = 0.05 + params["num_gpu_start"] = 0 + params["disable_blending"] = False + params["default_model_folder"] = "models/" + openpose = OpenPose(params) + + img = cv2.imread("examples/media/COCO_val2014_000000000192.jpg") + arr, output_image = openpose.forward(img, True) + print arr + + while 1: + cv2.imshow("output", output_image) + cv2.waitKey(15) + diff --git a/src/openpose/gpu/opencl.cpp b/src/openpose/gpu/opencl.cpp index 2e11177e..f7f81425 100644 --- a/src/openpose/gpu/opencl.cpp +++ b/src/openpose/gpu/opencl.cpp @@ -272,6 +272,28 @@ namespace op #endif } + cl::Device& OpenCL::getDevice() + { + #ifdef USE_OPENCL + return upImpl->mDevice; + #else + error("OpenPose must be compiled with the `USE_OPENCL` macro definition in order to use this" + " functionality.", __LINE__, __FUNCTION__, __FILE__); + throw std::runtime_error(""); + #endif + } + + cl::Context& OpenCL::getContext() + { + #ifdef USE_OPENCL + return upImpl->mContext; + #else + error("OpenPose must be compiled with the `USE_OPENCL` macro definition in order to use this" + " functionality.", __LINE__, __FUNCTION__, __FILE__); + throw std::runtime_error(""); + #endif + } + template bool OpenCL::buildKernelIntoManager(const std::string& kernelName, const std::string& src, bool isFile) { @@ -474,6 +496,20 @@ namespace op #endif } + int OpenCL::getAlignment() + { + #ifdef USE_OPENCL + cl::Device& device = this->getDevice(); + cl_uint mem_align; + clGetDeviceInfo(device.get(), CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof(mem_align), &mem_align, nullptr); + return mem_align; + #else + error("OpenPose must be compiled with the `USE_OPENCL` macro definition in order to use this" + " functionality.", __LINE__, __FUNCTION__, __FILE__); + return 0; + #endif + } + template void OpenCL::getBufferRegion(cl_buffer_region& region, const int origin, const int size); template void OpenCL::getBufferRegion(cl_buffer_region& region, const int origin, const int size); template cl::Kernel& OpenCL::getKernelFromManager(const std::string& kernelName, const std::string& src, bool isFile); diff --git a/src/openpose/net/nmsBaseCL.cpp b/src/openpose/net/nmsBaseCL.cpp index 185937bd..1e060bce 100644 --- a/src/openpose/net/nmsBaseCL.cpp +++ b/src/openpose/net/nmsBaseCL.cpp @@ -22,7 +22,7 @@ namespace op Type scoreAcc = 0.f; const int dWidth = 3; const int dHeight = 3; - for (auto dy = -dHeight ; dy <= dHeight ; dy++) + for (int dy = -dHeight ; dy <= dHeight ; dy++) { const int y = peakLocY + dy; if (0 <= y && y < height) // Default height = 368 diff --git a/src/openpose/net/resizeAndMergeBaseCL.cpp b/src/openpose/net/resizeAndMergeBaseCL.cpp index 3784967c..e06e3517 100644 --- a/src/openpose/net/resizeAndMergeBaseCL.cpp +++ b/src/openpose/net/resizeAndMergeBaseCL.cpp @@ -4,6 +4,7 @@ #endif #include #include +#include namespace op { @@ -67,7 +68,7 @@ namespace op Type temp[4]; for (unsigned char i = 0; i < 4; i++) { - const auto offset = yIntArray[i]*widthSourcePtr; + const int offset = yIntArray[i]*widthSourcePtr; temp[i] = cubicInterpolate(sourcePtr[offset + xIntArray[0]], sourcePtr[offset + xIntArray[1]], sourcePtr[offset + xIntArray[2]], sourcePtr[offset + xIntArray[3]], dx); } @@ -86,7 +87,7 @@ namespace op int x = get_global_id(2); Type* targetPtrC = &targetPtr[c*targetWidth*targetHeight]; - const Type* sourcePtrC = &sourcePtr[c*sourceWidth*sourceHeight]; + const __global Type* sourcePtrC = &sourcePtr[c*sourceWidth*sourceHeight]; if (x < targetWidth && y < targetHeight) { @@ -98,31 +99,33 @@ namespace op } ); - typedef cl::KernelFunctor ResizeAndMergeFunctor; + typedef cl::KernelFunctor ResizeAndMergeFunctor; const std::string resizeAndMergeKernel = MULTI_LINE_STRING( __kernel void resizeAndMergeKernel(__global Type* targetPtr, __global const Type* sourcePtr, const int sourceWidth, const int sourceHeight, - const int targetWidth, const int targetHeight) + const int targetWidth, const int targetHeight, + const int widthPadding, const int heightPadding) { int x = get_global_id(0); int y = get_global_id(1); if (x < targetWidth && y < targetHeight) { - const Type xSource = (x + 0.5f) * sourceWidth / (Type)targetWidth - 0.5f; - const Type ySource = (y + 0.5f) * sourceHeight / (Type)targetHeight - 0.5f; + const Type xSource = (x + 0.5f) * (sourceWidth-widthPadding) / (Type)targetWidth - 0.5f; + const Type ySource = (y + 0.5f) * (sourceHeight-heightPadding) / (Type)targetHeight - 0.5f; targetPtr[y*targetWidth+x] = bicubicInterpolate(sourcePtr, xSource, ySource, sourceWidth, sourceHeight, sourceWidth); } } ); - typedef cl::KernelFunctor ResizeAndAddFunctor; + typedef cl::KernelFunctor ResizeAndAddFunctor; const std::string resizeAndAddKernel = MULTI_LINE_STRING( __kernel void resizeAndAddKernel(__global Type* targetPtr, __global const Type* sourcePtr, const Type scaleWidth, const Type scaleHeight, const int sourceWidth, const int sourceHeight, - const int targetWidth, const int targetHeight) + const int targetWidth, const int targetHeight, + const int widthPadding, const int heightPadding) { int x = get_global_id(0); int y = get_global_id(1); @@ -137,12 +140,14 @@ namespace op } ); - typedef cl::KernelFunctor ResizeAndAverageFunctor; + typedef cl::KernelFunctor ResizeAndAverageFunctor; const std::string resizeAndAverageKernel = MULTI_LINE_STRING( __kernel void resizeAndAverageKernel(__global Type* targetPtr, __global const Type* sourcePtr, const Type scaleWidth, const Type scaleHeight, const int sourceWidth, const int sourceHeight, - const int targetWidth, const int targetHeight, const int counter) + const int targetWidth, const int targetHeight, + const int widthPadding, const int heightPadding, + const int counter) { int x = get_global_id(0); int y = get_global_id(1); @@ -170,10 +175,43 @@ namespace op targetPtrC[y*targetWidth+x] = 0; } ); + + typedef cl::KernelFunctor CopyBufferFunctor; + const std::string copyBufferKernel = MULTI_LINE_STRING( + __kernel void copyBufferKernel(__global Type* targetPtr, __global const Type* sourcePtr, + const int sourceWidth, const int sourceHeight, + const int targetWidth, const int targetHeight) + { + int x = get_global_id(0); + int y = get_global_id(1); + int c = get_global_id(2); + + __global Type* targetPtrC = &targetPtr[c*targetWidth*targetHeight]; + __global const Type* sourcePtrC = &sourcePtr[c*sourceWidth*sourceHeight]; + + if(x < sourceWidth && y < sourceHeight) + targetPtrC[y*targetWidth+x] = sourcePtrC[y*sourceWidth+x]; + else + targetPtrC[y*targetWidth+x] = 0; + } + ); #endif + int roundUps(int numToRound, int multiple) + { + if (multiple == 0) + return numToRound; + + int remainder = numToRound % multiple; + if (remainder == 0) + return numToRound; + + return numToRound + multiple - remainder; + } + template void resizeAndMergeOcl(T* targetPtr, const std::vector& sourcePtrs, + std::vector& sourceTempPtrs, const std::array& targetSize, const std::vector>& sourceSizes, const std::vector& scaleInputToNetInputs, @@ -192,18 +230,24 @@ namespace op // Get Kernels cl::Buffer targetPtrBuffer = cl::Buffer((cl_mem)(targetPtr), true); - auto resizeAndMergeKernel = OpenCL::getInstance(gpuID)->getKernelFunctorFromManager - ( - "resizeAndMergeKernel", resizeAndMergeOclCommonFunctions+resizeAndMergeKernel); - auto resizeAndAddKernel = OpenCL::getInstance(gpuID)->getKernelFunctorFromManager - ( - "resizeAndAddKernel", resizeAndMergeOclCommonFunctions+resizeAndAddKernel); - auto resizeAndAverageKernel = OpenCL::getInstance(gpuID)->getKernelFunctorFromManager - ( - "resizeAndAverageKernel", resizeAndMergeOclCommonFunctions+resizeAndAverageKernel); - auto zeroBufferKernel = OpenCL::getInstance(gpuID)->getKernelFunctorFromManager - ( - "zeroBufferKernel", zeroBufferKernel); + auto resizeAndMergeFullKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager + ( + "resizeAndMergeFullKernel",op::resizeAndMergeOclCommonFunctions+op::resizeAndMergeFullKernel); + auto resizeAndMergeKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager + ( + "resizeAndMergeKernel",op::resizeAndMergeOclCommonFunctions+op::resizeAndMergeKernel); + auto resizeAndAddKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager + ( + "resizeAndAddKernel",op::resizeAndMergeOclCommonFunctions+op::resizeAndAddKernel); + auto resizeAndAverageKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager + ( + "resizeAndAverageKernel",op::resizeAndMergeOclCommonFunctions+op::resizeAndAverageKernel); + auto zeroBufferKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager + ( + "zeroBufferKernel",op::zeroBufferKernel); + auto copyBufferKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager + ( + "copyBufferKernel",op::copyBufferKernel); // Parameters const auto channels = targetSize[1]; @@ -212,18 +256,36 @@ namespace op const auto& sourceSize = sourceSizes[0]; const auto sourceHeight = sourceSize[2]; const auto sourceWidth = sourceSize[3]; + //int gpuAlign = (op::OpenCL::getInstance(gpuID)->getAlignment() / 8) / sizeof(T); // No multi-scale merging or no merging required - if (sourceSizes.size() == 10) + if (sourceSizes.size() == 1) { const auto num = sourceSize[0]; if (targetSize[0] > 1 || num == 1) { cl::Buffer sourcePtrBuffer = cl::Buffer((cl_mem)(sourcePtrs.at(0)), true); const auto sourceChannelOffset = sourceHeight * sourceWidth; + const auto sourceWidthIdeal = roundUps(sourceWidth, 16); + const auto sourceHeightIdeal = roundUps(sourceHeight, 16); + const auto sourceChannelOffsetIdeal = sourceHeightIdeal * sourceWidthIdeal; const auto targetChannelOffset = targetWidth * targetHeight; for (auto n = 0; n < num; n++) { + // Allocate memory on GPU once + if(sourceTempPtrs[0] == nullptr){ + cl::Buffer* sourcePtrBufferIdealX = new cl::Buffer(op::OpenCL::getInstance(gpuID)->getContext(), CL_MEM_READ_WRITE, + sizeof(float) * channels * sourceWidthIdeal * sourceHeightIdeal); + sourceTempPtrs[0] = (T*)sourcePtrBufferIdealX->get(); + } + cl::Buffer sourcePtrBufferIdeal = cl::Buffer((cl_mem)(sourceTempPtrs[0]), true); + + // Copy to Buffer + copyBufferKernel(cl::EnqueueArgs(op::OpenCL::getInstance(gpuID)->getQueue(), + cl::NDRange(sourceWidthIdeal, sourceHeightIdeal, channels)), + sourcePtrBufferIdeal, sourcePtrBuffer, + sourceWidth, sourceHeight, sourceWidthIdeal, sourceHeightIdeal); + const auto offsetBase = n*channels; for (auto c = 0 ; c < channels ; c++) { @@ -233,14 +295,15 @@ namespace op cl::Buffer targetBuffer = targetPtrBuffer.createSubBuffer(CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &targerRegion); - OpenCL::getBufferRegion(sourceRegion, offset * sourceChannelOffset, sourceChannelOffset); - cl::Buffer sourceBuffer = sourcePtrBuffer.createSubBuffer(CL_MEM_READ_ONLY, - CL_BUFFER_CREATE_TYPE_REGION, - &sourceRegion); - resizeAndMergeKernel(cl::EnqueueArgs(OpenCL::getInstance(gpuID)->getQueue(), + op::OpenCL::getBufferRegion(sourceRegion, offset * sourceChannelOffsetIdeal, sourceChannelOffsetIdeal); + cl::Buffer sourceBuffer = sourcePtrBufferIdeal.createSubBuffer(CL_MEM_READ_ONLY, + CL_BUFFER_CREATE_TYPE_REGION, + &sourceRegion); + resizeAndMergeKernel(cl::EnqueueArgs(op::OpenCL::getInstance(gpuID)->getQueue(), cl::NDRange(targetWidth, targetHeight)), targetBuffer, sourceBuffer, - sourceWidth, sourceHeight, targetWidth, targetHeight); + sourceWidthIdeal, sourceHeightIdeal, targetWidth, targetHeight, + (sourceWidthIdeal-sourceWidth), (sourceHeightIdeal-sourceHeight)); } } } @@ -269,26 +332,46 @@ namespace op const auto scaleWidth = scaleToMainScaleWidth / scaleInputToNet; const auto scaleHeight = scaleToMainScaleHeight / scaleInputToNet; cl::Buffer sourcePtrBuffer = cl::Buffer((cl_mem)(sourcePtrs.at(i)), true); + + const auto currentHeightIdeal = roundUps(currentHeight, 16); + const auto currentWidthIdeal = roundUps(currentWidth, 16); + const auto sourceChannelOffsetIdeal = currentHeightIdeal * currentWidthIdeal; + + // Allocate memory on GPU once + if(sourceTempPtrs[i] == nullptr){ + cl::Buffer* sourcePtrBufferIdealX = new cl::Buffer(op::OpenCL::getInstance(gpuID)->getContext(), CL_MEM_READ_WRITE, + sizeof(float) * channels * currentWidthIdeal * currentHeightIdeal); + sourceTempPtrs[i] = (T*)sourcePtrBufferIdealX->get(); + } + cl::Buffer sourcePtrBufferIdeal = cl::Buffer((cl_mem)(sourceTempPtrs[i]), true); + + // Copy to Buffer + copyBufferKernel(cl::EnqueueArgs(op::OpenCL::getInstance(gpuID)->getQueue(), + cl::NDRange(currentWidthIdeal, currentHeightIdeal, channels)), + sourcePtrBufferIdeal, sourcePtrBuffer, + currentWidth, currentHeight, currentWidthIdeal, currentHeightIdeal); + // All but last image --> add if (i < sourceSizes.size() - 1) { for (auto c = 0 ; c < channels ; c++) { cl_buffer_region targerRegion, sourceRegion; - OpenCL::getBufferRegion(targerRegion, c * targetChannelOffset, targetChannelOffset); - OpenCL::getBufferRegion(sourceRegion, c * sourceChannelOffset, sourceChannelOffset); + op::OpenCL::getBufferRegion(targerRegion, c * targetChannelOffset, targetChannelOffset); + op::OpenCL::getBufferRegion(sourceRegion, c * sourceChannelOffsetIdeal, sourceChannelOffsetIdeal); cl::Buffer targetBuffer = targetPtrBuffer.createSubBuffer(CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &targerRegion); - cl::Buffer sourceBuffer = sourcePtrBuffer.createSubBuffer(CL_MEM_READ_WRITE, + cl::Buffer sourceBuffer = sourcePtrBufferIdeal.createSubBuffer(CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &sourceRegion); resizeAndAddKernel(cl::EnqueueArgs(OpenCL::getInstance(gpuID)->getQueue(), cl::NDRange(targetWidth, targetHeight)), targetBuffer, sourceBuffer, scaleWidth, scaleHeight, - currentWidth, currentHeight, - targetWidth, targetHeight); + currentWidthIdeal, currentHeightIdeal, + targetWidth, targetHeight, + (currentWidthIdeal-currentWidth), (currentHeightIdeal-currentHeight)); } } // Last image --> average all @@ -297,20 +380,22 @@ namespace op for (auto c = 0 ; c < channels ; c++) { cl_buffer_region targerRegion, sourceRegion; - OpenCL::getBufferRegion(targerRegion, c * targetChannelOffset, targetChannelOffset); - OpenCL::getBufferRegion(sourceRegion, c * sourceChannelOffset, sourceChannelOffset); + op::OpenCL::getBufferRegion(targerRegion, c * targetChannelOffset, targetChannelOffset); + op::OpenCL::getBufferRegion(sourceRegion, c * sourceChannelOffsetIdeal, sourceChannelOffsetIdeal); cl::Buffer targetBuffer = targetPtrBuffer.createSubBuffer(CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &targerRegion); - cl::Buffer sourceBuffer = sourcePtrBuffer.createSubBuffer(CL_MEM_READ_WRITE, + cl::Buffer sourceBuffer = sourcePtrBufferIdeal.createSubBuffer(CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &sourceRegion); resizeAndAverageKernel(cl::EnqueueArgs(OpenCL::getInstance(gpuID)->getQueue(), cl::NDRange(targetWidth, targetHeight)), targetBuffer, sourceBuffer, scaleWidth, scaleHeight, - currentWidth, currentHeight, - targetWidth, targetHeight, (int)sourceSizes.size()); + currentWidthIdeal, currentHeightIdeal, + targetWidth, targetHeight, + (currentWidthIdeal-currentWidth), (currentHeightIdeal-currentHeight), + (int)sourceSizes.size()); } } } @@ -322,6 +407,7 @@ namespace op UNUSED(sourceSizes); UNUSED(scaleInputToNetInputs); UNUSED(gpuID); + UNUSED(sourceTempPtrs); error("OpenPose must be compiled with the `USE_OPENCL` macro definition in order to use this" " functionality.", __LINE__, __FUNCTION__, __FILE__); #endif @@ -340,12 +426,14 @@ namespace op } template void resizeAndMergeOcl(float* targetPtr, const std::vector& sourcePtrs, + std::vector& sourceTempPtrs, const std::array& targetSize, const std::vector>& sourceSizes, const std::vector& scaleInputToNetInputs, int gpuID); template void resizeAndMergeOcl(double* targetPtr, const std::vector& sourcePtrs, + std::vector& sourceTempPtrs, const std::array& targetSize, const std::vector>& sourceSizes, const std::vector& scaleInputToNetInputs, diff --git a/src/openpose/net/resizeAndMergeCaffe.cpp b/src/openpose/net/resizeAndMergeCaffe.cpp index 01f79103..f17eaa79 100644 --- a/src/openpose/net/resizeAndMergeCaffe.cpp +++ b/src/openpose/net/resizeAndMergeCaffe.cpp @@ -4,6 +4,10 @@ #include #include #include +#ifdef USE_OPENCL + #include + #include +#endif namespace op { @@ -85,6 +89,7 @@ namespace op #ifdef USE_OPENCL // GPU ID mGpuID = gpuID; + mTempGPUData.resize(mBottomSizes.size(), nullptr); #else UNUSED(gpuID); #endif @@ -173,7 +178,7 @@ namespace op std::vector sourcePtrs(bottom.size()); for (auto i = 0u ; i < sourcePtrs.size() ; i++) sourcePtrs[i] = bottom[i]->gpu_data(); - resizeAndMergeOcl(top.at(0)->mutable_gpu_data(), sourcePtrs, mTopSize, mBottomSizes, + resizeAndMergeOcl(top.at(0)->mutable_gpu_data(), sourcePtrs, mTempGPUData, mTopSize, mBottomSizes, mScaleRatios, mGpuID); #else UNUSED(bottom);