add mobilenetv3 small

mobilenetv3/CMakeLists.txt

@@ -0,0 +1,28 @@
+cmake_minimum_required(VERSION 2.6)
+
+project(mobilenetv3)
+
+add_definitions(-std=c++11)
+
+option(CUDA_USE_STATIC_CUDA_RUNTIME OFF)
+set(CMAKE_CXX_STANDARD 11)
+set(CMAKE_BUILD_TYPE Debug)
+
+find_package(CUDA REQUIRED)
+
+set(CUDA_NVCC_PLAGS ${CUDA_NVCC_PLAGS};-std=c++11;-g;-G;-gencode;arch=compute_30;code=sm_30;)
+
+include_directories(${PROJECT_SOURCE_DIR}/include)
+include_directories(/usr/local/cuda/targets/aarch64-linux/include)
+link_directories(/usr/local/cuda/targets/aarch64-linux/lib)
+cuda_add_library(h_sigmoid ${PROJECT_SOURCE_DIR}/h_sigmoid.cu)
+#cuda_add_library(leaky ${PROJECT_SOURCE_DIR}/leaky.cu)
+
+add_executable(mobilenetv3  ${PROJECT_SOURCE_DIR}/h_sigmoidplugin.cpp ${PROJECT_SOURCE_DIR}/mobilenet_v3.cpp)
+target_link_libraries(mobilenetv3 nvinfer)
+target_link_libraries(mobilenetv3 cudart)
+#target_link_libraries(mobilenetv3 leaky)
+target_link_libraries(mobilenetv3 h_sigmoid)
+
+add_definitions(-O2 -pthread)
+

mobilenetv3/README.md

@@ -0,0 +1,37 @@
+# mobilenet v3
+
+MobileNetV3 architecture from
+     "Searching for MobileNetV3" <https://arxiv.org/abs/1905.02244?context=cs>.
+
+For the Pytorch implementation, you can refer to [mobilenetv3.pytorch](https://github.com/d-li14/mobilenetv3.pytorch)
+
+Following tricks are used in this mobilenet,
+
+- Hsigmoid is used in mobilenet v3. We create a plugin in tensorrt.
+- Batchnorm layer, implemented by scale layer.
+
+```
+// 1. generate mbv3_small.wts from pytorch implementation
+
+// 2. put mbv3_small.wts into tensorrtx/mobilenet
+
+// 3. build and run
+
+cd tensorrtx/mobilenetv3
+
+mkdir build
+
+cd build
+
+cmake ..
+
+make
+
+sudo ./mobilenetv3 -s   // serialize model to plan file i.e. 'mobilenetv3_small.engine'
+
+sudo ./mobilenetv3 -d   // deserialize plan file and run inference
+
+// 4. see if the output is same as pytorch implementation
+```
+
+

mobilenetv3/common.h

@@ -0,0 +1,356 @@
+#ifndef _TRT_COMMON_H_
+#define _TRT_COMMON_H_
+#include "NvInfer.h"
+#include "NvOnnxConfig.h"
+#include "NvOnnxParser.h"
+#include <cuda_runtime_api.h>
+#include <algorithm>
+#include <cassert>
+#include <fstream>
+#include <iostream>
+#include <iterator>
+#include <map>
+#include <memory>
+#include <numeric>
+#include <string>
+#include <vector>
+#include <cstring>
+#include <cmath>
+
+using namespace std;
+
+#define CHECK(status)                             \
+    do                                            \
+    {                                             \
+        auto ret = (status);                      \
+        if (ret != 0)                             \
+        {                                         \
+            std::cout << "Cuda failure: " << ret; \
+            abort();                              \
+        }                                         \
+    } while (0)
+
+constexpr long double operator"" _GB(long double val) { return val * (1 << 30); }
+constexpr long double operator"" _MB(long double val) { return val * (1 << 20); }
+constexpr long double operator"" _KB(long double val) { return val * (1 << 10); }
+
+// These is necessary if we want to be able to write 1_GB instead of 1.0_GB.
+// Since the return type is signed, -1_GB will work as expected.
+constexpr long long int operator"" _GB(long long unsigned int val) { return val * (1 << 30); }
+constexpr long long int operator"" _MB(long long unsigned int val) { return val * (1 << 20); }
+constexpr long long int operator"" _KB(long long unsigned int val) { return val * (1 << 10); }
+
+// Logger for TensorRT info/warning/errors
+class Logger : public nvinfer1::ILogger
+{
+public:
+
+    Logger(): Logger(Severity::kWARNING) {}
+
+    Logger(Severity severity): reportableSeverity(severity) {}
+
+    void log(Severity severity, const char* msg) override
+    {
+        // suppress messages with severity enum value greater than the reportable
+        if (severity > reportableSeverity) return;
+
+        switch (severity)
+        {
+        case Severity::kINTERNAL_ERROR: std::cerr << "INTERNAL_ERROR: "; break;
+        case Severity::kERROR: std::cerr << "ERROR: "; break;
+        case Severity::kWARNING: std::cerr << "WARNING: "; break;
+        case Severity::kINFO: std::cerr << "INFO: "; break;
+        default: std::cerr << "UNKNOWN: "; break;
+        }
+        std::cerr << msg << std::endl;
+    }
+
+    Severity reportableSeverity{Severity::kWARNING};
+};
+
+// Locate path to file, given its filename or filepath suffix and possible dirs it might lie in
+// Function will also walk back MAX_DEPTH dirs from CWD to check for such a file path
+inline std::string locateFile(const std::string& filepathSuffix, const std::vector<std::string>& directories)
+{
+    const int MAX_DEPTH{10};
+    bool found{false};
+    std::string filepath;
+
+    for (auto& dir : directories)
+    {
+        filepath = dir + filepathSuffix;
+
+        for (int i = 0; i < MAX_DEPTH && !found; i++)
+        {
+            std::ifstream checkFile(filepath);
+            found = checkFile.is_open();
+            if (found) break;
+            filepath = "../" + filepath; // Try again in parent dir
+        }
+
+        if (found)
+        {
+            break;
+        }
+
+        filepath.clear();
+    }
+
+    if (filepath.empty()) {
+        std::string directoryList = std::accumulate(directories.begin() + 1, directories.end(), directories.front(),
+            [](const std::string& a, const std::string& b) { return a + "\n\t" + b; });
+        throw std::runtime_error("Could not find " + filepathSuffix + " in data directories:\n\t" + directoryList);
+    }
+    return filepath;
+}
+
+inline void readPGMFile(const std::string& fileName, uint8_t* buffer, int inH, int inW)
+{
+    std::ifstream infile(fileName, std::ifstream::binary);
+    assert(infile.is_open() && "Attempting to read from a file that is not open.");
+    std::string magic, h, w, max;
+    infile >> magic >> h >> w >> max;
+    infile.seekg(1, infile.cur);
+    infile.read(reinterpret_cast<char*>(buffer), inH * inW);
+}
+
+namespace samples_common
+{
+
+inline void* safeCudaMalloc(size_t memSize)
+{
+    void* deviceMem;
+    CHECK(cudaMalloc(&deviceMem, memSize));
+    if (deviceMem == nullptr)
+    {
+        std::cerr << "Out of memory" << std::endl;
+        exit(1);
+    }
+    return deviceMem;
+}
+
+inline bool isDebug()
+{
+    return (std::getenv("TENSORRT_DEBUG") ? true : false);
+}
+
+struct InferDeleter
+{
+    template <typename T>
+    void operator()(T* obj) const
+    {
+        if (obj) {
+            obj->destroy();
+        }
+    }
+};
+
+template <typename T>
+inline std::shared_ptr<T> infer_object(T* obj)
+{
+    if (!obj) {
+        throw std::runtime_error("Failed to create object");
+    }
+    return std::shared_ptr<T>(obj, InferDeleter());
+}
+
+template <class Iter>
+inline std::vector<size_t> argsort(Iter begin, Iter end, bool reverse = false)
+{
+    std::vector<size_t> inds(end - begin);
+    std::iota(inds.begin(), inds.end(), 0);
+    if (reverse) {
+        std::sort(inds.begin(), inds.end(), [&begin](size_t i1, size_t i2) {
+            return begin[i2] < begin[i1];
+        });
+    }
+    else
+    {
+        std::sort(inds.begin(), inds.end(), [&begin](size_t i1, size_t i2) {
+            return begin[i1] < begin[i2];
+        });
+    }
+    return inds;
+}
+
+inline bool readReferenceFile(const std::string& fileName, std::vector<std::string>& refVector)
+{
+    std::ifstream infile(fileName);
+    if (!infile.is_open()) {
+        cout << "ERROR: readReferenceFile: Attempting to read from a file that is not open." << endl;
+        return false;
+    }
+    std::string line;
+    while (std::getline(infile, line)) {
+        if (line.empty()) continue;
+        refVector.push_back(line);
+    }
+    infile.close();
+    return true;
+}
+
+template <typename result_vector_t>
+inline std::vector<std::string> classify(const vector<string>& refVector, const result_vector_t& output, const size_t topK)
+{
+    auto inds = samples_common::argsort(output.cbegin(), output.cend(), true);
+    std::vector<std::string> result;
+    for (size_t k = 0; k < topK; ++k) {
+        result.push_back(refVector[inds[k]]);
+    }
+    return result;
+}
+
+//...LG returns top K indices, not values.
+template <typename T>
+inline vector<size_t> topK(const vector<T> inp, const size_t k)
+{
+    vector<size_t> result;
+    std::vector<size_t> inds = samples_common::argsort(inp.cbegin(), inp.cend(), true);
+    result.assign(inds.begin(), inds.begin()+k);
+    return result;
+}
+
+template <typename T>
+inline bool readASCIIFile(const string& fileName, const size_t size, vector<T>& out)
+{
+    std::ifstream infile(fileName);
+    if (!infile.is_open()) {
+        cout << "ERROR readASCIIFile: Attempting to read from a file that is not open." << endl;
+        return false;
+    }
+    out.clear();
+    out.reserve(size);
+    out.assign(std::istream_iterator<T>(infile), std::istream_iterator<T>());
+    infile.close();
+    return true;
+}
+
+template <typename T>
+inline bool writeASCIIFile(const string& fileName, const vector<T>& in)
+{
+    std::ofstream outfile(fileName);
+    if (!outfile.is_open()) {
+        cout << "ERROR: writeASCIIFile: Attempting to write to a file that is not open." << endl;
+        return false;
+    }
+    for (auto fn : in) {
+        outfile << fn << " ";
+    }
+    outfile.close();
+    return true;
+}
+
+inline void print_version()
+{
+//... This can be only done after statically linking this support into parserONNX.library
+#if 0
+    std::cout << "Parser built against:" << std::endl;
+    std::cout << "  ONNX IR version:  " << nvonnxparser::onnx_ir_version_string(onnx::IR_VERSION) << std::endl;
+#endif
+    std::cout << "  TensorRT version: "
+              << NV_TENSORRT_MAJOR << "."
+              << NV_TENSORRT_MINOR << "."
+              << NV_TENSORRT_PATCH << "."
+              << NV_TENSORRT_BUILD << std::endl;
+}
+
+inline string getFileType(const string& filepath)
+{
+    return filepath.substr(filepath.find_last_of(".") + 1);
+}
+
+inline string toLower(const string& inp)
+{
+    string out = inp;
+    std::transform(out.begin(), out.end(), out.begin(), ::tolower);
+    return out;
+}
+
+inline unsigned int getElementSize(nvinfer1::DataType t)
+{
+    switch (t)
+    {
+        case nvinfer1::DataType::kINT32: return 4;
+        case nvinfer1::DataType::kFLOAT: return 4;
+        case nvinfer1::DataType::kHALF: return 2;
+        case nvinfer1::DataType::kINT8: return 1;
+    }
+    throw std::runtime_error("Invalid DataType.");
+    return 0;
+}
+
+inline int64_t volume(const nvinfer1::Dims& d)
+{
+    return std::accumulate(d.d, d.d + d.nbDims, 1, std::multiplies<int64_t>());
+}
+
+// Struct to maintain command-line arguments.
+struct Args
+{
+    bool runInInt8 = false;
+};
+
+// Populates the Args struct with the provided command-line parameters.
+inline void parseArgs(Args& args, int argc, char* argv[])
+{
+    if (argc >= 1)
+    {
+        for (int i = 1; i < argc; ++i)
+        {
+            if (!strcmp(argv[i], "--int8")) args.runInInt8 = true;
+        }
+    }
+}
+
+template <int C, int H, int W>
+struct PPM
+{
+    std::string magic, fileName;
+    int h, w, max;
+    uint8_t buffer[C * H * W];
+};
+
+struct BBox
+{
+    float x1, y1, x2, y2;
+};
+
+template <int C, int H, int W>
+inline void writePPMFileWithBBox(const std::string& filename, PPM<C, H, W>& ppm, const BBox& bbox)
+{
+    std::ofstream outfile("./" + filename, std::ofstream::binary);
+    assert(!outfile.fail());
+    outfile << "P6" << "\n" << ppm.w << " " << ppm.h << "\n" << ppm.max << "\n";
+    auto round = [](float x) -> int { return int(std::floor(x + 0.5f)); };
+    const int x1 = std::min(std::max(0, round(int(bbox.x1))), W - 1);
+    const int x2 = std::min(std::max(0, round(int(bbox.x2))), W - 1);
+    const int y1 = std::min(std::max(0, round(int(bbox.y1))), H - 1);
+    const int y2 = std::min(std::max(0, round(int(bbox.y2))), H - 1);
+    for (int x = x1; x <= x2; ++x)
+    {
+        // bbox top border
+        ppm.buffer[(y1 * ppm.w + x) * 3] = 255;
+        ppm.buffer[(y1 * ppm.w + x) * 3 + 1] = 0;
+        ppm.buffer[(y1 * ppm.w + x) * 3 + 2] = 0;
+        // bbox bottom border
+        ppm.buffer[(y2 * ppm.w + x) * 3] = 255;
+        ppm.buffer[(y2 * ppm.w + x) * 3 + 1] = 0;
+        ppm.buffer[(y2 * ppm.w + x) * 3 + 2] = 0;
+    }
+    for (int y = y1; y <= y2; ++y)
+    {
+        // bbox left border
+        ppm.buffer[(y * ppm.w + x1) * 3] = 255;
+        ppm.buffer[(y * ppm.w + x1) * 3 + 1] = 0;
+        ppm.buffer[(y * ppm.w + x1) * 3 + 2] = 0;
+        // bbox right border
+        ppm.buffer[(y * ppm.w + x2) * 3] = 255;
+        ppm.buffer[(y * ppm.w + x2) * 3 + 1] = 0;
+        ppm.buffer[(y * ppm.w + x2) * 3 + 2] = 0;
+    }
+    outfile.write(reinterpret_cast<char*>(ppm.buffer), ppm.w * ppm.h * 3);
+}
+
+} // namespace samples_common
+
+#endif // _TRT_COMMON_H_

mobilenetv3/h_sigmoid.cu

@@ -0,0 +1,27 @@
+#include <cuda_runtime.h>
+#include <stdio.h>
+#include "h_sigmoid.cuh"
+
+
+__global__ void _hSigmoidKer(float const *in, float *out, int size) {
+    int index = threadIdx.x + blockIdx.x * blockDim.x;
+    if (index >= size)
+        return ;
+
+    if (in[index] > 3 )
+        out[index] = 1;
+    else if (in[index] < -3)
+        out[index] = 0;
+    else
+        out[index] = (in[index] + 3)/6;
+}
+
+extern "C" void cuh_sigmoid(float const *in, float *out, int size) {
+    int block_size = 256;
+    int grid_size = (size + block_size - 1) / block_size;
+    _hSigmoidKer<<<grid_size, block_size>>>(in, out, size);
+    cudaError_t err = cudaGetLastError();
+    if (err != cudaSuccess) {
+        fprintf(stderr, "Failed to launch _leakyReluKer kernel (error code %s)!\n", cudaGetErrorString(err));
+    }
+}

mobilenetv3/h_sigmoid.cuh

@@ -0,0 +1,8 @@
+#ifndef HSIGMOID_H
+#define HSIGMOID_H
+
+extern "C"
+
+void cuh_sigmoid(float const *in, float *out, int size);
+
+#endif

mobilenetv3/h_sigmoidplugin.cpp

@@ -0,0 +1,62 @@
+#include "common.h"
+#include "h_sigmoid.cuh"
+#include "h_sigmoidplugin.h"
+
+using namespace nvinfer1;
+using nvinfer1::HSigmoidPlugin;
+using nvinfer1::PluginFactory;
+
+HSigmoidPlugin::HSigmoidPlugin() {
+}
+
+HSigmoidPlugin::HSigmoidPlugin(const void* buffer, size_t size) {
+    assert(size == sizeof(input_size_));
+    input_size_ = *reinterpret_cast<const int*>(buffer);
+}
+
+int HSigmoidPlugin::getNbOutputs() const {
+    return 1;
+}
+
+Dims HSigmoidPlugin::getOutputDimensions(int index, const Dims* inputs, int nbInputDims) {
+    assert(nbInputDims == 1);
+    assert(index == 0);
+    // Output dimensions
+    return DimsCHW(inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]);
+}
+
+void HSigmoidPlugin::configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) {
+    input_size_ = inputDims[0].d[0] * inputDims[0].d[1] * inputDims[0].d[2];
+}
+
+int HSigmoidPlugin::initialize() {
+    return 0;
+}
+
+void HSigmoidPlugin::terminate() {}
+
+size_t HSigmoidPlugin::getWorkspaceSize(int maxBatchSize) const {
+    return 0;
+}
+
+int HSigmoidPlugin::enqueue(int batchSize, const void* const* inputs, void** outputs, void* workspace, cudaStream_t stream) {
+    cuh_sigmoid(reinterpret_cast<float const*>(inputs[0]), reinterpret_cast<float*>(outputs[0]), input_size_);
+    return 0;
+}
+
+size_t HSigmoidPlugin::getSerializationSize() {
+    return sizeof(input_size_);
+}
+
+void HSigmoidPlugin::serialize(void* buffer) {
+    *reinterpret_cast<int*>(buffer) = input_size_;
+}
+
+IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) {
+    IPlugin *plugin = nullptr;
+    if (strstr(layerName, "h_sigmoid") != NULL) {
+        plugin = new HSigmoidPlugin(serialData, serialLength);
+    }
+    return plugin;
+}
+

mobilenetv3/h_sigmoidplugin.h

@@ -0,0 +1,32 @@
+#ifndef HSIGMOID_PLUGIN_H
+#define HSIGMOID_PLUGIN_H
+#include <NvInfer.h>
+
+namespace nvinfer1 {
+class HSigmoidPlugin : public IPlugin {
+    public:
+        HSigmoidPlugin();
+        HSigmoidPlugin(const void* buffer, size_t size);
+        ~HSigmoidPlugin() override = default;
+        int getNbOutputs() const override;
+        Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override;
+        void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override;
+        int initialize() override;
+        void terminate() override;
+        size_t getWorkspaceSize(int maxBatchSize) const override;
+        int enqueue(
+                int batchSize, const void* const* inputs, void** outputs, void* workspace, cudaStream_t stream) override;
+        size_t getSerializationSize() override;
+        void serialize(void* buffer) override;
+
+    private:
+        int input_size_;
+};
+
+class PluginFactory : public IPluginFactory {
+    public:
+        IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength) override;
+};
+
+}
+#endif

mobilenetv3/mobilenet_v3.cpp

@@ -0,0 +1,433 @@
+#include "NvInfer.h"
+#include "cuda_runtime_api.h"
+#include "common.h"
+#include <fstream>
+#include <iostream>
+#include <map>
+#include <sstream>
+#include <vector>
+#include <chrono>
+//#include "plugin_factory.h"
+#include "h_sigmoidplugin.h"
+//#include "leakyplugin.h"
+
+// stuff we know about the network and the input/output blobs
+static const int INPUT_H = 224;
+static const int INPUT_W = 224;
+static const int OUTPUT_SIZE = 1000;
+static const int BS = 1;
+
+const char* INPUT_BLOB_NAME = "data";
+const char* OUTPUT_BLOB_NAME = "prob";
+
+using namespace nvinfer1;
+
+static Logger gLogger;
+
+// Load weights from files shared with TensorRT samples.
+// TensorRT weight files have a simple space delimited format:
+// [type] [size] <data x size in hex>
+std::map<std::string, Weights> loadWeights(const std::string file)
+{
+    std::cout << "Loading weights: " << file << std::endl;
+    std::map<std::string, Weights> weightMap;
+
+    // Open weights file
+    std::ifstream input(file);
+    assert(input.is_open() && "Unable to load weight file.");
+
+    // Read number of weight blobs
+    int32_t count;
+    input >> count;
+    assert(count > 0 && "Invalid weight map file.");
+
+    while (count--)
+    {
+        Weights wt{DataType::kFLOAT, nullptr, 0};
+        uint32_t size;
+
+        // Read name and type of blob
+        std::string name;
+        input >> name >> std::dec >> size;
+        wt.type = DataType::kFLOAT;
+
+        // Load blob
+        uint32_t* val = reinterpret_cast<uint32_t*>(malloc(sizeof(val) * size));
+        for (uint32_t x = 0, y = size; x < y; ++x)
+        {
+            input >> std::hex >> val[x];
+        }
+        wt.values = val;
+        
+        wt.count = size;
+        weightMap[name] = wt;
+    }
+
+    return weightMap;
+}
+
+IScaleLayer* addBatchNorm(INetworkDefinition *network, std::map<std::string, Weights>& weightMap, ITensor& input, std::string lname, float eps) {
+    float *gamma = (float*)weightMap[lname + ".weight"].values;
+    float *beta = (float*)weightMap[lname + ".bias"].values;
+    float *mean = (float*)weightMap[lname + ".running_mean"].values;
+    float *var = (float*)weightMap[lname + ".running_var"].values;
+    int len = weightMap[lname + ".running_var"].count;
+    std::cout << "len " << len << std::endl;
+
+    float *scval = reinterpret_cast<float*>(malloc(sizeof(float) * len));
+    for (int i = 0; i < len; i++) {
+        scval[i] = gamma[i] / sqrt(var[i] + eps);
+    }
+    Weights scale{DataType::kFLOAT, scval, len};
+    
+    float *shval = reinterpret_cast<float*>(malloc(sizeof(float) * len));
+    for (int i = 0; i < len; i++) {
+        shval[i] = beta[i] - mean[i] * gamma[i] / sqrt(var[i] + eps);
+    }
+    Weights shift{DataType::kFLOAT, shval, len};
+
+    float *pval = reinterpret_cast<float*>(malloc(sizeof(float) * len));
+    for (int i = 0; i < len; i++) {
+        pval[i] = 1.0;
+    }
+    Weights power{DataType::kFLOAT, pval, len};
+
+    weightMap[lname + ".scale"] = scale;
+    weightMap[lname + ".shift"] = shift;
+    weightMap[lname + ".power"] = power;
+    IScaleLayer* scale_1 = network->addScale(input, ScaleMode::kCHANNEL, shift, scale, power);
+    assert(scale_1);
+    return scale_1;
+}
+
+ILayer* hSwish(INetworkDefinition *network, ITensor& input, std::string name) {
+    //auto hsg = new LeakyPlugin();
+    auto hsg = new HSigmoidPlugin();
+    ITensor* inputTensors[] = {&input};
+    auto hs1 = network->addPlugin(inputTensors,1,*hsg);
+    assert(hs1);
+    hs1->setName(("h_sigmoid"+name).c_str());
+    ILayer* hsw = network->addElementWise(input, *hs1->getOutput(0),ElementWiseOperation::kPROD);
+    assert(hsw);
+    return hsw;
+}
+
+
+ILayer* convBnHswish(INetworkDefinition *network, std::map<std::string, Weights>& weightMap, ITensor& input,  int outch, int ksize, int s, int g, std::string lname) {
+    Weights emptywts{DataType::kFLOAT, nullptr, 0};
+    int p = (ksize - 1) / 2;
+    IConvolutionLayer* conv1 = network->addConvolution(input, outch, DimsHW{ksize, ksize}, weightMap[lname + "0.weight"], emptywts);
+    assert(conv1);
+    conv1->setStride(DimsHW{s, s});
+    conv1->setPadding(DimsHW{p, p});
+    conv1->setNbGroups(g);
+
+    IScaleLayer* bn1 = addBatchNorm(network, weightMap, *conv1->getOutput(0), lname + "1", 1e-5);
+    ILayer* hsw = hSwish(network, *bn1->getOutput(0), lname+"2");
+    assert(hsw);
+    return hsw;
+}
+
+ILayer* seLayer(INetworkDefinition *network, std::map<std::string, Weights>& weightMap, ITensor& input, int c, int w, std::string lname) {
+    int h = w;
+    IPoolingLayer* l1 = network->addPooling(input,PoolingType::kAVERAGE,DimsHW(w, h));
+    assert(l1);
+    l1->setStride(DimsHW{w, h});
+    IFullyConnectedLayer* l2 = network->addFullyConnected(*l1->getOutput(0), BS*c/4,weightMap[lname+"fc.0.weight"],weightMap[lname+"fc.0.bias"]);
+    IActivationLayer* relu1 = network->addActivation(*l2->getOutput(0),ActivationType::kRELU);
+    IFullyConnectedLayer* l4 = network->addFullyConnected(*relu1->getOutput(0), BS*c,weightMap[lname+"fc.2.weight"],weightMap[lname+"fc.2.bias"]);
+    auto hsg = new HSigmoidPlugin();
+    ITensor* inputTensors[] = {l4->getOutput(0)};
+    auto hs1 = network->addPlugin(inputTensors,1,*hsg);
+    assert(hs1);
+    hs1->setName(("h_sigmoid"+lname + "seLayer").c_str());
+    ILayer* se = network->addElementWise(input, *hs1->getOutput(0), ElementWiseOperation::kPROD);
+    assert(se);
+    return se;
+}
+
+ILayer* convSeq1(INetworkDefinition *network, std::map<std::string, Weights>& weightMap, ITensor& input, int output, int hdim, int k, int s, bool use_se, bool use_hs, int w, std::string lname) {
+    Weights emptywts{DataType::kFLOAT, nullptr, 0};
+    int p = (k - 1) / 2;
+    IConvolutionLayer* conv1 = network->addConvolution(input, hdim, DimsHW{k, k}, weightMap[lname + "0.weight"], emptywts);
+    conv1->setStride(DimsHW{s, s});
+    conv1->setPadding(DimsHW{p, p});
+    conv1->setNbGroups(hdim);
+
+    IScaleLayer* bn1 = addBatchNorm(network, weightMap, *conv1->getOutput(0), lname + "1", 1e-5);
+    ITensor *tensor3, *tensor4;
+    tensor3 = nullptr;
+    tensor4 = nullptr;
+    if (use_hs) {
+        ILayer* hsw = hSwish(network, *bn1->getOutput(0), lname+"2");
+        tensor3 = hsw->getOutput(0);
+    }
+    else {
+        IActivationLayer* relu1 = network->addActivation(*bn1->getOutput(0),ActivationType::kRELU);
+        tensor3 = relu1->getOutput(0);
+    }
+    if (use_se) {
+         ILayer* se1 = seLayer(network, weightMap, *tensor3, hdim, w, lname + "3.");
+         tensor4 = se1->getOutput(0);
+    }
+    else {
+         tensor4 = tensor3;
+    }
+    IConvolutionLayer* conv2 = network->addConvolution(*tensor4, output, DimsHW{1, 1}, weightMap[lname + "4.weight"], emptywts);
+    IScaleLayer* bn2 = addBatchNorm(network, weightMap, *conv2->getOutput(0), lname + "5", 1e-5);
+    assert(bn2);
+    return bn2;
+}
+ILayer* convSeq2(INetworkDefinition *network, std::map<std::string, Weights>& weightMap, ITensor& input, int output, int hdim, int k, int s, bool use_se, bool use_hs, int w, std::string lname) {
+    Weights emptywts{DataType::kFLOAT, nullptr, 0};
+    int p = (k - 1) / 2;
+    IConvolutionLayer* conv1 = network->addConvolution(input, hdim, DimsHW{1, 1}, weightMap[lname + "0.weight"], emptywts);
+    IScaleLayer* bn1 = addBatchNorm(network, weightMap, *conv1->getOutput(0), lname + "1", 1e-5);
+    ITensor *tensor3, *tensor6, *tensor7;
+    tensor3 = nullptr;
+    tensor6 = nullptr;
+    tensor7 = nullptr;
+    if (use_hs) {
+        ILayer* hsw1 = hSwish(network, *bn1->getOutput(0), lname + "2");
+        tensor3 = hsw1->getOutput(0);
+    }
+    else {
+        IActivationLayer* relu1 = network->addActivation(*bn1->getOutput(0),ActivationType::kRELU);
+        tensor3 = relu1->getOutput(0);
+    }
+    IConvolutionLayer* conv2 = network->addConvolution(*tensor3, hdim, DimsHW{k, k}, weightMap[lname + "3.weight"], emptywts);
+    conv2->setStride(DimsHW{s, s});
+    conv2->setPadding(DimsHW{p, p});
+    conv2->setNbGroups(hdim);
+    IScaleLayer* bn2 = addBatchNorm(network, weightMap, *conv2->getOutput(0), lname + "4", 1e-5);
+    if (use_se) {
+         ILayer* se1 = seLayer(network, weightMap, *bn2->getOutput(0), hdim, w, lname + "5.");
+         tensor6 = se1->getOutput(0);
+    }
+    else {
+         tensor6 = bn2->getOutput(0);
+    }
+    if (use_hs) {
+        ILayer* hsw2 = hSwish(network, *tensor6, lname + "6");
+        tensor7 = hsw2->getOutput(0);
+    }
+    else {
+        IActivationLayer* relu2 = network->addActivation(*tensor6, ActivationType::kRELU);
+        tensor7 = relu2->getOutput(0);
+    }
+    IConvolutionLayer* conv3 = network->addConvolution(*tensor7, output, DimsHW{1, 1}, weightMap[lname + "7.weight"], emptywts);
+    IScaleLayer* bn3 = addBatchNorm(network, weightMap, *conv3->getOutput(0), lname + "8", 1e-5);
+    assert(bn3);
+    return bn3;
+}
+ILayer* invertedRes(INetworkDefinition *network, std::map<std::string, Weights>& weightMap, ITensor& input, std::string lname,
+    int inch, int outch, int s, int hidden, int k, bool use_se, bool use_hs, int w) {
+    bool use_res_connect = (s == 1 && inch == outch);
+    ILayer *conv = nullptr;
+    if (inch == hidden) {
+        conv = convSeq1(network, weightMap, input, outch, hidden, k, s, use_se, use_hs, w, lname + "conv.");
+    }
+    else {
+        conv = convSeq2(network, weightMap, input, outch, hidden, k, s, use_se, use_hs, w, lname + "conv.");
+    }
+
+    if (!use_res_connect) return conv;
+    IElementWiseLayer* ew3 = network->addElementWise(input, *conv->getOutput(0), ElementWiseOperation::kSUM);
+    assert(ew3);
+    return ew3;
+}
+
+// Creat the engine using only the API and not any parser.
+ICudaEngine* createEngine(unsigned int maxBatchSize, IBuilder* builder, DataType dt)
+{
+    INetworkDefinition* network = builder->createNetwork();
+
+    // Create input tensor of shape { 1, 1, 32, 32 } with name INPUT_BLOB_NAME
+    ITensor* data = network->addInput(INPUT_BLOB_NAME, dt, Dims3{3, INPUT_H, INPUT_W});
+    assert(data);
+
+    std::map<std::string, Weights> weightMap = loadWeights("../mbv3_small.wts");
+    Weights emptywts{DataType::kFLOAT, nullptr, 0};
+
+    //auto test1 = network->addActivation(*data, ActivationType::kRELU);
+    auto ew1 = convBnHswish(network, weightMap, *data, 16, 3, 2, 1, "features.0.");
+    auto ir1 = invertedRes(network, weightMap, *ew1->getOutput(0), "features.1.", 16, 16, 2, 16, 3, 1, 0, 56);
+    auto ir2 = invertedRes(network, weightMap, *ir1->getOutput(0), "features.2.", 16, 24, 2, 72, 3, 0, 0, 28);
+    auto ir3 = invertedRes(network, weightMap, *ir2->getOutput(0), "features.3.", 24, 24, 1, 88, 3, 0, 0, 28);
+    auto ir4 = invertedRes(network, weightMap, *ir3->getOutput(0), "features.4.", 24, 40, 2, 96, 5, 1, 1, 14);
+    auto ir5 = invertedRes(network, weightMap, *ir4->getOutput(0), "features.5.", 40, 40, 1, 240, 5, 1, 1, 14);
+    auto ir6 = invertedRes(network, weightMap, *ir5->getOutput(0), "features.6.", 40, 40, 1, 240, 5, 1, 1, 14);
+    auto ir7 = invertedRes(network, weightMap, *ir6->getOutput(0), "features.7.", 40, 48, 1, 120, 5, 1, 1, 14);
+    auto ir8 = invertedRes(network, weightMap, *ir7->getOutput(0), "features.8.", 48, 48, 1, 144, 5, 1, 1, 14);
+    auto ir9 = invertedRes(network, weightMap, *ir8->getOutput(0), "features.9.", 48, 96, 2, 288, 5, 1, 1, 7);
+    auto ir10 = invertedRes(network, weightMap, *ir9->getOutput(0), "features.10.", 96, 96, 1, 576, 5, 1, 1, 7);
+    auto ir11 = invertedRes(network, weightMap, *ir10->getOutput(0), "features.11.", 96, 96, 1, 576, 5, 1, 1, 7);
+    ILayer* ew2 = convBnHswish(network, weightMap, *ir11->getOutput(0), 576, 1, 1, 1, "conv.0.");
+    ILayer* se1 = seLayer(network, weightMap, *ew2->getOutput(0), 576, 7, "conv.1.");
+
+    IPoolingLayer* pool1 = network->addPooling(*se1->getOutput(0), PoolingType::kAVERAGE, DimsHW{7, 7});
+    assert(pool1);
+    pool1->setStride(DimsHW{7, 7});
+    ILayer* sw1 = hSwish(network, *pool1->getOutput(0), "hSwish.0");
+
+    IFullyConnectedLayer* fc1 = network->addFullyConnected(*sw1->getOutput(0), 1280, weightMap["classifier.0.weight"], weightMap["classifier.0.bias"]);
+    assert(fc1);
+    ILayer* bn1 = addBatchNorm(network, weightMap, *fc1->getOutput(0), "classifier.1", 1e-5);
+    ILayer* sw2 = hSwish(network, *bn1->getOutput(0), "hSwish.1");
+    IFullyConnectedLayer* fc2 = network->addFullyConnected(*sw2->getOutput(0), 1000, weightMap["classifier.3.weight"], weightMap["classifier.3.bias"]);
+    ILayer* bn2 = addBatchNorm(network, weightMap, *fc2->getOutput(0), "classifier.4", 1e-5);
+    ILayer* sw3 = hSwish(network, *bn2->getOutput(0), "hSwish.2");
+
+    sw3->getOutput(0)->setName(OUTPUT_BLOB_NAME);
+    std::cout << "set name out" << std::endl;
+    network->markOutput(*sw3->getOutput(0));
+
+    // Build engine
+    builder->setMaxBatchSize(maxBatchSize);
+    builder->setMaxWorkspaceSize(1 << 20);
+    ICudaEngine* engine = builder->buildCudaEngine(*network);
+    std::cout << "build out" << std::endl;
+
+    // Don't need the network any more
+    network->destroy();
+
+    // Release host memory
+    for (auto& mem : weightMap)
+    {
+        free((void*) (mem.second.values));
+    }
+
+    return engine;
+}
+
+void APIToModel(unsigned int maxBatchSize, IHostMemory** modelStream)
+{
+    // Create builder
+    IBuilder* builder = createInferBuilder(gLogger);
+
+    // Create model to populate the network, then set the outputs and create an engine
+    ICudaEngine* engine = createEngine(maxBatchSize, builder, DataType::kFLOAT);
+    assert(engine != nullptr);
+
+    // Serialize the engine
+    (*modelStream) = engine->serialize();
+
+    // Close everything down
+    engine->destroy();
+    builder->destroy();
+}
+
+void doInference(IExecutionContext& context, float* input, float* output, int batchSize)
+{
+    const ICudaEngine& engine = context.getEngine();
+
+    // Pointers to input and output device buffers to pass to engine.
+    // Engine requires exactly IEngine::getNbBindings() number of buffers.
+    assert(engine.getNbBindings() == 2);
+    void* buffers[2];
+
+    // 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()
+    const int inputIndex = engine.getBindingIndex(INPUT_BLOB_NAME);
+    const int outputIndex = engine.getBindingIndex(OUTPUT_BLOB_NAME);
+
+    // Create GPU buffers on device
+    CHECK(cudaMalloc(&buffers[inputIndex], batchSize * 3 * INPUT_H * INPUT_W * sizeof(float)));
+    CHECK(cudaMalloc(&buffers[outputIndex], batchSize * OUTPUT_SIZE * sizeof(float)));
+
+    // Create stream
+    cudaStream_t stream;
+    CHECK(cudaStreamCreate(&stream));
+
+    // DMA input batch data to device, infer on the batch asynchronously, and DMA output back to host
+    CHECK(cudaMemcpyAsync(buffers[inputIndex], input, batchSize * 3 * INPUT_H * INPUT_W * sizeof(float), cudaMemcpyHostToDevice, stream));
+    context.enqueue(batchSize, buffers, stream, nullptr);
+    CHECK(cudaMemcpyAsync(output, buffers[outputIndex], batchSize * OUTPUT_SIZE * sizeof(float), cudaMemcpyDeviceToHost, stream));
+    cudaStreamSynchronize(stream);
+
+    // Release stream and buffers
+    cudaStreamDestroy(stream);
+    CHECK(cudaFree(buffers[inputIndex]));
+    CHECK(cudaFree(buffers[outputIndex]));
+}
+
+int main(int argc, char** argv)
+{
+    if (argc != 2) {
+        std::cerr << "arguments not right!" << std::endl;
+        std::cerr << "./mobilenet -s   // serialize model to plan file" << std::endl;
+        std::cerr << "./mobilenet -d   // deserialize plan file and run inference" << std::endl;
+        return -1;
+    }
+
+    // create a model using the API directly and serialize it to a stream
+    char *trtModelStream{nullptr};
+    size_t size{0};
+
+    if (std::string(argv[1]) == "-s") {
+        IHostMemory* modelStream{nullptr};
+        APIToModel(1, &modelStream);
+        assert(modelStream != nullptr);
+
+        std::ofstream p("mobilenetv3_small.engine");
+        if (!p)
+        {
+            std::cerr << "could not open plan output file" << std::endl;
+            return -1;
+        }
+        p.write(reinterpret_cast<const char*>(modelStream->data()), modelStream->size());
+        modelStream->destroy();
+        return 1;
+    } else if (std::string(argv[1]) == "-d") {
+        std::ifstream file("mobilenetv3_small.engine", std::ios::binary);
+        if (file.good()) {
+            file.seekg(0, file.end);
+            size = file.tellg();
+            file.seekg(0, file.beg);
+            trtModelStream = new char[size];
+            assert(trtModelStream);
+            file.read(trtModelStream, size);
+            file.close();
+        }
+    } else {
+        return -1;
+    }
+
+
+    // Subtract mean from image
+    float data[3 * INPUT_H * INPUT_W];
+    for (int i = 0; i < 3 * INPUT_H * INPUT_W; i++)
+        data[i] = 1.0;
+    PluginFactory pf;
+    IRuntime* runtime = createInferRuntime(gLogger);
+    assert(runtime != nullptr);
+    ICudaEngine* engine = runtime->deserializeCudaEngine(trtModelStream, size, &pf);
+    assert(engine != nullptr);
+    IExecutionContext* context = engine->createExecutionContext();
+    assert(context != nullptr);
+
+    // Run inference
+    float prob[OUTPUT_SIZE];
+    for (int i = 0; i < 100; i++) {
+        auto start = std::chrono::system_clock::now();
+        doInference(*context, data, prob, 1);
+        auto end = std::chrono::system_clock::now();
+        std::cout << std::chrono::duration_cast<std::chrono::milliseconds>(end - start).count() << "ms" << std::endl;
+    }
+
+    // Destroy the engine
+    context->destroy();
+    engine->destroy();
+    runtime->destroy();
+
+    // Print histogram of the output distribution
+    std::cout << "\nOutput:\n\n";
+    for (unsigned int i = 0; i < 20; i++)
+    {
+        std::cout << prob[i] << ", ";
+        //if (i % 10 == 0) std::cout << i / 10 << std::endl;
+    }
+    std::cout << std::endl;
+
+    return 0;
+}