[TensorRT-8] Add independent padding and depth NN #278
+12
-2
@@ -85,12 +85,14 @@ endif()
|
||||
find_package(CUDNN REQUIRED)
|
||||
include_directories(${CUDNN_INCLUDE_DIR})
|
||||
|
||||
find_package(yaml-cpp REQUIRED)
|
||||
|
||||
|
||||
# compile
|
||||
file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/sorting.cu" "src/pluginsRT/*.cpp")
|
||||
cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS})
|
||||
cuda_add_library(kernels SHARED ${tkdnn_CUSRC})
|
||||
target_link_libraries(kernels ${CUDA_CUBLAS_LIBRARIES} ${CUDA_LIBRARIES} ${CUDNN_LIBRARIES})
|
||||
target_link_libraries(kernels ${CUDA_CUBLAS_LIBRARIES} ${CUDA_LIBRARIES} ${CUDNN_LIBRARIES} yaml-cpp)
|
||||
|
||||
|
||||
|
||||
@@ -120,7 +122,6 @@ endif()
|
||||
# endif()
|
||||
|
||||
# gives problems in cross-compiling, probably malformed cmake config
|
||||
find_package(yaml-cpp REQUIRED)
|
||||
|
||||
#-------------------------------------------------------------------------------
|
||||
# Build Libraries
|
||||
@@ -202,6 +203,12 @@ target_link_libraries(test_shelfnet_berkeley tkDNN)
|
||||
add_executable(test_shelfnet_mapillary tests/shelfnet/shelfnet_mapillary.cpp)
|
||||
target_link_libraries(test_shelfnet_mapillary tkDNN)
|
||||
|
||||
# MONODEPTH2
|
||||
add_executable(test_monodepth2_640 tests/monodepth2/monodepth2_640.cpp)
|
||||
target_link_libraries(test_monodepth2_640 tkDNN)
|
||||
|
||||
add_executable(test_monodepth2_1024 tests/monodepth2/monodepth2_1024.cpp)
|
||||
target_link_libraries(test_monodepth2_1024 tkDNN)
|
||||
# DEMOS
|
||||
add_executable(test_rtinference tests/test_rtinference/rtinference.cpp)
|
||||
target_link_libraries(test_rtinference tkDNN)
|
||||
@@ -221,6 +228,9 @@ target_link_libraries(demoTracker tkDNN)
|
||||
add_executable(seg_demo demo/demo/seg_demo.cpp)
|
||||
target_link_libraries(seg_demo tkDNN)
|
||||
|
||||
add_executable(demoDepth demo/demo/demoDepth.cpp)
|
||||
target_link_libraries(demoDepth tkDNN)
|
||||
|
||||
#-------------------------------------------------------------------------------
|
||||
# Install
|
||||
#-------------------------------------------------------------------------------
|
||||
|
||||
+3
-3
@@ -46,9 +46,9 @@ int main(int argc, char *argv[]) {
|
||||
std::string cfgPath = YAMLgetConf<std::string>(conf,"cfg_input", "../tests/darknet/cfg/yolo4tiny.cfg");
|
||||
std::string namePath = YAMLgetConf<std::string>(conf,"name_input","../tests/darknet/names/coco.names");
|
||||
#elif _WIN32
|
||||
std::string input = YAMLgetConf(conf, "win_input", "..\\..\\..\\demo\\yolo_test.mp4");
|
||||
std::string cfgPath = YAMLgetConf(conf,"cfg_win_input","..\\..\\..\\tests\\darknet\\cfg\\yolo4tiny.cfg");
|
||||
std::string namePath = YAMLgetConf(conf,"name_win_input","..\\..\\..\\tests\\darknet\\names\\coco.names");
|
||||
std::string input = YAMLgetConf<std::string>(conf, "win_input", "..\\..\\..\\demo\\yolo_test.mp4");
|
||||
std::string cfgPath = YAMLgetConf<std::string>(conf,"cfg_win_input","..\\..\\..\\tests\\darknet\\cfg\\yolo4tiny.cfg");
|
||||
std::string namePath = YAMLgetConf<std::string>(conf,"name_win_input","..\\..\\..\\tests\\darknet\\names\\coco.names");
|
||||
#endif
|
||||
if(!fileExist(input.c_str()))
|
||||
FatalError("The given input video does not exist.");
|
||||
|
||||
@@ -0,0 +1,106 @@
|
||||
#include <iostream>
|
||||
#include <signal.h>
|
||||
#include <stdlib.h> /* srand, rand */
|
||||
//#include <unistd.h>
|
||||
#include <mutex>
|
||||
|
||||
#include "tkDNN/DepthNN.h"
|
||||
|
||||
bool gRun;
|
||||
|
||||
void sig_handler(int signo) {
|
||||
std::cout<<"request gateway stop\n";
|
||||
gRun = false;
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[]) {
|
||||
|
||||
signal(SIGINT, sig_handler);
|
||||
|
||||
std::string net = "monodepth2_fp32.rt";
|
||||
if(argc > 1)
|
||||
net = argv[1];
|
||||
#ifdef __linux__
|
||||
std::string input = "../demo/yolo_test.mp4";
|
||||
#elif _WIN32
|
||||
std::string input = "..\\..\\..\\demo\\yolo_test.mp4";
|
||||
#endif
|
||||
if(argc > 2)
|
||||
input = argv[2];
|
||||
bool show = true;
|
||||
if(argc > 3)
|
||||
show = atoi(argv[3]);
|
||||
bool save = true;
|
||||
if(argc > 4)
|
||||
save = atoi(argv[4]);
|
||||
|
||||
std::cout <<"Net settings - net: "<< net
|
||||
<<"\n";
|
||||
std::cout <<"Demo settings - input: "<< input
|
||||
<<", show: "<< show
|
||||
<<", save: "<< save<<"\n\n";
|
||||
|
||||
tk::dnn::DepthNN depthNN;
|
||||
|
||||
// create depth network
|
||||
int n_batch = 1;
|
||||
depthNN.init(net, n_batch);
|
||||
|
||||
// open video stream
|
||||
cv::VideoCapture cap(input);
|
||||
if(!cap.isOpened())
|
||||
gRun = false;
|
||||
else
|
||||
std::cout<<"camera started\n";
|
||||
|
||||
cv::VideoWriter resultVideo;
|
||||
if(save) {
|
||||
int w = depthNN.output_w;
|
||||
int h = depthNN.output_h;
|
||||
resultVideo.open("result.mp4", cv::VideoWriter::fourcc('M','J','P','G'), 30, cv::Size(w, h));
|
||||
}
|
||||
|
||||
if(show)
|
||||
cv::namedWindow("depth", cv::WINDOW_NORMAL);
|
||||
|
||||
cv::Mat frame;
|
||||
std::vector<cv::Mat> batch_frame;
|
||||
std::vector<cv::Mat> batch_dnn_input;
|
||||
|
||||
// start detection loop
|
||||
gRun = true;
|
||||
while(gRun) {
|
||||
batch_dnn_input.clear();
|
||||
batch_frame.clear();
|
||||
|
||||
//read frame
|
||||
cap >> frame;
|
||||
if(!frame.data)
|
||||
break;
|
||||
batch_frame.push_back(frame);
|
||||
batch_dnn_input.push_back(frame.clone());
|
||||
|
||||
//inference
|
||||
depthNN.update(batch_dnn_input, 1);
|
||||
if(show){
|
||||
cv::imshow("depth", depthNN.depthMats[0]);
|
||||
cv::waitKey(1);
|
||||
|
||||
}
|
||||
|
||||
if(save)
|
||||
resultVideo << depthNN.depthMats[0];
|
||||
}
|
||||
|
||||
std::cout<<"detection end\n";
|
||||
|
||||
double mean = 0;
|
||||
std::cout<<COL_GREENB<<"\n\nTime stats depth:\n";
|
||||
std::cout<<"Min: "<<*std::min_element(depthNN.stats.begin(), depthNN.stats.end())<<" ms\n";
|
||||
std::cout<<"Max: "<<*std::max_element(depthNN.stats.begin(), depthNN.stats.end())<<" ms\n";
|
||||
for(int i=0; i<depthNN.stats.size(); i++) mean += depthNN.stats[i]; mean /= depthNN.stats.size();
|
||||
std::cout<<"Avg: "<<mean<<" ms\t"<<1000/(mean)<<" FPS\n";
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,180 @@
|
||||
#ifndef DEPTHNN_H
|
||||
#define DEPTHNN_H
|
||||
|
||||
#include <iostream>
|
||||
#include <signal.h>
|
||||
#include <stdlib.h>
|
||||
#ifdef __linux__
|
||||
#include <unistd.h>
|
||||
#endif
|
||||
|
||||
#include <mutex>
|
||||
|
||||
#include <opencv2/core/core.hpp>
|
||||
#include <opencv2/highgui/highgui.hpp>
|
||||
#include <opencv2/imgproc/imgproc.hpp>
|
||||
|
||||
#include "tkDNN/utils.h"
|
||||
#include "tkDNN/tkdnn.h"
|
||||
|
||||
#include "NetworkViz.h"
|
||||
|
||||
|
||||
namespace tk { namespace dnn {
|
||||
|
||||
class DepthNN {
|
||||
|
||||
public:
|
||||
tk::dnn::NetworkRT *netRT = nullptr;
|
||||
dnnType *input_h;
|
||||
dnnType *input_d;
|
||||
float* depth_h;
|
||||
|
||||
int output_w;
|
||||
int output_h;
|
||||
|
||||
int nBatches = 1;
|
||||
|
||||
cv::Mat bgr[3];
|
||||
cv::Mat imagePreproc;
|
||||
|
||||
std::vector<double> stats; /*keeps track of inference times (ms)*/
|
||||
std::vector<std::vector<float>> depths;
|
||||
std::vector<cv::Mat> depthMats;
|
||||
|
||||
DepthNN() {};
|
||||
~DepthNN(){};
|
||||
|
||||
/**
|
||||
* Method used to initialize the class, allocate memory and compute
|
||||
* needed data.
|
||||
*
|
||||
* @param tensor_path path to the rt file of the NN.
|
||||
* @param n_batches maximum number of batches to use in inference
|
||||
* @return true if everything is correct, false otherwise.
|
||||
*/
|
||||
void init(const std::string& tensor_path, const int n_batches=1){
|
||||
//create net
|
||||
|
||||
std::cout<<(tensor_path).c_str()<<"\n";
|
||||
nBatches = n_batches;
|
||||
netRT = new tk::dnn::NetworkRT(NULL, (tensor_path).c_str());
|
||||
|
||||
//allocate memory for NN input
|
||||
checkCuda(cudaMallocHost(&input_h, sizeof(dnnType) * netRT->input_dim.tot() * nBatches));
|
||||
checkCuda(cudaMalloc(&input_d, sizeof(dnnType) * netRT->input_dim.tot() * nBatches));
|
||||
|
||||
//allocate memory for NN output
|
||||
depthMats.resize(nBatches);
|
||||
depths.resize(nBatches);
|
||||
for(int i=0; i< depths.size();++i)
|
||||
depths[i].resize(netRT->buffersDIM[1].tot());
|
||||
|
||||
depth_h = (float *)malloc(netRT->buffersDIM[1].tot() * sizeof(float));
|
||||
|
||||
output_h = netRT->buffersDIM[1].h;
|
||||
output_w = netRT->buffersDIM[1].w;
|
||||
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* This method preprocess the image, before feeding it to the NN.
|
||||
*
|
||||
* @param frame original frame to adapt for inference.
|
||||
* @param bi batch index
|
||||
*/
|
||||
void preprocess(cv::Mat &frame, const int bi=0) {
|
||||
//resize image, remove mean, divide by std
|
||||
cv::Mat frame_nomean;
|
||||
resize(frame, frame, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
|
||||
frame.convertTo(frame_nomean, CV_32FC3);
|
||||
frame_nomean.convertTo(imagePreproc, CV_32FC3, 1 / 255.0, 0);
|
||||
|
||||
//copy image into tensor and copy it into GPU
|
||||
cv::split(imagePreproc, bgr);
|
||||
for (int i = 0; i < netRT->input_dim.c; i++){
|
||||
int idx = i * imagePreproc.rows * imagePreproc.cols;
|
||||
int ch = netRT->input_dim.c-1 -i;
|
||||
memcpy((void *)&input_h[idx + netRT->input_dim.tot()*bi], (void *)bgr[ch].data, imagePreproc.rows * imagePreproc.cols * sizeof(dnnType));
|
||||
}
|
||||
checkCuda(cudaMemcpyAsync(input_d+ netRT->input_dim.tot()*bi, input_h + netRT->input_dim.tot()*bi, netRT->input_dim.tot() * sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
|
||||
}
|
||||
|
||||
/**
|
||||
* This method postprocess the output of the NN to obtain the correct
|
||||
* boundig boxes.
|
||||
*
|
||||
* @param bi batch index
|
||||
* @param mAP set to true only if all the probabilities for a bounding
|
||||
* box are needed, as in some cases for the mAP calculation
|
||||
*/
|
||||
void postprocess(const int bi=0) {
|
||||
|
||||
dnnType *rt_out[1];
|
||||
rt_out[0] = (dnnType *)netRT->buffersRT[1]+ netRT->buffersDIM[1].tot()*bi;
|
||||
checkCuda(cudaMemcpy(depth_h, rt_out[0], netRT->buffersDIM[1].tot()* sizeof(float), cudaMemcpyDeviceToHost));
|
||||
memcpy(&depths[bi][0], &depth_h[0], netRT->buffersDIM[1].tot()* sizeof(float));
|
||||
|
||||
// cv::Mat d(netRT->buffersDIM[1].h, netRT->buffersDIM[1].w, CV_8UC1, depth_h);
|
||||
// depthMats[bi] = d.clone();
|
||||
|
||||
cv::Mat depth_mat = vizData2Mat(rt_out[0], netRT->buffersDIM[1], netRT->buffersDIM[1].h, netRT->buffersDIM[1].w);
|
||||
// cv::Mat depth_mat = vizData2Mat((dnnType *)netRT->buffersRT[0], netRT->buffersDIM[0], netRT->buffersDIM[0].h, netRT->buffersDIM[0].w);
|
||||
depthMats[bi] = depth_mat.clone();
|
||||
|
||||
}
|
||||
|
||||
/**
|
||||
* This method performs the inference of the NN.
|
||||
*
|
||||
* @param frames frames to build the embedding from.
|
||||
* @param cur_batches number of batches to use in inference
|
||||
*/
|
||||
void update(std::vector<cv::Mat>& frames, const int cur_batches=1){
|
||||
if(cur_batches > nBatches)
|
||||
FatalError("A batch size greater than nBatches cannot be used");
|
||||
|
||||
if(TKDNN_VERBOSE) printCenteredTitle(" TENSORRT feature extraction ", '=', 30);
|
||||
{
|
||||
TKDNN_TSTART
|
||||
for(int bi=0; bi<cur_batches;++bi){
|
||||
if(!frames[bi].data)
|
||||
FatalError("No image data feed to extract features");
|
||||
preprocess(frames[bi], bi);
|
||||
}
|
||||
TKDNN_TSTOP
|
||||
}
|
||||
|
||||
//do inference
|
||||
tk::dnn::dataDim_t dim = netRT->input_dim;
|
||||
dim.n = cur_batches;
|
||||
{
|
||||
if(TKDNN_VERBOSE) dim.print();
|
||||
TKDNN_TSTART
|
||||
netRT->infer(dim, input_d);
|
||||
TKDNN_TSTOP
|
||||
if(TKDNN_VERBOSE) dim.print();
|
||||
stats.push_back(t_ns);
|
||||
}
|
||||
|
||||
{
|
||||
TKDNN_TSTART
|
||||
for(int bi=0; bi<cur_batches;++bi)
|
||||
postprocess(bi);
|
||||
TKDNN_TSTOP
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Method to draw the result.
|
||||
*
|
||||
*/
|
||||
void draw() { }
|
||||
|
||||
};
|
||||
|
||||
|
||||
}}
|
||||
|
||||
#endif /* DEPTHNN_H*/
|
||||
+29
-2
@@ -31,7 +31,8 @@ enum layerType_t {
|
||||
LAYER_SHORTCUT,
|
||||
LAYER_UPSAMPLE,
|
||||
LAYER_REGION,
|
||||
LAYER_YOLO
|
||||
LAYER_YOLO,
|
||||
LAYER_PADDING,
|
||||
};
|
||||
|
||||
#define TKDNN_BN_MIN_EPSILON 1e-5
|
||||
@@ -87,6 +88,7 @@ public:
|
||||
case LAYER_UPSAMPLE: return "Upsample";
|
||||
case LAYER_REGION: return "Region";
|
||||
case LAYER_YOLO: return "Yolo";
|
||||
case LAYER_PADDING: return "Padding";
|
||||
default: return "unknown";
|
||||
}
|
||||
}
|
||||
@@ -473,7 +475,6 @@ public:
|
||||
|
||||
virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
|
||||
|
||||
protected:
|
||||
dnnType mul, add;
|
||||
dnnType *add_vector;
|
||||
};
|
||||
@@ -520,9 +521,35 @@ protected:
|
||||
bool poolOn3d;
|
||||
};
|
||||
|
||||
/**
|
||||
* Padding Layers
|
||||
* tkDNN supports reflection,constant and zero padding
|
||||
*/
|
||||
|
||||
typedef enum {
|
||||
PADDING_MODE_CONSTANT = 0,
|
||||
PADDING_MODE_ZERO = 1,
|
||||
PADDING_MODE_REFLECTION = 2
|
||||
} tkdnnPaddingMode_t;
|
||||
|
||||
class Padding : public Layer {
|
||||
public:
|
||||
Padding(Network *net,int32_t pad_h,int32_t pad_w,tkdnnPaddingMode_t padding_mode,float constant = 0.0);
|
||||
virtual ~Padding();
|
||||
virtual layerType_t getLayerType(){return LAYER_PADDING ;};
|
||||
virtual dnnType* infer(dataDim_t& dim,dnnType* srcData);
|
||||
int32_t paddingH,paddingW;
|
||||
tkdnnPaddingMode_t padding_mode;
|
||||
float constant;
|
||||
|
||||
};
|
||||
|
||||
|
||||
|
||||
/**
|
||||
Softmax layer
|
||||
*/
|
||||
|
||||
class Softmax : public Layer {
|
||||
|
||||
public:
|
||||
|
||||
@@ -23,6 +23,8 @@
|
||||
#include <pluginsRT/ShortcutRT.h>
|
||||
#include <pluginsRT/UpsampleRT.h>
|
||||
#include <pluginsRT/YoloRT.h>
|
||||
#include <pluginsRT/ConstantPaddingRT.h>
|
||||
#include <pluginsRT/ReflectionPadding.h>
|
||||
|
||||
|
||||
|
||||
@@ -93,8 +95,10 @@ public:
|
||||
nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Region *l);
|
||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Shortcut *l);
|
||||
nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Yolo *l);
|
||||
nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Upsample *l);
|
||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Upsample *l);
|
||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, DeformConv2d *l);
|
||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input,Padding *l);
|
||||
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor* input,MulAdd *l);
|
||||
|
||||
#if NV_TENSORRT_MAJOR > 5 && NV_TENSORRT_MAJOR < 8
|
||||
bool serialize(const char *filename);
|
||||
|
||||
@@ -6,7 +6,7 @@
|
||||
namespace tk { namespace dnn {
|
||||
|
||||
cv::Mat vizFloat2colorMap(cv::Mat map, double min=0, double max=0, int classes=19);
|
||||
cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int img_h, int img_w, double min=0, double max=0, int classes=19);
|
||||
cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int img_h, int img_w, double min=0, double max=0, int classes=0);
|
||||
cv::Mat vizLayer2Mat(tk::dnn::Network *net, int layer, int imgdim = 1000);
|
||||
|
||||
}}
|
||||
|
||||
@@ -48,4 +48,11 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
|
||||
const int dst_dim, cudaStream_t stream = cudaStream_t(0));
|
||||
|
||||
void scalAdd(dnnType* dstData, int size, float alpha, float beta, int inc, cudaStream_t stream = cudaStream_t(0));
|
||||
|
||||
void reflection_pad2d_out_forward(int32_t pad_h,int32_t pad_w,float *srcData,float *dstData,int32_t input_h,int32_t input_w,int32_t plane_dim,int32_t n_batch,cudaStream_t cudaStream = cudaStream_t(0));
|
||||
|
||||
void constant_pad2d_forward(dnnType *srcData,dnnType *dstData,int32_t input_h,int32_t input_w,int32_t output_h,
|
||||
int32_t output_w,int32_t c,int32_t n,int32_t padT,int32_t padL,dnnType constant,cudaStream_t cudaStream = cudaStream_t(0));
|
||||
|
||||
|
||||
#endif //KERNELS_H
|
||||
|
||||
@@ -0,0 +1,109 @@
|
||||
//
|
||||
// Created by perseusdg on 1/7/22.
|
||||
//
|
||||
|
||||
#ifndef _CONSTANTPADDINGRT_PLUGIN_H
|
||||
#define _CONSTANTPADDINGRT_PLUGIN_H
|
||||
|
||||
#include<cassert>
|
||||
#include <NvInfer.h>
|
||||
#include <vector>
|
||||
#include <utils.h>
|
||||
#include <kernels.h>
|
||||
|
||||
namespace nvinfer1{
|
||||
class ConstantPaddingRT : public IPluginV2Ext {
|
||||
public:
|
||||
ConstantPaddingRT(int32_t padH,int32_t padW,int32_t n,int32_t c,int32_t i_h,int32_t i_w,int32_t o_h,int32_t o_w,float constant);
|
||||
|
||||
ConstantPaddingRT(const void *data,size_t length);
|
||||
|
||||
~ConstantPaddingRT();
|
||||
|
||||
int getNbOutputs() const NOEXCEPT override;
|
||||
|
||||
Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
|
||||
|
||||
int initialize() NOEXCEPT override ;
|
||||
|
||||
void terminate() NOEXCEPT override ;
|
||||
|
||||
size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
|
||||
|
||||
|
||||
#if NV_TENSORRT_MAJOR > 7
|
||||
int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ;
|
||||
#elif NV_TENSORRT_MAJOR <= 7
|
||||
int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
|
||||
#endif
|
||||
|
||||
size_t getSerializationSize() const NOEXCEPT override ;
|
||||
|
||||
void serialize(void *buffer) const NOEXCEPT override ;
|
||||
|
||||
void destroy() NOEXCEPT override ;
|
||||
|
||||
const char *getPluginType() const NOEXCEPT override ;
|
||||
|
||||
const char *getPluginVersion() const NOEXCEPT override;
|
||||
|
||||
const char *getPluginNamespace() const NOEXCEPT override ;
|
||||
|
||||
void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
|
||||
|
||||
IPluginV2Ext *clone() const NOEXCEPT override ;
|
||||
|
||||
DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
|
||||
|
||||
void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
|
||||
|
||||
bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
|
||||
|
||||
bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
|
||||
|
||||
void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
|
||||
int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
|
||||
bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
|
||||
int32_t maxBatchSize) NOEXCEPT override;
|
||||
|
||||
void detachFromContext() NOEXCEPT override;
|
||||
|
||||
bool supportsFormat (DataType type, PluginFormat format) const NOEXCEPT override;
|
||||
|
||||
int32_t i_h,i_w,o_h,o_w,n,c,padH,padW;
|
||||
float constant;
|
||||
private:
|
||||
std::string mPluginNamespace;
|
||||
|
||||
};
|
||||
|
||||
class ConstantPaddingRTPluginCreator : public IPluginCreator {
|
||||
public:
|
||||
ConstantPaddingRTPluginCreator();
|
||||
|
||||
void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override;
|
||||
|
||||
const char *getPluginNamespace() const NOEXCEPT override;
|
||||
|
||||
IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
|
||||
|
||||
IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
|
||||
|
||||
const char *getPluginName() const NOEXCEPT override ;
|
||||
|
||||
const char *getPluginVersion() const NOEXCEPT override;
|
||||
|
||||
const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
|
||||
|
||||
private:
|
||||
static PluginFieldCollection mFC;
|
||||
static std::vector<PluginField> mPluginAttributes;
|
||||
std::string mPluginNamespace;
|
||||
|
||||
};
|
||||
|
||||
REGISTER_TENSORRT_PLUGIN(ConstantPaddingRTPluginCreator);
|
||||
};
|
||||
|
||||
|
||||
#endif //TKDNN_CONSTANTPADDINGRT_H
|
||||
@@ -1,3 +1,6 @@
|
||||
#ifndef _FLATTENCONCATRT_PLUGIN_H
|
||||
#define _FLATTENCONCATRT_PLUGIN_H
|
||||
|
||||
#include<cassert>
|
||||
#include <NvInfer.h>
|
||||
#include <vector>
|
||||
@@ -93,4 +96,5 @@ namespace nvinfer1 {
|
||||
};
|
||||
|
||||
REGISTER_TENSORRT_PLUGIN(FlattenConcatRTPluginCreator);
|
||||
};
|
||||
};
|
||||
#endif
|
||||
@@ -0,0 +1,101 @@
|
||||
#ifndef _REFLECTIONPADDINGRT_PLUGIN_H
|
||||
#define _REFLECTIONPADDINGRT_PLUGIN_H
|
||||
|
||||
#include<cassert>
|
||||
#include <NvInfer.h>
|
||||
#include <vector>
|
||||
#include <utils.h>
|
||||
#include <kernels.h>
|
||||
|
||||
namespace nvinfer1{
|
||||
class ReflectionPaddingRT : public IPluginV2Ext {
|
||||
public:
|
||||
ReflectionPaddingRT(int32_t padH,int32_t padW,int32_t input_h,int32_t input_w,int32_t output_h,int32_t output_w,int32_t c,int32_t n);
|
||||
|
||||
ReflectionPaddingRT(const void *data,size_t length);
|
||||
|
||||
~ReflectionPaddingRT();
|
||||
|
||||
int getNbOutputs() const NOEXCEPT override;
|
||||
|
||||
Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
|
||||
|
||||
int initialize() NOEXCEPT override ;
|
||||
|
||||
void terminate() NOEXCEPT override ;
|
||||
|
||||
size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
|
||||
|
||||
#if NV_TENSORRT_MAJOR > 7
|
||||
int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ;
|
||||
#elif NV_TENSORRT_MAJOR <= 7
|
||||
int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
|
||||
#endif
|
||||
|
||||
size_t getSerializationSize() const NOEXCEPT override ;
|
||||
|
||||
void serialize(void *buffer) const NOEXCEPT override ;
|
||||
|
||||
void destroy() NOEXCEPT override ;
|
||||
|
||||
const char *getPluginType() const NOEXCEPT override ;
|
||||
|
||||
const char *getPluginVersion() const NOEXCEPT override;
|
||||
|
||||
const char *getPluginNamespace() const NOEXCEPT override ;
|
||||
|
||||
void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
|
||||
|
||||
IPluginV2Ext *clone() const NOEXCEPT override ;
|
||||
|
||||
DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
|
||||
|
||||
void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
|
||||
|
||||
bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
|
||||
|
||||
bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
|
||||
|
||||
void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
|
||||
int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
|
||||
bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
|
||||
int32_t maxBatchSize) NOEXCEPT override;
|
||||
|
||||
void detachFromContext() NOEXCEPT override;
|
||||
|
||||
bool supportsFormat (DataType type, PluginFormat format) const NOEXCEPT override;
|
||||
|
||||
int32_t padH,padW,input_h,input_w,output_h,output_w,n,c;
|
||||
private:
|
||||
std::string mPluginNamespace;
|
||||
|
||||
};
|
||||
|
||||
class ReflectionPaddingRTPluginCreator : public IPluginCreator {
|
||||
public:
|
||||
ReflectionPaddingRTPluginCreator();
|
||||
|
||||
void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
|
||||
|
||||
const char *getPluginNamespace() const NOEXCEPT override ;
|
||||
|
||||
IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
|
||||
|
||||
IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
|
||||
|
||||
const char *getPluginName() const NOEXCEPT override ;
|
||||
|
||||
const char *getPluginVersion() const NOEXCEPT override;
|
||||
|
||||
const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
|
||||
|
||||
private:
|
||||
static PluginFieldCollection mFC;
|
||||
static std::vector<PluginField> mPluginAttributes;
|
||||
std::string mPluginNamespace;
|
||||
};
|
||||
|
||||
REGISTER_TENSORRT_PLUGIN(ReflectionPaddingRTPluginCreator);
|
||||
};
|
||||
#endif
|
||||
|
||||
@@ -56,6 +56,9 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
|
||||
else if(act_mode == ACTIVATION_LOGISTIC) {
|
||||
activationLOGISTICForward(srcData, dstData, dim.tot());
|
||||
|
||||
} else if(act_mode == ACTIVATION_ELU) {
|
||||
activationELUForward(srcData, dstData, dim.tot());
|
||||
|
||||
} else {
|
||||
dnnType alpha = dnnType(1);
|
||||
dnnType beta = dnnType(0);
|
||||
|
||||
+202
-43
@@ -26,15 +26,15 @@ class Logger : public ILogger {
|
||||
|
||||
namespace tk { namespace dnn {
|
||||
|
||||
std::map<Layer*, nvinfer1::ITensor*>tensors;
|
||||
std::map<Layer*, nvinfer1::ITensor*>tensors;
|
||||
|
||||
NetworkRT::NetworkRT(Network *net, const char *name) {
|
||||
|
||||
float rt_ver = float(NV_TENSORRT_MAJOR) +
|
||||
float(NV_TENSORRT_MINOR)/10 +
|
||||
float rt_ver = float(NV_TENSORRT_MAJOR) +
|
||||
float(NV_TENSORRT_MINOR)/10 +
|
||||
float(NV_TENSORRT_PATCH)/100;
|
||||
std::cout<<"New NetworkRT (TensorRT v"<<rt_ver<<")\n";
|
||||
|
||||
|
||||
builderRT = createInferBuilder(loggerRT);
|
||||
std::cout<<"Float16 support: "<<builderRT->platformHasFastFp16()<<"\n";
|
||||
std::cout<<"Int8 support: "<<builderRT->platformHasFastInt8()<<"\n";
|
||||
@@ -42,12 +42,12 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
|
||||
std::cout<<"DLAs: "<<builderRT->getNbDLACores()<<"\n";
|
||||
#endif
|
||||
networkRT = builderRT->createNetworkV2(0U);
|
||||
#if NV_TENSORRT_MAJOR >= 6
|
||||
#if NV_TENSORRT_MAJOR >= 6
|
||||
configRT = builderRT->createBuilderConfig();
|
||||
#endif
|
||||
|
||||
|
||||
if(!fileExist(name)) {
|
||||
#if NV_TENSORRT_MAJOR >= 6
|
||||
#if NV_TENSORRT_MAJOR >= 6
|
||||
// Calibrator life time needs to last until after the engine is built.
|
||||
std::unique_ptr<IInt8EntropyCalibrator> calibrator;
|
||||
|
||||
@@ -78,14 +78,14 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
|
||||
configRT->setDLACore(0);
|
||||
}
|
||||
#endif
|
||||
#if NV_TENSORRT_MAJOR >= 6
|
||||
#if NV_TENSORRT_MAJOR >= 6
|
||||
if(net->int8 && builderRT->platformHasFastInt8()){
|
||||
// dtRT = DataType::kINT8;
|
||||
// builderRT->setInt8Mode(true);
|
||||
configRT->setFlag(BuilderFlag::kINT8);
|
||||
BatchStream calibrationStream(dim, 1, 100, //TODO: check if 100 images are sufficient to the calibration (or 4951)
|
||||
BatchStream calibrationStream(dim, 1, 100, //TODO: check if 100 images are sufficient to the calibration (or 4951)
|
||||
net->fileImgList, net->fileLabelList);
|
||||
|
||||
|
||||
/* The calibTableFilePath contains the path+filename of the calibration table.
|
||||
* Each calibration table can be found in the corresponding network folder (../Test/*).
|
||||
* Each network is located in a folder with the same name as the network.
|
||||
@@ -96,15 +96,15 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
|
||||
if(!fileExist((const char *)calib_table_path.c_str()))
|
||||
calib_table_name = "./" + net->networkNameRT.substr(0, net->networkNameRT.find('.')) + "-calibration.table";
|
||||
|
||||
calibrator.reset(new Int8EntropyCalibrator(calibrationStream, 1,
|
||||
calib_table_name,
|
||||
calibrator.reset(new Int8EntropyCalibrator(calibrationStream, 1,
|
||||
calib_table_name,
|
||||
"data"));
|
||||
configRT->setInt8Calibrator(calibrator.get());
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
// add input layer
|
||||
ITensor *input = networkRT->addInput("data", DataType::kFLOAT,
|
||||
ITensor *input = networkRT->addInput("data", DataType::kFLOAT,
|
||||
Dims3{ dim.c, dim.h, dim.w});
|
||||
checkNULL(input);
|
||||
|
||||
@@ -112,17 +112,17 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
|
||||
for(int i=0; i<net->num_layers; i++) {
|
||||
Layer *l = net->layers[i];
|
||||
ILayer *Ilay = convert_layer(input, l);
|
||||
#if NV_TENSORRT_MAJOR >= 6
|
||||
#if NV_TENSORRT_MAJOR >= 6
|
||||
if(net->int8 && builderRT->platformHasFastInt8())
|
||||
{
|
||||
Ilay->setPrecision(DataType::kINT8);
|
||||
}
|
||||
#endif
|
||||
Ilay->setName( (l->getLayerName() + std::to_string(i)).c_str() );
|
||||
|
||||
|
||||
input = Ilay->getOutput(0);
|
||||
input->setName( (l->getLayerName() + std::to_string(i) + "_out").c_str() );
|
||||
|
||||
|
||||
if(l->final)
|
||||
networkRT->markOutput(*input);
|
||||
tensors[l] = input;
|
||||
@@ -182,7 +182,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
|
||||
|
||||
// In order to bind the buffers, we need to know the names of the input and output tensors.
|
||||
// note that indices are guaranteed to be less than IEngine::getNbBindings()
|
||||
buf_input_idx = engineRT->getBindingIndex("data");
|
||||
buf_input_idx = engineRT->getBindingIndex("data");
|
||||
buf_output_idx = engineRT->getBindingIndex("out");
|
||||
std::cout<<"input index = "<<buf_input_idx<<" -> output index = "<<buf_output_idx<<"\n";
|
||||
|
||||
@@ -275,6 +275,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
|
||||
return convert_layer(input, (Upsample*) l);
|
||||
if(type == LAYER_DEFORMCONV2D)
|
||||
return convert_layer(input, (DeformConv2d*) l);
|
||||
if(type == LAYER_PADDING)
|
||||
return convert_layer(input, (Padding*) l);
|
||||
if(type == LAYER_MULADD)
|
||||
return convert_layer(input,(MulAdd*) l);
|
||||
|
||||
std::cout<<l->getLayerName()<<"\n";
|
||||
FatalError("Layer not implemented in tensorRT");
|
||||
@@ -285,10 +289,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Dense *l) {
|
||||
//std::cout<<"convert Dense\n";
|
||||
void *data_b, *bias_b;
|
||||
if(dtRT == DataType::kHALF) {
|
||||
data_b = l->data16_h;
|
||||
data_b = l->data16_h;
|
||||
bias_b = l->bias16_h;
|
||||
} else {
|
||||
data_b = l->data_h;
|
||||
data_b = l->data_h;
|
||||
bias_b = l->bias_h;
|
||||
}
|
||||
|
||||
@@ -308,7 +312,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
|
||||
|
||||
void *data_b, *bias_b, *bias2_b, *power_b, *mean_b, *variance_b, *scales_b;
|
||||
if(dtRT == DataType::kHALF) {
|
||||
data_b = l->data16_h;
|
||||
data_b = l->data16_h;
|
||||
bias_b = l->bias16_h;
|
||||
bias2_b = l->bias216_h;
|
||||
power_b = l->power16_h;
|
||||
@@ -316,7 +320,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
|
||||
variance_b = l->variance16_h;
|
||||
scales_b = l->scales16_h;
|
||||
} else {
|
||||
data_b = l->data_h;
|
||||
data_b = l->data_h;
|
||||
bias_b = l->bias_h;
|
||||
bias2_b = l->bias2_h;
|
||||
power_b = l->power_h;
|
||||
@@ -332,7 +336,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
|
||||
b = { dtRT, bias_b, l->outputs};
|
||||
else{
|
||||
if (l->additional_bias)
|
||||
b = { dtRT, bias2_b, l->outputs};
|
||||
b = { dtRT, bias2_b, l->outputs};
|
||||
else
|
||||
b = { dtRT, nullptr, 0}; //on batchnorm bias are added later
|
||||
}
|
||||
@@ -340,7 +344,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
|
||||
ILayer *lRT = nullptr;
|
||||
#if NV_TENSORRT_MAJOR < 8
|
||||
if(!l->deConv) {
|
||||
IConvolutionLayer *lRTconv = networkRT->addConvolution(*input,
|
||||
IConvolutionLayer *lRTconv = networkRT->addConvolution(*input,
|
||||
l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b);
|
||||
checkNULL(lRTconv);
|
||||
lRTconv->setStride(DimsHW{l->strideH, l->strideW});
|
||||
@@ -348,14 +352,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
|
||||
lRTconv->setNbGroups(l->groups);
|
||||
lRT = (ILayer*) lRTconv;
|
||||
} else {
|
||||
IDeconvolutionLayer *lRTconv = networkRT->addDeconvolution(*input,
|
||||
IDeconvolutionLayer *lRTconv = networkRT->addDeconvolution(*input,
|
||||
l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b);
|
||||
checkNULL(lRTconv);
|
||||
lRTconv->setStride(DimsHW{l->strideH, l->strideW});
|
||||
lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW});
|
||||
lRTconv->setNbGroups(l->groups);
|
||||
lRT = (ILayer*) lRTconv;
|
||||
|
||||
|
||||
Dims d = lRTconv->getOutput(0)->getDimensions();
|
||||
//std::cout<<"DECONV: "<<d.d[0]<<" "<<d.d[1]<<" "<<d.d[2]<<" "<<d.d[3]<<"\n";
|
||||
}
|
||||
@@ -388,14 +392,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
|
||||
Weights shift{dtRT, mean_b, l->outputs};
|
||||
Weights scale{dtRT, variance_b, l->outputs};
|
||||
// std::cout<<lRT->getNbOutputs()<<std::endl;
|
||||
IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL,
|
||||
IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL,
|
||||
shift, scale, power);
|
||||
|
||||
|
||||
checkNULL(lRT2);
|
||||
|
||||
Weights shift2{dtRT, bias_b, l->outputs};
|
||||
Weights scale2{dtRT, scales_b, l->outputs};
|
||||
IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL,
|
||||
IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL,
|
||||
shift2, scale2, power);
|
||||
checkNULL(lRT3);
|
||||
|
||||
@@ -405,6 +409,80 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
|
||||
return lRT;
|
||||
}
|
||||
|
||||
ILayer* NetworkRT::convert_layer(ITensor *input,MulAdd *l){
|
||||
|
||||
void *power_b, *shift_b, *scales_b;
|
||||
int size = l->input_dim.tot();
|
||||
|
||||
power_b = new dnnType[size];
|
||||
shift_b = new dnnType[size];
|
||||
scales_b = new dnnType[size];
|
||||
|
||||
for(int i=0; i<size; i++) {
|
||||
((dnnType*) power_b)[i] = 1.0;
|
||||
((dnnType*) shift_b)[i] = l->add;
|
||||
((dnnType*) scales_b)[i] = l->mul;
|
||||
}
|
||||
|
||||
if(dtRT == DataType::kHALF) {
|
||||
|
||||
__half *power16_h = nullptr, *power16_d = nullptr;
|
||||
__half *scales16_h = nullptr, *scales16_d = nullptr;
|
||||
__half *shift16_h = nullptr, *shift16_d = nullptr;
|
||||
|
||||
dnnType * power_d = nullptr;
|
||||
dnnType * scales_d = nullptr;
|
||||
dnnType * shift_d = nullptr;
|
||||
|
||||
cudaMalloc(&power_d, size*sizeof(dnnType));
|
||||
cudaMemcpy(power_d, power_b, size*sizeof(dnnType), cudaMemcpyHostToDevice);
|
||||
|
||||
cudaMalloc(&shift_d, size*sizeof(dnnType));
|
||||
cudaMemcpy(shift_d, shift_b, size*sizeof(dnnType), cudaMemcpyHostToDevice);
|
||||
|
||||
cudaMalloc(&scales_d, size*sizeof(dnnType));
|
||||
cudaMemcpy(scales_d, scales_b, size*sizeof(dnnType), cudaMemcpyHostToDevice);
|
||||
|
||||
//convert to fp16
|
||||
power16_h = new __half[size];
|
||||
cudaMalloc(&power16_d, size*sizeof(__half));
|
||||
float2half(power_d, power16_d, size);
|
||||
cudaMemcpy(power16_h, power16_d, size*sizeof(__half), cudaMemcpyDeviceToHost);
|
||||
|
||||
shift16_h = new __half[size];
|
||||
cudaMalloc(&shift16_d, size*sizeof(__half));
|
||||
float2half(shift_d, shift16_d, size);
|
||||
cudaMemcpy(shift16_h, shift16_d, size*sizeof(__half), cudaMemcpyDeviceToHost);
|
||||
|
||||
scales16_h = new __half[size];
|
||||
cudaMalloc(&scales16_d, size*sizeof(__half));
|
||||
float2half(scales_d, scales16_d, size);
|
||||
cudaMemcpy(scales16_h, scales16_d, size*sizeof(__half), cudaMemcpyDeviceToHost);
|
||||
|
||||
power_b = power16_h;
|
||||
shift_b = shift16_h;
|
||||
scales_b = scales16_h;
|
||||
|
||||
|
||||
cudaFree(power16_d);
|
||||
cudaFree(shift16_d);
|
||||
cudaFree(scales16_d);
|
||||
|
||||
cudaFree(power_d);
|
||||
cudaFree(shift_d);
|
||||
cudaFree(scales_d);
|
||||
}
|
||||
|
||||
Weights power{dtRT, power_b, size};
|
||||
Weights shift{dtRT, shift_b, size};
|
||||
Weights scale{dtRT, scales_b, size};
|
||||
IScaleLayer *lRT = networkRT->addScale(*input, ScaleMode::kELEMENTWISE,
|
||||
shift, scale, power);
|
||||
checkNULL(lRT);
|
||||
return lRT;
|
||||
}
|
||||
|
||||
|
||||
ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) {
|
||||
// std::cout<<"convert Pooling\n";
|
||||
|
||||
@@ -450,7 +528,68 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) {
|
||||
lRT->setStrideNd(Dims2{l->strideH,l->strideW});
|
||||
return lRT;
|
||||
#endif
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
ILayer* NetworkRT::convert_layer(ITensor *input,Padding *l){
|
||||
|
||||
float rt_ver = float(NV_TENSORRT_MAJOR) +
|
||||
float(NV_TENSORRT_MINOR)/10 +
|
||||
float(NV_TENSORRT_PATCH)/100;
|
||||
|
||||
/*#if ((NV_TENSORRT_MAJOR == 8 && NV_TENSORRT_MINOR >= 2) || NV_TENSORRT_MAJOR > 8)
|
||||
auto *lRT = networkRT->addSlice(*input,Dims3{0,0,0},Dims3{l->output_dim.c,l->output_dim.h,l->output_dim.w},Dims3{0,0,0});
|
||||
if(l->padding_mode == PADDING_MODE_REFLECTION){
|
||||
lRT->setMode(SliceMode::kREFLECT);
|
||||
}else if(l->padding_mode == PADDING_MODE_CONSTANT || l->padding_mode == PADDING_MODE_ZERO){
|
||||
lRT->setMode(SliceMode::kFILL);
|
||||
lRT->setInput(4, reinterpret_cast<ITensor &>(l->constant));
|
||||
}
|
||||
checkNULL(lRT);
|
||||
return lRT;
|
||||
#else*/
|
||||
//todo use ISliceLayer for padding,currently using ISliceLayer for reflection padding generates an error with monodepth2
|
||||
if(l->padding_mode == PADDING_MODE_REFLECTION){
|
||||
auto creator = getPluginRegistry()->getPluginCreator("ReflectionPaddingRT_tkDNN","1");
|
||||
std::vector<PluginField> mPluginAttributes;
|
||||
PluginFieldCollection mFC{};
|
||||
mPluginAttributes.emplace_back(PluginField("padH",&l->paddingH,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("padW",&l->paddingW,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("inputH",&l->input_dim.h,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("inputW",&l->input_dim.w,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("outputH",&l->output_dim.h,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("outputW",&l->output_dim.w,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("n",&l->input_dim.n,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("c",&l->input_dim.c,PluginFieldType::kINT32,1));
|
||||
mFC.nbFields = mPluginAttributes.size();
|
||||
mFC.fields = mPluginAttributes.data();
|
||||
auto *plugin = creator->createPlugin(l->getLayerName().c_str(),&mFC);
|
||||
auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
|
||||
checkNULL(lRT);
|
||||
return lRT;
|
||||
}else if(l->padding_mode == PADDING_MODE_CONSTANT || l->padding_mode == PADDING_MODE_ZERO){
|
||||
auto creator = getPluginRegistry()->getPluginCreator("ConstantPaddingRT_tkDNN","1");
|
||||
std::vector<PluginField> mPluginAttributes;
|
||||
PluginFieldCollection mFC{};
|
||||
mPluginAttributes.emplace_back(PluginField("padH",&l->paddingH,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("padW",&l->paddingW,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("inputH",&l->input_dim.h,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("inputW",&l->input_dim.w,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("outputH",&l->output_dim.h,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("outputW",&l->output_dim.w,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("n",&l->input_dim.n,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("c",&l->input_dim.c,PluginFieldType::kINT32,1));
|
||||
mPluginAttributes.emplace_back(PluginField("constant",&l->constant,PluginFieldType::kFLOAT32,1));
|
||||
mFC.nbFields = mPluginAttributes.size();
|
||||
mFC.fields = mPluginAttributes.data();
|
||||
auto *plugin = creator->createPlugin(l->getLayerName().c_str(),&mFC);
|
||||
auto *lRT = networkRT->addPluginV2(&input,1,*plugin);
|
||||
checkNULL(lRT);
|
||||
return lRT;
|
||||
}
|
||||
|
||||
return nullptr;
|
||||
|
||||
}
|
||||
|
||||
ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
|
||||
@@ -458,14 +597,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
|
||||
|
||||
if(l->act_mode == ACTIVATION_LEAKY) {
|
||||
//std::cout<<"New plugin LEAKY\n";
|
||||
|
||||
#if NV_TENSORRT_MAJOR < 6
|
||||
|
||||
#if NV_TENSORRT_MAJOR < 6
|
||||
// plugin version
|
||||
IPlugin *plugin = new ActivationLeakyRT(l->slope);
|
||||
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
|
||||
checkNULL(lRT);
|
||||
return lRT;
|
||||
#else
|
||||
#else
|
||||
IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kLEAKY_RELU);
|
||||
lRT->setAlpha(l->slope);
|
||||
checkNULL(lRT);
|
||||
@@ -490,7 +629,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
|
||||
//IPluginV2Layer *lRT = networkRT->addPluginV2(&input, 1, *plugin);
|
||||
//checkNULL(lRT);
|
||||
return lRT;
|
||||
}
|
||||
}
|
||||
else if(l->act_mode == ACTIVATION_MISH) {
|
||||
IActivationLayer *lRT1 = networkRT->addActivation(*input, ActivationType::kSOFTPLUS);
|
||||
lRT1->setAlpha(1);
|
||||
@@ -504,6 +643,11 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
|
||||
checkNULL(lRT);
|
||||
return lRT;
|
||||
}
|
||||
else if(l->act_mode == CUDNN_ACTIVATION_ELU || l->act_mode == ACTIVATION_ELU){
|
||||
IActivationLayer *lRT = networkRT->addActivation(*input,ActivationType::kELU);
|
||||
checkNULL(lRT);
|
||||
return lRT;
|
||||
}
|
||||
else {
|
||||
FatalError("this Activation mode is not yet implemented");
|
||||
return NULL;
|
||||
@@ -521,7 +665,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Softmax *l) {
|
||||
|
||||
ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) {
|
||||
// std::cout<<"convert route\n";
|
||||
|
||||
|
||||
|
||||
|
||||
ITensor **tens = new ITensor*[l->layers_n];
|
||||
@@ -633,10 +777,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) {
|
||||
//std::cout<<"convert Shortcut\n";
|
||||
|
||||
//std::cout<<"New plugin Shortcut\n";
|
||||
|
||||
|
||||
ITensor *back_tens = tensors[l->backLayer];
|
||||
|
||||
if(l->backLayer->output_dim.c == l->output_dim.c && !l->mul)
|
||||
if(l->backLayer->output_dim.c == l->output_dim.c && !l->mul)
|
||||
{
|
||||
IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM);
|
||||
checkNULL(lRT);
|
||||
@@ -660,7 +804,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) {
|
||||
auto *plugin = creator->createPlugin(l->getLayerName().c_str(),&mFC);
|
||||
auto **inputs = new ITensor*[2];
|
||||
inputs[0] = input;
|
||||
inputs[1] = back_tens;
|
||||
inputs[1] = back_tens;
|
||||
auto *lRT = networkRT->addPluginV2(inputs, 2, *plugin);
|
||||
checkNULL(lRT);
|
||||
return lRT;
|
||||
@@ -690,8 +834,9 @@ IPluginV2Layer* NetworkRT::convert_layer(ITensor *input, Yolo *l) {
|
||||
return lRT;
|
||||
}
|
||||
|
||||
IPluginV2Layer* NetworkRT::convert_layer(ITensor *input, Upsample *l) {
|
||||
//std::cout<<"convert Upsample\n";
|
||||
ILayer* NetworkRT::convert_layer(ITensor *input, Upsample *l) {
|
||||
|
||||
#if NV_TENSORRT_MAJOR < 8
|
||||
auto creator = getPluginRegistry()->getPluginCreator("UpSample_tkDNN","1");
|
||||
std::vector<PluginField> mPluginAttributes;
|
||||
PluginFieldCollection mFC{};
|
||||
@@ -705,6 +850,13 @@ IPluginV2Layer* NetworkRT::convert_layer(ITensor *input, Upsample *l) {
|
||||
auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
|
||||
checkNULL(lRT);
|
||||
return lRT;
|
||||
#else
|
||||
auto *lRT = networkRT->addResize(*input);
|
||||
lRT->setResizeMode(ResizeMode::kNEAREST);
|
||||
lRT->setOutputDimensions(Dims3{l->output_dim.c, l->output_dim.h, l->output_dim.w});
|
||||
checkNULL(lRT);
|
||||
return lRT;
|
||||
#endif
|
||||
}
|
||||
|
||||
ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {
|
||||
@@ -787,14 +939,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {
|
||||
Weights shift{dtRT, mean_b, l->outputs};
|
||||
Weights scale{dtRT, variance_b, l->outputs};
|
||||
//std::cout<<lRT->getNbOutputs()<<std::endl;
|
||||
IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL,
|
||||
IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL,
|
||||
shift, scale, power);
|
||||
|
||||
|
||||
checkNULL(lRT2);
|
||||
|
||||
Weights shift2{dtRT, bias_b, l->outputs};
|
||||
Weights scale2{dtRT, scales_b, l->outputs};
|
||||
IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL,
|
||||
IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL,
|
||||
shift2, scale2, power);
|
||||
checkNULL(lRT3);
|
||||
|
||||
@@ -856,6 +1008,7 @@ bool NetworkRT::deserialize(const char *filename) {
|
||||
return true;
|
||||
}
|
||||
|
||||
#if NV_TENSORRT_MAJOR > 7
|
||||
void NetworkRT::destroy() {
|
||||
delete contextRT;
|
||||
if(builderActive) {
|
||||
@@ -863,5 +1016,11 @@ void NetworkRT::destroy() {
|
||||
delete builderRT;
|
||||
}
|
||||
}
|
||||
#elif NV_TENSORRT_MAJOR <=7
|
||||
void NetworkRT::destroy() {
|
||||
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
}}
|
||||
|
||||
+1
-1
@@ -383,7 +383,7 @@ cv::Mat vizFloat2colorMap(cv::Mat map,double min, double max, int classes) {
|
||||
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);
|
||||
applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_PARULA);
|
||||
}
|
||||
return falseColorsMap;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,45 @@
|
||||
//
|
||||
// Created by perseusdg on 03/01/22.
|
||||
//
|
||||
|
||||
#include <iostream>
|
||||
#include "Layer.h"
|
||||
#include "kernels.h"
|
||||
|
||||
namespace tk{ namespace dnn {
|
||||
Padding::Padding(Network *net, int32_t pad_h, int32_t pad_w, tkdnnPaddingMode_t padding_mode,float constant) : Layer(net) {
|
||||
this->paddingH = pad_h;
|
||||
this->paddingW = pad_w;
|
||||
this->padding_mode = padding_mode;
|
||||
output_dim.c = input_dim.c;
|
||||
output_dim.n = input_dim.n;
|
||||
output_dim.h = input_dim.h + 2 * (this->paddingH);
|
||||
output_dim.w = input_dim.w + 2 * (this->paddingW);
|
||||
if(padding_mode == tkdnnPaddingMode_t::PADDING_MODE_CONSTANT){
|
||||
this->constant = constant;
|
||||
}else{
|
||||
this->constant = 0;
|
||||
}
|
||||
checkCuda(cudaMalloc(&dstData,output_dim.tot()*sizeof(dnnType)));
|
||||
}
|
||||
|
||||
Padding::~Padding() {
|
||||
checkCuda(cudaFree(dstData));
|
||||
}
|
||||
dnnType* Padding::infer(dataDim_t &dim, float *srcData) {
|
||||
fill(dstData,output_dim.tot(),0.0);
|
||||
if(padding_mode == tkdnnPaddingMode_t::PADDING_MODE_REFLECTION)
|
||||
{
|
||||
reflection_pad2d_out_forward(paddingH, paddingW, srcData, dstData, input_dim.h, input_dim.w, input_dim.c,
|
||||
input_dim.n);
|
||||
}
|
||||
else if(padding_mode == tkdnnPaddingMode_t::PADDING_MODE_CONSTANT){
|
||||
constant_pad2d_forward(srcData,dstData,input_dim.h,input_dim.w,output_dim.h,output_dim.w,input_dim.c,
|
||||
input_dim.n,paddingH,paddingW,constant);
|
||||
}
|
||||
|
||||
dim = output_dim;
|
||||
return dstData;
|
||||
}
|
||||
|
||||
}}
|
||||
@@ -0,0 +1,110 @@
|
||||
#include "kernels.h"
|
||||
#include <thrust/pair.h>
|
||||
#include <stdio.h>
|
||||
|
||||
/*
|
||||
* Reflection padding is from https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/native/cuda/ReflectionPad.cu
|
||||
*/
|
||||
__device__
|
||||
inline thrust::pair<int32_t,int32_t> get_index_mapping2d(
|
||||
int32_t input_dim_x,int32_t input_dim_y,int32_t output_dim_x,
|
||||
int32_t output_dim_y,int32_t pad_l,int32_t pad_t,int32_t output_xy,
|
||||
int32_t y_shift,int32_t z_shift,int32_t n_plane){
|
||||
auto input_offset = ((blockIdx.y + y_shift) + (blockIdx.z + z_shift)*n_plane)*input_dim_x*input_dim_y;
|
||||
auto output_offset = ((blockIdx.y + y_shift) + (blockIdx.z + z_shift)*n_plane)*output_dim_x*output_dim_y;
|
||||
auto output_x = output_xy % output_dim_x;
|
||||
auto output_y = output_xy/output_dim_x;
|
||||
|
||||
auto i_start_x = ::max(int32_t(0),-pad_l);
|
||||
auto i_start_y = ::max(int32_t(0),-pad_t);
|
||||
auto o_start_x = ::max(int32_t(0),pad_l);
|
||||
auto o_start_y = ::max(int32_t(0),pad_t);
|
||||
|
||||
auto input_x = ::abs(output_x - pad_l) - ::abs(output_x - (input_dim_x + pad_l -1)) -output_x + 2*pad_l + input_dim_x -1 -o_start_x + i_start_x;
|
||||
auto input_y = ::abs(output_y - pad_t) - ::abs(output_y - (input_dim_y + pad_t -1)) -output_y + 2*pad_t + input_dim_y -1 -o_start_y + i_start_y;
|
||||
|
||||
return thrust::make_pair<int32_t,int32_t>(input_offset + input_y*input_dim_x + input_x,output_offset + output_y*output_dim_x+output_x);
|
||||
}
|
||||
|
||||
__global__
|
||||
void reflection_pad2d_out_kernel(
|
||||
float* input,float* output,int32_t input_dim_x,
|
||||
int32_t input_dim_y,int32_t pad_t,int32_t pad_b,int32_t pad_l,
|
||||
int32_t pad_r,int32_t y_shift,int32_t z_shift,int32_t n_plane){
|
||||
auto output_xy = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
auto output_dim_x = input_dim_x + pad_l + pad_r;
|
||||
auto output_dim_y = input_dim_y + pad_t + pad_b;
|
||||
|
||||
if(output_xy < output_dim_x*output_dim_y){
|
||||
auto index_pair = get_index_mapping2d(input_dim_x,input_dim_y,output_dim_x,output_dim_y,pad_l,pad_t,output_xy,y_shift,z_shift,n_plane);
|
||||
output[index_pair.second] = input[index_pair.first];
|
||||
}
|
||||
}
|
||||
|
||||
int32_t ceilDiv(int32_t a,int32_t b){
|
||||
return (a+b-1)/b;
|
||||
}
|
||||
|
||||
|
||||
void reflection_pad2d_out_forward(int32_t pad_h,int32_t pad_w,float *srcData,float *dstData,int32_t input_h,int32_t input_w,int32_t plane_dim,int32_t n_batch,cudaStream_t cudaStream){
|
||||
int32_t pad_l = pad_w;
|
||||
int32_t pad_r = pad_w;
|
||||
int32_t pad_t = pad_h;
|
||||
int32_t pad_b = pad_w;
|
||||
int32_t output_h = input_h + pad_t + pad_b;
|
||||
int32_t output_w = input_w + pad_l + pad_r;
|
||||
int32_t size_y = plane_dim;
|
||||
int32_t size_z = n_batch;
|
||||
int32_t output_plane_size = output_h*output_w;
|
||||
dim3 block_size(output_plane_size>256 ?256:output_plane_size);
|
||||
for(int32_t block_y=0;block_y<size_y;block_y += 65535){
|
||||
int32_t block_y_size = std::min(size_y - block_y,static_cast<int32_t>(65535));
|
||||
for(int32_t block_z=0;block_z<size_z;block_z += 65535){
|
||||
int32_t block_z_size = std::min(size_z -block_z,static_cast<int32_t>(65535));
|
||||
|
||||
dim3 grid_size(ceilDiv(output_plane_size,static_cast<int32_t>(256)),block_y_size,block_z_size);
|
||||
reflection_pad2d_out_kernel<<<grid_size,block_size,0,cudaStream>>>(srcData,dstData,input_w,input_h,pad_t,pad_b,pad_l,pad_r,block_y,block_z,plane_dim);
|
||||
}
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
/*
|
||||
* constant padding is inspired from https://github.com/apache/incubator-mxnet/blob/master/src/operator/pad.cu
|
||||
*/
|
||||
|
||||
__global__
|
||||
void constant_pad2d_kernel(dnnType *srcData,dnnType *dstData,const int32_t padT,const int32_t padL,float constant,int32_t n,int32_t c,int32_t i_h,int32_t i_w,int32_t o_h,int32_t o_w){
|
||||
int outputPointId = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
if(outputPointId >= o_h*o_w){
|
||||
return ;
|
||||
}
|
||||
|
||||
int Ny = i_h;
|
||||
int Nx = i_w;
|
||||
|
||||
int plane = blockIdx.y;
|
||||
int batch = blockIdx.z;
|
||||
int outputPointX = outputPointId % o_w;
|
||||
int outputPointY = outputPointId / o_w;
|
||||
int checkT = max(0, outputPointY - padT + 1);
|
||||
int checkB = max(0, padT + Ny - outputPointY);
|
||||
int checkL = max(0, outputPointX - padL + 1);
|
||||
int checkR = max(0, padL + Nx - outputPointX);
|
||||
int inputPointX = min(max(outputPointX - padL, 0), Nx - 1);
|
||||
int inputPointY = min(max(outputPointY - padT, 0), Ny - 1);
|
||||
int need_pad = !(checkT * checkB * checkL * checkR);
|
||||
float value_to_copy = srcData[batch*c*i_h*i_w + plane*i_h*i_w + inputPointY*i_w + inputPointX];
|
||||
dstData[batch*c*o_w*o_h + plane*o_h*o_w + outputPointY*o_w + outputPointX] = value_to_copy * (!need_pad) + need_pad*constant;
|
||||
|
||||
}
|
||||
|
||||
void constant_pad2d_forward(dnnType *srcData,dnnType *dstData,int32_t input_h,int32_t input_w,int32_t output_h,
|
||||
int32_t output_w,int32_t c,int32_t n,int32_t padT,int32_t padL,dnnType constant,cudaStream_t cudaStream){
|
||||
int32_t output_plane_size = output_h*output_w;
|
||||
dim3 block_size(output_plane_size>256 ?256:output_plane_size);
|
||||
dim3 grid_size(ceilDiv(output_plane_size,static_cast<int32_t>(256)),c,n);
|
||||
constant_pad2d_kernel<<<grid_size,block_size,0,cudaStream>>>(srcData,dstData,padT,padL,constant,n,c,input_h,input_w,output_h,output_w);
|
||||
|
||||
}
|
||||
|
||||
@@ -0,0 +1,204 @@
|
||||
#include <tkDNN/pluginsRT/ConstantPaddingRT.h>
|
||||
|
||||
using namespace nvinfer1;
|
||||
|
||||
std::vector<PluginField> ConstantPaddingRTPluginCreator::mPluginAttributes;
|
||||
PluginFieldCollection ConstantPaddingRTPluginCreator::mFC{};
|
||||
|
||||
static const char* CONSTANTPADDINGRT_PLUGIN_VERSION{"1"};
|
||||
static const char* CONSTANTPADDINGRT_PLUGIN_NAME{"ConstantPaddingRT_tkDNN"};
|
||||
|
||||
ConstantPaddingRT::ConstantPaddingRT(int32_t padH, int32_t padW, int32_t n, int32_t c, int32_t i_h, int32_t i_w,
|
||||
int32_t o_h, int32_t o_w, float constant) {
|
||||
this->padH = padH;
|
||||
this->padW = padW;
|
||||
this->n = n;
|
||||
this->c = c;
|
||||
this->i_h = i_h;
|
||||
this->i_w = i_w;
|
||||
this->o_h = o_h;
|
||||
this->o_w = o_w;
|
||||
this->constant = constant;
|
||||
|
||||
}
|
||||
|
||||
ConstantPaddingRT::ConstantPaddingRT(const void *data, size_t length) {
|
||||
const char* buf = reinterpret_cast<const char*>(data),*bufcheck=buf;
|
||||
padH = readBUF<int32_t>(buf);
|
||||
padW = readBUF<int32_t>(buf);
|
||||
i_h = readBUF<int32_t>(buf);
|
||||
i_w = readBUF<int32_t>(buf);
|
||||
o_h = readBUF<int32_t>(buf);
|
||||
o_w = readBUF<int32_t>(buf);
|
||||
n = readBUF<int32_t>(buf);
|
||||
c = readBUF<int32_t>(buf);
|
||||
constant = readBUF<float>(buf);
|
||||
assert(buf = bufcheck + length);
|
||||
}
|
||||
|
||||
ConstantPaddingRT::~ConstantPaddingRT() {}
|
||||
|
||||
int ConstantPaddingRT::getNbOutputs() const NOEXCEPT{
|
||||
return 1;
|
||||
}
|
||||
|
||||
Dims ConstantPaddingRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
|
||||
return Dims3{c,o_h,o_w};
|
||||
}
|
||||
|
||||
int ConstantPaddingRT::initialize() NOEXCEPT {
|
||||
return 0;
|
||||
}
|
||||
|
||||
void ConstantPaddingRT::terminate() NOEXCEPT {
|
||||
|
||||
}
|
||||
|
||||
size_t ConstantPaddingRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
|
||||
return 0;
|
||||
}
|
||||
|
||||
#if NV_TENSORRT_MAJOR > 7
|
||||
int ConstantPaddingRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT {
|
||||
dnnType* srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
|
||||
dnnType* dstData = reinterpret_cast<dnnType*>(outputs[0]);
|
||||
constant_pad2d_forward(srcData,dstData,i_h,i_w,o_h,o_w,c,n,padH,padW,constant,stream);
|
||||
return 0;
|
||||
}
|
||||
#elif NV_TENSORRT_MAJOR <= 7
|
||||
int32_t ConstantPaddingRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace,
|
||||
cudaStream_t stream) {
|
||||
dnnType* srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
|
||||
dnnType* dstData = reinterpret_cast<dnnType*>(outputs[0]);
|
||||
constant_pad2d_forward(srcData,dstData,i_h,i_w,o_h,o_w,c,n,padH,padW,constant,stream);
|
||||
return 0;
|
||||
}
|
||||
#endif
|
||||
|
||||
size_t ConstantPaddingRT::getSerializationSize() const NOEXCEPT {
|
||||
return (8*sizeof(int32_t) + 1*sizeof(float));
|
||||
}
|
||||
|
||||
void ConstantPaddingRT::serialize(void *buffer) const NOEXCEPT {
|
||||
char *buf = reinterpret_cast<char*>(buffer),*a=buf;
|
||||
writeBUF(buf,padH);
|
||||
writeBUF(buf,padW);
|
||||
writeBUF(buf,i_h);
|
||||
writeBUF(buf,i_w);
|
||||
writeBUF(buf,o_h);
|
||||
writeBUF(buf,o_w);
|
||||
writeBUF(buf,n);
|
||||
writeBUF(buf,c);
|
||||
writeBUF(buf,constant);
|
||||
}
|
||||
|
||||
void ConstantPaddingRT::destroy() NOEXCEPT {
|
||||
delete this;
|
||||
}
|
||||
|
||||
const char* ConstantPaddingRT::getPluginType() const NOEXCEPT {
|
||||
return CONSTANTPADDINGRT_PLUGIN_NAME;
|
||||
}
|
||||
|
||||
const char* ConstantPaddingRT::getPluginVersion() const NOEXCEPT {
|
||||
return CONSTANTPADDINGRT_PLUGIN_VERSION;
|
||||
}
|
||||
|
||||
const char* ConstantPaddingRT::getPluginNamespace() const NOEXCEPT {
|
||||
return mPluginNamespace.c_str();
|
||||
}
|
||||
|
||||
void ConstantPaddingRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
|
||||
mPluginNamespace = pluginNamespace;
|
||||
}
|
||||
|
||||
IPluginV2Ext *ConstantPaddingRT::clone() const NOEXCEPT {
|
||||
auto *p = new ConstantPaddingRT(padH,padW,n,c,i_h,i_w,o_h,o_w,constant);
|
||||
p->setPluginNamespace(mPluginNamespace.c_str());
|
||||
return p;
|
||||
}
|
||||
|
||||
DataType ConstantPaddingRT::getOutputDataType(int index, const nvinfer1::DataType *inputTypes,
|
||||
int nbInputs) const NOEXCEPT {
|
||||
return DataType::kFLOAT;
|
||||
}
|
||||
|
||||
void ConstantPaddingRT::attachToContext(cudnnContext *cudnnContext, cublasContext *cublasContext,
|
||||
IGpuAllocator *gpuAllocator) NOEXCEPT {
|
||||
|
||||
}
|
||||
|
||||
bool ConstantPaddingRT::isOutputBroadcastAcrossBatch(int outputIndex, const bool *inputIsBroadcasted,
|
||||
int nbInputs) const NOEXCEPT {
|
||||
return false;
|
||||
}
|
||||
|
||||
bool ConstantPaddingRT::canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT {
|
||||
return false;
|
||||
}
|
||||
|
||||
void ConstantPaddingRT::configurePlugin(const Dims *inputDims, int32_t nbInputs, const Dims *outputDims,
|
||||
int32_t nbOutputs, const DataType *inputTypes, const DataType *outputTypes,
|
||||
const bool *inputIsBroadcast, const bool *outputIsBroadcast,
|
||||
PluginFormat floatFormat, int32_t maxBatchSize) NOEXCEPT {
|
||||
|
||||
}
|
||||
|
||||
void ConstantPaddingRT::detachFromContext() NOEXCEPT {
|
||||
|
||||
}
|
||||
|
||||
bool ConstantPaddingRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
|
||||
return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR);
|
||||
}
|
||||
|
||||
|
||||
|
||||
ConstantPaddingRTPluginCreator::ConstantPaddingRTPluginCreator() {
|
||||
mPluginAttributes.clear();
|
||||
mFC.nbFields = mPluginAttributes.size();
|
||||
mFC.fields = mPluginAttributes.data();
|
||||
}
|
||||
|
||||
void ConstantPaddingRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
|
||||
mPluginNamespace = pluginNamespace;
|
||||
}
|
||||
|
||||
const char *ConstantPaddingRTPluginCreator::getPluginNamespace() const NOEXCEPT {
|
||||
return mPluginNamespace.c_str();
|
||||
}
|
||||
|
||||
IPluginV2Ext *ConstantPaddingRTPluginCreator::deserializePlugin(const char *name, const void *serialData,
|
||||
size_t serialLength) NOEXCEPT {
|
||||
auto *pluginObj = new ConstantPaddingRT(serialData,serialLength);
|
||||
pluginObj->setPluginNamespace(mPluginNamespace.c_str());
|
||||
return pluginObj;
|
||||
}
|
||||
|
||||
IPluginV2Ext *ConstantPaddingRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
|
||||
const PluginField *fields = fc->fields;
|
||||
int padH = *(static_cast<const int32_t*>(fields[0].data));
|
||||
int padW = *(static_cast<const int32_t*>(fields[1].data));
|
||||
int inputH = *(static_cast<const int32_t*>(fields[2].data));
|
||||
int inputW = *(static_cast<const int32_t*>(fields[3].data));
|
||||
int outputH = *(static_cast<const int32_t*>(fields[4].data));
|
||||
int outputW = *(static_cast<const int32_t*>(fields[5].data));
|
||||
int n = *(static_cast<const int32_t*>(fields[6].data));
|
||||
int c = *(static_cast<const int32_t*>(fields[7].data));
|
||||
float constant = *(static_cast<const float*>(fields[8].data));
|
||||
auto *pluginObj = new ConstantPaddingRT(padH,padW,n,c,inputH,inputW,outputH,outputW,constant);
|
||||
pluginObj->setPluginNamespace(mPluginNamespace.c_str());
|
||||
return pluginObj;
|
||||
}
|
||||
|
||||
const char *ConstantPaddingRTPluginCreator::getPluginName() const NOEXCEPT {
|
||||
return CONSTANTPADDINGRT_PLUGIN_NAME;
|
||||
}
|
||||
|
||||
const char *ConstantPaddingRTPluginCreator::getPluginVersion() const NOEXCEPT {
|
||||
return CONSTANTPADDINGRT_PLUGIN_VERSION;
|
||||
}
|
||||
|
||||
const PluginFieldCollection *ConstantPaddingRTPluginCreator::getFieldNames() NOEXCEPT {
|
||||
return &mFC;
|
||||
}
|
||||
@@ -0,0 +1,202 @@
|
||||
#include <tkDNN/pluginsRT/ReflectionPadding.h>
|
||||
using namespace nvinfer1;
|
||||
|
||||
std::vector<PluginField> ReflectionPaddingRTPluginCreator::mPluginAttributes;
|
||||
PluginFieldCollection ReflectionPaddingRTPluginCreator::mFC{};
|
||||
|
||||
static const char* REFLECTIONPADDINGRT_PLUGIN_VERSION{"1"};
|
||||
static const char* REFLECTIONPADDINGRT_PLUGIN_NAME{"ReflectionPaddingRT_tkDNN"};
|
||||
|
||||
ReflectionPaddingRT::ReflectionPaddingRT(int32_t padH, int32_t padW, int32_t input_h, int32_t input_w, int32_t output_h,
|
||||
int32_t output_w, int32_t c, int32_t n) {
|
||||
this->padH = padH;
|
||||
this->padW = padW;
|
||||
this->input_h = input_h;
|
||||
this->input_w = input_w;
|
||||
this->output_h = output_h;
|
||||
this->output_w = output_w;
|
||||
this->n = n;
|
||||
this->c = c;
|
||||
}
|
||||
|
||||
ReflectionPaddingRT::ReflectionPaddingRT(const void *data, size_t length) {
|
||||
const char* buf = reinterpret_cast<const char*>(data),*bufcheck=buf;
|
||||
padH = readBUF<int32_t>(buf);
|
||||
padW = readBUF<int32_t>(buf);
|
||||
input_h = readBUF<int32_t>(buf);
|
||||
input_w = readBUF<int32_t>(buf);
|
||||
output_h = readBUF<int32_t>(buf);
|
||||
output_w = readBUF<int32_t>(buf);
|
||||
n = readBUF<int32_t>(buf);
|
||||
c = readBUF<int32_t>(buf);
|
||||
assert(buf = bufcheck + length);
|
||||
}
|
||||
|
||||
ReflectionPaddingRT::~ReflectionPaddingRT() {}
|
||||
|
||||
int ReflectionPaddingRT::getNbOutputs() const NOEXCEPT {
|
||||
return 1;
|
||||
}
|
||||
|
||||
Dims ReflectionPaddingRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
|
||||
return Dims3{c,output_h,output_w};
|
||||
}
|
||||
|
||||
int ReflectionPaddingRT::initialize() NOEXCEPT {
|
||||
return 0;
|
||||
}
|
||||
|
||||
void ReflectionPaddingRT::terminate() NOEXCEPT {
|
||||
|
||||
}
|
||||
|
||||
size_t ReflectionPaddingRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
|
||||
return 0;
|
||||
}
|
||||
|
||||
#if NV_TENSORRT_MAJOR > 7
|
||||
int ReflectionPaddingRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
|
||||
cudaStream_t stream) NOEXCEPT {
|
||||
dnnType* srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
|
||||
dnnType* dstData = reinterpret_cast<dnnType*>(outputs[0]);
|
||||
reflection_pad2d_out_forward(padH,padW,srcData,dstData,input_h,input_w,c,n,stream);
|
||||
return 0;
|
||||
}
|
||||
|
||||
#elif NV_TENSORRT_MAJOR <= 7
|
||||
int32_t ReflectionPaddingRT::enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream){
|
||||
dnnType* srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
|
||||
dnnType* dstData = reinterpret_cast<dnnType*>(outputs[0]);
|
||||
reflection_pad2d_out_forward(padH,padW,srcData,dstData,input_h,input_w,c,n,stream);
|
||||
return 0;
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
size_t ReflectionPaddingRT::getSerializationSize() const NOEXCEPT {
|
||||
return 8*sizeof(int32_t);
|
||||
}
|
||||
|
||||
void ReflectionPaddingRT::serialize(void *buffer) const NOEXCEPT {
|
||||
char *buf = reinterpret_cast<char*>(buffer),*a=buf;
|
||||
writeBUF(buf,padH);
|
||||
writeBUF(buf,padW);
|
||||
writeBUF(buf,input_h);
|
||||
writeBUF(buf,input_w);
|
||||
writeBUF(buf,output_h);
|
||||
writeBUF(buf,output_w);
|
||||
writeBUF(buf,n);
|
||||
writeBUF(buf,c);
|
||||
}
|
||||
|
||||
void ReflectionPaddingRT::destroy() NOEXCEPT {
|
||||
delete this;
|
||||
}
|
||||
|
||||
const char *ReflectionPaddingRT::getPluginType() const NOEXCEPT {
|
||||
return REFLECTIONPADDINGRT_PLUGIN_NAME;
|
||||
}
|
||||
|
||||
const char *ReflectionPaddingRT::getPluginVersion() const NOEXCEPT {
|
||||
return REFLECTIONPADDINGRT_PLUGIN_VERSION;
|
||||
}
|
||||
|
||||
const char *ReflectionPaddingRT::getPluginNamespace() const NOEXCEPT {
|
||||
return mPluginNamespace.c_str();
|
||||
}
|
||||
|
||||
void ReflectionPaddingRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
|
||||
mPluginNamespace = pluginNamespace;
|
||||
}
|
||||
|
||||
IPluginV2Ext *ReflectionPaddingRT::clone() const NOEXCEPT {
|
||||
auto *p = new ReflectionPaddingRT(padH,padW,input_h,input_w,output_h,output_w,c,n);
|
||||
p->setPluginNamespace(mPluginNamespace.c_str());
|
||||
return p;
|
||||
}
|
||||
|
||||
DataType
|
||||
ReflectionPaddingRT::getOutputDataType(int index, const nvinfer1::DataType *inputTypes, int nbInputs) const NOEXCEPT {
|
||||
return DataType::kFLOAT;
|
||||
}
|
||||
|
||||
void ReflectionPaddingRT::attachToContext(cudnnContext *cudnnContext, cublasContext *cublasContext,
|
||||
IGpuAllocator *gpuAllocator) NOEXCEPT {
|
||||
}
|
||||
|
||||
bool ReflectionPaddingRT::isOutputBroadcastAcrossBatch(int outputIndex, const bool *inputIsBroadcasted,
|
||||
int nbInputs) const NOEXCEPT {
|
||||
return false;
|
||||
}
|
||||
|
||||
bool ReflectionPaddingRT::canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT {
|
||||
return false;
|
||||
}
|
||||
|
||||
void
|
||||
ReflectionPaddingRT::configurePlugin(const Dims *inputDims, int32_t nbInputs, const Dims *outputDims, int32_t nbOutputs,
|
||||
const DataType *inputTypes, const DataType *outputTypes,
|
||||
const bool *inputIsBroadcast, const bool *outputIsBroadcast,
|
||||
PluginFormat floatFormat, int32_t maxBatchSize) NOEXCEPT {
|
||||
|
||||
}
|
||||
|
||||
void ReflectionPaddingRT::detachFromContext() NOEXCEPT {
|
||||
|
||||
}
|
||||
|
||||
bool ReflectionPaddingRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
|
||||
return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR);
|
||||
}
|
||||
|
||||
|
||||
ReflectionPaddingRTPluginCreator::ReflectionPaddingRTPluginCreator() {
|
||||
mPluginAttributes.clear();
|
||||
mFC.nbFields = mPluginAttributes.size();
|
||||
mFC.fields = mPluginAttributes.data();
|
||||
}
|
||||
|
||||
void ReflectionPaddingRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
|
||||
mPluginNamespace = pluginNamespace;
|
||||
}
|
||||
|
||||
const char *ReflectionPaddingRTPluginCreator::getPluginNamespace() const NOEXCEPT {
|
||||
return mPluginNamespace.c_str();
|
||||
}
|
||||
|
||||
IPluginV2Ext *ReflectionPaddingRTPluginCreator::deserializePlugin(const char *name, const void *serialData,
|
||||
size_t serialLength) NOEXCEPT {
|
||||
auto *pluginObj = new ReflectionPaddingRT(serialData,serialLength);
|
||||
pluginObj->setPluginNamespace(mPluginNamespace.c_str());
|
||||
return pluginObj;
|
||||
}
|
||||
|
||||
IPluginV2Ext *
|
||||
ReflectionPaddingRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
|
||||
const PluginField *fields = fc->fields;
|
||||
int padH = *(static_cast<const int32_t*>(fields[0].data));
|
||||
int padW = *(static_cast<const int32_t*>(fields[1].data));
|
||||
int inputH = *(static_cast<const int32_t*>(fields[2].data));
|
||||
int inputW = *(static_cast<const int32_t*>(fields[3].data));
|
||||
int outputH = *(static_cast<const int32_t*>(fields[4].data));
|
||||
int outputW = *(static_cast<const int32_t*>(fields[5].data));
|
||||
int n = *(static_cast<const int32_t*>(fields[6].data));
|
||||
int c = *(static_cast<const int32_t*>(fields[7].data));
|
||||
auto *pluginObj = new ReflectionPaddingRT(padH,padW,inputH,inputW,outputH,outputW,c,n);
|
||||
pluginObj->setPluginNamespace(mPluginNamespace.c_str());
|
||||
return pluginObj;
|
||||
}
|
||||
|
||||
const char *ReflectionPaddingRTPluginCreator::getPluginName() const NOEXCEPT {
|
||||
return REFLECTIONPADDINGRT_PLUGIN_NAME;
|
||||
}
|
||||
|
||||
const char *ReflectionPaddingRTPluginCreator::getPluginVersion() const NOEXCEPT {
|
||||
return REFLECTIONPADDINGRT_PLUGIN_VERSION;
|
||||
}
|
||||
|
||||
const PluginFieldCollection *ReflectionPaddingRTPluginCreator::getFieldNames() NOEXCEPT {
|
||||
return &mFC;
|
||||
}
|
||||
|
||||
|
||||
@@ -0,0 +1,266 @@
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <opencv2/imgproc/imgproc.hpp>
|
||||
#include <opencv2/highgui/highgui.hpp>
|
||||
#include <tkdnn.h>
|
||||
#include "tkDNN/NetworkViz.h"
|
||||
|
||||
const char* encoder_conv1_bin = "monodepth2_1024/layers/encoder/encoder-conv1.bin";
|
||||
const char* encoder_layer1_bin[] = {
|
||||
"monodepth2_1024/layers/encoder/encoder-layer1-0-conv1.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer1-0-conv2.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer1-1-conv1.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer1-1-conv2.bin",
|
||||
};
|
||||
|
||||
const char* encoder_layer2_bin[] = {
|
||||
"monodepth2_1024/layers/encoder/encoder-layer2-0-conv1.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer2-0-conv2.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer2-0-downsample-0.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer2-1-conv1.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer2-1-conv2.bin"
|
||||
};
|
||||
|
||||
const char* encoder_layer3_bin[]={
|
||||
"monodepth2_1024/layers/encoder/encoder-layer3-0-conv1.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer3-0-conv2.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer3-0-downsample-0.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer3-1-conv1.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer3-1-conv2.bin"
|
||||
};
|
||||
|
||||
const char* encoder_layer4_bin[] = {
|
||||
"monodepth2_1024/layers/encoder/encoder-layer4-0-conv1.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer4-0-conv2.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer4-0-downsample-0.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer4-1-conv1.bin",
|
||||
"monodepth2_1024/layers/encoder/encoder-layer4-1-conv2.bin"
|
||||
};
|
||||
|
||||
const char *encoder_fc_bin = "monodepth2_1024/layers/encoder/encoder-fc.bin";
|
||||
|
||||
const char* decoder_layer_bin[] = {
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-0-conv-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-1-conv-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-2-conv-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-3-conv-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-4-conv-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-5-conv-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-6-conv-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-7-conv-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-8-conv-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-9-conv-conv.bin"
|
||||
};
|
||||
|
||||
const char* decoder_dispconv_layer_bin[] = {
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-10-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-11-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-12-conv.bin",
|
||||
"monodepth2_1024/layers/depth_decoder/decoder-13-conv.bin"
|
||||
};
|
||||
|
||||
const char* output_bin[] = {
|
||||
"monodepth2_1024/debug/outputs/output-disp-0.bin",
|
||||
"monodepth2_1024/debug/outputs/output-disp-1.bin",
|
||||
"monodepth2_1024/debug/outputs/output-disp-2.bin",
|
||||
"monodepth2_1024/debug/outputs/output-disp-3.bin"
|
||||
};
|
||||
|
||||
const char* input_bin = "monodepth2_1024/debug/input.bin";
|
||||
|
||||
|
||||
int main(){
|
||||
|
||||
//downloadWeightsifDoNotExist(input_bin, "monodepth2_1024", "https://cloud.hipert.unimore.it/s/iYw9QwgP6CsqxLR/download");
|
||||
|
||||
tk::dnn::dataDim_t dim(1,3,320,1024,1);
|
||||
tk::dnn::Network net(dim);
|
||||
|
||||
tk::dnn::Layer* muladd_sub = new tk::dnn::MulAdd(&net, 1.0f, -0.45f);
|
||||
tk::dnn::Layer* muladd_mul = new tk::dnn::MulAdd(&net, 1.0f / 0.225f, 0.0f);
|
||||
tk::dnn::Layer* encoder_conv = new tk::dnn::Conv2d(&net,64,7,7,2,2,3,3,encoder_conv1_bin,true);
|
||||
tk::dnn::Layer* encoder_relu = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_maxpool = new tk::dnn::Pooling(&net,3,3,2,2,1,1,tk::dnn::POOLING_MAX);
|
||||
|
||||
//layer-1
|
||||
tk::dnn::Layer* encoder_layer_1_0_convbn_1 = new tk::dnn::Conv2d(&net,64,3,3,1,1,1,1,encoder_layer1_bin[0],true);
|
||||
tk::dnn::Layer* encoder_relu_1 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_1_0_convbn_2 = new tk::dnn::Conv2d(&net,64,3,3,1,1,1,1,encoder_layer1_bin[1],true);
|
||||
tk::dnn::Layer* encoder_layer_1_0_shortcut_1 = new tk::dnn::Shortcut(&net,encoder_maxpool);
|
||||
tk::dnn::Layer* encoder_relu_2 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_1_1_convbn_1 = new tk::dnn::Conv2d(&net,64,3,3,1,1,1,1,encoder_layer1_bin[2],true);
|
||||
tk::dnn::Layer* encoder_relu_3 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_1_1_convbn_2 = new tk::dnn::Conv2d(&net,64,3,3,1,1,1,1,encoder_layer1_bin[3],true);
|
||||
tk::dnn::Layer* encoder_layer_1_1_shortcut_1 = new tk::dnn::Shortcut(&net,encoder_relu_2);
|
||||
tk::dnn::Layer* encoder_relu_4 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
|
||||
//layer-2
|
||||
tk::dnn::Layer* encoder_layer_2_0_convbn_1 = new tk::dnn::Conv2d(&net,128,3,3,2,2,1,1,encoder_layer2_bin[0],true);
|
||||
tk::dnn::Layer* encoder_relu_5 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_2_0_convbn_2 = new tk::dnn::Conv2d(&net,128,3,3,1,1,1,1,encoder_layer2_bin[1],true);
|
||||
tk::dnn::Layer* encoder_layer_2_0_route = new tk::dnn::Route(&net,&encoder_relu_4,1);
|
||||
tk::dnn::Layer* encoder_layer_2_0_downsample_convbn = new tk::dnn::Conv2d(&net,128,1,1,2,2,0,0,encoder_layer2_bin[2],true);
|
||||
tk::dnn::Layer* encoder_layer_2_0_shortcut = new tk::dnn::Shortcut(&net,encoder_layer_2_0_convbn_2);
|
||||
tk::dnn::Layer* encoder_relu_6 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_2_1_convbn_1 = new tk::dnn::Conv2d(&net,128,3,3,1,1,1,1,encoder_layer2_bin[3],true);
|
||||
tk::dnn::Layer* encoder_relu_7 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_2_1_convbn_2 = new tk::dnn::Conv2d(&net,128,3,3,1,1,1,1,encoder_layer2_bin[4],true);
|
||||
tk::dnn::Layer* encoder_layer_2_1shortcut = new tk::dnn::Shortcut(&net,encoder_relu_6);
|
||||
tk::dnn::Layer* encoder_relu_8 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
|
||||
//layer-3
|
||||
tk::dnn::Layer* encoder_layer_3_0_convbn_1 = new tk::dnn::Conv2d(&net,256,3,3,2,2,1,1,encoder_layer3_bin[0],true);
|
||||
tk::dnn::Layer* encoder_relu_9 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_3_0_convbn_2 = new tk::dnn::Conv2d(&net,256,3,3,1,1,1,1,encoder_layer3_bin[1],true);
|
||||
tk::dnn::Layer* encoder_layer_3_0_route = new tk::dnn::Route(&net,&encoder_relu_8,1);
|
||||
tk::dnn::Layer* encoder_layer_3_0_downsample_convbn = new tk::dnn::Conv2d(&net,256,1,1,2,2,0,0,encoder_layer3_bin[2],true);
|
||||
tk::dnn::Layer* encoder_layer_3_0_shortcut = new tk::dnn::Shortcut(&net,encoder_layer_3_0_convbn_2);
|
||||
tk::dnn::Layer* encoder_relu_10 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_3_1_convbn_1 = new tk::dnn::Conv2d(&net,256,3,3,1,1,1,1,encoder_layer3_bin[3],true);
|
||||
tk::dnn::Layer* encoder_relu_11 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_3_1_convbn_2 = new tk::dnn::Conv2d(&net,256,3,3,1,1,1,1,encoder_layer3_bin[4],true);
|
||||
tk::dnn::Layer* encoder_layer_3_1shortcut = new tk::dnn::Shortcut(&net,encoder_relu_10);
|
||||
tk::dnn::Layer* encoder_relu_12 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
|
||||
//layer-4
|
||||
tk::dnn::Layer* encoder_layer_4_0_convbn_1 = new tk::dnn::Conv2d(&net,512,3,3,2,2,1,1,encoder_layer4_bin[0],true);
|
||||
tk::dnn::Layer* encoder_relu_13 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_4_0_convbn_2 = new tk::dnn::Conv2d(&net,512,3,3,1,1,1,1,encoder_layer4_bin[1],true);
|
||||
tk::dnn::Layer* encoder_layer_4_0_route = new tk::dnn::Route(&net,&encoder_relu_12,1);
|
||||
tk::dnn::Layer* encoder_layer_4_0_downsample_convbn = new tk::dnn::Conv2d(&net,512,1,1,2,2,0,0,encoder_layer4_bin[2],true);
|
||||
tk::dnn::Layer* encoder_layer_4_0_shortcut = new tk::dnn::Shortcut(&net,encoder_layer_4_0_convbn_2);
|
||||
tk::dnn::Layer* encoder_relu_14 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_4_1_convbn_1 = new tk::dnn::Conv2d(&net,512,3,3,1,1,1,1,encoder_layer4_bin[3],true);
|
||||
tk::dnn::Layer* encoder_relu_15 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_4_1_convbn_2 = new tk::dnn::Conv2d(&net,512,3,3,1,1,1,1,encoder_layer4_bin[4],true);
|
||||
tk::dnn::Layer* encoder_layer_4_1shortcut = new tk::dnn::Shortcut(&net,encoder_relu_14);
|
||||
tk::dnn::Layer* encoder_relu_16 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
|
||||
//decoder
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_4_0 = new tk::dnn::Conv2d(&net,256,3,3,1,1,0,0,decoder_layer_bin[0]);
|
||||
tk::dnn::Layer* decoder_elu = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_upsampling_2d = new tk::dnn::Upsample(&net,2);
|
||||
tk::dnn::Layer* concatenate_layer[2] = {decoder_upsampling_2d,encoder_relu_12};
|
||||
tk::dnn::Layer* decoder_concatenate = new tk::dnn::Route(&net,concatenate_layer,2);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_1 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_4_1 = new tk::dnn::Conv2d(&net,256,3,3,1,1,0,0,decoder_layer_bin[1]);
|
||||
tk::dnn::Layer* decoder_elu_1 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_2 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_3_0 = new tk::dnn::Conv2d(&net,128,3,3,1,1,0,0,decoder_layer_bin[2]);
|
||||
tk::dnn::Layer* decoder_elu_2 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_upsampling_2d_1 = new tk::dnn::Upsample(&net,2);
|
||||
tk::dnn::Layer* concatenate_layer_1[2] = {decoder_upsampling_2d_1,encoder_relu_8};
|
||||
tk::dnn::Layer* decoder_concatenate_layer_1 = new tk::dnn::Route{&net,concatenate_layer_1,2};
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_3 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_3_1 = new tk::dnn::Conv2d(&net,128,3,3,1,1,0,0,decoder_layer_bin[3]);
|
||||
tk::dnn::Layer* decoder_elu_3 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_5 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_2_0 = new tk::dnn::Conv2d(&net,64,3,3,1,1,0,0,decoder_layer_bin[4]);
|
||||
tk::dnn::Layer* decoder_elu_4 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_upsampling_2d_2 = new tk::dnn::Upsample(&net,2);
|
||||
tk::dnn::Layer* concatenate_layer_2[2] = {decoder_upsampling_2d_2,encoder_relu_4};
|
||||
tk::dnn::Layer* decoder_concatenate_layer_2 = new tk::dnn::Route(&net,concatenate_layer_2,2);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_6 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_2_1 = new tk::dnn::Conv2d(&net,64,3,3,1,1,0,0,decoder_layer_bin[5]);
|
||||
tk::dnn::Layer* decoder_elu_5 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_8 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_1_0 = new tk::dnn::Conv2d(&net,32,3,3,1,1,0,0,decoder_layer_bin[6]);
|
||||
tk::dnn::Layer* decoder_elu_6 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_upsampling_2d_3 = new tk::dnn::Upsample(&net,2);
|
||||
tk::dnn::Layer* concatenate_layer_3[2] = {decoder_upsampling_2d_3,encoder_relu};
|
||||
tk::dnn::Layer* decoder_concatenate_layer_3 = new tk::dnn::Route(&net,concatenate_layer_3,2);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_9 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_1_1 = new tk::dnn::Conv2d(&net,32,3,3,1,1,0,0,decoder_layer_bin[7]);
|
||||
tk::dnn::Layer* decoder_elu_7 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_11 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_0_0 = new tk::dnn::Conv2d(&net,16,3,3,1,1,0,0,decoder_layer_bin[8]);
|
||||
tk::dnn::Layer* decoder_elu_8 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_upsampling_2d_4 = new tk::dnn::Upsample(&net,2);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_12 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_0_1 = new tk::dnn::Conv2d(&net,16,3,3,1,1,0,0,decoder_layer_bin[9]);
|
||||
tk::dnn::Layer* decoder_elu_9 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_13 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_dispconv_0 = new tk::dnn::Conv2d(&net,1,3,3,1,1,0,0,decoder_dispconv_layer_bin[0]);
|
||||
tk::dnn::Layer* disp0 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_SIGMOID);
|
||||
disp0->setFinal();
|
||||
|
||||
tk::dnn::Layer* route_elu_7 = new tk::dnn::Route(&net,&decoder_elu_7,1);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_10 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_dispconv_1 = new tk::dnn::Conv2d(&net,1,3,3,1,1,0,0,decoder_dispconv_layer_bin[1]);
|
||||
tk::dnn::Layer* disp1 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_SIGMOID);
|
||||
disp1->setFinal();
|
||||
|
||||
tk::dnn::Layer* route_elu_5 = new tk::dnn::Route(&net,&decoder_elu_5,1);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_7 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_dispconv_2 = new tk::dnn::Conv2d(&net,1,3,3,1,1,0,0,decoder_dispconv_layer_bin[2]);
|
||||
tk::dnn::Layer* disp2 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_SIGMOID);
|
||||
disp2->setFinal();
|
||||
|
||||
tk::dnn::Layer* route_elu_3 = new tk::dnn::Route(&net,&decoder_elu_3,1);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_4 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_dispconv_3 = new tk::dnn::Conv2d(&net,1,3,3,1,1,0,0,decoder_dispconv_layer_bin[3]);
|
||||
tk::dnn::Layer* disp3 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_SIGMOID);
|
||||
disp3->setFinal();
|
||||
|
||||
|
||||
dnnType *data;
|
||||
dnnType *input_H;
|
||||
readBinaryFile(input_bin, dim.tot(),&input_H,&data);
|
||||
std::cout<<"INPUT DIMENSIONS : "<<dim.tot()<<std::endl;
|
||||
|
||||
net.print();
|
||||
|
||||
tk::dnn::NetworkRT netRT(&net, net.getNetworkRTName("monodepth2_1024"));
|
||||
tk::dnn::dataDim_t dim1 = dim;
|
||||
dnnType *cudnn_out = nullptr;
|
||||
printCenteredTitle(" CUDNN inference ", '=', 30);
|
||||
{
|
||||
dim1.print();
|
||||
TKDNN_TSTART
|
||||
net.infer(dim1, data);
|
||||
TKDNN_TSTOP
|
||||
dim1.print();
|
||||
}
|
||||
|
||||
tk::dnn::dataDim_t dim2 = dim;
|
||||
printCenteredTitle(" TENSORRT inference ", '=', 30);
|
||||
{
|
||||
dim2.print();
|
||||
TKDNN_TSTART
|
||||
netRT.infer(dim2, data);
|
||||
TKDNN_TSTOP
|
||||
dim2.print();
|
||||
}
|
||||
tk::dnn::Layer *outs[4] = {disp0,disp1,disp2,disp3};
|
||||
std::cout<<std::endl<<std::endl;
|
||||
disp3->output_dim.print();
|
||||
int ret_cudnn = 0, ret_tensorrt = 0, ret_cudnn_tensorrt = 0;
|
||||
for(int i=0;i<4;i++){
|
||||
printCenteredTitle((std::string("MONODEPTH2 CHECK RESULTS ") + std::to_string(i) + " ").c_str(), '=', 30);
|
||||
outs[i]->output_dim.print();
|
||||
|
||||
dnnType *out, *out_h;
|
||||
int odim = outs[i]->output_dim.tot();
|
||||
readBinaryFile(output_bin[i], odim, &out_h, &out);
|
||||
|
||||
dnnType *cudnn_out, *rt_out;
|
||||
cudnn_out = outs[i]->dstData;
|
||||
rt_out = (dnnType *)netRT.buffersRT[1+i];
|
||||
std::cout<<"CUDNN vs correct";
|
||||
ret_cudnn |= checkResult(odim, cudnn_out, out) == 0 ? 0: ERROR_CUDNN;
|
||||
std::cout<<"TRT vs correct";
|
||||
ret_tensorrt |= checkResult(odim, rt_out, out) == 0 ? 0 : ERROR_TENSORRT;
|
||||
std::cout<<"CUDNN vs TRT ";
|
||||
ret_cudnn_tensorrt |= checkResult(odim, cudnn_out, rt_out) == 0 ? 0 : ERROR_CUDNNvsTENSORRT;
|
||||
|
||||
cv::Mat depth_mat = vizData2Mat(outs[i]->dstData, outs[i]->output_dim, outs[i]->output_dim.h, outs[i]->output_dim.w);
|
||||
cv::imshow("depth", depth_mat);
|
||||
cv::waitKey(0);
|
||||
}
|
||||
|
||||
|
||||
return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt;
|
||||
|
||||
|
||||
}
|
||||
@@ -0,0 +1,266 @@
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <opencv2/imgproc/imgproc.hpp>
|
||||
#include <opencv2/highgui/highgui.hpp>
|
||||
#include <tkdnn.h>
|
||||
#include "tkDNN/NetworkViz.h"
|
||||
|
||||
const char* encoder_conv1_bin = "monodepth2_640/layers/encoder/encoder-conv1.bin";
|
||||
const char* encoder_layer1_bin[] = {
|
||||
"monodepth2_640/layers/encoder/encoder-layer1-0-conv1.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer1-0-conv2.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer1-1-conv1.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer1-1-conv2.bin",
|
||||
};
|
||||
|
||||
const char* encoder_layer2_bin[] = {
|
||||
"monodepth2_640/layers/encoder/encoder-layer2-0-conv1.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer2-0-conv2.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer2-0-downsample-0.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer2-1-conv1.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer2-1-conv2.bin"
|
||||
};
|
||||
|
||||
const char* encoder_layer3_bin[]={
|
||||
"monodepth2_640/layers/encoder/encoder-layer3-0-conv1.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer3-0-conv2.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer3-0-downsample-0.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer3-1-conv1.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer3-1-conv2.bin"
|
||||
};
|
||||
|
||||
const char* encoder_layer4_bin[] = {
|
||||
"monodepth2_640/layers/encoder/encoder-layer4-0-conv1.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer4-0-conv2.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer4-0-downsample-0.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer4-1-conv1.bin",
|
||||
"monodepth2_640/layers/encoder/encoder-layer4-1-conv2.bin"
|
||||
};
|
||||
|
||||
const char *encoder_fc_bin = "monodepth2_640/layers/encoder/encoder-fc.bin";
|
||||
|
||||
const char* decoder_layer_bin[] = {
|
||||
"monodepth2_640/layers/depth_decoder/decoder-0-conv-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-1-conv-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-2-conv-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-3-conv-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-4-conv-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-5-conv-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-6-conv-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-7-conv-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-8-conv-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-9-conv-conv.bin"
|
||||
};
|
||||
|
||||
const char* decoder_dispconv_layer_bin[] = {
|
||||
"monodepth2_640/layers/depth_decoder/decoder-10-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-11-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-12-conv.bin",
|
||||
"monodepth2_640/layers/depth_decoder/decoder-13-conv.bin"
|
||||
};
|
||||
|
||||
const char* output_bin[] = {
|
||||
"monodepth2_640/debug/outputs/output-disp-0.bin",
|
||||
"monodepth2_640/debug/outputs/output-disp-1.bin",
|
||||
"monodepth2_640/debug/outputs/output-disp-2.bin",
|
||||
"monodepth2_640/debug/outputs/output-disp-3.bin"
|
||||
};
|
||||
|
||||
const char* input_bin = "monodepth2_640/debug/input.bin";
|
||||
|
||||
|
||||
int main(){
|
||||
|
||||
downloadWeightsifDoNotExist(input_bin, "monodepth2_640", "https://cloud.hipert.unimore.it/s/iYw9QwgP6CsqxLR/download");
|
||||
|
||||
tk::dnn::dataDim_t dim(1,3,192,640,1);
|
||||
tk::dnn::Network net(dim);
|
||||
|
||||
tk::dnn::Layer* muladd_sub = new tk::dnn::MulAdd(&net, 1.0f, -0.45f);
|
||||
tk::dnn::Layer* muladd_mul = new tk::dnn::MulAdd(&net, 1.0f / 0.225f, 0.0f);
|
||||
tk::dnn::Layer* encoder_conv = new tk::dnn::Conv2d(&net,64,7,7,2,2,3,3,encoder_conv1_bin,true);
|
||||
tk::dnn::Layer* encoder_relu = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_maxpool = new tk::dnn::Pooling(&net,3,3,2,2,1,1,tk::dnn::POOLING_MAX);
|
||||
|
||||
//layer-1
|
||||
tk::dnn::Layer* encoder_layer_1_0_convbn_1 = new tk::dnn::Conv2d(&net,64,3,3,1,1,1,1,encoder_layer1_bin[0],true);
|
||||
tk::dnn::Layer* encoder_relu_1 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_1_0_convbn_2 = new tk::dnn::Conv2d(&net,64,3,3,1,1,1,1,encoder_layer1_bin[1],true);
|
||||
tk::dnn::Layer* encoder_layer_1_0_shortcut_1 = new tk::dnn::Shortcut(&net,encoder_maxpool);
|
||||
tk::dnn::Layer* encoder_relu_2 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_1_1_convbn_1 = new tk::dnn::Conv2d(&net,64,3,3,1,1,1,1,encoder_layer1_bin[2],true);
|
||||
tk::dnn::Layer* encoder_relu_3 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_1_1_convbn_2 = new tk::dnn::Conv2d(&net,64,3,3,1,1,1,1,encoder_layer1_bin[3],true);
|
||||
tk::dnn::Layer* encoder_layer_1_1_shortcut_1 = new tk::dnn::Shortcut(&net,encoder_relu_2);
|
||||
tk::dnn::Layer* encoder_relu_4 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
|
||||
//layer-2
|
||||
tk::dnn::Layer* encoder_layer_2_0_convbn_1 = new tk::dnn::Conv2d(&net,128,3,3,2,2,1,1,encoder_layer2_bin[0],true);
|
||||
tk::dnn::Layer* encoder_relu_5 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_2_0_convbn_2 = new tk::dnn::Conv2d(&net,128,3,3,1,1,1,1,encoder_layer2_bin[1],true);
|
||||
tk::dnn::Layer* encoder_layer_2_0_route = new tk::dnn::Route(&net,&encoder_relu_4,1);
|
||||
tk::dnn::Layer* encoder_layer_2_0_downsample_convbn = new tk::dnn::Conv2d(&net,128,1,1,2,2,0,0,encoder_layer2_bin[2],true);
|
||||
tk::dnn::Layer* encoder_layer_2_0_shortcut = new tk::dnn::Shortcut(&net,encoder_layer_2_0_convbn_2);
|
||||
tk::dnn::Layer* encoder_relu_6 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_2_1_convbn_1 = new tk::dnn::Conv2d(&net,128,3,3,1,1,1,1,encoder_layer2_bin[3],true);
|
||||
tk::dnn::Layer* encoder_relu_7 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_2_1_convbn_2 = new tk::dnn::Conv2d(&net,128,3,3,1,1,1,1,encoder_layer2_bin[4],true);
|
||||
tk::dnn::Layer* encoder_layer_2_1shortcut = new tk::dnn::Shortcut(&net,encoder_relu_6);
|
||||
tk::dnn::Layer* encoder_relu_8 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
|
||||
//layer-3
|
||||
tk::dnn::Layer* encoder_layer_3_0_convbn_1 = new tk::dnn::Conv2d(&net,256,3,3,2,2,1,1,encoder_layer3_bin[0],true);
|
||||
tk::dnn::Layer* encoder_relu_9 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_3_0_convbn_2 = new tk::dnn::Conv2d(&net,256,3,3,1,1,1,1,encoder_layer3_bin[1],true);
|
||||
tk::dnn::Layer* encoder_layer_3_0_route = new tk::dnn::Route(&net,&encoder_relu_8,1);
|
||||
tk::dnn::Layer* encoder_layer_3_0_downsample_convbn = new tk::dnn::Conv2d(&net,256,1,1,2,2,0,0,encoder_layer3_bin[2],true);
|
||||
tk::dnn::Layer* encoder_layer_3_0_shortcut = new tk::dnn::Shortcut(&net,encoder_layer_3_0_convbn_2);
|
||||
tk::dnn::Layer* encoder_relu_10 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_3_1_convbn_1 = new tk::dnn::Conv2d(&net,256,3,3,1,1,1,1,encoder_layer3_bin[3],true);
|
||||
tk::dnn::Layer* encoder_relu_11 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_3_1_convbn_2 = new tk::dnn::Conv2d(&net,256,3,3,1,1,1,1,encoder_layer3_bin[4],true);
|
||||
tk::dnn::Layer* encoder_layer_3_1shortcut = new tk::dnn::Shortcut(&net,encoder_relu_10);
|
||||
tk::dnn::Layer* encoder_relu_12 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
|
||||
//layer-4
|
||||
tk::dnn::Layer* encoder_layer_4_0_convbn_1 = new tk::dnn::Conv2d(&net,512,3,3,2,2,1,1,encoder_layer4_bin[0],true);
|
||||
tk::dnn::Layer* encoder_relu_13 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_4_0_convbn_2 = new tk::dnn::Conv2d(&net,512,3,3,1,1,1,1,encoder_layer4_bin[1],true);
|
||||
tk::dnn::Layer* encoder_layer_4_0_route = new tk::dnn::Route(&net,&encoder_relu_12,1);
|
||||
tk::dnn::Layer* encoder_layer_4_0_downsample_convbn = new tk::dnn::Conv2d(&net,512,1,1,2,2,0,0,encoder_layer4_bin[2],true);
|
||||
tk::dnn::Layer* encoder_layer_4_0_shortcut = new tk::dnn::Shortcut(&net,encoder_layer_4_0_convbn_2);
|
||||
tk::dnn::Layer* encoder_relu_14 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_4_1_convbn_1 = new tk::dnn::Conv2d(&net,512,3,3,1,1,1,1,encoder_layer4_bin[3],true);
|
||||
tk::dnn::Layer* encoder_relu_15 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
tk::dnn::Layer* encoder_layer_4_1_convbn_2 = new tk::dnn::Conv2d(&net,512,3,3,1,1,1,1,encoder_layer4_bin[4],true);
|
||||
tk::dnn::Layer* encoder_layer_4_1shortcut = new tk::dnn::Shortcut(&net,encoder_relu_14);
|
||||
tk::dnn::Layer* encoder_relu_16 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_RELU);
|
||||
|
||||
//decoder
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_4_0 = new tk::dnn::Conv2d(&net,256,3,3,1,1,0,0,decoder_layer_bin[0]);
|
||||
tk::dnn::Layer* decoder_elu = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_upsampling_2d = new tk::dnn::Upsample(&net,2);
|
||||
tk::dnn::Layer* concatenate_layer[2] = {decoder_upsampling_2d,encoder_relu_12};
|
||||
tk::dnn::Layer* decoder_concatenate = new tk::dnn::Route(&net,concatenate_layer,2);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_1 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_4_1 = new tk::dnn::Conv2d(&net,256,3,3,1,1,0,0,decoder_layer_bin[1]);
|
||||
tk::dnn::Layer* decoder_elu_1 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_2 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_3_0 = new tk::dnn::Conv2d(&net,128,3,3,1,1,0,0,decoder_layer_bin[2]);
|
||||
tk::dnn::Layer* decoder_elu_2 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_upsampling_2d_1 = new tk::dnn::Upsample(&net,2);
|
||||
tk::dnn::Layer* concatenate_layer_1[2] = {decoder_upsampling_2d_1,encoder_relu_8};
|
||||
tk::dnn::Layer* decoder_concatenate_layer_1 = new tk::dnn::Route{&net,concatenate_layer_1,2};
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_3 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_3_1 = new tk::dnn::Conv2d(&net,128,3,3,1,1,0,0,decoder_layer_bin[3]);
|
||||
tk::dnn::Layer* decoder_elu_3 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_5 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_2_0 = new tk::dnn::Conv2d(&net,64,3,3,1,1,0,0,decoder_layer_bin[4]);
|
||||
tk::dnn::Layer* decoder_elu_4 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_upsampling_2d_2 = new tk::dnn::Upsample(&net,2);
|
||||
tk::dnn::Layer* concatenate_layer_2[2] = {decoder_upsampling_2d_2,encoder_relu_4};
|
||||
tk::dnn::Layer* decoder_concatenate_layer_2 = new tk::dnn::Route(&net,concatenate_layer_2,2);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_6 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_2_1 = new tk::dnn::Conv2d(&net,64,3,3,1,1,0,0,decoder_layer_bin[5]);
|
||||
tk::dnn::Layer* decoder_elu_5 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_8 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_1_0 = new tk::dnn::Conv2d(&net,32,3,3,1,1,0,0,decoder_layer_bin[6]);
|
||||
tk::dnn::Layer* decoder_elu_6 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_upsampling_2d_3 = new tk::dnn::Upsample(&net,2);
|
||||
tk::dnn::Layer* concatenate_layer_3[2] = {decoder_upsampling_2d_3,encoder_relu};
|
||||
tk::dnn::Layer* decoder_concatenate_layer_3 = new tk::dnn::Route(&net,concatenate_layer_3,2);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_9 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_1_1 = new tk::dnn::Conv2d(&net,32,3,3,1,1,0,0,decoder_layer_bin[7]);
|
||||
tk::dnn::Layer* decoder_elu_7 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_11 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_0_0 = new tk::dnn::Conv2d(&net,16,3,3,1,1,0,0,decoder_layer_bin[8]);
|
||||
tk::dnn::Layer* decoder_elu_8 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_upsampling_2d_4 = new tk::dnn::Upsample(&net,2);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_12 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_upconv_0_1 = new tk::dnn::Conv2d(&net,16,3,3,1,1,0,0,decoder_layer_bin[9]);
|
||||
tk::dnn::Layer* decoder_elu_9 = new tk::dnn::Activation(&net,tk::dnn::ACTIVATION_ELU);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_13 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_dispconv_0 = new tk::dnn::Conv2d(&net,1,3,3,1,1,0,0,decoder_dispconv_layer_bin[0]);
|
||||
tk::dnn::Layer* disp0 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_SIGMOID);
|
||||
disp0->setFinal();
|
||||
|
||||
tk::dnn::Layer* route_elu_7 = new tk::dnn::Route(&net,&decoder_elu_7,1);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_10 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_dispconv_1 = new tk::dnn::Conv2d(&net,1,3,3,1,1,0,0,decoder_dispconv_layer_bin[1]);
|
||||
tk::dnn::Layer* disp1 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_SIGMOID);
|
||||
disp1->setFinal();
|
||||
|
||||
tk::dnn::Layer* route_elu_5 = new tk::dnn::Route(&net,&decoder_elu_5,1);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_7 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_dispconv_2 = new tk::dnn::Conv2d(&net,1,3,3,1,1,0,0,decoder_dispconv_layer_bin[2]);
|
||||
tk::dnn::Layer* disp2 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_SIGMOID);
|
||||
disp2->setFinal();
|
||||
|
||||
tk::dnn::Layer* route_elu_3 = new tk::dnn::Route(&net,&decoder_elu_3,1);
|
||||
tk::dnn::Layer* decoder_reflection_padding_2d_4 = new tk::dnn::Padding(&net,1,1,tk::dnn::PADDING_MODE_REFLECTION);
|
||||
tk::dnn::Layer* decoder_dispconv_3 = new tk::dnn::Conv2d(&net,1,3,3,1,1,0,0,decoder_dispconv_layer_bin[3]);
|
||||
tk::dnn::Layer* disp3 = new tk::dnn::Activation(&net,CUDNN_ACTIVATION_SIGMOID);
|
||||
disp3->setFinal();
|
||||
|
||||
|
||||
dnnType *data;
|
||||
dnnType *input_H;
|
||||
readBinaryFile(input_bin, dim.tot(),&input_H,&data);
|
||||
std::cout<<"INPUT DIMENSIONS : "<<dim.tot()<<std::endl;
|
||||
|
||||
net.print();
|
||||
|
||||
tk::dnn::NetworkRT netRT(&net, net.getNetworkRTName("monodepth2_640"));
|
||||
tk::dnn::dataDim_t dim1 = dim;
|
||||
dnnType *cudnn_out = nullptr;
|
||||
printCenteredTitle(" CUDNN inference ", '=', 30);
|
||||
{
|
||||
dim1.print();
|
||||
TKDNN_TSTART
|
||||
net.infer(dim1, data);
|
||||
TKDNN_TSTOP
|
||||
dim1.print();
|
||||
}
|
||||
|
||||
tk::dnn::dataDim_t dim2 = dim;
|
||||
printCenteredTitle(" TENSORRT inference ", '=', 30);
|
||||
{
|
||||
dim2.print();
|
||||
TKDNN_TSTART
|
||||
netRT.infer(dim2, data);
|
||||
TKDNN_TSTOP
|
||||
dim2.print();
|
||||
}
|
||||
tk::dnn::Layer *outs[4] = {disp0,disp1,disp2,disp3};
|
||||
std::cout<<std::endl<<std::endl;
|
||||
disp3->output_dim.print();
|
||||
int ret_cudnn = 0, ret_tensorrt = 0, ret_cudnn_tensorrt = 0;
|
||||
for(int i=0;i<4;i++){
|
||||
printCenteredTitle((std::string("MONODEPTH2 CHECK RESULTS ") + std::to_string(i) + " ").c_str(), '=', 30);
|
||||
outs[i]->output_dim.print();
|
||||
|
||||
dnnType *out, *out_h;
|
||||
int odim = outs[i]->output_dim.tot();
|
||||
readBinaryFile(output_bin[i], odim, &out_h, &out);
|
||||
|
||||
dnnType *cudnn_out, *rt_out;
|
||||
cudnn_out = outs[i]->dstData;
|
||||
rt_out = (dnnType *)netRT.buffersRT[1+i];
|
||||
std::cout<<"CUDNN vs correct";
|
||||
ret_cudnn |= checkResult(odim, cudnn_out, out) == 0 ? 0: ERROR_CUDNN;
|
||||
std::cout<<"TRT vs correct";
|
||||
ret_tensorrt |= checkResult(odim, rt_out, out) == 0 ? 0 : ERROR_TENSORRT;
|
||||
std::cout<<"CUDNN vs TRT ";
|
||||
ret_cudnn_tensorrt |= checkResult(odim, cudnn_out, rt_out) == 0 ? 0 : ERROR_CUDNNvsTENSORRT;
|
||||
|
||||
cv::Mat depth_mat = vizData2Mat(outs[i]->dstData, outs[i]->output_dim, outs[i]->output_dim.h, outs[i]->output_dim.w);
|
||||
cv::imshow("depth", depth_mat);
|
||||
cv::waitKey(0);
|
||||
}
|
||||
|
||||
|
||||
return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt;
|
||||
|
||||
|
||||
}
|
||||
Reference in New Issue
Block a user