mmr/tkDNN
Archived
1
0
Fork 0
Code Issues 46 Pull Requests 12 Releases 4 Activity

Merge branch 'cnet' of https://github.com/ceccocats/tkDNN into cnet

Browse Source
This commit is contained in:
xavier
2020-03-12 20:18:11 +01:00
parent 5c347b1e11 ef1ebea5e9
commit a65c8aa91e
15 changed files with 1816 additions and 65 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/CenternetDetection.cpp
+3 -3
View File
@@ -234,7 +234,7 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
sz_old = sz;
cv::cuda::GpuMat im_Orig;
im_Orig = cv::cuda::GpuMat(imageORIG);
cv::cuda::resize (im_Orig, imageF1_d, cv::Size(new_width, new_height));
// cv::cuda::resize (im_Orig, imageF1_d, cv::Size(new_width, new_height));
checkCuda( cudaDeviceSynchronize() );
sz = imageF1_d.size();
@@ -243,7 +243,7 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
std::cout << " TIME resize: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
step_t = end_t;
cv::cuda::warpAffine(imageF1_d, imageF2_d, trans, cv::Size(inp_width, inp_height), cv::INTER_LINEAR );
// cv::cuda::warpAffine(imageF1_d, imageF2_d, trans, cv::Size(inp_width, inp_height), cv::INTER_LINEAR );
checkCuda( cudaDeviceSynchronize() );
end_t = std::chrono::steady_clock::now();
std::cout << " TIME warpAffine: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
@@ -256,7 +256,7 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
step_t = end_t;
dim2 = dim;
cv::cuda::split(imageF1_d,bgr);//split source
// cv::cuda::split(imageF1_d,bgr);//split source
end_t = std::chrono::steady_clock::now();
std::cout << " TIME split: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
step_t = end_t;
src/NetworkRT.cpp
+50 -19
View File
@@ -304,23 +304,36 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) {
if(l->pool_mode == tkdnnPoolingMode_t::POOLING_AVERAGE) ptype = PoolingType::kAVERAGE;
if(l->pool_mode == tkdnnPoolingMode_t::POOLING_AVERAGE_EXCLUDE_PADDING) ptype = PoolingType::kMAX_AVERAGE_BLEND;
if(l->paddingH == 0 && l->paddingW == 0 && l->input_dim.h == l->output_dim.h && l->input_dim.w == l->output_dim.w)
if(l->maxpoolfixedsize)
{
IPlugin *plugin = new ResizeLayerRT( l->output_dim.c,l->output_dim.h+1,l->output_dim.w+1 );
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
lRT->setName( "Resize" );
input = lRT->getOutput(0);
IPlugin *plugin = new MaxPoolFixedSizeRT(l->output_dim.c, l->output_dim.h, l->output_dim.w, l->output_dim.n, l->strideH, l->strideW, l->winH, l->winH-1);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
lRT->setName( "MaxPoolingFixedSize" );
return lRT;
}
else
{
if(l->paddingH == 0 && l->paddingW == 0 && l->input_dim.h == l->output_dim.h && l->input_dim.w == l->output_dim.w)
{
IPlugin *plugin = new ResizeLayerRT( l->output_dim.c,l->output_dim.h+1,l->output_dim.w+1 );
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
lRT->setName( "Resize" );
IPoolingLayer *lRT = networkRT->addPooling(*input, ptype, DimsHW{l->winH, l->winW});
checkNULL(lRT);
input = lRT->getOutput(0);
}
lRT->setPadding(DimsHW{l->paddingH, l->paddingW});
lRT->setStride(DimsHW{l->strideH, l->strideW});
return lRT;
IPoolingLayer *lRT = networkRT->addPooling(*input, ptype, DimsHW{l->winH, l->winW});
checkNULL(lRT);
lRT->setPadding(DimsHW{l->paddingH, l->paddingW});
lRT->setStride(DimsHW{l->strideH, l->strideW});
return lRT;
}
}
ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
@@ -437,18 +450,18 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) {
//std::cout<<"New plugin Shortcut\n";
ITensor *back_tens = tensors[l->backLayer];
/*
// plugin version
IPlugin *plugin = new ShortcutRT();
IPlugin *plugin = new ShortcutRT(l->backLayer->output_dim);
ITensor **inputs = new ITensor*[2];
inputs[0] = input;
inputs[1] = back_tens;
IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin);
checkNULL(lRT);
*/
IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM);
checkNULL(lRT);
// IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM);
// checkNULL(lRT);
return lRT;
}
@@ -602,13 +615,31 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
}
if(name.find("Shortcut") == 0) {
ShortcutRT *r = new ShortcutRT();
tk::dnn::dataDim_t bdim;
bdim.c = readBUF<int>(buf);
bdim.h = readBUF<int>(buf);
bdim.w = readBUF<int>(buf);
bdim.l = 1;
ShortcutRT *r = new ShortcutRT(bdim);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
return r;
}
if(name.find("Pooling") == 0) {
MaxPoolFixedSizeRT *r = new MaxPoolFixedSizeRT( readBUF<int>(buf), //c
readBUF<int>(buf), //h
readBUF<int>(buf), //w
readBUF<int>(buf), //n
readBUF<int>(buf), //strideH
readBUF<int>(buf), //strideW
readBUF<int>(buf), //winSize
readBUF<int>(buf)); //padding
return r;
}
if(name.find("Resize") == 0) {
ResizeLayerRT *r = new ResizeLayerRT(readBUF<int>(buf), //o_c
readBUF<int>(buf), //o_h
src/Pooling.cpp
+18 -6
View File
@@ -7,7 +7,7 @@ namespace tk { namespace dnn {
Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW,
int paddingH, int paddingW,
tkdnnPoolingMode_t pool_mode, bool final) :
tkdnnPoolingMode_t pool_mode, bool final, bool maxpoolfixedsize) :
Layer(net, final) {
this->winH = winH;
@@ -17,6 +17,7 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW,
this->pool_mode = pool_mode;
this->paddingH = paddingH;
this->paddingW = paddingW;
this->maxpoolfixedsize = maxpoolfixedsize;
checkCUDNN( cudnnCreatePoolingDescriptor(&poolingDesc) );
@@ -111,11 +112,22 @@ dnnType* Pooling::infer(dataDim_t &dim, dnnType* srcData) {
poolDst = tmpOutputData;
}
dnnType alpha = dnnType(1);
dnnType beta = dnnType(0);
checkCUDNN( cudnnPoolingForward(net->cudnnHandle, poolingDesc,
&alpha, srcTensorDesc, poolSrc,
&beta, dstTensorDesc, poolDst) );
if(this->maxpoolfixedsize)
{
MaxPoolingForward(poolSrc, poolDst, dim.n, dim.c, dim.h, dim.w, this->strideH, this->strideW, this->winH, this->winH-1);
}
else
{
dnnType alpha = dnnType(1);
dnnType beta = dnnType(0);
checkCUDNN( cudnnPoolingForward(net->cudnnHandle, poolingDesc,
&alpha, srcTensorDesc, poolSrc,
&beta, dstTensorDesc, poolDst) );
}
//update dim
dim = output_dim;
src/Shortcut.cpp
+1 -1
View File
@@ -10,7 +10,7 @@ Shortcut::Shortcut(Network *net, Layer *backLayer) : Layer(net) {
this->backLayer = backLayer;
checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
if( backLayer->output_dim.c != input_dim.c ||
if( /*backLayer->output_dim.c != input_dim.c ||*/
backLayer->output_dim.w != input_dim.w ||
backLayer->output_dim.h != input_dim.h )
FatalError("Shortcut dim missmatch");
src/kernels/pooling.cu
+52
View File
@@ -0,0 +1,52 @@
#include "kernels.h"
__global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c, int stride_x, int stride_y, int size, int pad, float *input, float *output)
{
int h = (in_h + pad - size) / stride_y + 1;
int w = (in_w + pad - size) / stride_x + 1;
int c = in_c;
int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
if(id >= n) return;
int j = id % w;
id /= w;
int i = id % h;
id /= h;
int k = id % c;
id /= c;
int b = id;
int w_offset = -pad / 2;
int h_offset = -pad / 2;
int out_index = j + w*(i + h*(k + c*b));
float max = -9999999;
int max_i = -1;
int l, m;
for(l = 0; l < size; ++l){
for(m = 0; m < size; ++m){
int cur_h = h_offset + i*stride_y + l;
int cur_w = w_offset + j*stride_x + m;
int index = cur_w + in_w*(cur_h + in_h*(k + b*in_c));
int valid = (cur_h >= 0 && cur_h < in_h &&
cur_w >= 0 && cur_w < in_w);
float val = (valid != 0) ? input[index] : -9999999;
max_i = (val > max) ? index : max_i;
max = (val > max) ? val : max;
}
}
output[out_index] = max;
}
void MaxPoolingForward(dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, int stride_x, int stride_y, int size, int padding, cudaStream_t stream)
{
int tot_size = n*c*h*w;
int blocks = (tot_size+255)/256;
int threads = 256;
forward_maxpool_layer_kernel<<<blocks, threads, 0, stream>>>(tot_size, h, w, c, stride_x, stride_y, size, padding, srcData, dstData);
}