Amd fix and Python API (#648)

2026-08-12 04:10:28 +00:00 · 2018-06-18 00:55:30 -06:00
parent b1305573c4
commit 3cf6fb0b89
21 changed files with 1209 additions and 55 deletions
@@ -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)
@@ -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:

@@ -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)
+
+```
+
+
+
@@ -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)
@@ -0,0 +1,233 @@
+// ------------------------- OpenPose Resize Layer Testing -------------------------
+
+#include <openpose/headers.hpp>
+#ifdef USE_OPENCL
+#include <openpose/gpu/opencl.hcl>
+#include <openpose/gpu/cl2.hpp>
+#include <chrono> // `std::chrono::` functions and classes, e.g. std::chrono::milliseconds
+// GFlags: DEFINE_bool, _int32, _int64, _uint64, _double, _string
+#include <gflags/gflags.h>
+// Allow Google Flags in Ubuntu 14
+#ifndef GFLAGS_GFLAGS_H_
+namespace gflags = google;
+#endif
+#ifdef USE_CAFFE
+#include <caffe/net.hpp>
+#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<float>(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<float>(0);
+//            float* newGpuPtr;
+//            cudaMallocHost((void **)&newGpuPtr, newSize.width * newSize.height * sizeof(float));
+//            cudaMemcpy(newGpuPtr, newCpuPtr, newSize.width * newSize.height * sizeof(float),
+//                       cudaMemcpyHostToDevice);
+
+//            std::vector<const float*> sourcePtrs;
+//            sourcePtrs.emplace_back(gpuPtr);
+//            std::array<int, 4> targetSize = {1,1,newImg.size().height,newImg.size().width};
+//            std::array<int, 4> sourceSize = {1,1,img.size().height,img.size().width};
+//            std::vector<std::array<int, 4>> 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<float>(0);
+
+//        // Upload to Dest to GPU
+//        cv::Mat newImg = cv::Mat(newSize,CV_32FC1,cv::Scalar(0));
+
+//        std::vector<const float*> sourcePtrs;
+//        sourcePtrs.emplace_back(cpuPtr);
+//        std::array<int, 4> targetSize = {1,1,newImg.size().height,newImg.size().width};
+//        std::array<int, 4> sourceSize = {1,1,img.size().height,img.size().width};
+//        std::vector<std::array<int, 4>> sourceSizes;
+//        sourceSizes.emplace_back(sourceSize);
+//        op::resizeAndMergeCpu(&newImg.at<float>(0), sourcePtrs, targetSize, sourceSizes);
+
+//        return newImg;
+//    }
+
+typedef cl::KernelFunctor<cl::Buffer, int, int, float> 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<int> 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<caffe::Net<float>> upCaffeNet;
+        caffe::Caffe::set_mode(caffe::Caffe::GPU);
+        caffe::Caffe::SelectDevice(0, true);
+        upCaffeNet.reset(new caffe::Net<float>{"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<cv::Vec3f>(h, w)[c]);
+                }
+            }
+        }
+        blob_proto.set_num(1);
+        caffe::Blob<float>* input_layer = upCaffeNet->input_blobs()[0];
+        input_layer->FromProto(blob_proto);
+        upCaffeNet->Forward(0);
+
+        boost::shared_ptr<caffe::Blob<float> > 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; i<heatmapChannels; i++){
+                cv::Mat hm(cv::Size(output_blob->shape()[2], output_blob->shape()[3]), CV_32FC1);
+                // Read subbuffer
+                cl_buffer_region sourceRegion;
+                op::OpenCL::getBufferRegion<float>(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
+}
@@ -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)
@@ -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)
@@ -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)
@@ -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 <typename T>
        bool buildKernelIntoManager(const std::string& kernelName, const std::string& src = "", bool isFile = false);

@@ -53,6 +59,8 @@ namespace op
      
        template <typename T> static void getBufferRegion(cl_buffer_region& region, const int origin, const int size);

+        int getAlignment();
+
        static std::string clErrorToString(int err);

        static int getTotalGPU();
@@ -19,6 +19,7 @@ namespace op

    template <typename T>
    OP_API void resizeAndMergeOcl(T* targetPtr, const std::vector<const T*>& sourcePtrs,
+                                  std::vector<T*>& sourceTempPtrs,
                                  const std::array<int, 4>& targetSize,
                                  const std::vector<std::array<int, 4>>& sourceSizes,
                                  const std::vector<T>& scaleInputToNetInputs = {1.f},
@@ -44,6 +44,7 @@ namespace op
                                  const std::vector<caffe::Blob<T>*>& bottom);

    private:
+        std::vector<T*> mTempGPUData;
        std::vector<T> mScaleRatios;
        std::vector<std::array<int, 4>> mBottomSizes;
        std::array<int, 4> mTopSize;
@@ -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<unsigned int, std::string>& getPoseBodyPartMapping(const PoseModel poseModel);
@@ -0,0 +1,2 @@
+add_subdirectory(openpose)
+
@@ -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")
@@ -0,0 +1 @@
+from openpose import *
@@ -0,0 +1,343 @@
+#ifndef OPENPOSE_PYTHON_HPP
+#define OPENPOSE_PYTHON_HPP
+
+// OpenPose dependencies
+#include <openpose/core/headers.hpp>
+#include <openpose/filestream/headers.hpp>
+#include <openpose/gui/headers.hpp>
+#include <openpose/pose/headers.hpp>
+#include <openpose/utilities/headers.hpp>
+#include <caffe/caffe.hpp>
+#include <stdlib.h>
+
+#include <openpose/net/nmsCaffe.hpp>
+#include <openpose/net/resizeAndMergeCaffe.hpp>
+#include <openpose/pose/bodyPartConnectorCaffe.hpp>
+#include <boost/make_shared.hpp>
+#include <openpose/pose/poseParameters.hpp>
+#include <openpose/pose/enumClasses.hpp>
+#include <openpose/pose/poseExtractor.hpp>
+#include <openpose/gpu/cuda.hpp>
+#include <openpose/gpu/opencl.hcl>
+
+#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<op::PoseExtractorCaffe> poseExtractorCaffe;
+    std::unique_ptr<op::PoseCpuRenderer> poseRenderer;
+    std::unique_ptr<op::FrameDisplayer> frameDisplayer;
+    std::unique_ptr<op::ScaleAndSizeExtractor> scaleAndSizeExtractor;
+
+    std::unique_ptr<op::ResizeAndMergeCaffe<float>> resizeAndMergeCaffe;
+    std::unique_ptr<op::NmsCaffe<float>> nmsCaffe;
+    std::unique_ptr<op::BodyPartConnectorCaffe<float>> bodyPartConnectorCaffe;
+    std::shared_ptr<caffe::Blob<float>> heatMapsBlob;
+    std::shared_ptr<caffe::Blob<float>> peaksBlob;
+    op::Array<float> mPoseKeypoints;
+    op::Array<float> 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<int> 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<op::ScaleAndSizeExtractor>(new op::ScaleAndSizeExtractor(netInputSize, outputSize, FLAGS_scale_number, FLAGS_scale_gap));
+
+        poseExtractorCaffe = std::unique_ptr<op::PoseExtractorCaffe>(new op::PoseExtractorCaffe{poseModel, FLAGS_model_folder, FLAGS_num_gpu_start});
+
+        poseRenderer = std::unique_ptr<op::PoseCpuRenderer>(new op::PoseCpuRenderer{poseModel, (float)FLAGS_render_threshold, !FLAGS_disable_blending,
+                                                                                                    (float)FLAGS_alpha_pose});
+        frameDisplayer = std::unique_ptr<op::FrameDisplayer>(new op::FrameDisplayer{"OpenPose Tutorial - Example 1", outputSize});
+
+        // Custom
+        resizeAndMergeCaffe = std::unique_ptr<op::ResizeAndMergeCaffe<float>>(new op::ResizeAndMergeCaffe<float>{});
+        nmsCaffe = std::unique_ptr<op::NmsCaffe<float>>(new op::NmsCaffe<float>{});
+        bodyPartConnectorCaffe = std::unique_ptr<op::BodyPartConnectorCaffe<float>>(new op::BodyPartConnectorCaffe<float>{});
+        heatMapsBlob = {std::make_shared<caffe::Blob<float>>(1,1,1,1)};
+        peaksBlob = {std::make_shared<caffe::Blob<float>>(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<caffe::Blob<float>*> caffeNetSharedToPtr(
+        std::vector<boost::shared_ptr<caffe::Blob<float>>>& caffeNetOutputBlob)
+    {
+        try
+        {
+            // Prepare spCaffeNetOutputBlobss
+            std::vector<caffe::Blob<float>*> 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<float>& 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<int> imageSize{inputImage.cols, inputImage.rows};
+        // Step 2 - Get desired scale sizes
+        std::vector<double> scaleInputToNetInputs;
+        std::vector<op::Point<int>> netInputSizes;
+        double scaleInputToOutput;
+        op::Point<int> 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<boost::shared_ptr<caffe::Blob<float>>>& caffeNetOutputBlob, op::Array<float>& poseKeypoints, cv::Mat& displayImage, std::vector<op::Point<int>>& imageSizes){
+        // Get Scale
+        const op::Point<int> inputDataSize{inputImage.cols, inputImage.rows};
+
+        // Convert to Ptr
+        //std::vector<boost::shared_ptr<caffe::Blob<float>>> 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<int> imageSize{inputImage.cols, inputImage.rows};
+        // Step 2 - Get desired scale sizes
+        std::vector<double> scaleInputToNetInputs;
+        std::vector<op::Point<int>> netInputSizes;
+        double scaleInputToOutput;
+        op::Point<int> outputResolution;
+
+        std::tie(scaleInputToNetInputs, netInputSizes, scaleInputToOutput, outputResolution)
+            = scaleAndSizeExtractor->extract(imageSize);
+
+        const auto netInputArray = cvMatToOpInput.createArray(inputImage, scaleInputToNetInputs, netInputSizes);
+
+        // Run the modes
+        const std::vector<float> floatScaleRatios(scaleInputToNetInputs.begin(), scaleInputToNetInputs.end());
+        resizeAndMergeCaffe->setScaleRatios(floatScaleRatios);
+        std::vector<caffe::Blob<float>*> heatMapsBlobs{heatMapsBlob.get()};
+        std::vector<caffe::Blob<float>*> 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<float> 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<boost::shared_ptr<caffe::Blob<float>>> caffeNetOutputBlob;
+
+    for(int i=0; i<size[0]; i++){
+        boost::shared_ptr<caffe::Blob<float>> caffeHmPtr(new caffe::Blob<float>());
+        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; d<caffeHmPtr->shape()[1]; d++){
+            for(int r=0; r<caffeHmPtr->shape()[2]; r++){
+                for(int c=0; c<caffeHmPtr->shape()[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<op::Point<int>> imageSizes;
+    for(int i=0; i<size[0]; i++){
+        op::Point<int> 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
@@ -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)
+
@@ -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 <typename T>
    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<float>(cl_buffer_region& region, const int origin, const int size);
    template void OpenCL::getBufferRegion<double>(cl_buffer_region& region, const int origin, const int size);
    template cl::Kernel&  OpenCL::getKernelFromManager<float>(const std::string& kernelName, const std::string& src, bool isFile);
@@ -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
@@ -4,6 +4,7 @@
 #endif
 #include <openpose/core/common.hpp>
 #include <openpose/net/resizeAndMergeBase.hpp>
+#include <iostream>

 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<cl::Buffer, cl::Buffer, int, int, int, int> ResizeAndMergeFunctor;
+        typedef cl::KernelFunctor<cl::Buffer, cl::Buffer, int, int, int, int, int, int> 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<cl::Buffer, cl::Buffer, float, float, int, int, int, int> ResizeAndAddFunctor;
+        typedef cl::KernelFunctor<cl::Buffer, cl::Buffer, float, float, int, int, int, int, int, int> 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<cl::Buffer, cl::Buffer, float, float, int, int, int, int, int> ResizeAndAverageFunctor;
+        typedef cl::KernelFunctor<cl::Buffer, cl::Buffer, float, float, int, int, int, int, int, int, int> 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<cl::Buffer, cl::Buffer, int, int, int, int> 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 <typename T>
    void resizeAndMergeOcl(T* targetPtr, const std::vector<const T*>& sourcePtrs,
+                           std::vector<T*>& sourceTempPtrs,
                           const std::array<int, 4>& targetSize,
                           const std::vector<std::array<int, 4>>& sourceSizes,
                           const std::vector<T>& scaleInputToNetInputs,
@@ -192,18 +230,24 @@ namespace op

                // Get Kernels
                cl::Buffer targetPtrBuffer = cl::Buffer((cl_mem)(targetPtr), true);
-                auto resizeAndMergeKernel = OpenCL::getInstance(gpuID)->getKernelFunctorFromManager
-                        <ResizeAndMergeFunctor, T>(
-                         "resizeAndMergeKernel", resizeAndMergeOclCommonFunctions+resizeAndMergeKernel);
-                auto resizeAndAddKernel = OpenCL::getInstance(gpuID)->getKernelFunctorFromManager
-                        <ResizeAndAddFunctor, T>(
-                         "resizeAndAddKernel", resizeAndMergeOclCommonFunctions+resizeAndAddKernel);
-                auto resizeAndAverageKernel = OpenCL::getInstance(gpuID)->getKernelFunctorFromManager
-                        <ResizeAndAverageFunctor, T>(
-                         "resizeAndAverageKernel", resizeAndMergeOclCommonFunctions+resizeAndAverageKernel);
-                auto zeroBufferKernel = OpenCL::getInstance(gpuID)->getKernelFunctorFromManager
-                        <ZeroBufferFunctor, T>(
-                         "zeroBufferKernel", zeroBufferKernel);
+                auto resizeAndMergeFullKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager
+                        <op::ResizeAndMergeFullFunctor, T>(
+                            "resizeAndMergeFullKernel",op::resizeAndMergeOclCommonFunctions+op::resizeAndMergeFullKernel);
+                auto resizeAndMergeKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager
+                        <op::ResizeAndMergeFunctor, T>(
+                            "resizeAndMergeKernel",op::resizeAndMergeOclCommonFunctions+op::resizeAndMergeKernel);
+                auto resizeAndAddKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager
+                        <op::ResizeAndAddFunctor, T>(
+                            "resizeAndAddKernel",op::resizeAndMergeOclCommonFunctions+op::resizeAndAddKernel);
+                auto resizeAndAverageKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager
+                        <op::ResizeAndAverageFunctor, T>(
+                            "resizeAndAverageKernel",op::resizeAndMergeOclCommonFunctions+op::resizeAndAverageKernel);
+                auto zeroBufferKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager
+                        <op::ZeroBufferFunctor, T>(
+                            "zeroBufferKernel",op::zeroBufferKernel);
+                auto copyBufferKernel = op::OpenCL::getInstance(gpuID)->getKernelFunctorFromManager
+                        <op::CopyBufferFunctor, T>(
+                            "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<T>(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<T>(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<T>(targerRegion, c * targetChannelOffset, targetChannelOffset);
-                                OpenCL::getBufferRegion<T>(sourceRegion, c * sourceChannelOffset, sourceChannelOffset);
+                                op::OpenCL::getBufferRegion<T>(targerRegion, c * targetChannelOffset, targetChannelOffset);
+                                op::OpenCL::getBufferRegion<T>(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<T>(targerRegion, c * targetChannelOffset, targetChannelOffset);
-                                OpenCL::getBufferRegion<T>(sourceRegion, c * sourceChannelOffset, sourceChannelOffset);
+                                op::OpenCL::getBufferRegion<T>(targerRegion, c * targetChannelOffset, targetChannelOffset);
+                                op::OpenCL::getBufferRegion<T>(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<const float*>& sourcePtrs,
+                                    std::vector<float*>& sourceTempPtrs,
                                    const std::array<int, 4>& targetSize,
                                    const std::vector<std::array<int, 4>>& sourceSizes,
                                    const std::vector<float>& scaleInputToNetInputs,
                                    int gpuID);

    template void resizeAndMergeOcl(double* targetPtr, const std::vector<const double*>& sourcePtrs,
+                                    std::vector<double*>& sourceTempPtrs,
                                    const std::array<int, 4>& targetSize,
                                    const std::vector<std::array<int, 4>>& sourceSizes,
                                    const std::vector<double>& scaleInputToNetInputs,
@@ -4,6 +4,10 @@
 #include <openpose/net/resizeAndMergeBase.hpp>
 #include <openpose/utilities/fastMath.hpp>
 #include <openpose/net/resizeAndMergeCaffe.hpp>
+#ifdef USE_OPENCL
+    #include <openpose/gpu/opencl.hcl>
+    #include <openpose/gpu/cl2.hpp>
+#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<const T*> 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);