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

merge with master

Browse Source 
Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>
This commit is contained in:
Micaela Verucchi
2021-07-20 18:38:46 +02:00
parent 4b3731928c 946b1fe5f5
commit 318dfe4b55
68 changed files with 5577 additions and 276 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/Activation.cpp
+9 -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,12 +47,15 @@ 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());
}
else if(act_mode == ACTIVATION_LOGISTIC) {
activationLOGISTICForward(srcData, dstData, dim.tot());
} else {
dnnType alpha = dnnType(1);
dnnType beta = dnnType(0);
src/DarknetParser.cpp
+13 -2
View File
@@ -37,7 +37,10 @@ namespace tk { namespace dnn {
std::string name,value;
if(!divideNameAndValue(line, name, value))
return false;
if(name.find("width") != std::string::npos)
if(name.find("new_coords") != std::string::npos)
fields.new_coords = std::stoi(value);
else if(name.find("width") != std::string::npos)
fields.width = std::stoi(value);
else if(name.find("height") != std::string::npos)
fields.height = std::stoi(value);
@@ -79,6 +82,13 @@ namespace tk { namespace dnn {
fields.group_id = std::stoi(value);
else if(name.find("scale_x_y") != std::string::npos)
fields.scale_xy = std::stof(value);
else if(name.find("beta_nms") != std::string::npos)
fields.nms_thresh = std::stof(value);
else if(name.find("nms_kind") != std::string::npos){
if(value == "greedynms") fields.nms_kind = 0;
else if(value == "diounms") fields.nms_kind = 1;
else std::cout<<"Not supported nms_kind "<<value<<", setting to greedynms"<<std::endl;
}
else if(name.find("from") != std::string::npos)
fields.layers.push_back(std::stof(value));
else if(name.find("mask") != std::string::npos){
@@ -161,7 +171,7 @@ namespace tk { namespace dnn {
} else if(f.type == "yolo") {
std::string wgs = wgs_path + "/g" + std::to_string(netLayers.size()) + ".bin";
//printf("%d %d %s %d %f\n", f.classes, f.num/f.n_mask, wgs.c_str(), f.n_mask, f.scale_xy);
tk::dnn::Yolo *l = new tk::dnn::Yolo(net, f.classes, f.num/f.n_mask, wgs, f.n_mask, f.scale_xy);
tk::dnn::Yolo *l = new tk::dnn::Yolo(net, f.classes, f.num/f.n_mask, wgs, f.n_mask, f.scale_xy, f.nms_thresh, (tk::dnn::Yolo::nmsKind_t) f.nms_kind, f.new_coords);
if(names.size() != f.classes)
FatalError("Mismatch between number of classes and names");
l->classesNames = names;
@@ -177,6 +187,7 @@ namespace tk { namespace dnn {
if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU);
else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY;
else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH;
else if(f.activation == "logistic") act = tk::dnn::ACTIVATION_LOGISTIC;
else { FatalError("activation not supported: " + f.activation); }
netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act);
};
src/LSTM.cpp
+9 -4
View File
@@ -87,17 +87,22 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig
checkCUDNN(cudnnCreateRNNDescriptor(&rnnDesc));
#if CUDNN_MAJOR > 7
checkCUDNN(cudnnSetRNNDescriptor_v6(net->cudnnHandle,
checkCUDNN(cudnnSetRNNDescriptor_v6(net->cudnnHandle,rnnDesc, stateSize, numLayers, dropoutDesc,
cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT,
//(bidirectional ? cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL : cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL),
cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL,
cudnnRNNMode_t::CUDNN_LSTM,
cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD,
net->dataType));
#else
checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle,
#endif
rnnDesc, stateSize, numLayers, dropoutDesc,
checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle,rnnDesc, stateSize, numLayers, dropoutDesc,
cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT,
//(bidirectional ? cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL : cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL),
cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL,
cudnnRNNMode_t::CUDNN_LSTM,
cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD,
net->dataType));
#endif
// Get temp space sizes
src/LayerWgs.cpp
+1
View File
@@ -32,6 +32,7 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs,
this->batchnorm = batchnorm;
if(batchnorm) {
readBinaryFile(weights_path.c_str(), outputs, &scales_h, &scales_d, seek);
seek += outputs;
readBinaryFile(weights_path.c_str(), outputs, &mean_h, &mean_d, seek);
src/NetworkRT.cpp
+113 -40
View File
@@ -140,6 +140,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT);
#else
engineRT = builderRT->buildCudaEngine(*networkRT);
//engineRT = std::shared_ptr<nvinfer1::ICudaEngine>(builderRT->buildCudaEngine(*networkRT));
#endif
if(engineRT == nullptr)
FatalError("cloud not build cuda engine")
@@ -226,7 +227,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
return convert_layer(input, (Conv2d*) l);
if(type == LAYER_POOLING)
return convert_layer(input, (Pooling*) l);
if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH)
if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_LOGISTIC)
return convert_layer(input, (Activation*) l);
if(type == LAYER_SOFTMAX)
return convert_layer(input, (Softmax*) l);
@@ -236,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)
@@ -389,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
@@ -421,6 +424,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
checkNULL(lRT);
return lRT;
}
else if(l->act_mode == ACTIVATION_LOGISTIC) {
IPlugin *plugin = new ActivationLogisticRT();
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
return lRT;
}
else {
FatalError("this Activation mode is not yet implemented");
return NULL;
@@ -472,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";
@@ -506,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);
@@ -515,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;
@@ -529,7 +548,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Yolo *l) {
//std::cout<<"convert Yolo\n";
//std::cout<<"New plugin YOLO\n";
IPlugin *plugin = new YoloRT(l->classes, l->num, l, l->n_masks, l->scaleXY);
IPlugin *plugin = new YoloRT(l->classes, l->num, l, l->n_masks, l->scaleXY, l->nms_thresh, l->nsm_kind, l->new_coords);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
return lRT;
@@ -561,7 +580,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {
IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin);
checkNULL(lRT);
lRT->setName( ("Deformable" + std::to_string(l->id)).c_str() );
delete(inputs);
delete[](inputs);
// batchnorm
void *bias_b, *power_b, *mean_b, *variance_b, *scales_b;
if(dtRT == DataType::kHALF) {
@@ -638,43 +657,61 @@ bool NetworkRT::deserialize(const char *filename) {
IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) {
const char * buf = reinterpret_cast<const char*>(serialData);
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
std::string name(layerName);
//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;
}
if(name.find("ActivationMish") == 0) {
ActivationMishRT *a = new ActivationMishRT();
a->size = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return a;
}
if(name.find("ActivationLogistic") == 0) {
ActivationLogisticRT *a = new ActivationLogisticRT();
a->size = readBUF<int>(buf);
return a;
}
if(name.find("ActivationLogistic") == 0) {
ActivationLogisticRT *a = new ActivationLogisticRT();
a->size = readBUF<int>(buf);
return a;
}
if(name.find("ActivationCReLU") == 0) {
ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF<float>(buf));
float activationReluTemp = readBUF<float>(buf);
ActivationReLUCeiling* a = new ActivationReLUCeiling(activationReluTemp);
a->size = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return a;
}
if(name.find("Region") == 0) {
RegionRT *r = new RegionRT(readBUF<int>(buf), //classes
readBUF<int>(buf), //coords
readBUF<int>(buf)); //num
int classesTemp = readBUF<int>(buf);
int coordsTemp = readBUF<int>(buf);
int numTemp = readBUF<int>(buf);
RegionRT* r = new RegionRT(classesTemp, coordsTemp, numTemp);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
if(name.find("Reorg") == 0) {
ReorgRT *r = new ReorgRT(readBUF<int>(buf)); //stride
int strideTemp = readBUF<int>(buf);
ReorgRT *r = new ReorgRT(strideTemp);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
@@ -685,32 +722,39 @@ 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);
return r;
assert(buf == bufCheck + serialLength);
}
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
int cTemp = readBUF<int>(buf);
int hTemp = readBUF<int>(buf);
int wTemp = readBUF<int>(buf);
int nTemp = readBUF<int>(buf);
int strideHTemp = readBUF<int>(buf);
int strideWTemp = readBUF<int>(buf);
int winSizeTemp = readBUF<int>(buf);
int paddingTemp = readBUF<int>(buf);
MaxPoolFixedSizeRT* r = new MaxPoolFixedSizeRT(cTemp, hTemp, wTemp, nTemp, strideHTemp, strideWTemp, winSizeTemp, paddingTemp);
assert(buf == bufCheck + serialLength);
return r;
}
if(name.find("Resize") == 0) {
ResizeLayerRT *r = new ResizeLayerRT(readBUF<int>(buf), //o_c
readBUF<int>(buf), //o_h
readBUF<int>(buf)); //o_w
int o_cTemp = readBUF<int>(buf);
int o_hTemp = readBUF<int>(buf);
int o_wTemp = readBUF<int>(buf);
ResizeLayerRT* r = new ResizeLayerRT(o_cTemp, o_hTemp, o_wTemp);
r->i_c = readBUF<int>(buf);
r->i_h = readBUF<int>(buf);
r->i_w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
@@ -721,6 +765,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
r->w = readBUF<int>(buf);
r->rows = readBUF<int>(buf);
r->cols = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
@@ -732,19 +777,28 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
new_dim.h = readBUF<int>(buf);
new_dim.w = readBUF<int>(buf);
ReshapeRT *r = new ReshapeRT(new_dim);
assert(buf == bufCheck + serialLength);
return r;
}
if(name.find("Yolo") == 0) {
YoloRT *r = new YoloRT(readBUF<int>(buf), //classes
readBUF<int>(buf), //num
nullptr,
readBUF<int>(buf)); //n_masks
int classes_temp = readBUF<int>(buf);
int num_temp = readBUF<int>(buf);
int n_masks_temp = readBUF<int>(buf);
float scale_xy_temp = readBUF<float>(buf);
float nms_thresh_temp = readBUF<float>(buf);
int nms_kind_temp = readBUF<int>(buf);
int new_coords_temp = readBUF<int>(buf);
YoloRT *r = new YoloRT(classes_temp,num_temp,nullptr,n_masks_temp,scale_xy_temp,nms_thresh_temp,nms_kind_temp,new_coords_temp);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
r->scaleXY = readBUF<float>(buf);
for(int i=0; i<r->n_masks; i++)
r->mask[i] = readBUF<dnnType>(buf);
for(int i=0; i<r->n_masks*2*r->num; i++)
@@ -758,36 +812,54 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
tmp[j] = readBUF<char>(buf);
r->classesNames[i] = std::string(tmp);
}
assert(buf == bufCheck + serialLength);
yolos[n_yolos++] = r;
return r;
}
if(name.find("Upsample") == 0) {
UpsampleRT *r = new UpsampleRT(readBUF<int>(buf)); //stride
int strideTemp = readBUF<int>(buf);
UpsampleRT* r = new UpsampleRT(strideTemp);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
if(name.find("Route") == 0) {
RouteRT *r = new RouteRT(readBUF<int>(buf),readBUF<int>(buf));
int groupsTemp = readBUF<int>(buf);
int group_idTemp = readBUF<int>(buf);
RouteRT* r = new RouteRT(groupsTemp, group_idTemp);
r->in = readBUF<int>(buf);
for(int i=0; i<RouteRT::MAX_INPUTS; i++)
r->c_in[i] = readBUF<int>(buf);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
if(name.find("Deformable") == 0) {
DeformableConvRT *r = new DeformableConvRT(readBUF<int>(buf), readBUF<int>(buf), readBUF<int>(buf),
readBUF<int>(buf), readBUF<int>(buf), readBUF<int>(buf),
readBUF<int>(buf), readBUF<int>(buf),
readBUF<int>(buf),readBUF<int>(buf),readBUF<int>(buf),readBUF<int>(buf),
readBUF<int>(buf),readBUF<int>(buf),readBUF<int>(buf),readBUF<int>(buf),
nullptr);
int chuck_dimTemp = readBUF<int>(buf);
int khTemp = readBUF<int>(buf);
int kwTemp = readBUF<int>(buf);
int shTemp = readBUF<int>(buf);
int swTemp = readBUF<int>(buf);
int phTemp = readBUF<int>(buf);
int pwTemp = readBUF<int>(buf);
int deformableGroupTemp = readBUF<int>(buf);
int i_nTemp = readBUF<int>(buf);
int i_cTemp = readBUF<int>(buf);
int i_hTemp = readBUF<int>(buf);
int i_wTemp = readBUF<int>(buf);
int o_nTemp = readBUF<int>(buf);
int o_cTemp = readBUF<int>(buf);
int o_hTemp = readBUF<int>(buf);
int o_wTemp = readBUF<int>(buf);
DeformableConvRT* r = new DeformableConvRT(chuck_dimTemp, khTemp, kwTemp, shTemp, swTemp, phTemp, pwTemp, deformableGroupTemp, i_nTemp, i_cTemp, i_hTemp, i_wTemp, o_nTemp, o_cTemp, o_hTemp, o_wTemp, nullptr);
dnnType *aus = new dnnType[r->chunk_dim*2];
for(int i=0; i<r->chunk_dim*2; i++)
aus[i] = readBUF<dnnType>(buf);
@@ -818,6 +890,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->ones_d2, aus, sizeof(dnnType)*r->dim_ones, cudaMemcpyHostToDevice) );
free(aus);
assert(buf == bufCheck + serialLength);
return r;
}
src/NetworkViz.cpp
+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);
src/Reshape.cpp
+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() {
src/Resize.cpp
+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;
}
}}
src/Shortcut.cpp
+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;
src/Yolo.cpp
+60 -18
View File
@@ -9,9 +9,10 @@
#include "Layer.h"
#include "kernels.h"
namespace tk { namespace dnn {
Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights, int n_masks, float scale_xy) :
Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights, int n_masks, float scale_xy, double nms_thresh, nmsKind_t nsm_kind, int new_coords) :
Layer(net) {
this->final = true;
@@ -19,6 +20,9 @@ Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights, int n_
this->num = num;
this->n_masks = n_masks;
this->scaleXY = scale_xy;
this->nms_thresh = nms_thresh;
this->nsm_kind = nsm_kind;
this->new_coords = new_coords;
// load anchors
if(fname_weights != "") {
@@ -59,12 +63,21 @@ int entry_index(int batch, int location, int entry,
entry*input_dim.w*input_dim.h + loc;
}
Yolo::box get_yolo_box(float *x, float *biases, int n, int index, int i, int j, int lw, int lh, int w, int h, int stride) {
Yolo::box get_yolo_box(float *x, float *biases, int n, int index, int i, int j, int lw, int lh, int w, int h, int stride, int new_coords) {
Yolo::box b;
b.x = (i + x[index + 0*stride]) / lw;
b.y = (j + x[index + 1*stride]) / lh;
b.w = exp(x[index + 2*stride]) * biases[2*n] / w;
b.h = exp(x[index + 3*stride]) * biases[2*n+1] / h;
if(new_coords == 0){
b.x = (i + x[index + 0*stride]) / lw;
b.y = (j + x[index + 1*stride]) / lh;
b.w = exp(x[index + 2*stride]) * biases[2*n] / w;
b.h = exp(x[index + 3*stride]) * biases[2*n+1] / h;
}
else{
b.x = (i + x[index + 0 * stride] ) / lw;
b.y = (j + x[index + 1 * stride] ) / lh;
b.w = x[index + 2 * stride] * x[index + 2 * stride] * 4 * biases[2 * n] / w;
b.h = x[index + 3 * stride] * x[index + 3 * stride] * 4 * biases[2 * n + 1] / h;
}
return b;
}
@@ -75,12 +88,16 @@ 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);
activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h);
if (new_coords == 1){
if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
}
else{
activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h);
if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim);
activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h);
if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim);
activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h);
}
}
}
@@ -116,7 +133,7 @@ void correct_yolo_boxes(Yolo::detection *dets, int n, int w, int h, int netw, in
}
}
int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh) {
int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh, int new_coords) {
if(predictions == nullptr)
predictions = new dnnType[output_dim.tot()];
@@ -140,7 +157,7 @@ int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int net
if(objectness <= thresh) continue;
int box_index = entry_index(0, n*lw*lh + i, 0, classes, input_dim, output_dim);
dets[count].bbox = get_yolo_box(predictions, bias_h, mask_h[n], box_index, col, row, lw, lh, netw, neth, lw*lh);
dets[count].bbox = get_yolo_box(predictions, bias_h, mask_h[n], box_index, col, row, lw, lh, netw, neth, lw*lh, new_coords);
dets[count].objectness = objectness;
dets[count].classes = classes;
for(j = 0; j < classes; ++j){
@@ -193,6 +210,32 @@ float yolo_box_iou(Yolo::box a, Yolo::box b)
return yolo_box_intersection(a, b)/yolo_box_union(a, b);
}
void box_c(const Yolo::box a, const Yolo::box b, float& top, float& bot, float& left, float& right) {
top = (std::min)(a.y - a.h / 2, b.y - b.h / 2);
bot = (std::max)(a.y + a.h / 2, b.y + b.h / 2);
left = (std::min)(a.x - a.w / 2, b.x - b.w / 2);
right = (std::max)(a.x + a.w / 2, b.x + b.w / 2);
}
// https://github.com/Zzh-tju/DIoU-darknet
// https://arxiv.org/abs/1911.08287
float yolo_box_diou(const Yolo::box a, const Yolo::box b, const float nms_thresh=0.6)
{
float top, bot, left, right;
box_c(a, b, top, bot, left, right);
float w = right - left;
float h = bot - top;
float c = w * w + h * h;
float iou = yolo_box_iou(a, b);
if (c == 0)
return iou;
float d = (a.x - b.x) * (a.x - b.x) + (a.y - b.y) * (a.y - b.y);
float u = pow(d / c, nms_thresh);
float diou_term = u;
return iou - diou_term;
}
int yolo_nms_comparator(const void *pa, const void *pb)
{
Yolo::detection a = *(Yolo::detection *)pa;
@@ -219,8 +262,7 @@ Yolo::detection *Yolo::allocateDetections(int nboxes, int classes) {
return dets;
}
void Yolo::mergeDetections(Yolo::detection *dets, int ndets, int classes) {
double nms_thresh = 0.45;
void Yolo::mergeDetections(Yolo::detection *dets, int ndets, int classes, double nms_thresh, nmsKind_t nsm_kind) {
int total = ndets;
int i, j, k;
@@ -246,13 +288,13 @@ void Yolo::mergeDetections(Yolo::detection *dets, int ndets, int classes) {
box a = dets[i].bbox;
for(j = i+1; j < total; ++j){
box b = dets[j].bbox;
if (yolo_box_iou(a, b) > nms_thresh){
if (nsm_kind == GREEDY_NMS && yolo_box_iou(a, b) > nms_thresh)
dets[j].prob[k] = 0;
else if (nsm_kind == DIOU_NMS && yolo_box_diou(a, b, nms_thresh) > nms_thresh)
dets[j].prob[k] = 0;
}
}
}
}
}
}}
src/Yolo3Detection.cpp
+9 -5
View File
@@ -32,6 +32,9 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c
memcpy(yolo[i]->bias_h, yRT->bias, sizeof(dnnType)*num*nMasks*2);
yolo[i]->input_dim = yolo[i]->output_dim = tk::dnn::dataDim_t(1, yRT->c, yRT->h, yRT->w);
yolo[i]->classesNames = yRT->classesNames;
yolo[i]->nms_thresh = yRT->nms_thresh;
yolo[i]->nsm_kind = (tk::dnn::Yolo::nmsKind_t) yRT->nms_kind;
yolo[i]->new_coords = yRT->new_coords;
}
dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes);
@@ -91,9 +94,10 @@ void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
void Yolo3Detection::postprocess(const int bi, const bool mAP){
//get yolo outputs
dnnType *rt_out[netRT->pluginFactory->n_yolos];
for(int i=0; i<netRT->pluginFactory->n_yolos; i++)
rt_out[i] = (dnnType*)netRT->buffersRT[i+1] + netRT->buffersDIM[i+1].tot()*bi;
std::vector<float *> rt_out;
//dnnType *rt_out[netRT->pluginFactory->n_yolos];
for(int i=0; i<netRT->pluginFactory->n_yolos; i++)
rt_out.push_back((dnnType*)netRT->buffersRT[i+1] + netRT->buffersDIM[i+1].tot()*bi);
float x_ratio = float(originalSize[bi].width) / float(netRT->input_dim.w);
float y_ratio = float(originalSize[bi].height) / float(netRT->input_dim.h);
@@ -102,9 +106,9 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){
nDets = 0;
for(int i=0; i<netRT->pluginFactory->n_yolos; i++) {
yolo[i]->dstData = rt_out[i];
yolo[i]->computeDetections(dets, nDets, netRT->input_dim.w, netRT->input_dim.h, confThreshold);
yolo[i]->computeDetections(dets, nDets, netRT->input_dim.w, netRT->input_dim.h, confThreshold, yolo[i]->new_coords);
}
tk::dnn::Yolo::mergeDetections(dets, nDets, classes);
tk::dnn::Yolo::mergeDetections(dets, nDets, classes, yolo[0]->nms_thresh, yolo[0]->nsm_kind);
// fill detected
detected.clear();
src/kernels/activation_leaky.cu
+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);
}
src/kernels/deformable_conv.cu
+1 -1
View File
@@ -18,7 +18,7 @@ inline int GET_BLOCKS(const int N)
}
__device__ float dmcn_im2col_bilinear(const float *bottom_data, const int data_width,
__device__ __host__ float dmcn_im2col_bilinear(const float *bottom_data, const int data_width,
const int height, const int width, float h, float w) {
int h_low = floor(h);
int w_low = floor(w);
src/kernels/postprocessing.cu
+19
View File
@@ -34,6 +34,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>());
src/kernels/resize.cu
+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);
}
src/kernels/shortcut.cu
+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);
}
}
src/utils.cpp
+16 -2
View File
@@ -23,14 +23,23 @@ bool fileExist(const char *fname) {
void downloadWeightsifDoNotExist(const std::string& input_bin, const std::string& test_folder, const std::string& weights_url){
if(!fileExist(input_bin.c_str())){
std::string mkdir_cmd = "mkdir " + test_folder;
std::string wget_cmd = "wget " + weights_url + " -O " + test_folder + "/weights.zip";
std::string wget_cmd = "curl " + weights_url + " --output " + test_folder + "/weights.zip";
#ifdef __linux__
std::string unzip_cmd = "unzip " + test_folder + "/weights.zip -d" + test_folder;
std::string rm_cmd = "rm " + test_folder + "/weights.zip";
#elif _WIN32
std::string unzip_cmd = "7z x " + test_folder + "/weights.zip -o" + test_folder;
#endif
int err = 0;
err = system(mkdir_cmd.c_str());
err = system(wget_cmd.c_str());
err = system(unzip_cmd.c_str());
#ifdef __linux__
err = system(rm_cmd.c_str());
#endif
}
}
@@ -102,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;
@@ -193,8 +203,12 @@ void getMemUsage(double& vm_usage_kb, double& resident_set_kb){
>> O >> itrealvalue >> starttime >> vsize >> rss;
stat_stream.close();
#ifdef __linux__
long page_size_kb = sysconf(_SC_PAGE_SIZE) / 1024; // in case x86-64 is configured to use 2MB pages
#elif _WIN32
long page_size_kb = 4096/1024;
#endif
vm_usage_kb = vsize / 1024.0;
resident_set_kb = rss * page_size_kb;
}