YOLOv3处理图片优化——cuda bilinear resize
Posted laosan007
tags:
篇首语:本文由小常识网(cha138.com)小编为大家整理,主要介绍了YOLOv3处理图片优化——cuda bilinear resize相关的知识,希望对你有一定的参考价值。
YOLOv3中处理一张1080P的图片,resize到输入416*416尺寸,调用内部接口做cpu resize,可能80%~90%的时间耗在图像解码、resize上,对比推理时间耗时严重。尝试用cuda做外部resize。
修改下工程用于Ubuntu16.04,1080ti显卡,提供个包其中需要cmakelist修改下opencv路径。
下载:https://pan.baidu.com/s/10RC1Lvxt4FFg5bsbrtnX8w
resizeGPU.cu
#include "resizeGPU.cuh" //#define _DEBUG #define BLOCK_DIM 64 #define threadNum 1024 #define WARP_SIZE 32 #define elemsPerThread 1 int32_t* deviceDataResized; //отмасштабированное изображение в памяти GPU int32_t* deviceData; //оригинальное изображение в памяти GPU int32_t* hostOriginalImage; int32_t* hostResizedImage; void reAllocPinned(int w, int h, int w2, int h2, int32_t* dataSource) { cudaMallocHost((void**)&hostOriginalImage, w*h* sizeof(int32_t)); // host pinned cudaMallocHost((void**)&hostResizedImage, w2*h2 * sizeof(int32_t)); // host pinned memcpy(hostOriginalImage, dataSource, w*h * sizeof(int32_t)); return; } void freePinned() { cudaFreeHost(hostOriginalImage); cudaFreeHost(hostResizedImage); return; } void initGPU(const int maxResolutionX, const int maxResolutionY) { cudaMalloc((void**)&deviceDataResized, maxResolutionX*maxResolutionY * sizeof(int32_t)); cudaMalloc((void**)&deviceData, maxResolutionX*maxResolutionY * sizeof(int32_t)); return; } void deinitGPU() { cudaFree(deviceData); cudaFree(deviceDataResized); return; } __global__ void SomeKernel(int32_t* originalImage, int32_t* resizedImage, int w, int h, int w2, int h2/*, float x_ratio, float y_ratio*/) { __shared__ int32_t tile[1024]; const float x_ratio = ((float)(w - 1)) / w2; const float y_ratio = ((float)(h - 1)) / h2; //const int blockbx = blockIdx.y * w2 + blockIdx.x*BLOCK_DIM; //unsigned int threadId = blockIdx.x * threadNum*elemsPerThread + threadIdx.x; unsigned int threadId = blockIdx.x * threadNum*elemsPerThread + threadIdx.x*elemsPerThread; //__shared__ float result[threadNum*elemsPerThread]; unsigned int shift = 0; //int32_t a, b, c, d, x, y, index; while((threadId < w2*h2 && shift<elemsPerThread)) { const int32_t i = threadId / w2; const int32_t j = threadId - (i*w2); //float x_diff, y_diff, blue, red, green; const int32_t x = (int)(x_ratio * j); const int32_t y = (int)(y_ratio * i); const float x_diff = (x_ratio * j) - x; const float y_diff = (y_ratio * i) - y; const int32_t index = (y*w + x); const int32_t a = originalImage[index]; const int32_t b = originalImage[index + 1]; const int32_t c = originalImage[index + w]; const int32_t d = originalImage[index + w + 1]; // blue element // Yb = Ab(1-w)(1-h) + Bb(w)(1-h) + Cb(h)(1-w) + Db(wh) const float blue = (a & 0xff)*(1 - x_diff)*(1 - y_diff) + (b & 0xff)*(x_diff)*(1 - y_diff) + (c & 0xff)*(y_diff)*(1 - x_diff) + (d & 0xff)*(x_diff*y_diff); // green element // Yg = Ag(1-w)(1-h) + Bg(w)(1-h) + Cg(h)(1-w) + Dg(wh) const float green = ((a >> 8) & 0xff)*(1 - x_diff)*(1 - y_diff) + ((b >> 8) & 0xff)*(x_diff)*(1 - y_diff) + ((c >> 8) & 0xff)*(y_diff)*(1 - x_diff) + ((d >> 8) & 0xff)*(x_diff*y_diff); // red element // Yr = Ar(1-w)(1-h) + Br(w)(1-h) + Cr(h)(1-w) + Dr(wh) const float red = ((a >> 16) & 0xff)*(1 - x_diff)*(1 - y_diff) + ((b >> 16) & 0xff)*(x_diff)*(1 - y_diff) + ((c >> 16) & 0xff)*(y_diff)*(1 - x_diff) + ((d >> 16) & 0xff)*(x_diff*y_diff); /* resizedImage[threadId] = 0xff000000 | ((((int32_t)red) << 16) & 0xff0000) | ((((int32_t)green) << 8) & 0xff00) | ((int32_t)blue); */ tile[threadIdx.x] = 0xff000000 | ((((int32_t)red) << 16) & 0xff0000) | ((((int32_t)green) << 8) & 0xff00) | ((int32_t)blue); threadId++; //threadId+= WARP_SIZE; shift++; } __syncthreads(); threadId = blockIdx.x * threadNum*elemsPerThread + threadIdx.x*elemsPerThread; resizedImage[threadId] = tile[threadIdx.x]; /* shift--; threadId = blockIdx.x * threadNum*elemsPerThread + threadIdx.x*elemsPerThread+ shift; while (shift >= 0) { resizedImage[threadId] = tile[shift]; shift--; threadId--; } */ } int32_t* resizeBilinear_gpu(int w, int h, int w2, int h2) { #ifdef _DEBUG cudaError_t error; //store cuda error codes #endif int length = w2 * h2; // Копирование исходных данных в GPU для обработки cudaMemcpy(deviceData, hostOriginalImage, w*h * sizeof(int32_t), cudaMemcpyHostToDevice); //cudaMemcpy2D(deviceData, w * sizeof(int32_t), hostOriginalImage, w * sizeof(int32_t), w * sizeof(int32_t), h, cudaMemcpyHostToDevice); //error = cudaMemcpyToSymbol(deviceData, pixels, w*h * sizeof(int32_t),0, cudaMemcpyHostToDevice); #ifdef _DEBUG if (error != cudaSuccess) { printf("cudaMemcpy (pixels->deviceData), returned error %s (code %d), line(%d) ", cudaGetErrorString(error), error, __LINE__); exit(EXIT_FAILURE); } #endif dim3 threads = dim3(threadNum, 1,1); //block size 32,32,x dim3 blocks = dim3(w2*h2/ threadNum*elemsPerThread, 1,1); //printf("Blockdim.x %d ", blocks.x); //printf("thrdim.x %d ", threads.x); // Запуск ядра из (length / 256) блоков по 256 потоков, // предполагая, что length кратно 256 SomeKernel << <blocks, threads >> >(deviceData, deviceDataResized, w, h, w2, h2/*, x_ratio, y_ratio*/); cudaDeviceSynchronize(); // Считывание результата из GPU cudaMemcpy(hostResizedImage, deviceDataResized, length * sizeof(int32_t), cudaMemcpyDeviceToHost); return hostResizedImage; }
converter.cpp
#include "converter.hpp" int32_t* cvtMat2Int32(const cv::Mat& srcImage) { int32_t *result = new int32_t[srcImage.cols*srcImage.rows]; int offset = 0; for (int i = 0; i<srcImage.cols*srcImage.rows * 3; i += 3) { int32_t blue = srcImage.data[i]; int32_t green = srcImage.data[i + 1]; int32_t red = srcImage.data[i + 2]; result[offset++] = 0xff000000 | ((((int32_t)red) << 16) & 0xff0000) | ((((int32_t)green) << 8) & 0xff00) | ((int32_t)blue); } return result; } void cvtInt322Mat(int32_t *pxArray, cv::Mat& outImage) { int offset = 0; for (int i = 0; i<outImage.cols*outImage.rows * 3; i += 3) { int32_t a = pxArray[offset++]; int32_t blue = a & 0xff; int32_t green = ((a >> 8) & 0xff); int32_t red = ((a >> 16) & 0xff); outImage.data[i] = blue; outImage.data[i + 1] = green; outImage.data[i + 2] = red; } return; }
resizeCPU.cpp
#include "resizeCPU.hpp" int* resizeBilinear_cpu(int32_t* pixels, int w, int h, int w2, int h2) { int32_t* temp = new int32_t[w2*h2]; int32_t a, b, c, d, x, y, index; float x_ratio = ((float)(w - 1)) / w2; float y_ratio = ((float)(h - 1)) / h2; float x_diff, y_diff, blue, red, green; int offset = 0; for (int i = 0; i<h2; i++) { for (int j = 0; j<w2; j++) { x = (int)(x_ratio * j); y = (int)(y_ratio * i); x_diff = (x_ratio * j) - x; y_diff = (y_ratio * i) - y; index = (y*w + x); a = pixels[index]; b = pixels[index + 1]; c = pixels[index + w]; d = pixels[index + w + 1]; // blue element // Yb = Ab(1-w)(1-h) + Bb(w)(1-h) + Cb(h)(1-w) + Db(wh) blue = (a & 0xff)*(1 - x_diff)*(1 - y_diff) + (b & 0xff)*(x_diff)*(1 - y_diff) + (c & 0xff)*(y_diff)*(1 - x_diff) + (d & 0xff)*(x_diff*y_diff); // green element // Yg = Ag(1-w)(1-h) + Bg(w)(1-h) + Cg(h)(1-w) + Dg(wh) green = ((a >> 8) & 0xff)*(1 - x_diff)*(1 - y_diff) + ((b >> 8) & 0xff)*(x_diff)*(1 - y_diff) + ((c >> 8) & 0xff)*(y_diff)*(1 - x_diff) + ((d >> 8) & 0xff)*(x_diff*y_diff); // red element // Yr = Ar(1-w)(1-h) + Br(w)(1-h) + Cr(h)(1-w) + Dr(wh) red = ((a >> 16) & 0xff)*(1 - x_diff)*(1 - y_diff) + ((b >> 16) & 0xff)*(x_diff)*(1 - y_diff) + ((c >> 16) & 0xff)*(y_diff)*(1 - x_diff) + ((d >> 16) & 0xff)*(x_diff*y_diff); temp[offset++] = 0xff000000 | ((((int32_t)red) << 16) & 0xff0000) | ((((int32_t)green) << 8) & 0xff00) | ((int32_t)blue); } } return temp; }
对比下结果,在1080ti下,resize 1080P图片到416*416尺寸,cuda resize 1.6ms,cpu resize 3.8ms,
darknet内部接口cpu resize 8.0ms。
cpu resize相比darknet resize 接口主要是移位操作有提速,cuda resize处理时间减少很多,但是需要做数据类型Mat与Int32相互转换。
以上是关于YOLOv3处理图片优化——cuda bilinear resize的主要内容,如果未能解决你的问题,请参考以下文章
YOLOv3配置(win10+opencv3.40+cuda9.1+cudnn7.1+vs2015)