diff --git a/CMakeLists.txt b/CMakeLists.txt index b430195..4528544 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -22,7 +22,7 @@ set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32) find_package(CUDNN REQUIRED) # compile -file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/*.cu") +file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/sorting.cu") cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS}) cuda_add_library(kernels SHARED ${tkdnn_CUSRC}) @@ -44,7 +44,7 @@ set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDNN_LIBRAR set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -std=c++11") include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS} ${NVINFER_INCLUDES}) add_library(tkDNN SHARED ${tkdnn_SRC}) -target_link_libraries(tkDNN ${tkdnn_LIBS}) +target_link_libraries(tkDNN ${tkdnn_LIBS} nvinfer_plugin) #static #add_library(tkDNN_static STATIC ${tkdnn_SRC}) @@ -94,6 +94,8 @@ target_link_libraries(test_yolo3_berkeley tkDNN) add_executable(test_yolo3_flir tests/yolo3_flir/yolo3_flir.cpp) target_link_libraries(test_yolo3_flir tkDNN) +add_executable(test_mobilenetv2ssd tests/mobilenetv2ssd/mobilenetv2ssd.cpp) +target_link_libraries(test_mobilenetv2ssd tkDNN) add_executable(test_resnet101 tests/resnet101/resnet101.cpp) target_link_libraries(test_resnet101 tkDNN) diff --git a/demo/config.yaml b/demo/config.yaml index 3577abf..3c2d188 100644 --- a/demo/config.yaml +++ b/demo/config.yaml @@ -3,5 +3,5 @@ map_points : 101 #number of recall points (0 for all, 101 for COCO, 11 Pascal map_levels : 10 #number of IoU step for the AP map_step : 0.05 #step of IoU IoU_thresh : 0.5 #starting IoU threshold -conf_thresh : 0.3 #threshold on the condifence of the bbox +conf_thresh : 0.0 #threshold on the condifence of the bbox verbose : false #print on screen information \ No newline at end of file diff --git a/demo/demo/map.cpp b/demo/demo/map.cpp index 13636fb..88f879f 100644 --- a/demo/demo/map.cpp +++ b/demo/demo/map.cpp @@ -34,7 +34,7 @@ int main(int argc, char *argv[]) char * labels_path = "../demo/COCO_val2017/all_labels.txt"; bool show = false; bool write_dets = false; - int n_images = 1000; + int n_images = 5000; if(argc > 1) net = argv[1]; diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 8950d53..7ef9710 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -14,6 +14,8 @@ enum layerType_t { LAYER_DECONV2D, LAYER_DEFORMCONV2D, LAYER_ACTIVATION, + LAYER_ACTIVATION_CRELU, + LAYER_ACTIVATION_LEAKY, LAYER_FLATTEN, LAYER_MULADD, LAYER_POOLING, @@ -52,22 +54,24 @@ public: std::string getLayerName() { layerType_t type = getLayerType(); switch(type) { - case LAYER_DENSE: return "Dense"; - case LAYER_CONV2D: return "Conv2d"; - case LAYER_DECONV2D: return "DeConv2d"; - case LAYER_DEFORMCONV2D:return "DeformConv2d"; - case LAYER_ACTIVATION: return "Activation"; - case LAYER_FLATTEN: return "Flatten"; - case LAYER_MULADD: return "MulAdd"; - case LAYER_POOLING: return "Pooling"; - case LAYER_SOFTMAX: return "Softmax"; - case LAYER_ROUTE: return "Route"; - case LAYER_REORG: return "Reorg"; - case LAYER_SHORTCUT: return "Shortcut"; - case LAYER_UPSAMPLE: return "Upsample"; - case LAYER_REGION: return "Region"; - case LAYER_YOLO: return "Yolo"; - default: return "unknown"; + case LAYER_DENSE: return "Dense"; + case LAYER_CONV2D: return "Conv2d"; + case LAYER_DECONV2D: return "DeConv2d"; + case LAYER_DEFORMCONV2D: return "DeformConv2d"; + case LAYER_ACTIVATION: return "Activation"; + case LAYER_ACTIVATION_CRELU: return "ActivationCReLU"; + case LAYER_ACTIVATION_LEAKY: return "ActivationLeaky"; + case LAYER_FLATTEN: return "Flatten"; + case LAYER_MULADD: return "MulAdd"; + case LAYER_POOLING: return "Pooling"; + case LAYER_SOFTMAX: return "Softmax"; + case LAYER_ROUTE: return "Route"; + case LAYER_REORG: return "Reorg"; + case LAYER_SHORTCUT: return "Shortcut"; + case LAYER_UPSAMPLE: return "Upsample"; + case LAYER_REGION: return "Region"; + case LAYER_YOLO: return "Yolo"; + default: return "unknown"; } } @@ -145,10 +149,18 @@ class Activation : public Layer { public: int act_mode; + float ceiling; - Activation(Network *net, int act_mode); + Activation(Network *net, int act_mode, const float ceiling=0.0); virtual ~Activation(); - virtual layerType_t getLayerType() { return LAYER_ACTIVATION; }; + virtual layerType_t getLayerType() { + if(act_mode == CUDNN_ACTIVATION_CLIPPED_RELU) + return LAYER_ACTIVATION_CRELU; + else if (act_mode == ACTIVATION_LEAKY) + return LAYER_ACTIVATION_LEAKY; + else + return LAYER_ACTIVATION; + }; virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); @@ -165,14 +177,14 @@ class Conv2d : public LayerWgs { public: Conv2d( Network *net, int out_ch, int kernelH, int kernelW, int strideH, int strideW, int paddingH, int paddingW, - std::string fname_weights, bool batchnorm = false, bool deConv = false, bool final = false, int groups = 1); + std::string fname_weights, bool batchnorm = false, bool deConv = false, bool final = false, int groups = 1, bool additional_bias=false); virtual ~Conv2d(); virtual layerType_t getLayerType() { return LAYER_CONV2D; }; virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); int kernelH, kernelW, strideH, strideW, paddingH, paddingW; - bool deConv; + bool deConv, additional_bias; int groups; protected: diff --git a/include/tkDNN/Network.h b/include/tkDNN/Network.h index b234f71..26f9f53 100644 --- a/include/tkDNN/Network.h +++ b/include/tkDNN/Network.h @@ -59,7 +59,7 @@ public: dataDim_t input_dim; dataDim_t getOutputDim(); - bool fp16, dla; + bool fp16, dla, int8; bool dontLoadWeights; }; diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 82d082a..e8a2b42 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -24,6 +24,7 @@ template T readBUF(const char*& buffer) using namespace nvinfer1; #include "pluginsRT/ActivationLeakyRT.h" +#include "pluginsRT/ActivationReLUCeilingRT.h" #include "pluginsRT/ReorgRT.h" #include "pluginsRT/RegionRT.h" //#include "pluginsRT/RouteRT.h" diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index f514ffc..d3a5d0d 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -5,6 +5,7 @@ void activationELUForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0)); void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0)); +void activationReLUCeilingForward(dnnType* srcData, dnnType* dstData, int size, const float ceiling, cudaStream_t stream= cudaStream_t(0)); void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0)); void activationSIGMOIDForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0)); diff --git a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h new file mode 100644 index 0000000..0074520 --- /dev/null +++ b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h @@ -0,0 +1,62 @@ +#include +#include "../kernels.h" + +class ActivationReLUCeiling : public IPlugin { + +public: + ActivationReLUCeiling(const float ceiling) { + this->ceiling = ceiling; + } + + ~ActivationReLUCeiling(){ + + } + + int getNbOutputs() const override { + return 1; + } + + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return inputs[0]; + } + + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + size = 1; + for(int i=0; i(inputs[0]), + reinterpret_cast(outputs[0]), size, ceiling, stream); + return 0; + } + + + virtual size_t getSerializationSize() override { + return 1*sizeof(int) + 1*sizeof(float); + } + + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer); + tk::dnn::writeBUF(buf, ceiling); + tk::dnn::writeBUF(buf, size); + + } + + int size; + float ceiling; +}; diff --git a/src/Activation.cpp b/src/Activation.cpp index 1f6f0e7..d619331 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -5,10 +5,11 @@ namespace tk { namespace dnn { -Activation::Activation(Network *net, int act_mode) : +Activation::Activation(Network *net, int act_mode, const float ceiling) : Layer(net) { this->act_mode = act_mode; + this->ceiling = ceiling; checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) ); if(int(act_mode) < 100) { @@ -31,7 +32,7 @@ Activation::Activation(Network *net, int act_mode) : checkCUDNN( cudnnSetActivationDescriptor(activDesc, (cudnnActivationMode_t) act_mode, CUDNN_PROPAGATE_NAN, - 0.0) ); + ceiling) ); } } diff --git a/src/CenternetDetection.cpp b/src/CenternetDetection.cpp index c707c6c..65b9e99 100644 --- a/src/CenternetDetection.cpp +++ b/src/CenternetDetection.cpp @@ -294,8 +294,6 @@ void CenternetDetection::update(cv::Mat &imageORIG) { netRT->infer(dim2, input_d); TIMER_STOP dim2.print(); - - stats.push_back(t_ns); } // checkResult(dim2.tot(), input_h, input); step_t = std::chrono::steady_clock::now(); @@ -456,5 +454,6 @@ void CenternetDetection::update(cv::Mat &imageORIG) { std::cout<<"TOTAL: \n"; TIMER_STOP + stats.push_back(t_ns); } }} \ No newline at end of file diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index 18e907d..458cd43 100644 --- a/src/Conv2d.cpp +++ b/src/Conv2d.cpp @@ -100,7 +100,7 @@ void Conv2d::inferCUDNN(dnnType* srcData, bool back) { &beta, dstTensorDesc, dstData)); } - if(!batchnorm) { + if(!batchnorm && !additional_bias) { //CHECK WITH IF CORRECT // bias alpha = dnnType(1); beta = dnnType(1); @@ -108,23 +108,34 @@ void Conv2d::inferCUDNN(dnnType* srcData, bool back) { &alpha, biasTensorDesc, bias_d, &beta, dstTensorDesc, dstData) ); } else { - alpha = dnnType(1); - beta = dnnType(0); - checkCUDNN( cudnnBatchNormalizationForwardInference(net->cudnnHandle, - CUDNN_BATCHNORM_SPATIAL, &alpha, &beta, - dstTensorDesc, dstData, dstTensorDesc, - dstData, biasTensorDesc, //same tensor descriptor as bias - scales_d, bias_d, mean_d, variance_d, - TKDNN_BN_MIN_EPSILON) ); + if(additional_bias) + { + alpha = dnnType(1); + beta = dnnType(1); + checkCUDNN( cudnnAddTensor(net->cudnnHandle, + &alpha, biasTensorDesc, bias2_d, + &beta, dstTensorDesc, dstData) ); + } + if(batchnorm) + { + alpha = dnnType(1); + beta = dnnType(0); + checkCUDNN( cudnnBatchNormalizationForwardInference(net->cudnnHandle, + CUDNN_BATCHNORM_SPATIAL, &alpha, &beta, + dstTensorDesc, dstData, dstTensorDesc, + dstData, biasTensorDesc, //same tensor descriptor as bias + scales_d, bias_d, mean_d, variance_d, + TKDNN_BN_MIN_EPSILON) ); + } } } Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW, int strideH, int strideW, int paddingH, int paddingW, - std::string fname_weights, bool batchnorm, bool deConv, bool final, int groups) : + std::string fname_weights, bool batchnorm, bool deConv, bool final, int groups, bool additional_bias) : LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1, - fname_weights, batchnorm, false, final, deConv, groups) { + fname_weights, batchnorm, additional_bias, final, deConv, groups) { this->kernelH = kernelH; this->kernelW = kernelW; this->strideH = strideH; @@ -133,6 +144,7 @@ Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW, this->paddingW = paddingW; this->deConv = deConv; this->groups = groups; + this->additional_bias = additional_bias; if(!deConv) { output_dim.n = input_dim.n; diff --git a/src/LayerWgs.cpp b/src/LayerWgs.cpp index 41ca038..9dad5e9 100644 --- a/src/LayerWgs.cpp +++ b/src/LayerWgs.cpp @@ -9,11 +9,7 @@ namespace tk { namespace dnn { LayerWgs::LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kl, std::string fname_weights, bool batchnorm, bool additional_bias, bool final, bool deConv, int groups) : Layer(net, final) { - - if(deConv) - inputs = inputs/groups; - else - outputs = outputs/groups; + inputs = inputs/groups; this->inputs = inputs; this->outputs = outputs; diff --git a/src/Network.cpp b/src/Network.cpp index 3cfebd5..559855f 100644 --- a/src/Network.cpp +++ b/src/Network.cpp @@ -22,19 +22,25 @@ Network::Network(dataDim_t input_dim) { fp16 = false; dla = false; + int8 = false; if(const char* env_p = std::getenv("TKDNN_MODE")) { if(strcmp(env_p, "FP16") == 0) fp16 = true; - else if(strcmp(env_p, "DLA") == 0) { - dla = true; - fp16 = true; - } + else if(strcmp(env_p, "DLA") == 0) { + dla = true; + fp16 = true; + } + else if(strcmp(env_p, "INT8") == 0) { + int8 = true; + } } if(fp16) - std::cout<setDefaultDeviceType(DeviceType::kDLA); builderRT->setDLACore(0); } + // if(net->int8 && builderRT->platformHasFastInt8()) + // { + // dtRT = DataType::kINT8; + // builderRT->setInt8Mode(true); + // Int8EntropyCalibrator calibrator(1, "../demo/images.txt","../demo/yolov3-calibration.table", 416*416*3, 416, 416); + // builderRT->setInt8Calibrator((nvinfer1::IInt8Calibrator * )&calibrator); + // // builderRT->setStrictTypeConstraints(true); + // } //add input layer ITensor *input = networkRT->addInput("data", DataType::kFLOAT, @@ -164,7 +173,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) + if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY) return convert_layer(input, (Activation*) l); if(type == LAYER_SOFTMAX) return convert_layer(input, (Softmax*) l); @@ -337,7 +346,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kSIGMOID); checkNULL(lRT); return lRT; - + } + else if(l->act_mode == CUDNN_ACTIVATION_CLIPPED_RELU) { + IPlugin *plugin = new ActivationReLUCeiling(l->ceiling); + IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + checkNULL(lRT); + return lRT; } else { FatalError("this Activation mode is not yet implemented"); return NULL; @@ -535,11 +549,16 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa std::string name(layerName); std::cout<size = readBUF(buf); return a; } + if(name.find("ActivationCReLU") == 0) { + ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF(buf)); + a->size = readBUF(buf); + return a; + } if(name.find("Region") == 0) { RegionRT *r = new RegionRT(readBUF(buf), //classes diff --git a/src/Yolo3Detection.cpp b/src/Yolo3Detection.cpp index 03a1b31..f21efa7 100644 --- a/src/Yolo3Detection.cpp +++ b/src/Yolo3Detection.cpp @@ -63,6 +63,7 @@ bool Yolo3Detection::init(std::string tensor_path) { void Yolo3Detection::update(cv::Mat &imageORIG) { + TIMER_START if(!imageORIG.data) { std::cout<<"YOLO: NO IMAGE DATA\n"; return; @@ -100,7 +101,6 @@ void Yolo3Detection::update(cv::Mat &imageORIG) { stats.push_back(t_ns); } - TIMER_START // compute dets ndets = 0; for(int i=0; ipluginFactory->n_yolos; i++) { @@ -109,7 +109,6 @@ void Yolo3Detection::update(cv::Mat &imageORIG) { yolo[i]->computeDetections(dets, ndets, netRT->input_dim.w, netRT->input_dim.h, thresh); } tk::dnn::Yolo::mergeDetections(dets, ndets, classes); - TIMER_STOP // fill detected detected.clear(); @@ -148,6 +147,8 @@ void Yolo3Detection::update(cv::Mat &imageORIG) { detected.push_back(res); } } + TIMER_STOP + stats.push_back(t_ns); } diff --git a/src/kernels/activation_relu_ceiling.cu b/src/kernels/activation_relu_ceiling.cu new file mode 100644 index 0000000..72ccf4f --- /dev/null +++ b/src/kernels/activation_relu_ceiling.cu @@ -0,0 +1,33 @@ +#include "kernels.h" + +__global__ +void activation_relu_ceiling(dnnType *input, dnnType *output, int size, const float ceiling) { + + int i = blockDim.x*blockIdx.x + threadIdx.x; + + if(i0) + { + if (input[i]>ceiling) + output[i] = ceiling; + else + output[i] = input[i]; + } + else + output[i] = 0.0f; + } + } + + +/** + Relu ceiling activation function +*/ +void activationReLUCeilingForward(dnnType* srcData, dnnType* dstData, int size, const float ceiling, cudaStream_t stream) +{ + int blocks = (size+255)/256; + int threads = 256; + + activation_relu_ceiling<<>>(srcData, dstData, size, ceiling); +} + + diff --git a/tests/mobilenetv2ssd/mobilenetv2ssd.cpp b/tests/mobilenetv2ssd/mobilenetv2ssd.cpp new file mode 100644 index 0000000..8c91214 --- /dev/null +++ b/tests/mobilenetv2ssd/mobilenetv2ssd.cpp @@ -0,0 +1,448 @@ +#include +#include "tkdnn.h" + +const char *output_bin = "../tests/mobilenetv2ssd/debug/regression_headers-5.bin"; +const char *input_bin = "../tests/mobilenetv2ssd/debug/input.bin"; + +const char *conv0_bin = "../tests/mobilenetv2ssd/layers/base_net-0-0.bin"; +const char *inverted_residual1[]={ +"../tests/mobilenetv2ssd/layers/base_net-1-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-1-conv-3.bin"}; +const char *inverted_residual2[]={ +"../tests/mobilenetv2ssd/layers/base_net-2-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-2-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-2-conv-6.bin"}; +const char *inverted_residual3[]={ +"../tests/mobilenetv2ssd/layers/base_net-3-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-3-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-3-conv-6.bin"}; +const char *inverted_residual4[]={ +"../tests/mobilenetv2ssd/layers/base_net-4-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-4-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-4-conv-6.bin"}; +const char *inverted_residual5[]={ +"../tests/mobilenetv2ssd/layers/base_net-5-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-5-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-5-conv-6.bin"}; +const char *inverted_residual6[]={ +"../tests/mobilenetv2ssd/layers/base_net-6-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-6-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-6-conv-6.bin"}; +const char *inverted_residual7[]={ +"../tests/mobilenetv2ssd/layers/base_net-7-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-7-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-7-conv-6.bin"}; +const char *inverted_residual8[]={ +"../tests/mobilenetv2ssd/layers/base_net-8-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-8-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-8-conv-6.bin"}; +const char *inverted_residual9[]={ +"../tests/mobilenetv2ssd/layers/base_net-9-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-9-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-9-conv-6.bin"}; +const char *inverted_residual10[]={ +"../tests/mobilenetv2ssd/layers/base_net-10-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-10-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-10-conv-6.bin"}; +const char *inverted_residual11[]={ +"../tests/mobilenetv2ssd/layers/base_net-11-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-11-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-11-conv-6.bin"}; +const char *inverted_residual12[]={ +"../tests/mobilenetv2ssd/layers/base_net-12-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-12-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-12-conv-6.bin"}; +const char *inverted_residual13[]={ +"../tests/mobilenetv2ssd/layers/base_net-13-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-13-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-13-conv-6.bin"}; +const char *inverted_residual14[]={ +"../tests/mobilenetv2ssd/layers/base_net-14-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-14-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-14-conv-6.bin"}; +const char *inverted_residual15[]={ +"../tests/mobilenetv2ssd/layers/base_net-15-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-15-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-15-conv-6.bin"}; +const char *inverted_residual16[]={ +"../tests/mobilenetv2ssd/layers/base_net-16-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-16-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-16-conv-6.bin"}; +const char *inverted_residual17[]={ +"../tests/mobilenetv2ssd/layers/base_net-17-conv-0.bin", +"../tests/mobilenetv2ssd/layers/base_net-17-conv-3.bin", +"../tests/mobilenetv2ssd/layers/base_net-17-conv-6.bin"}; + +const char *conv18 = "../tests/mobilenetv2ssd/layers/base_net-18-0.bin"; + +const char *extras0[]={ +"../tests/mobilenetv2ssd/layers/extras-0-conv-0.bin", +"../tests/mobilenetv2ssd/layers/extras-0-conv-3.bin", +"../tests/mobilenetv2ssd/layers/extras-0-conv-6.bin"}; +const char *extras1[]={ +"../tests/mobilenetv2ssd/layers/extras-1-conv-0.bin", +"../tests/mobilenetv2ssd/layers/extras-1-conv-3.bin", +"../tests/mobilenetv2ssd/layers/extras-1-conv-6.bin"}; +const char *extras2[]={ +"../tests/mobilenetv2ssd/layers/extras-2-conv-0.bin", +"../tests/mobilenetv2ssd/layers/extras-2-conv-3.bin", +"../tests/mobilenetv2ssd/layers/extras-2-conv-6.bin"}; +const char *extras3[]={ +"../tests/mobilenetv2ssd/layers/extras-3-conv-0.bin", +"../tests/mobilenetv2ssd/layers/extras-3-conv-3.bin", +"../tests/mobilenetv2ssd/layers/extras-3-conv-6.bin"}; + +const char *classification_header0[]={ +"../tests/mobilenetv2ssd/layers/classification_headers-0-0.bin", +"../tests/mobilenetv2ssd/layers/classification_headers-0-3.bin"}; +const char *classification_header1[]={ +"../tests/mobilenetv2ssd/layers/classification_headers-1-0.bin", +"../tests/mobilenetv2ssd/layers/classification_headers-1-3.bin"}; +const char *classification_header2[]={ +"../tests/mobilenetv2ssd/layers/classification_headers-2-0.bin", +"../tests/mobilenetv2ssd/layers/classification_headers-2-3.bin"}; +const char *classification_header3[]={ +"../tests/mobilenetv2ssd/layers/classification_headers-3-0.bin", +"../tests/mobilenetv2ssd/layers/classification_headers-3-3.bin"}; +const char *classification_header4[]={ +"../tests/mobilenetv2ssd/layers/classification_headers-4-0.bin", +"../tests/mobilenetv2ssd/layers/classification_headers-4-3.bin"}; + +const char *classification_header5 = "../tests/mobilenetv2ssd/layers/classification_headers-5.bin"; + +const char *regression_header0[]={ +"../tests/mobilenetv2ssd/layers/regression_headers-0-0.bin", +"../tests/mobilenetv2ssd/layers/regression_headers-0-3.bin"}; +const char *regression_header1[]={ +"../tests/mobilenetv2ssd/layers/regression_headers-1-0.bin", +"../tests/mobilenetv2ssd/layers/regression_headers-1-3.bin"}; +const char *regression_header2[]={ +"../tests/mobilenetv2ssd/layers/regression_headers-2-0.bin", +"../tests/mobilenetv2ssd/layers/regression_headers-2-3.bin"}; +const char *regression_header3[]={ +"../tests/mobilenetv2ssd/layers/regression_headers-3-0.bin", +"../tests/mobilenetv2ssd/layers/regression_headers-3-3.bin"}; +const char *regression_header4[]={ +"../tests/mobilenetv2ssd/layers/regression_headers-4-0.bin", +"../tests/mobilenetv2ssd/layers/regression_headers-4-3.bin"}; + +const char *regression_header5 = "../tests/mobilenetv2ssd/layers/regression_headers-5.bin"; + + + + + +int main() +{ + + // Network layout + tk::dnn::dataDim_t dim(1, 3, 300, 300, 1); + tk::dnn::Network net(dim); + + tk::dnn::Conv2d conv1(&net, 32, 3, 3, 2, 2, 1, 1, conv0_bin, true); + tk::dnn::Activation relu3(&net, CUDNN_ACTIVATION_RELU); + + //Inverted Residual 1 + + tk::dnn::Conv2d conv2(&net, 32, 3, 3, 1, 1, 1, 1, inverted_residual1[0], true,false, false,32); + tk::dnn::Activation relu5(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d conv3(&net, 16, 1, 1, 1, 1, 0, 0, inverted_residual1[1], true); + + + //Inverted Residual 2 + tk::dnn::Conv2d ir_2_conv1(&net, 96, 1, 1, 1, 1, 0, 0, inverted_residual2[0], true); + tk::dnn::Activation relu_2_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_2_conv2(&net, 96, 3, 3, 2, 2, 1, 1, inverted_residual2[1], true, false, false, 96); + tk::dnn::Activation relu_2_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_2_conv3(&net, 24, 1, 1, 1, 1, 0, 0, inverted_residual2[2], true); + + //Inverted Residual 3 + tk::dnn::Layer *last = &ir_2_conv3; + tk::dnn::Conv2d ir_3_conv1(&net, 144, 1, 1, 1, 1, 0, 0, inverted_residual3[0], true); + tk::dnn::Activation relu_3_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_3_conv2(&net, 144, 3, 3, 1, 1, 1, 1, inverted_residual3[1], true, false, false, 144); + tk::dnn::Activation relu_3_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_3_conv3(&net, 24, 1, 1, 1, 1, 0, 0, inverted_residual3[2], true); + + tk::dnn::Shortcut s3_0 (&net, last); + // //Inverted Residual 4 + tk::dnn::Conv2d ir_4_conv1(&net, 144, 1, 1, 1, 1, 0, 0, inverted_residual4[0], true); + tk::dnn::Activation relu_4_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_4_conv2(&net, 144, 3, 3, 2, 2, 1, 1, inverted_residual4[1], true, false, false, 144); + tk::dnn::Activation relu_4_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_4_conv3(&net, 32, 1, 1, 1, 1, 0, 0, inverted_residual4[2], true); + + // // //Inverted Residual 5 + last = &ir_4_conv3; + tk::dnn::Conv2d ir_5_conv1(&net, 192, 1, 1, 1, 1, 0, 0, inverted_residual5[0], true); + tk::dnn::Activation relu_5_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_5_conv2(&net, 192, 3, 3, 1, 1, 1, 1, inverted_residual5[1], true, false, false, 192); + tk::dnn::Activation relu_5_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_5_conv3(&net, 32, 1, 1, 1, 1, 0, 0, inverted_residual5[2], true); + + tk::dnn::Shortcut s5_0 (&net, last); + // // // //Inverted Residual 6 + last = &s5_0; + tk::dnn::Conv2d ir_6_conv1(&net, 192, 1, 1, 1, 1, 0, 0, inverted_residual6[0], true); + tk::dnn::Activation relu_6_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_6_conv2(&net, 192, 3, 3, 1, 1, 1, 1, inverted_residual6[1], true, false, false, 192); + tk::dnn::Activation relu_6_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_6_conv3(&net, 32, 1, 1, 1, 1, 0, 0, inverted_residual6[2], true); + + tk::dnn::Shortcut s6_0 (&net, last); + //Inverted Residual 7 + tk::dnn::Conv2d ir_7_conv1(&net, 192, 1, 1, 1, 1, 0, 0, inverted_residual7[0], true); + tk::dnn::Activation relu_7_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_7_conv2(&net, 192, 3, 3, 2, 2, 1, 1, inverted_residual7[1], true, false, false, 192); + tk::dnn::Activation relu_7_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_7_conv3(&net, 64, 1, 1, 1, 1, 0, 0, inverted_residual7[2], true); + + // //Inverted Residual 8 + last = &ir_7_conv3; + tk::dnn::Conv2d ir_8_conv1(&net, 384, 1, 1, 1, 1, 0, 0, inverted_residual8[0], true); + tk::dnn::Activation relu_8_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_8_conv2(&net, 384, 3, 3, 1, 1, 1, 1, inverted_residual8[1], true, false, false, 384); + tk::dnn::Activation relu_8_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_8_conv3(&net, 64, 1, 1, 1, 1, 0, 0, inverted_residual8[2], true); + + tk::dnn::Shortcut s8_0 (&net, last); + //Inverted Residual 9 + last = &s8_0; + tk::dnn::Conv2d ir_9_conv1(&net, 384, 1, 1, 1, 1, 0, 0, inverted_residual9[0], true); + tk::dnn::Activation relu_9_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_9_conv2(&net, 384, 3, 3, 1, 1, 1, 1, inverted_residual9[1], true, false, false, 384); + tk::dnn::Activation relu_9_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_9_conv3(&net, 64, 1, 1, 1, 1, 0, 0, inverted_residual9[2], true); + + tk::dnn::Shortcut s9_0 (&net, last); + //Inverted Residual 10 + last = &s9_0; + tk::dnn::Conv2d ir_10_conv1(&net, 384, 1, 1, 1, 1, 0, 0, inverted_residual10[0], true); + tk::dnn::Activation relu_10_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_10_conv2(&net, 384, 3, 3, 1, 1, 1, 1, inverted_residual10[1], true, false, false, 384); + tk::dnn::Activation relu_10_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_10_conv3(&net, 64, 1, 1, 1, 1, 0, 0, inverted_residual10[2], true); + + tk::dnn::Shortcut s10_0 (&net, last); + //Inverted Residual 11 + tk::dnn::Conv2d ir_11_conv1(&net, 384, 1, 1, 1, 1, 0, 0, inverted_residual11[0], true); + tk::dnn::Activation relu_11_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_11_conv2(&net, 384, 3, 3, 1, 1, 1, 1, inverted_residual11[1], true, false, false, 384); + tk::dnn::Activation relu_11_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_11_conv3(&net, 96, 1, 1, 1, 1, 0, 0, inverted_residual11[2], true); + + last = &ir_11_conv3; + //Inverted Residual 12 + tk::dnn::Conv2d ir_12_conv1(&net, 576, 1, 1, 1, 1, 0, 0, inverted_residual12[0], true); + tk::dnn::Activation relu_12_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_12_conv2(&net, 576, 3, 3, 1, 1, 1, 1, inverted_residual12[1], true, false, false, 576); + tk::dnn::Activation relu_12_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_12_conv3(&net, 96, 1, 1, 1, 1, 0, 0, inverted_residual12[2], true); + + tk::dnn::Shortcut s12_0 (&net, last); + last = &s12_0; + //Inverted Residual 13 + tk::dnn::Conv2d ir_13_conv1(&net, 576, 1, 1, 1, 1, 0, 0, inverted_residual13[0], true); + tk::dnn::Activation relu_13_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_13_conv2(&net, 576, 3, 3, 1, 1, 1, 1, inverted_residual13[1], true, false, false, 576); + tk::dnn::Activation relu_13_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_13_conv3(&net, 96, 1, 1, 1, 1, 0, 0, inverted_residual13[2], true); + + tk::dnn::Shortcut s13_0 (&net, last); + // //Inverted Residual 14 + tk::dnn::Conv2d ir_14_conv1(&net, 576, 1, 1, 1, 1, 0, 0, inverted_residual14[0], true); + tk::dnn::Activation relu_14_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_14_conv2(&net, 576, 3, 3, 2, 2, 1, 1, inverted_residual14[1], true, false, false, 576); + tk::dnn::Activation relu_14_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_14_conv3(&net, 160, 1, 1, 1, 1, 0, 0, inverted_residual14[2], true); + + // //Inverted Residual 15 + last = &ir_14_conv3; + tk::dnn::Conv2d ir_15_conv1(&net, 960, 1, 1, 1, 1, 0, 0, inverted_residual15[0], true); + tk::dnn::Activation relu_15_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_15_conv2(&net, 960, 3, 3, 1, 1, 1, 1, inverted_residual15[1], true, false, false, 960); + tk::dnn::Activation relu_15_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_15_conv3(&net, 160, 1, 1, 1, 1, 0, 0, inverted_residual15[2], true); + + tk::dnn::Shortcut s15_0 (&net, last); + //Inverted Residual 16 + last = &s15_0; + tk::dnn::Conv2d ir_16_conv1(&net, 960, 1, 1, 1, 1, 0, 0, inverted_residual16[0], true); + tk::dnn::Activation relu_16_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_16_conv2(&net, 960, 3, 3, 1, 1, 1, 1, inverted_residual16[1], true, false, false, 960); + tk::dnn::Activation relu_16_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_16_conv3(&net, 160, 1, 1, 1, 1, 0, 0, inverted_residual16[2], true); + + tk::dnn::Shortcut s16_0 (&net, last); + //Inverted Residual 17 + tk::dnn::Conv2d ir_17_conv1(&net, 960, 1, 1, 1, 1, 0, 0, inverted_residual17[0], true); + tk::dnn::Activation relu_17_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_17_conv2(&net, 960, 3, 3, 1, 1, 1, 1, inverted_residual17[1], true, false, false, 960); + tk::dnn::Activation relu_17_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d ir_17_conv3(&net, 320, 1, 1, 1, 1, 0, 0, inverted_residual17[2], true); + + //Conv 18 + tk::dnn::Conv2d ir_18_conv1(&net, 1280, 1, 1, 1, 1, 0, 0, conv18, true); + tk::dnn::Activation relu_18_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Layer * header_1[1] = {&relu_18_1}; + + // //extras Inverted Residual 0 + tk::dnn::Conv2d e_0_conv1(&net, 256, 1, 1, 1, 1, 0, 0, extras0[0], true); + tk::dnn::Activation e_relu_0_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d e_0_conv2(&net, 256, 3, 3, 2, 2, 1, 1, extras0[1], true, false, false, 256); + tk::dnn::Activation e_relu_0_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d e_0_conv3(&net, 512, 1, 1, 1, 1, 0, 0, extras0[2], true); + tk::dnn::Layer * header_2[1] = {&e_0_conv3}; + + // //extras Inverted Residual 1 + tk::dnn::Conv2d e_1_conv1(&net, 128, 1, 1, 1, 1, 0, 0, extras1[0], true); + tk::dnn::Activation e_relu_1_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d e_1_conv2(&net, 128, 3, 3, 2, 2, 1, 1, extras1[1], true, false, false, 128); + tk::dnn::Activation e_relu_1_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d e_1_conv3(&net, 256, 1, 1, 1, 1, 0, 0, extras1[2], true); + tk::dnn::Layer * header_3[1] = {&e_1_conv3}; + + //extras Inverted Residual 2 + tk::dnn::Conv2d e_2_conv1(&net, 128, 1, 1, 1, 1, 0, 0, extras2[0], true); + tk::dnn::Activation e_relu_2_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d e_2_conv2(&net, 128, 3, 3, 2, 2, 1, 1, extras2[1], true, false, false, 128); + tk::dnn::Activation e_relu_2_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d e_2_conv3(&net, 256, 1, 1, 1, 1, 0, 0, extras2[2], true); + tk::dnn::Layer * header_4[1] = {&e_2_conv3}; + + //extras Inverted Residual 3 + tk::dnn::Conv2d e_3_conv1(&net, 64, 1, 1, 1, 1, 0, 0, extras3[0], true); + tk::dnn::Activation e_relu_3_1(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d e_3_conv2(&net, 64, 3, 3, 2, 2, 1, 1, extras3[1], true, false, false, 64); + tk::dnn::Activation e_relu_3_2(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d e_3_conv3(&net, 64, 1, 1, 1, 1, 0, 0, extras3[2], true); + tk::dnn::Layer * header_5[1] = {&e_3_conv3}; + + // classification header 0 + tk::dnn::Layer * header_0[1] = {&relu_14_1}; + tk::dnn::Route rout_ch_0(&net, header_0, 1); + tk::dnn::Conv2d ch_0_conv1(&net, 576, 3, 3, 1, 1, 1, 1, classification_header0[0], true, false, false, 576,true); + tk::dnn::Activation ch_relu_0_1(&net, CUDNN_ACTIVATION_CLIPPED_RELU,6); + tk::dnn::Conv2d ch_0_conv2(&net, 126, 1, 1, 1, 1, 0, 0, classification_header0[1], false); + + // // classification header 1 + tk::dnn::Route rout_ch_1(&net, header_1, 1); + tk::dnn::Conv2d ch_1_conv1(&net, 1280, 3, 3, 1, 1, 1, 1, classification_header1[0], true, false, false, 1280,true); + tk::dnn::Activation ch_relu_1_1(&net, CUDNN_ACTIVATION_CLIPPED_RELU,6); + tk::dnn::Conv2d ch_1_conv2(&net, 126, 1, 1, 1, 1, 0, 0, classification_header1[1], false); + + // //classification header 2 + tk::dnn::Route rout_ch_2(&net, header_2, 1); + tk::dnn::Conv2d ch_2_conv1(&net, 512, 3, 3, 1, 1, 1, 1, classification_header2[0], true, false, false, 512, true); + tk::dnn::Activation ch_relu_2_1(&net, CUDNN_ACTIVATION_CLIPPED_RELU,6); + tk::dnn::Conv2d ch_2_conv2(&net, 126, 1, 1, 1, 1, 0, 0, classification_header2[1], false); + + // //classification header 3 + tk::dnn::Route rout_ch_3(&net, header_3, 1); + tk::dnn::Conv2d ch_3_conv1(&net, 256, 3, 3, 1, 1, 1, 1, classification_header3[0], true, false, false, 256, true); + tk::dnn::Activation ch_relu_3_1(&net, CUDNN_ACTIVATION_CLIPPED_RELU,6); + tk::dnn::Conv2d ch_3_conv2(&net, 126, 1, 1, 1, 1, 0, 0, classification_header3[1], false); + + // //classification header 4 + tk::dnn::Route rout_ch_4(&net, header_4, 1); + tk::dnn::Conv2d ch_4_conv1(&net, 256, 3, 3, 1, 1, 1, 1, classification_header4[0], true, false, false, 256, true); + tk::dnn::Activation ch_relu_4_1(&net, CUDNN_ACTIVATION_CLIPPED_RELU,6); + tk::dnn::Conv2d ch_4_conv2(&net, 126, 1, 1, 1, 1, 0, 0, classification_header4[1], false); + + // //classification header 5 + tk::dnn::Route rout_ch_5(&net, header_5, 1); + tk::dnn::Conv2d ch_5_conv(&net, 126, 1, 1, 1, 1, 0, 0, classification_header5, false); + + //regression header 0 + tk::dnn::Route rout_rh_0(&net, header_0, 1); + tk::dnn::Conv2d rh_0_conv1(&net, 576, 3, 3, 1, 1, 1, 1, regression_header0[0], true, false, false, 576, true); + tk::dnn::Activation rh_relu_0_1(&net, CUDNN_ACTIVATION_CLIPPED_RELU,6); + tk::dnn::Conv2d rh_0_conv2(&net, 24, 1, 1, 1, 1, 0, 0, regression_header0[1], false); + + // //regression header 1 + tk::dnn::Route rout_rh_1(&net, header_1, 1); + tk::dnn::Conv2d rh_1_conv1(&net, 1280, 3, 3, 1, 1, 1, 1, regression_header1[0], true, false, false, 1280, true); + tk::dnn::Activation rh_relu_1_1(&net, CUDNN_ACTIVATION_CLIPPED_RELU,6); + tk::dnn::Conv2d rh_1_conv2(&net, 24, 1, 1, 1, 1, 0, 0, regression_header1[1], false); + + //regression header 2 + tk::dnn::Route rout_rh_2(&net, header_2, 1); + tk::dnn::Conv2d rh_2_conv1(&net, 512, 3, 3, 1, 1, 1, 1, regression_header2[0], true, false, false, 512, true); + tk::dnn::Activation rh_relu_2_1(&net, CUDNN_ACTIVATION_CLIPPED_RELU,6); + tk::dnn::Conv2d rh_2_conv2(&net, 24, 1, 1, 1, 1, 0, 0, regression_header2[1], false); + + //regression header 3 + tk::dnn::Route rout_rh_3(&net, header_3, 1); + tk::dnn::Conv2d rh_3_conv1(&net, 256, 3, 3, 1, 1, 1, 1, regression_header3[0], true, false, false, 256, true); + tk::dnn::Activation rh_relu_3_1(&net, CUDNN_ACTIVATION_CLIPPED_RELU,6); + tk::dnn::Conv2d rh_3_conv2(&net, 24, 1, 1, 1, 1, 0, 0, regression_header3[1], false); + + //regression header 4 + + tk::dnn::Route rout_rh_4(&net, header_4, 1); + tk::dnn::Conv2d rh_4_conv1(&net, 256, 3, 3, 1, 1, 1, 1, regression_header4[0], true, false, false, 256, true); + tk::dnn::Activation rh_relu_4_1(&net, CUDNN_ACTIVATION_CLIPPED_RELU,6); + tk::dnn::Conv2d rh_4_conv2(&net, 24, 1, 1, 1, 1, 0, 0, regression_header4[1], false); + + //regression header 5 + tk::dnn::Route rout_rh_5(&net, header_5, 1); + tk::dnn::Conv2d rh_5_conv(&net, 24, 1, 1, 1, 1, 0, 0, regression_header5, false); + + //flatten confidence and flatten locations + + // Load input + dnnType *data; + dnnType *input_h; + readBinaryFile(input_bin, dim.tot(), &input_h, &data); + //printDeviceVector(64, data, true); + + //print network model + net.print(); + + //convert network to tensorRT + tk::dnn::NetworkRT netRT(&net, "mobilenetv2ssd.rt"); + + + tk::dnn::dataDim_t out_dim; + out_dim = net.layers[net.num_layers-1]->output_dim; + dnnType *cudnn_out, *rt_out; + + tk::dnn::dataDim_t dim1 = dim; //input dim + printCenteredTitle(" CUDNN inference ", '=', 30); + { + dim1.print(); + TIMER_START + net.infer(dim1, data); + TIMER_STOP + dim1.print(); + } + cudnn_out = net.layers[net.num_layers-1]->dstData; + + printDeviceVector(64, cudnn_out, true); + + tk::dnn::dataDim_t dim2 = dim; + printCenteredTitle(" TENSORRT inference ", '=', 30); + { + dim2.print(); + TIMER_START + netRT.infer(dim2, data); + TIMER_STOP + dim2.print(); + } + rt_out = (dnnType *)netRT.buffersRT[1]; + + + printCenteredTitle(std::string(" RESNET CHECK RESULTS ").c_str(), '=', 30); + dnnType *out, *out_h; + int odim = out_dim.tot(); + readBinaryFile(output_bin, odim, &out_h, &out); + std::cout << "CUDNN vs correct"; + checkResult(odim, cudnn_out, out); + + std::cout << "TRT vs correct"; + checkResult(odim, rt_out, out); + std::cout << "CUDNN vs TRT "; + checkResult(odim, cudnn_out, rt_out); + + return 0; +}