From 00b6842a5e4dba5bc5cebdc99c856207bb50a8f3 Mon Sep 17 00:00:00 2001 From: Gines Hidalgo Date: Tue, 7 May 2019 08:30:28 -0400 Subject: [PATCH] Reduced latency with CUDA cvMatToOpOutput --- doc/release_notes.md | 1 + include/openpose/wrapper/wrapperAuxiliary.hpp | 3 +- src/openpose/core/cvMatToOpInput.cpp | 75 ++++++++++--------- src/openpose/core/cvMatToOpOutput.cpp | 61 ++++++++------- 4 files changed, 75 insertions(+), 65 deletions(-) diff --git a/doc/release_notes.md b/doc/release_notes.md index 2c144fb1..9d57aa51 100644 --- a/doc/release_notes.md +++ b/doc/release_notes.md @@ -267,6 +267,7 @@ OpenPose Library - Release Notes 3. ~2-4x speedup for NMS. 4. ~2x speedup for image resize. 5. +25-30% speedup for rendering. + 6. Reduced latency and increased speed by moving the resize in cvMatToOpOutput to CUDA. It generalizes better to higher number of GPUs. 3. Unity binding of OpenPose released. OpenPose adds the flag `BUILD_UNITY_SUPPORT` on CMake, which enables special Unity code so it can be built as a Unity plugin. 4. If camera is unplugged, OpenPose GUI and command line will display a warning and try to reconnect it. 5. Wrapper classes simplified and renamed. Wrapper renamed as WrapperT, and created Wrapper as the non-templated class equivalent. diff --git a/include/openpose/wrapper/wrapperAuxiliary.hpp b/include/openpose/wrapper/wrapperAuxiliary.hpp index 28abc5bc..c0bc622f 100644 --- a/include/openpose/wrapper/wrapperAuxiliary.hpp +++ b/include/openpose/wrapper/wrapperAuxiliary.hpp @@ -273,8 +273,7 @@ namespace op cvMatToOpInputW = std::make_shared>(cvMatToOpInput); } // Note: We realized that somehow doing it on GPU for any number of GPUs does speedup the whole OP - resizeOnCpu = true; - // resizeOnCpu = (numberGpuThreads < 3); + resizeOnCpu = false; if (addCvMatToOpOutput && (resizeOnCpu || !renderOutputGpu)) { const auto gpuResize = false; diff --git a/src/openpose/core/cvMatToOpInput.cpp b/src/openpose/core/cvMatToOpInput.cpp index 601ccad9..2cfff4d1 100644 --- a/src/openpose/core/cvMatToOpInput.cpp +++ b/src/openpose/core/cvMatToOpInput.cpp @@ -87,41 +87,46 @@ namespace op // CUDA version (if #Gpus > n) else { - // (Re)Allocate temporary memory - const unsigned int inputImageSize = 3 * cvInputData.rows * cvInputData.cols; - const unsigned int outputImageSize = 3 * netInputSizes[i].x * netInputSizes[i].y; - if (pInputMaxSize < inputImageSize) - { - pInputMaxSize = inputImageSize; - // Free temporary memory - cudaFree(pInputImageCuda); - cudaFree(pInputImageReorderedCuda); - // Re-allocate memory - cudaMalloc((void**)&pInputImageCuda, sizeof(unsigned char) * inputImageSize); - cudaMalloc((void**)&pInputImageReorderedCuda, sizeof(float) * inputImageSize); - } - if (pOutputMaxSize < outputImageSize) - { - pOutputMaxSize = outputImageSize; - // Free temporary memory - cudaFree(pOutputImageCuda); - // Re-allocate memory - cudaMalloc((void**)&pOutputImageCuda, sizeof(float) * outputImageSize); - } - // Copy image to GPU - cudaMemcpy( - pInputImageCuda, cvInputData.data, sizeof(unsigned char) * inputImageSize, - cudaMemcpyHostToDevice); - // Resize image on GPU - reorderAndCast(pInputImageReorderedCuda, pInputImageCuda, cvInputData.cols, cvInputData.rows, 3); - resizeAndMergeRGBGPU( - pOutputImageCuda, pInputImageReorderedCuda, cvInputData.cols, cvInputData.rows, - netInputSizes[i].x, netInputSizes[i].y, (float)scaleInputToNetInputs[i]); - // Copy back to CPU - inputNetData[i].reset({1, 3, netInputSizes.at(i).y, netInputSizes.at(i).x}); - cudaMemcpy( - inputNetData[i].getPtr(), pOutputImageCuda, sizeof(float) * outputImageSize, - cudaMemcpyDeviceToHost); + #ifdef USE_CUDA + // (Re)Allocate temporary memory + const unsigned int inputImageSize = 3 * cvInputData.rows * cvInputData.cols; + const unsigned int outputImageSize = 3 * netInputSizes[i].x * netInputSizes[i].y; + if (pInputMaxSize < inputImageSize) + { + pInputMaxSize = inputImageSize; + // Free temporary memory + cudaFree(pInputImageCuda); + cudaFree(pInputImageReorderedCuda); + // Re-allocate memory + cudaMalloc((void**)&pInputImageCuda, sizeof(unsigned char) * inputImageSize); + cudaMalloc((void**)&pInputImageReorderedCuda, sizeof(float) * inputImageSize); + } + if (pOutputMaxSize < outputImageSize) + { + pOutputMaxSize = outputImageSize; + // Free temporary memory + cudaFree(pOutputImageCuda); + // Re-allocate memory + cudaMalloc((void**)&pOutputImageCuda, sizeof(float) * outputImageSize); + } + // Copy image to GPU + cudaMemcpy( + pInputImageCuda, cvInputData.data, sizeof(unsigned char) * inputImageSize, + cudaMemcpyHostToDevice); + // Resize image on GPU + reorderAndCast(pInputImageReorderedCuda, pInputImageCuda, cvInputData.cols, cvInputData.rows, 3); + resizeAndMergeRGBGPU( + pOutputImageCuda, pInputImageReorderedCuda, cvInputData.cols, cvInputData.rows, + netInputSizes[i].x, netInputSizes[i].y, (float)scaleInputToNetInputs[i]); + // Copy back to CPU + inputNetData[i].reset({1, 3, netInputSizes.at(i).y, netInputSizes.at(i).x}); + cudaMemcpy( + inputNetData[i].getPtr(), pOutputImageCuda, sizeof(float) * outputImageSize, + cudaMemcpyDeviceToHost); + #else + error("You need to compile OpenPose with CUDA support in order to use GPU resize.", + __LINE__, __FUNCTION__, __FILE__); + #endif } } return inputNetData; diff --git a/src/openpose/core/cvMatToOpOutput.cpp b/src/openpose/core/cvMatToOpOutput.cpp index 07a7813d..fadbc5c4 100644 --- a/src/openpose/core/cvMatToOpOutput.cpp +++ b/src/openpose/core/cvMatToOpOutput.cpp @@ -89,35 +89,40 @@ namespace op // CUDA version (if #Gpus > 3) else { + #ifdef USE_CUDA // Input image can be shared between this one and cvMatToOpInput.hpp - // (Free and re-)Allocate temporary memory - const unsigned int inputImageSize = 3 * cvInputData.rows * cvInputData.cols; - if (pInputMaxSize < inputImageSize) - { - pInputMaxSize = inputImageSize; - cudaFree(pInputImageCuda); - cudaMalloc((void**)&pInputImageCuda, sizeof(unsigned char) * inputImageSize); - } - // (Free and re-)Allocate temporary memory - const unsigned int outputImageSize = 3 * outputResolution.x * outputResolution.y; - if (*spOutputMaxSize < outputImageSize) - { - *spOutputMaxSize = outputImageSize; - cudaFree(*spOutputImageCuda); - cudaMalloc((void**)spOutputImageCuda.get(), sizeof(float) * outputImageSize); - } - // Copy original image to GPU - cudaMemcpy( - pInputImageCuda, cvInputData.data, sizeof(unsigned char) * inputImageSize, cudaMemcpyHostToDevice); - // Resize output image on GPU - resizeAndMergeRGBGPU( - *spOutputImageCuda, pInputImageCuda, cvInputData.cols, cvInputData.rows, outputResolution.x, - outputResolution.y, (float)scaleInputToOutput); - *spGpuMemoryAllocated = true; - // // No need to copy output image back to CPU - // cudaMemcpy( - // outputData.getPtr(), *spOutputImageCuda, sizeof(float) * outputImageSize, - // cudaMemcpyDeviceToHost); + // (Free and re-)Allocate temporary memory + const unsigned int inputImageSize = 3 * cvInputData.rows * cvInputData.cols; + if (pInputMaxSize < inputImageSize) + { + pInputMaxSize = inputImageSize; + cudaFree(pInputImageCuda); + cudaMalloc((void**)&pInputImageCuda, sizeof(unsigned char) * inputImageSize); + } + // (Free and re-)Allocate temporary memory + const unsigned int outputImageSize = 3 * outputResolution.x * outputResolution.y; + if (*spOutputMaxSize < outputImageSize) + { + *spOutputMaxSize = outputImageSize; + cudaFree(*spOutputImageCuda); + cudaMalloc((void**)spOutputImageCuda.get(), sizeof(float) * outputImageSize); + } + // Copy original image to GPU + cudaMemcpy( + pInputImageCuda, cvInputData.data, sizeof(unsigned char) * inputImageSize, cudaMemcpyHostToDevice); + // Resize output image on GPU + resizeAndMergeRGBGPU( + *spOutputImageCuda, pInputImageCuda, cvInputData.cols, cvInputData.rows, outputResolution.x, + outputResolution.y, (float)scaleInputToOutput); + *spGpuMemoryAllocated = true; + // // No need to copy output image back to CPU + // cudaMemcpy( + // outputData.getPtr(), *spOutputImageCuda, sizeof(float) * outputImageSize, + // cudaMemcpyDeviceToHost); + #else + error("You need to compile OpenPose with CUDA support in order to use GPU resize.", + __LINE__, __FUNCTION__, __FILE__); + #endif } // Return result return outputData;