From 56bc77272c8be911ecaecc258d7a62cda5453f96 Mon Sep 17 00:00:00 2001 From: Gines Hidalgo Date: Thu, 9 May 2019 10:00:03 -0400 Subject: [PATCH] Faster multi-scale resize --- doc/release_notes.md | 6 +- src/openpose/net/resizeAndMergeBase.cu | 264 ++++++++++++++++++++----- 2 files changed, 220 insertions(+), 50 deletions(-) diff --git a/doc/release_notes.md b/doc/release_notes.md index aca99850..84aa502f 100644 --- a/doc/release_notes.md +++ b/doc/release_notes.md @@ -264,9 +264,9 @@ OpenPose Library - Release Notes 2. Speed up of the CUDA functions of OpenPose: 1. Greedy body part connector implemented in CUDA: +~30% speedup in Nvidia (CUDA) version with default flags and +~10% in maximum accuracy configuration. In addition, it provides a small 0.5% boost in accuracy (default flags). 2. +5-30% additional speedup for the body part connector of point 1. - 3. ~2-4x speedup for NMS. - 4. ~2x speedup for image resize. - 5. +25-30% speedup for rendering. + 3. About 2-4x speedup for NMS. + 4. About 2x speedup for image resize and about 2x speedup for multi-scale resize. + 5. About 25-30% speedup for rendering. 6. Reduced latency and increased speed by moving the resize in CvMatToOpOutput and OpOutputToCvMat to CUDA. The linear speedup 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. diff --git a/src/openpose/net/resizeAndMergeBase.cu b/src/openpose/net/resizeAndMergeBase.cu index b0f0f14e..914609b8 100644 --- a/src/openpose/net/resizeAndMergeBase.cu +++ b/src/openpose/net/resizeAndMergeBase.cu @@ -30,14 +30,13 @@ namespace op const auto x = (blockIdx.x * blockDim.x) + threadIdx.x; const auto y = (blockIdx.y * blockDim.y) + threadIdx.y; const auto channel = (blockIdx.z * blockDim.z) + threadIdx.z; - if (x < widthTarget && y < heightTarget) { const auto sourceArea = widthSource * heightSource; const auto targetArea = widthTarget * heightTarget; const T xSource = (x + T(0.5f)) * widthSource / T(widthTarget) - T(0.5f); const T ySource = (y + T(0.5f)) * heightSource / T(heightTarget) - T(0.5f); - const T* sourcePtrChannel = sourcePtr + channel * sourceArea; + const T* const sourcePtrChannel = sourcePtr + channel * sourceArea; targetPtr[channel * targetArea + y*widthTarget+x] = bicubicInterpolate( sourcePtrChannel, xSource, ySource, widthSource, heightSource, widthSource); } @@ -51,7 +50,6 @@ namespace op const auto x = (blockIdx.x * blockDim.x) + threadIdx.x; const auto y = (blockIdx.y * blockDim.y) + threadIdx.y; const auto channel = (blockIdx.z * blockDim.z) + threadIdx.z; - if (x < widthTarget && y < heightTarget) { const auto targetArea = widthTarget * heightTarget; @@ -60,7 +58,7 @@ namespace op const auto sourceArea = widthSource * heightSource; const T xSource = (x + T(0.5f)) / T(rescaleFactor) - T(0.5f); const T ySource = (y + T(0.5f)) / T(rescaleFactor) - T(0.5f); - const T* sourcePtrChannel = sourcePtr + channel * sourceArea; + const T* const sourcePtrChannel = sourcePtr + channel * sourceArea; targetPtr[channel * targetArea + y*widthTarget+x] = bicubicInterpolate( sourcePtrChannel, xSource, ySource, widthSource, heightSource, widthSource); } @@ -77,7 +75,6 @@ namespace op const auto x = (blockIdx.x * blockDim.x) + threadIdx.x; const auto y = (blockIdx.y * blockDim.y) + threadIdx.y; const auto channel = (blockIdx.z * blockDim.z) + threadIdx.z; - if (x < widthTarget && y < heightTarget) { const auto targetArea = widthTarget * heightTarget; @@ -116,7 +113,7 @@ namespace op const auto targetArea = widthTarget * heightTarget; const T xSource = (x + T(0.5f)) / T(rescaleFactor) - T(0.5f); const T ySource = (y + T(0.5f)) / T(rescaleFactor) - T(0.5f); - const T* sourcePtrChannel = sourcePtr + channel * sourceArea; + const T* const sourcePtrChannel = sourcePtr + channel * sourceArea; targetPtr[channel * targetArea + y*widthTarget+x] = bicubicInterpolate( sourcePtrChannel, xSource, ySource, widthSource, heightSource, widthSource); return; @@ -155,13 +152,88 @@ namespace op } template - __global__ void resizeKernelAndAdd( + __global__ void resizeAndAddKernel( T* targetPtr, const T* const sourcePtr, const T scaleWidth, const T scaleHeight, const int widthSource, const int heightSource, const int widthTarget, const int heightTarget) { const auto x = (blockIdx.x * blockDim.x) + threadIdx.x; const auto y = (blockIdx.y * blockDim.y) + threadIdx.y; + const auto channel = (blockIdx.z * blockDim.z) + threadIdx.z; + if (x < widthTarget && y < heightTarget) + { + const auto sourceArea = widthSource * heightSource; + const auto targetArea = widthTarget * heightTarget; + const T xSource = (x + T(0.5f)) * widthSource / T(widthTarget) - T(0.5f); + const T ySource = (y + T(0.5f)) * heightSource / T(heightTarget) - T(0.5f); + const T* const sourcePtrChannel = sourcePtr + channel * sourceArea; + targetPtr[channel * targetArea + y*widthTarget+x] += bicubicInterpolate( + sourcePtrChannel, xSource, ySource, widthSource, heightSource, widthSource); + } + } + template + __global__ void resizeAndAverageKernel( + T* targetPtr, const T* const sourcePtr, const T scaleWidth, const T scaleHeight, const int widthSource, + const int heightSource, const int widthTarget, const int heightTarget, const int counter) + { + const auto x = (blockIdx.x * blockDim.x) + threadIdx.x; + const auto y = (blockIdx.y * blockDim.y) + threadIdx.y; + const auto channel = (blockIdx.z * blockDim.z) + threadIdx.z; + if (x < widthTarget && y < heightTarget) + { + const auto sourceArea = widthSource * heightSource; + const auto targetArea = widthTarget * heightTarget; + const T xSource = (x + T(0.5f)) / scaleWidth - T(0.5f); + const T ySource = (y + T(0.5f)) / scaleHeight - T(0.5f); + const T* const sourcePtrChannel = sourcePtr + channel * sourceArea; + const auto interpolated = bicubicInterpolate( + sourcePtrChannel, xSource, ySource, widthSource, heightSource, widthSource); + auto& targetPixel = targetPtr[channel * targetArea + y*widthTarget+x]; + targetPixel = (targetPixel + interpolated) / T(counter); + } + } + + template + __global__ void resizeAndAddAndAverageKernel( + T* targetPtr, const int counter, const T* const scaleWidths, const T* const scaleHeights, + const int* const widthSources, const int* const heightSources, const int widthTarget, const int heightTarget, + const T* const sourcePtr0, const T* const sourcePtr1, const T* const sourcePtr2, const T* const sourcePtr3, + const T* const sourcePtr4, const T* const sourcePtr5, const T* const sourcePtr6, const T* const sourcePtr7) + { + const auto x = (blockIdx.x * blockDim.x) + threadIdx.x; + const auto y = (blockIdx.y * blockDim.y) + threadIdx.y; + const auto channel = (blockIdx.z * blockDim.z) + threadIdx.z; + // For each pixel + if (x < widthTarget && y < heightTarget) + { + // Local variable for higher speed + T interpolated = T(0.f); + // For each input source pointer + for (auto i = 0 ; i < counter ; ++i) + { + const auto sourceArea = widthSources[i] * heightSources[i]; + const T xSource = (x + T(0.5f)) / scaleWidths[i] - T(0.5f); + const T ySource = (y + T(0.5f)) / scaleHeights[i] - T(0.5f); + const T* const sourcePtr = ( + i == 0 ? sourcePtr0 : i == 1 ? sourcePtr1 : i == 2 ? sourcePtr2 : i == 3 ? sourcePtr3 + : i == 4 ? sourcePtr4 : i == 5 ? sourcePtr5 : i == 6 ? sourcePtr6 : sourcePtr7); + const T* const sourcePtrChannel = sourcePtr + channel * sourceArea; + interpolated += bicubicInterpolate( + sourcePtrChannel, xSource, ySource, widthSources[i], heightSources[i], widthSources[i]); + } + // Save into memory + const auto targetArea = widthTarget * heightTarget; + targetPtr[channel * targetArea + y*widthTarget+x] = interpolated / T(counter); + } + } + + template + __global__ void resizeAndAddKernelOld( + T* targetPtr, const T* const sourcePtr, const T scaleWidth, const T scaleHeight, const int widthSource, + const int heightSource, const int widthTarget, const int heightTarget) + { + const auto x = (blockIdx.x * blockDim.x) + threadIdx.x; + const auto y = (blockIdx.y * blockDim.y) + threadIdx.y; if (x < widthTarget && y < heightTarget) { const T xSource = (x + T(0.5f)) / scaleWidth - T(0.5f); @@ -172,13 +244,12 @@ namespace op } template - __global__ void resizeKernelAndAverage( + __global__ void resizeAndAverageKernelOld( T* targetPtr, const T* const sourcePtr, const T scaleWidth, const T scaleHeight, const int widthSource, const int heightSource, const int widthTarget, const int heightTarget, const int counter) { const auto x = (blockIdx.x * blockDim.x) + threadIdx.x; const auto y = (blockIdx.y * blockDim.y) + threadIdx.y; - if (x < widthTarget && y < heightTarget) { const T xSource = (x + T(0.5f)) / scaleWidth - T(0.5f); @@ -210,8 +281,9 @@ namespace op const auto heightTarget = targetSize[2]; const auto widthTarget = targetSize[3]; const dim3 threadsPerBlock{THREADS_PER_BLOCK_1D, THREADS_PER_BLOCK_1D}; - const dim3 numBlocks{getNumberCudaBlocks(widthTarget, threadsPerBlock.x), - getNumberCudaBlocks(heightTarget, threadsPerBlock.y)}; + const dim3 numBlocks{ + getNumberCudaBlocks(widthTarget, threadsPerBlock.x), + getNumberCudaBlocks(heightTarget, threadsPerBlock.y)}; const auto& sourceSize = sourceSizes[0]; const auto heightSource = sourceSize[2]; const auto widthSource = sourceSize[3]; @@ -247,9 +319,10 @@ namespace op // // Optimized function for any resize size (suboptimal for 8x resize) // OP_CUDA_PROFILE_INIT(REPS); // const dim3 threadsPerBlock{THREADS_PER_BLOCK_1D, THREADS_PER_BLOCK_1D, 1}; - // const dim3 numBlocks{getNumberCudaBlocks(widthTarget, threadsPerBlock.x), - // getNumberCudaBlocks(heightTarget, threadsPerBlock.y), - // getNumberCudaBlocks(num * channels, threadsPerBlock.z)}; + // const dim3 numBlocks{ + // getNumberCudaBlocks(widthTarget, threadsPerBlock.x), + // getNumberCudaBlocks(heightTarget, threadsPerBlock.y), + // getNumberCudaBlocks(num * channels, threadsPerBlock.z)}; // resizeKernel<<>>( // targetPtr, sourcePtrs.at(0), widthSource, heightSource, widthTarget, heightTarget); // OP_CUDA_PROFILE_END(timeNormalize2, 1e3, REPS); @@ -282,45 +355,143 @@ namespace op // Multi-scaling merging else { - const auto targetChannelOffset = widthTarget * heightTarget; - cudaMemset(targetPtr, 0, channels*targetChannelOffset * sizeof(T)); const auto scaleToMainScaleWidth = widthTarget / T(widthSource); const auto scaleToMainScaleHeight = heightTarget / T(heightSource); - for (auto i = 0u ; i < sourceSizes.size(); i++) + // // Profiling code + // const auto REPS = 10; + // // const auto REPS = 100; + // double timeNormalize1 = 0.; + // double timeNormalize2 = 0.; + // double timeNormalize3 = 0.; + // // Non-optimized function + // OP_CUDA_PROFILE_INIT(REPS); + // const auto targetChannelOffset = widthTarget * heightTarget; + // cudaMemset(targetPtr, 0, channels*targetChannelOffset * sizeof(T)); + // for (auto i = 0u ; i < sourceSizes.size(); ++i) + // { + // const auto& currentSize = sourceSizes.at(i); + // const auto currentHeight = currentSize[2]; + // const auto currentWidth = currentSize[3]; + // const auto sourceChannelOffset = currentHeight * currentWidth; + // const auto scaleInputToNet = scaleInputToNetInputs[i] / scaleInputToNetInputs[0]; + // const auto scaleWidth = scaleToMainScaleWidth / scaleInputToNet; + // const auto scaleHeight = scaleToMainScaleHeight / scaleInputToNet; + // // All but last image --> add + // if (i < sourceSizes.size() - 1) + // { + // for (auto c = 0 ; c < channels ; c++) + // { + // resizeAndAddKernelOld<<>>( + // targetPtr + c * targetChannelOffset, sourcePtrs[i] + c * sourceChannelOffset, + // scaleWidth, scaleHeight, currentWidth, currentHeight, widthTarget, + // heightTarget); + // } + // } + // // Last image --> average all + // else + // { + // for (auto c = 0 ; c < channels ; c++) + // { + // resizeAndAverageKernelOld<<>>( + // targetPtr + c * targetChannelOffset, sourcePtrs[i] + c * sourceChannelOffset, + // scaleWidth, scaleHeight, currentWidth, currentHeight, widthTarget, + // heightTarget, (int)sourceSizes.size()); + // } + // } + // } + // OP_CUDA_PROFILE_END(timeNormalize1, 1e3, REPS); + // // Optimized function for any resize size (suboptimal for 8x resize) + // OP_CUDA_PROFILE_INIT(REPS); + // const auto targetChannelOffset = widthTarget * heightTarget; + // cudaMemset(targetPtr, 0, channels*targetChannelOffset * sizeof(T)); + // const dim3 threadsPerBlock{THREADS_PER_BLOCK_1D, THREADS_PER_BLOCK_1D, 1}; + // const dim3 numBlocks{ + // getNumberCudaBlocks(widthTarget, threadsPerBlock.x), + // getNumberCudaBlocks(heightTarget, threadsPerBlock.y), + // getNumberCudaBlocks(channels, threadsPerBlock.z)}; + // for (auto i = 0u ; i < sourceSizes.size(); ++i) + // { + // const auto& currentSize = sourceSizes.at(i); + // const auto currentHeight = currentSize[2]; + // const auto currentWidth = currentSize[3]; + // const auto scaleInputToNet = scaleInputToNetInputs[i] / scaleInputToNetInputs[0]; + // const auto scaleWidth = scaleToMainScaleWidth / scaleInputToNet; + // const auto scaleHeight = scaleToMainScaleHeight / scaleInputToNet; + // // All but last image --> add + // if (i < sourceSizes.size() - 1) + // resizeAndAddKernel<<>>( + // targetPtr, sourcePtrs[i], scaleWidth, scaleHeight, currentWidth, currentHeight, + // widthTarget, heightTarget); + // // Last image --> average all + // else + // resizeAndAverageKernelOld<<>>( + // targetPtr, sourcePtrs[i], scaleWidth, scaleHeight, currentWidth, currentHeight, + // widthTarget, heightTarget, (int)sourceSizes.size()); + // } + // OP_CUDA_PROFILE_END(timeNormalize2, 1e3, REPS); + + // Super optimized function + // OP_CUDA_PROFILE_INIT(REPS); + if (sourcePtrs.size() > 8) + error("More than 8 scales are not implemented (yet). Notify us to implement it.", + __LINE__, __FUNCTION__, __FILE__); + const dim3 threadsPerBlock{THREADS_PER_BLOCK_1D, THREADS_PER_BLOCK_1D, 1}; + const dim3 numBlocks{ + getNumberCudaBlocks(widthTarget, threadsPerBlock.x), + getNumberCudaBlocks(heightTarget, threadsPerBlock.y), + getNumberCudaBlocks(channels, threadsPerBlock.z)}; + // Fill auxiliary params + std::vector widthSourcesCpu(sourceSizes.size()); + std::vector heightSourcesCpu(sourceSizes.size()); + std::vector scaleWidthsCpu(sourceSizes.size()); + std::vector scaleHeightsCpu(sourceSizes.size()); + for (auto i = 0u ; i < sourceSizes.size(); ++i) { const auto& currentSize = sourceSizes.at(i); - const auto currentHeight = currentSize[2]; - const auto currentWidth = currentSize[3]; - const auto sourceChannelOffset = currentHeight * currentWidth; + heightSourcesCpu[i] = currentSize[2]; + widthSourcesCpu[i] = currentSize[3]; const auto scaleInputToNet = scaleInputToNetInputs[i] / scaleInputToNetInputs[0]; - const auto scaleWidth = scaleToMainScaleWidth / scaleInputToNet; - const auto scaleHeight = scaleToMainScaleHeight / scaleInputToNet; - // All but last image --> add - if (i < sourceSizes.size() - 1) - { - for (auto c = 0 ; c < channels ; c++) - { - resizeKernelAndAdd<<>>( - targetPtr + c * targetChannelOffset, sourcePtrs[i] + c * sourceChannelOffset, - scaleWidth, scaleHeight, currentWidth, currentHeight, widthTarget, - heightTarget - ); - } - } - // Last image --> average all - else - { - for (auto c = 0 ; c < channels ; c++) - { - resizeKernelAndAverage<<>>( - targetPtr + c * targetChannelOffset, sourcePtrs[i] + c * sourceChannelOffset, - scaleWidth, scaleHeight, currentWidth, currentHeight, widthTarget, - heightTarget, (int)sourceSizes.size() - ); - } - } + scaleWidthsCpu[i] = scaleToMainScaleWidth / scaleInputToNet; + scaleHeightsCpu[i] = scaleToMainScaleHeight / scaleInputToNet; } + // GPU params + int* widthSources; + cudaMalloc((void**)&widthSources, sizeof(int) * sourceSizes.size()); + cudaMemcpy( + widthSources, widthSourcesCpu.data(), sizeof(int) * sourceSizes.size(), + cudaMemcpyHostToDevice); + int* heightSources; + cudaMalloc((void**)&heightSources, sizeof(int) * sourceSizes.size()); + cudaMemcpy( + heightSources, heightSourcesCpu.data(), sizeof(int) * sourceSizes.size(), + cudaMemcpyHostToDevice); + T* scaleWidths; + cudaMalloc((void**)&scaleWidths, sizeof(T) * sourceSizes.size()); + cudaMemcpy( + scaleWidths, scaleWidthsCpu.data(), sizeof(T) * sourceSizes.size(), + cudaMemcpyHostToDevice); + T* scaleHeights; + cudaMalloc((void**)&scaleHeights, sizeof(T) * sourceSizes.size()); + cudaMemcpy( + scaleHeights, scaleHeightsCpu.data(), sizeof(T) * sourceSizes.size(), + cudaMemcpyHostToDevice); + // Resize each channel, add all, and get average + resizeAndAddAndAverageKernel<<>>( + targetPtr, (int)sourceSizes.size(), scaleWidths, scaleHeights, widthSources, heightSources, + widthTarget, heightTarget, sourcePtrs[0], sourcePtrs[1], sourcePtrs[2], sourcePtrs[3], + sourcePtrs[4], sourcePtrs[5], sourcePtrs[6], sourcePtrs[7]); + // Free memory + cudaFree(widthSources); + cudaFree(heightSources); + cudaFree(scaleWidths); + cudaFree(scaleHeights); + // OP_CUDA_PROFILE_END(timeNormalize3, 1e3, REPS); + + // // Profiling code + // log(" Res(orig)=" + std::to_string(timeNormalize1) + "ms"); + // log(" Res(new4)=" + std::to_string(timeNormalize2) + "ms"); + // log(" Res(new1)=" + std::to_string(timeNormalize3) + "ms"); } cudaCheck(__LINE__, __FUNCTION__, __FILE__); @@ -335,7 +506,6 @@ namespace op void resizeAndPadRbgGpu( T* targetPtr, const T* const srcPtr, const int widthSource, const int heightSource, const int widthTarget, const int heightTarget, const T scaleFactor) - { try {