Merge remote-tracking branch 'origin/master' into cnet

This commit is contained in:
Davide Sapienza
2021-07-22 17:03:23 +02:00
59 changed files with 7640 additions and 209 deletions
+5 -5
View File
@@ -5,11 +5,12 @@
namespace tk { namespace dnn {
Activation::Activation(Network *net, int act_mode, const float ceiling) :
Activation::Activation(Network *net, int act_mode, const float ceiling, const float slope) :
Layer(net) {
this->act_mode = act_mode;
this->ceiling = ceiling;
this->act_mode = act_mode;
this->ceiling = ceiling;
this->slope = slope;
checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) );
if(int(act_mode) < 100) {
@@ -46,8 +47,7 @@ Activation::~Activation() {
dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
if(act_mode == ACTIVATION_LEAKY) {
activationLEAKYForward(srcData, dstData, dim.tot());
activationLEAKYForward(srcData, dstData, dim.tot(), this->slope);
}
else if(act_mode == ACTIVATION_MISH) {
activationMishForward(srcData, dstData, dim.tot());
+5
View File
@@ -166,6 +166,11 @@ Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW,
}
initCUDNN(deConv);
if(this->groups != 1)
MACC = kernelH*kernelW*output_dim.c*output_dim.w*output_dim.h;
else
MACC = input_dim.c*kernelH*kernelW*output_dim.c*output_dim.w*output_dim.h;
// allocate warkspace
if (ws_sizeInBytes!=0) {
checkCuda( cudaMalloc(&workSpace, ws_sizeInBytes) );
+6
View File
@@ -73,6 +73,12 @@ DeformConv2d::DeformConv2d( Network *net, int out_ch, int deformable_group, int
output_dim.c = out_ch;
initCUDNN();
if(this->deformableGroup != 1)
MACC = kernelH*kernelW*output_dim.c*output_dim.w*output_dim.h;
else
MACC = input_dim.c*kernelH*kernelW*output_dim.c*output_dim.w*output_dim.h;
//allocate data for infer result
checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
}
+2
View File
@@ -18,6 +18,8 @@ Layer::Layer(Network *net) {
if(!net->addLayer(this))
FatalError("Net reached max number of layers");
}
feature_map_size = input_dim.tot() + output_dim.tot();
}
Layer::~Layer() {
+5 -1
View File
@@ -19,6 +19,8 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs,
int seek = 0;
readBinaryFile(weights_path.c_str(), inputs*outputs*kh*kw*kl, &data_h, &data_d, seek);
seek += inputs*outputs*kh*kw*kl;
n_params = seek;
this->additional_bias = additional_bias;
if(additional_bias) {
readBinaryFile(weights_path.c_str(), outputs, &bias2_h, &bias2_d, seek);
@@ -26,15 +28,17 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs,
}
readBinaryFile(weights_path.c_str(), outputs, &bias_h, &bias_d, seek);
seek += outputs;
this->batchnorm = batchnorm;
if(batchnorm) {
seek += outputs;
readBinaryFile(weights_path.c_str(), outputs, &scales_h, &scales_d, seek);
seek += outputs;
readBinaryFile(weights_path.c_str(), outputs, &mean_h, &mean_d, seek);
seek += outputs;
readBinaryFile(weights_path.c_str(), outputs, &variance_h, &variance_d, seek);
seek += outputs;
float eps = TKDNN_BN_MIN_EPSILON;
+36
View File
@@ -96,6 +96,28 @@ dataDim_t Network::getOutputDim() {
return layers[num_layers-1]->output_dim;
}
void Network::adjustFeatureMapSizeWithShortcuts(){
layerType_t layer_type;
int shortcutted_idx;
for(int i=0; i<num_layers; i++) {
layer_type = layers[i]->getLayerType();
if(layer_type == LAYER_SHORTCUT){
shortcutted_idx = -1;
for(int j=0; j<num_layers; j++) {
if(static_cast<tk::dnn::Shortcut*>(layers[i])->backLayer == layers[j]){
shortcutted_idx = j;
break;
}
}
if(shortcutted_idx == -1)
FatalError("Problem when computing featuer_map_size with shortcuts");
for(int j=shortcutted_idx+1; j<i; ++j)
layers[j]->feature_map_size += layers[shortcutted_idx]->output_dim.tot();
}
}
}
void Network::print() {
printCenteredTitle(" NETWORK MODEL ", '=', 60);
@@ -106,10 +128,21 @@ void Network::print() {
std::cout.width(16); std::cout<<std::left<<"output (H*W,CH)";
std::cout<<"\n";
adjustFeatureMapSizeWithShortcuts();
long long unsigned int tot_params = 0;
long long unsigned int max_feature_map_size = 0;
long long unsigned int tot_MACC = 0;
for(int i=0; i<num_layers; i++) {
dataDim_t in = layers[i]->input_dim;
dataDim_t out = layers[i]->output_dim;
tot_params += layers[i]->n_params;
tot_MACC += layers[i]->MACC;
if(layers[i]->feature_map_size> max_feature_map_size)
max_feature_map_size = layers[i]->feature_map_size;
std::cout.width(3); std::cout<<std::right<<i;
std::cout<<" ";
std::cout.width(16); std::cout<<std::left<<layers[i]->getLayerName();
@@ -128,6 +161,9 @@ void Network::print() {
}
printCenteredTitle("", '=', 60);
std::cout<<"\n";
std::cout<<"N params: "<<tot_params<<std::endl;
std::cout<<"Max feature map size: "<<max_feature_map_size<<std::endl;
std::cout<<"N MACC: "<<tot_MACC<<std::endl<<std::endl;
printCudaMemUsage();
}
const char *Network::getNetworkRTName(const char *network_name){
+20 -8
View File
@@ -139,8 +139,8 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
#if NV_TENSORRT_MAJOR >= 6
engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT);
#else
//engineRT = builderRT->buildCudaEngine(*networkRT);
engineRT = builderRT->buildCudaEngine(*networkRT);
//engineRT = std::shared_ptr<nvinfer1::ICudaEngine>(builderRT->buildCudaEngine(*networkRT));
#endif
if(engineRT == nullptr)
FatalError("cloud not build cuda engine")
@@ -237,6 +237,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
return convert_layer(input, (Flatten*) l);
if(type == LAYER_RESHAPE)
return convert_layer(input, (Reshape*) l);
if(type == LAYER_RESIZE)
return convert_layer(input, (Resize*) l);
if(type == LAYER_REORG)
return convert_layer(input, (Reorg*) l);
if(type == LAYER_REGION)
@@ -390,13 +392,13 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
#if NV_TENSORRT_MAJOR < 6
// plugin version
IPlugin *plugin = new ActivationLeakyRT();
IPlugin *plugin = new ActivationLeakyRT(l->slope);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
return lRT;
#else
IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kLEAKY_RELU);
lRT->setAlpha(0.1);
lRT->setAlpha(l->slope);
checkNULL(lRT);
return lRT;
#endif
@@ -479,13 +481,23 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Flatten *l) {
ILayer* NetworkRT::convert_layer(ITensor *input, Reshape *l) {
// std::cout<<"convert Reshape\n";
l->output_dim.print();
IPlugin *plugin = new ReshapeRT(l->output_dim);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
return lRT;
}
ILayer* NetworkRT::convert_layer(ITensor *input, Resize *l) {
// std::cout<<"convert Resize\n";
IResizeLayer *lRT = networkRT->addResize(*input); //default is kNEAREST
checkNULL(lRT);
Dims d{};
lRT->setResizeMode(ResizeMode(l->mode));
lRT->setOutputDimensions(DimsCHW{l->output_dim.c, l->output_dim.h, l->output_dim.w});
return lRT;
}
ILayer* NetworkRT::convert_layer(ITensor *input, Reorg *l) {
//std::cout<<"convert Reorg\n";
@@ -513,7 +525,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) {
ITensor *back_tens = tensors[l->backLayer];
if(l->backLayer->output_dim.c == l->output_dim.c)
if(l->backLayer->output_dim.c == l->output_dim.c && !l->mul)
{
IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM);
checkNULL(lRT);
@@ -522,7 +534,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) {
else
{
// plugin version
IPlugin *plugin = new ShortcutRT(l->backLayer->output_dim);
IPlugin *plugin = new ShortcutRT(l->backLayer->output_dim, l->mul);
ITensor **inputs = new ITensor*[2];
inputs[0] = input;
inputs[1] = back_tens;
@@ -651,7 +663,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
//std::cout<<name<<std::endl;
if(name.find("ActivationLeaky") == 0) {
ActivationLeakyRT *a = new ActivationLeakyRT();
ActivationLeakyRT *a = new ActivationLeakyRT(readBUF<float>(buf));
a->size = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return a;
@@ -710,7 +722,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
bdim.w = readBUF<int>(buf);
bdim.l = 1;
ShortcutRT *r = new ShortcutRT(bdim);
ShortcutRT *r = new ShortcutRT(bdim, readBUF<bool>(buf));
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
+380 -15
View File
@@ -6,23 +6,389 @@
namespace tk { namespace dnn {
cv::Mat vizFloat2colorMap(cv::Mat map) {
cv::Mat mapillary_15_map(cv::Mat adjMap){
// cv::imshow("test", adjMap);
// cv::waitKey(0);
cv::Mat M1(1, 256, CV_8UC1), M2(1, 256, CV_8UC1), M3(1, 256, CV_8UC1);
//animal
M3.at<uchar>(0)=165;
M2.at<uchar>(0)=42;
M1.at<uchar>(0)=45;
//curb
M3.at<uchar>(1)=196;
M2.at<uchar>(1)=196;
M1.at<uchar>(1)=196;
//barrier
M3.at<uchar>(2)=90;
M2.at<uchar>(2)=120;
M1.at<uchar>(2)=150;
//road
M3.at<uchar>(3)=128;
M2.at<uchar>(3)=64;
M1.at<uchar>(3)=128;
//building
M3.at<uchar>(4)=70;
M2.at<uchar>(4)=70;
M1.at<uchar>(4)=70;
//person
M3.at<uchar>(5)=220;
M2.at<uchar>(5)=20;
M1.at<uchar>(5)=60;
//roadmark
M3.at<uchar>(6)=255;
M2.at<uchar>(6)=255;
M1.at<uchar>(6)=255;
//nature
M3.at<uchar>(7)=107;
M2.at<uchar>(7)=142;
M1.at<uchar>(7)=35;
//sky
M3.at<uchar>(8)=70;
M2.at<uchar>(8)=130;
M1.at<uchar>(8)=180;
//billboard
M3.at<uchar>(9)=220;
M2.at<uchar>(9)=220;
M1.at<uchar>(9)=220;
//pole
M3.at<uchar>(10)=153;
M2.at<uchar>(10)=153;
M1.at<uchar>(10)=153;
//traffic sign
M3.at<uchar>(11)=128;
M2.at<uchar>(11)=128;
M1.at<uchar>(11)=128;
//bike
M3.at<uchar>(12)=119;
M2.at<uchar>(12)=11;
M1.at<uchar>(12)=32;
//vehicle
M3.at<uchar>(13)=0;
M2.at<uchar>(13)=0;
M1.at<uchar>(13)=142;
//void
for(int i=14;i<256;i++)
{
M1.at<uchar>(i)=0;
M2.at<uchar>(i)=0;
M3.at<uchar>(i)=0;
}
cv::Mat r1,r2,r3;
cv::LUT(adjMap,M1,r1);
cv::LUT(adjMap,M2,r2);
cv::LUT(adjMap,M3,r3);
std::vector<cv::Mat> planes;
planes.push_back(r1);
planes.push_back(r2);
planes.push_back(r3);
cv::Mat dst;
cv::merge(planes,dst);
return dst;
}
cv::Mat berkeley_20_map(cv::Mat adjMap){
cv::Mat M1(1, 256, CV_8UC1), M2(1, 256, CV_8UC1), M3(1, 256, CV_8UC1);
//road
M3.at<uchar>(0)=128;
M2.at<uchar>(0)=64;
M1.at<uchar>(0)=128;
//sidewalk
M3.at<uchar>(1)=244;
M2.at<uchar>(1)=35;
M1.at<uchar>(1)=232;
//building
M3.at<uchar>(2)=70;
M2.at<uchar>(2)=70;
M1.at<uchar>(2)=70;
//wall
M3.at<uchar>(3)=102;
M2.at<uchar>(3)=102;
M1.at<uchar>(3)=156;
//fence
M3.at<uchar>(4)=90;
M2.at<uchar>(4)=120;
M1.at<uchar>(4)=150;
//pole
M3.at<uchar>(5)=153;
M2.at<uchar>(5)=153;
M1.at<uchar>(5)=153;
//traffic light
M3.at<uchar>(6)=250;
M2.at<uchar>(6)=170;
M1.at<uchar>(6)=30;
//traffic sign
M3.at<uchar>(7)=128;
M2.at<uchar>(7)=128;
M1.at<uchar>(7)=128;
//nature
M3.at<uchar>(8)=107;
M2.at<uchar>(8)=142;
M1.at<uchar>(8)=35;
//ground
M3.at<uchar>(9)=0;
M2.at<uchar>(9)=192;
M1.at<uchar>(9)=0;
//sky
M3.at<uchar>(10)=70;
M2.at<uchar>(10)=130;
M1.at<uchar>(10)=180;
//person
M3.at<uchar>(11)=220;
M2.at<uchar>(11)=20;
M1.at<uchar>(11)=60;
//rider
M3.at<uchar>(12)=255;
M2.at<uchar>(12)=0;
M1.at<uchar>(12)=100;
//car
M3.at<uchar>(13)=0;
M2.at<uchar>(13)=0;
M1.at<uchar>(13)=142;
//truck
M3.at<uchar>(14)=0;
M2.at<uchar>(14)=0;
M1.at<uchar>(14)=70;
//bus
M3.at<uchar>(15)=0;
M2.at<uchar>(15)=60;
M1.at<uchar>(15)=100;
//train
M3.at<uchar>(16)=0;
M2.at<uchar>(16)=0;
M1.at<uchar>(16)=192;
//motorbike
M3.at<uchar>(17)=0;
M2.at<uchar>(17)=0;
M1.at<uchar>(17)=230;
//bike
M3.at<uchar>(18)=119;
M2.at<uchar>(18)=11;
M1.at<uchar>(18)=32;
//void
for(int i=19;i<256;i++)
{
M1.at<uchar>(i)=0;
M2.at<uchar>(i)=0;
M3.at<uchar>(i)=0;
}
cv::Mat r1,r2,r3;
cv::LUT(adjMap,M1,r1);
cv::LUT(adjMap,M2,r2);
cv::LUT(adjMap,M3,r3);
std::vector<cv::Mat> planes;
planes.push_back(r1);
planes.push_back(r2);
planes.push_back(r3);
cv::Mat dst;
cv::merge(planes,dst);
return dst;
}
cv::Mat cityscapes_19_map(cv::Mat adjMap){
cv::Mat M1(1, 256, CV_8UC1), M2(1, 256, CV_8UC1), M3(1, 256, CV_8UC1);
//road
M3.at<uchar>(0)=128;
M2.at<uchar>(0)=64;
M1.at<uchar>(0)=128;
//sidewalk
M3.at<uchar>(1)=244;
M2.at<uchar>(1)=35;
M1.at<uchar>(1)=232;
//building
M3.at<uchar>(2)=70;
M2.at<uchar>(2)=70;
M1.at<uchar>(2)=70;
//wall
M3.at<uchar>(3)=102;
M2.at<uchar>(3)=102;
M1.at<uchar>(3)=156;
//fence
M3.at<uchar>(4)=190;
M2.at<uchar>(4)=153;
M1.at<uchar>(4)=153;
//pole
M3.at<uchar>(5)=153;
M2.at<uchar>(5)=153;
M1.at<uchar>(5)=153;
//traffic light
M3.at<uchar>(6)=250;
M2.at<uchar>(6)=170;
M1.at<uchar>(6)=30;
//traffic sign
M3.at<uchar>(7)=220;
M2.at<uchar>(7)=220;
M1.at<uchar>(7)=0;
//vegetation
M3.at<uchar>(8)=107;
M2.at<uchar>(8)=142;
M1.at<uchar>(8)=35;
//terrain
M3.at<uchar>(9)=152;
M2.at<uchar>(9)=251;
M1.at<uchar>(9)=152;
//sky
M3.at<uchar>(10)=70;
M2.at<uchar>(10)=130;
M1.at<uchar>(10)=180;
//person
M3.at<uchar>(11)=220;
M2.at<uchar>(11)=20;
M1.at<uchar>(11)=60;
//rider
M3.at<uchar>(12)=255;
M2.at<uchar>(12)=0;
M1.at<uchar>(12)=0;
//car
M3.at<uchar>(13)=0;
M2.at<uchar>(13)=0;
M1.at<uchar>(13)=142;
//truck
M3.at<uchar>(14)=0;
M2.at<uchar>(14)=0;
M1.at<uchar>(14)=70;
//bus
M3.at<uchar>(15)=0;
M2.at<uchar>(15)=60;
M1.at<uchar>(15)=100;
//train
M3.at<uchar>(16)=0;
M2.at<uchar>(16)=80;
M1.at<uchar>(16)=100;
//motorcycle
M3.at<uchar>(17)=0;
M2.at<uchar>(17)=0;
M1.at<uchar>(17)=230;
//bicycle
M3.at<uchar>(18)=119;
M2.at<uchar>(18)=11;
M1.at<uchar>(18)=32;
//void
for(int i=19;i<256;i++)
{
M1.at<uchar>(i)=0;
M2.at<uchar>(i)=0;
M3.at<uchar>(i)=0;
}
cv::Mat r1,r2,r3;
cv::LUT(adjMap,M1,r1);
cv::LUT(adjMap,M2,r2);
cv::LUT(adjMap,M3,r3);
std::vector<cv::Mat> planes;
planes.push_back(r1);
planes.push_back(r2);
planes.push_back(r3);
cv::Mat dst;
cv::merge(planes,dst);
return dst;
}
cv::Mat vizFloat2colorMap(cv::Mat map,double min, double max, int classes) {
if(min == 0 && max == 0)
cv::minMaxIdx(map, &min, &max);
double min;
double max;
cv::minMaxIdx(map, &min, &max);
cv::Mat adjMap;
// expand your range to 0..255. Similar to histEq();
map.convertTo(adjMap,CV_8UC1, 255 / (max-min), -min);
//return adjMap;
cv::Mat falseColorsMap;
applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_HOT);
switch (classes)
{
case 15:
map.convertTo(adjMap,CV_8UC1);
falseColorsMap = mapillary_15_map(adjMap);
break;
case 20:
map.convertTo(adjMap,CV_8UC1);
falseColorsMap = berkeley_20_map(adjMap);
break;
case 19:
map.convertTo(adjMap,CV_8UC1);
falseColorsMap = cityscapes_19_map(adjMap);
break;
default:
// expand your range to 0..255. Similar to histEq();
map.convertTo(adjMap,CV_8UC1, 255 / (max-min), -min);
applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_JET);
}
return falseColorsMap;
}
cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim) {
cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int img_h, int img_w, double min, double max, int classes) {
dnnType *data = nullptr;
// copy to CPU
@@ -38,14 +404,13 @@ cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim) {
cv::Mat grid = cv::Mat(gridSize, CV_8UC3, cv::Scalar(0));
for(int i=0; i<dim.c;i++) {
cv::Mat raw = vizFloat2colorMap(cv::Mat(cv::Size(dim.w, dim.h),CV_32FC1, data + dim.w*dim.h*i));
cv::Mat raw = vizFloat2colorMap(cv::Mat(cv::Size(dim.w, dim.h),CV_32FC1, data + dim.w*dim.h*i), min, max, classes);
int r = i / gridDim;
int c = i - r * gridDim;
raw.copyTo(grid.rowRange(r*dim.h, r*dim.h + dim.h).colRange(c*dim.w, c*dim.w + dim.w));
}
float ar = float(dim.w)/dim.h;
cv::Size vdim(ar*imgdim, imgdim);
cv::Size vdim(img_w, img_h);
cv::Mat viz;
cv::resize(grid, viz, vdim, 0, 0, 0);
@@ -59,7 +424,7 @@ cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim) {
cv::Mat vizLayer2Mat(tk::dnn::Network *net, int layer, int imgdim) {
if(layer >= net->num_layers)
FatalError("Could not viz layer\n");
return vizData2Mat(net->layers[layer]->dstData, net->layers[layer]->output_dim, imgdim);
return vizData2Mat(net->layers[layer]->dstData, net->layers[layer]->output_dim, imgdim, imgdim);
//cv::imwrite("viz/layer" + std::to_string(layer) + ".png", viz);
//cv::imshow("layer", viz);
+5
View File
@@ -15,6 +15,11 @@ Reshape::Reshape(Network *net, dataDim_t new_dim) : Layer(net) {
output_dim.w = new_dim.w;
output_dim.l = new_dim.l;
output_dim = new_dim;
if(input_dim.tot() != output_dim.tot())
FatalError("Reshape dimension mismatch");
}
Reshape::~Reshape() {
+39
View File
@@ -0,0 +1,39 @@
#include <iostream>
#include "Layer.h"
#include "kernels.h"
namespace tk { namespace dnn {
Resize::Resize(Network *net, int scale_c, int scale_h, int scale_w, bool fixed, ResizeMode_t mode) : Layer(net) {
this->mode = mode;
if(fixed){
output_dim.c = scale_c;
output_dim.h = scale_h;
output_dim.w = scale_w;
}
else{
output_dim.c *= scale_c;
output_dim.h *= scale_h;
output_dim.w *= scale_w;
}
checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
}
Resize::~Resize() {
checkCuda( cudaFree(dstData) );
}
dnnType* Resize::infer(dataDim_t &dim, dnnType* srcData) {
resizeForward(srcData, dstData, dim.n, dim.c, dim.h, dim.w,
output_dim.c, output_dim.h, output_dim.w);
dim = output_dim;
return dstData;
}
}}
+7 -6
View File
@@ -5,15 +5,16 @@
namespace tk { namespace dnn {
Shortcut::Shortcut(Network *net, Layer *backLayer) : Layer(net) {
Shortcut::Shortcut(Network *net, Layer *backLayer, bool mul) : Layer(net) {
this->backLayer = backLayer;
this->mul = mul;
checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
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 mismatch");
if( ( backLayer->output_dim.c != input_dim.c && mul ) ||
(( backLayer->output_dim.w != input_dim.w || backLayer->output_dim.h != input_dim.h ) && !mul ) )
FatalError("Shortcut dim missmatch");
}
Shortcut::~Shortcut() {
@@ -26,7 +27,7 @@ dnnType* Shortcut::infer(dataDim_t &dim, dnnType* srcData) {
dataDim_t bdim = this->backLayer->output_dim;
checkCuda(cudaMemcpy(dstData, srcData, dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice));
shortcutForward(this->backLayer->dstData, dstData, dim.n, dim.c, dim.h, dim.w, 1, bdim.n, bdim.c, bdim.h, bdim.w, 1);
shortcutForward(this->backLayer->dstData, dstData, dim.n, dim.c, dim.h, dim.w, 1, bdim.n, bdim.c, bdim.h, bdim.w, 1, mul);
//update data dimensions
dim = output_dim;
-1
View File
@@ -88,7 +88,6 @@ dnnType* Yolo::infer(dataDim_t &dim, dnnType* srcData) {
for (int b = 0; b < dim.n; ++b){
for(int n = 0; n < n_masks; ++n){
int index = entry_index(b, n*dim.w*dim.h, 0, classes, input_dim, output_dim);
std::cout<<"new_coords"<<new_coords<<std::endl;
if (new_coords == 1){
if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
}
+4 -4
View File
@@ -1,7 +1,7 @@
#include "kernels.h"
__global__
void activation_leaky(dnnType *input, dnnType *output, int size) {
void activation_leaky(dnnType *input, dnnType *output, int size, float slope) {
int i = blockDim.x*blockIdx.x + threadIdx.x;
@@ -9,7 +9,7 @@ void activation_leaky(dnnType *input, dnnType *output, int size) {
if (input[i]>0)
output[i] = input[i];
else
output[i] = 0.1f*input[i];
output[i] = slope*input[i];
}
}
@@ -17,12 +17,12 @@ void activation_leaky(dnnType *input, dnnType *output, int size) {
/**
ELU activation function
*/
void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream)
void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, float slope, cudaStream_t stream)
{
int blocks = (size+255)/256;
int threads = 256;
activation_leaky<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
activation_leaky<<<blocks, threads, 0, stream>>>(srcData, dstData, size, slope);
}
+19
View File
@@ -40,6 +40,25 @@ void sortAndTopKonDevice(dnnType *src_begin, int *idsrc, float *topk_scores, int
sortAndTopK_kernel<<<blocks, threads, 0>>>(src_begin, idsrc, topk_scores, topk_inds, topk_ys, topk_xs, size, K);
}
__global__
void maxElem_kernel(float *src_begin, float *dst_begin, const int n_classes, const int size){
int i = blockDim.x*blockIdx.x + threadIdx.x;
if (i > size)
return;
thrust::device_ptr<float> dPbeg ( &src_begin[i*n_classes] ) ;
thrust::device_ptr<float> dPend = dPbeg + n_classes;
thrust::device_ptr<float> result = thrust::max_element(thrust::device,dPbeg, dPend);
dst_begin[i] = result - dPbeg;
}
void maxElem(dnnType *src_begin, dnnType *dst_begin, const int c, const int h, const int w){
int blocks = (h*w)/32+1;
int threads = 32;
maxElem_kernel<<<blocks, threads, 0>>>(src_begin, dst_begin, c, h*w);
}
void topKxyclasses(int *ids_begin, int *ids_end, const int K, const int size, const int wh, int *clses, int *xs, int *ys){
thrust::transform(thrust::device, ids_begin, ids_end, thrust::make_constant_iterator(wh), clses, thrust::divides<int>());
thrust::transform(thrust::device, ids_begin, ids_end, thrust::make_constant_iterator(wh), ids_begin, thrust::modulus<int>());
+14 -27
View File
@@ -1,46 +1,33 @@
#include "kernels.h"
#include <stdio.h>
#define MIN(a,b) (((a)<(b))?(a):(b))
#define MAX(a,b) (((a)>(b))?(a):(b))
__global__ void resize_kernel( int i_N,float *x, int i_w, int i_h, int i_c,
__global__ void resize_kernel( int size,float *x, int i_w, int i_h, int i_c,
int o_w, int o_h, int o_c, int batch, float *out)
{
int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
if(i >= i_N) return;
int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
if(id >= size) return;
int out_index = i;
int out_w = i%o_w;
i = i/o_w;
int out_h = i%o_h;
i = i/o_h;
int out_c = i%o_c;
i = i/o_c;
int i = id % o_w;
id /= o_w;
int j = id % o_h;
id /= o_h;
int k = id % o_c;
id /= o_c;
int b = id % batch;
//copying last column/last row as padding
int in_index = ((i*i_c + MIN(out_c,i_c-1))*i_h + MIN(out_h,i_h-1))*i_w + MIN(out_w, i_w-1);
out[out_index] = x[in_index];
int out_index = i + o_w*(j + o_h*(k + o_c*b));
int add_index = i/(o_w/i_w) + i_w*(j/(o_h/i_h) + i_h*(k + i_c*b));
out[out_index] = x[add_index];
}
void resizeForward( dnnType* srcData, dnnType* dstData, int n, int i_c, int i_h, int i_w,
int o_c, int o_h, int o_w, cudaStream_t stream )
{
int i_size = n*i_c*i_h*i_w;
int o_size = n*o_c*o_h*o_w;
int blocks = (o_size+255)/256;
int threads = 256;
if(i_c == o_c && i_h == o_h && i_w == o_w )
{
checkCuda(cudaMemcpy(dstData, srcData, i_size*sizeof(dnnType), cudaMemcpyDeviceToDevice));
}
else
{
checkCuda(cudaMemset(dstData, 0, o_size*sizeof(dnnType)));
resize_kernel<<<blocks, threads, 0, stream>>>(o_size, srcData, i_w, i_h, i_c, o_w, o_h, o_c, n, dstData);
// printDeviceVector(i_size, srcData);
// printDeviceVector(o_size, dstData);
}
resize_kernel<<<blocks, threads, 0, stream>>>(o_size, srcData, i_w, i_h, i_c, o_w, o_h, o_c, n, dstData);
}
+48 -15
View File
@@ -21,27 +21,60 @@ __global__ void shortcut_kernel(int size, int minw, int minh, int minc, int stri
//out[out_index] += add[add_index];
}
__global__ void shortcut_mul_kernel(int size, int minw, int minh, int minc, int sample, int batch,
int w1, int h1, int c1, dnnType *mul,
int w2, int h2, int c2, float s1, float s2, dnnType *out)
{
int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
if (id >= size) return;
int i = id % minw;
id /= minw;
int j = id % minh;
id /= minh;
int k = id % minc;
id /= minc;
int b = id % batch;
int out_index = i*sample + w1*(j*sample + h1*(k + c1*b));
out[out_index] = out[out_index] * mul[k + c2*b];
}
void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1, int w1, int s1,
int n2, int c2, int h2, int w2, int s2,
cudaStream_t stream)
bool mul, cudaStream_t stream)
{
assert(n1 == n2);
int batch = n1;
int minw = (w1 < w2) ? w1 : w2;
int minh = (h1 < h2) ? h1 : h2;
int minc = (c1 < c2) ? c1 : c2;
if(!mul){
int minw = (w1 < w2) ? w1 : w2;
int minh = (h1 < h2) ? h1 : h2;
int minc = (c1 < c2) ? c1 : c2;
int stride = w1/w2;
int sample = w2/w1;
assert(stride == h1/h2);
assert(sample == h2/h1);
if(stride < 1) stride = 1;
if(sample < 1) sample = 1;
int stride = w1/w2;
int sample = w2/w1;
assert(stride == h1/h2);
assert(sample == h2/h1);
if(stride < 1) stride = 1;
if(sample < 1) sample = 1;
int size = batch * minw * minh * minc;
int blocks = (size+255)/256;
int threads = 256;
shortcut_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, stride, sample, batch,
w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);
}
else{
int minw = w1;
int minh = h1;
int minc = c1;
int sample = 1;
int size = batch * minw * minh * minc;
int blocks = (size+255)/256;
int threads = 256;
shortcut_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, stride, sample, batch,
w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);
int size = batch * minw * minh * minc;
int blocks = (size+255)/256;
int threads = 256;
shortcut_mul_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, sample, batch,
w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);
}
}
+10 -7
View File
@@ -92,7 +92,7 @@ void printDeviceVector(int size, dnnType* vec_d, bool device){
delete [] vec;
}
int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device, int limit) {
int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device, int limit, bool verbose) {
dnnType *data_h, *correct_h;
const float eps = 0.02f;
@@ -111,6 +111,7 @@ int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device, int
}
int diffs = 0;
for(int i=0; i<size; i++) {
// data_h[i] = data_h[i]*1e-2;
if(data_h[i] != data_h[i] || correct_h[i] != correct_h[i] || //nan control
fabs(data_h[i] - correct_h[i]) > eps) {
diffs += 1;
@@ -126,13 +127,15 @@ int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device, int
delete [] correct_h;
}
std::cout<<" | ";
if(diffs == 0)
std::cout<<COL_GREENB<<"OK";
else
std::cout<<COL_REDB<<"Wrongs: "<<diffs;
if(verbose){
std::cout<<" | ";
if(diffs == 0)
std::cout<<COL_GREENB<<"OK";
else
std::cout<<COL_REDB<<"Wrongs: "<<diffs;
std::cout<<COL_END<<" ~"<<eps<<"\n";
std::cout<<COL_END<<" ~"<<eps<<"\n";
}
return diffs;
}