add swin-transformer (#934)

* add swin-transformer * change dir * reformat * fix gen wts and polish readme Co-authored-by: wdhao <873665182@qq.com>
2022-03-15 12:10:49 +08:00 · 2022-03-15 12:10:49 +08:00 · e7ec22d8b8
commit e7ec22d8b8
parent 3f31d76ac4
20 changed files with 4495 additions and 0 deletions
--- a/swin-transformer/semantic-segmentation/CMakeLists.txt
+++ b/swin-transformer/semantic-segmentation/CMakeLists.txt
@ -0,0 +1,75 @@
+cmake_minimum_required(VERSION 3.4)
+
+project(swintransformer)
+
+set(OpenCV_DIR "D:\\opencv\\opencv346\\build")
+set(TENSORRT_DIR "D:\\TensorRT-7.0.0.11.Windows10.x86_64.cuda-10.2.cudnn7.6\\TensorRT-7.0.0.11")
+
+add_definitions(-std=c++11)
+add_definitions(-DAPI_EXPORTS)
+option(CUDA_USE_STATIC_CUDA_RUNTIME OFF)
+set(CMAKE_CXX_STANDARD 11)
+set(CMAKE_BUILD_TYPE Debug)
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -Wall -Ofast -D_MWAITXINTRIN_H_INCLUDED")
+if(WIN32)
+include_directories(${PROJECT_SOURCE_DIR}/include)
+endif(WIN32)
+
+
+find_package(CUDA REQUIRED)
+message(STATUS "    libraries: ${CUDA_LIBRARIES}")
+message(STATUS "    include path: ${CUDA_INCLUDE_DIRS}")
+include_directories(${CUDA_INCLUDE_DIRS})
+set(CUDA_NVCC_PLAGS ${CUDA_NVCC_PLAGS};-std=c++11; -g; -G;-gencode; arch=compute_75;code=sm_75)
+enable_language(CUDA)  # 这一句添加后 ，就会在vs中不需要再手动设置cuda 
+include_directories(${TENSORRT_DIR}\\include)
+link_directories(${TENSORRT_DIR}\\lib)
+
+# file(GLOB SOURCE_FILES "*.cu")
+# cuda_add_library(myplugins SHARED ${PROJECT_SOURCE_DIR}/yololayer.cu ${PROJECT_SOURCE_DIR}/API.h)
+# target_link_libraries(myplugins nvinfer cudart)
+
+# 设置opencv的信息
+find_package(OpenCV QUIET
+    NO_MODULE
+    NO_DEFAULT_PATH
+    NO_CMAKE_PATH
+    NO_CMAKE_ENVIRONMENT_PATH
+    NO_SYSTEM_ENVIRONMENT_PATH
+    NO_CMAKE_PACKAGE_REGISTRY
+    NO_CMAKE_BUILDS_PATH
+    NO_CMAKE_SYSTEM_PATH
+    NO_CMAKE_SYSTEM_PACKAGE_REGISTRY
+)
+
+message(STATUS "OpenCV library status:")
+message(STATUS "    version: ${OpenCV_VERSION}")
+message(STATUS "    libraries: ${OpenCV_LIBS}")
+message(STATUS "    include path: ${OpenCV_INCLUDE_DIRS}")
+
+include_directories(${OpenCV_INCLUDE_DIRS})
+
+
+file(GLOB SOURCE_FILES "*.h" "*.cpp" "*.cu")
+add_executable(swintransformer ${SOURCE_FILES})
+
+target_link_libraries(swintransformer nvinfer nvonnxparser)
+target_link_libraries(swintransformer cudart)
+target_link_libraries(swintransformer ${OpenCV_LIBS})
+
+# if (WIN32)
+    # message(STATUS "copy dll......: ${CMAKE_COMMAND} ${TENSORRT_DIR}")
+    # add_custom_command(TARGET swintransformer POST_BUILD
+        # COMMAND ${CMAKE_COMMAND} -E copy_if_different ${TENSORRT_DIR}/lib/myelin64_1.dll ./${CMAKE_BUILD_TYPE}/myelin64_1.dll
+        # COMMAND ${CMAKE_COMMAND} -E copy_if_different ${TENSORRT_DIR}/lib/nvinfer.dll ./${CMAKE_BUILD_TYPE}/nvinfer.dll
+        # COMMAND ${CMAKE_COMMAND} -E copy_if_different ${TENSORRT_DIR}/lib/nvinfer_plugin.dll ./${CMAKE_BUILD_TYPE}/nvinfer_plugin.dll
+        # COMMAND ${CMAKE_COMMAND} -E copy_if_different ${TENSORRT_DIR}/lib/nvonnxparser.dll ./${CMAKE_BUILD_TYPE}/nvonnxparser.dll
+        # COMMAND ${CMAKE_COMMAND} -E copy_if_different ${TENSORRT_DIR}/lib/nvparsers.dll ./${CMAKE_BUILD_TYPE}/nvparsers.dll
+        # COMMAND ${CMAKE_COMMAND} -E copy_if_different ${TENSORRT_DIR}/lib/nvserialize.dll ./${CMAKE_BUILD_TYPE}/nvserialize.dll
+        # COMMAND ${CMAKE_COMMAND} -E copy_if_different ${CUDA_TOOLKIT_ROOT_DIR}/bin/cublas64_10.dll ./${CMAKE_BUILD_TYPE}/cublas64_10.dll
+    # )
+# endif(WIN32)
+
+if(UNIX)
+add_definitions(-O2 -pthread)
+endif(UNIX)
--- a/swin-transformer/semantic-segmentation/README.md
+++ b/swin-transformer/semantic-segmentation/README.md
@ -0,0 +1,36 @@
+# swin_transform
+
+The Pytorch implementation is [microsoft/Swin-Transformer](https://github.com/microsoft/Swin-Transformer.git).
+
+Only support Swin-T, welcome the PR for other backbones.
+
+## How to Run
+
+1. generate .wts from pytorch with .pt, or download .wts from model zoo
+
+```
+git clone https://github.com/microsoft/Swin-Transformer.git
+git clone https://github.com/wang-xinyu/tensorrtx.git
+
+python gen_wts.py Swin-Transform.pt
+// a file 'Swin-Transform.wts' will be generated.
+```
+
+2. build tensorrtx/swin-transform and run
+
+```
+cd {tensorrtx}/swin-transform/semantic-segmentation/
+mkdir build
+cd build
+cp {microsoft}/Swin-Transformer/Swin-Transform.wts {tensorrtx}/swin-transformer/semantic-segmentation/build
+cmake ..
+make
+sudo ./swintransformer -s [.wts] [.engine]   // serialize model to plan file
+sudo ./swintransformer -d [.engine] [image folder]  // deserialize and run inference, the images in [image folder] will be processed.
+
+```
+
+## More Information
+
+See the readme in [home page.](https://github.com/wang-xinyu/tensorrtx)
+
--- a/swin-transformer/semantic-segmentation/UpsampleKernel.cu
+++ b/swin-transformer/semantic-segmentation/UpsampleKernel.cu
@ -0,0 +1,146 @@
+#include "UpsmapleKernel.h"
+
+
+/**
+ * @brief caculate the number of cuda kernel for upsample. (Cite from: 《GPU高性能编程CUDA实战》P46,P47)
+ * 
+ * @param total_thread_num: the number of cuda thread of you want to used for upsample
+ * @param max_thread_num: the gpu device property
+ * @return int  the number of cuda kernel for upsample
+ */
+int get_kernel_num(int total_thread_num, int max_thread_num)
+{
+    return (total_thread_num + max_thread_num - 1)/max_thread_num;
+}
+
+int get_max_thread_num()
+{
+    cudaDeviceProp prop;
+    cudaGetDeviceProperties(&prop, 0);
+    return prop.maxThreadsPerBlock;
+}
+
+__host__ __forceinline__ float linear_upsampling_compute_scale(int input_size, int output_size)
+{
+    return float(input_size)/float(output_size) ;
+}
+
+__device__ __forceinline__ float linear_upsampling_compute_source_index(float scale, int dst_index, int intput_size)
+{
+    float src_idx = scale * (dst_index + 0.5)-0.5;
+    return (src_idx>=0) ? src_idx : 0;
+}
+
+
+__device__ __forceinline__ int get_index(const int batch_idx, const int channel_idx, const int height_idx, const int width_idx, 
+                const int batch_total, const int channel_total, const int width)
+{
+    int ret_idx = batch_idx * batch_total
+                    + channel_idx * channel_total
+                    + height_idx * width
+                    + width_idx;
+    return ret_idx;
+}
+
+/**
+ * @brief 
+ * 
+ * @tparam T 
+ * @param n 
+ * @param input_shape: input data shape. such as [batch, channel, height, width] 
+ * @param rate_h 
+ * @param rate_w 
+ * @param inputs 
+ * @param outputs 
+ * @return __global__ BilinearKernel 
+ * @TODO: 
+ *  
+ */
+
+
+template <typename T>
+__global__ void BilinearKernel(
+        const int n,
+        int input_b,
+        int input_c,
+        int input_h,
+        int input_w,
+        int output_h,
+        int output_w,
+        const float rate_h,
+        const float rate_w,
+        const T* inputs,
+        T* outputs)
+{
+
+    int index = threadIdx.x + blockIdx.x * blockDim.x;
+    if(index < n)
+    {
+        const int w2 = index % output_w;
+        const int h2 = index / output_w;
+
+
+        const float h1r = linear_upsampling_compute_source_index(rate_h, h2, input_h);
+        const int h1 = int(h1r);
+        const int h1p = (h1 < input_h - 1) ? 1 : 0;
+        const float h1lambda = h1r - h1;
+        const float h0lambda = 1 - h1lambda;
+
+        const float w1r = linear_upsampling_compute_source_index(rate_w, w2, input_w);
+        const int w1 = int(w1r);
+        const int w1p = (w1 < input_w - 1) ? 1 : 0;
+        const float w1lambda = w1r - w1;
+        const float w0lambda = 1 - w1lambda;
+
+        int s_batch_total_1 = input_c * input_h * input_w;
+        int s_channel_total_1 = input_h * input_w;
+
+        int s_batch_total_2 = input_c * output_h * output_w;
+        int s_channel_total_2 = output_h * output_w;
+
+
+        const int batch_size = input_b;
+        const int channel_size = input_c;
+
+        for(int b_idx=0; b_idx<batch_size; b_idx++)
+        {
+            for(int c=0; c<channel_size; c++)
+            {
+                const T val = h0lambda * (w0lambda * inputs[get_index(b_idx, c, h1, w1, s_batch_total_1, s_channel_total_1, input_w)]
+                                    + w1lambda * inputs[get_index(b_idx, c, h1, w1+w1p, s_batch_total_1, s_channel_total_1, input_w)])
+                                    + h1lambda * (w0lambda * inputs[get_index(b_idx, c, h1+h1p, w1, s_batch_total_1, s_channel_total_1, input_w)]
+                                    + w1lambda * inputs[get_index(b_idx, c, h1+h1p, w1+w1p, s_batch_total_1, s_channel_total_1, input_w)]);
+                outputs[get_index(b_idx, c, h2, w2, s_batch_total_2, s_channel_total_2, output_w)] = val;
+                
+            }
+        }
+    }
+}
+
+
+int UpsampleInference(
+    cudaStream_t stream,
+    int n,
+    int input_b,
+    int input_c,
+    int input_h,
+    int input_w,
+    float scale_h,
+    float scale_w,
+    const void* inputs,
+    void* outputs)
+{
+    int output_h = int(input_h * scale_h);
+    int output_w = int(input_w * scale_w);
+    int max_threads = get_max_thread_num();
+    int kernel_num = get_kernel_num(n, max_threads);
+    float rate_h = linear_upsampling_compute_scale(input_h, output_h);
+    float rate_w = linear_upsampling_compute_scale(input_w, output_w);
+
+    BilinearKernel<float><<< kernel_num, max_threads, 0, stream>>>(n,input_b,input_c,input_h,input_w,
+                                                                                    output_h, output_w, 
+                                                                                    rate_h, rate_w,
+                                                                                    static_cast<const float*>(inputs),
+                                                                                    static_cast<float*>(outputs));
+    return 0;
+}
--- a/swin-transformer/semantic-segmentation/UpsamplePlugin.cpp
+++ b/swin-transformer/semantic-segmentation/UpsamplePlugin.cpp
@ -0,0 +1,235 @@
+#include <iostream>
+#include "UpsmapleKernel.h"
+#include "UpsamplePlugin.h"
+
+#include <cassert>
+#include <cstring>
+
+using namespace nvinfer1;
+
+// Upsample plugin specific constants
+namespace {
+    static const char* UPSAMPLE_PLUGIN_VERSION{"1"};
+    static const char* UPSAMPLE_PLUGIN_NAME{"UpsamplePlugin"};
+}
+
+// Static class fields initialization
+PluginFieldCollection UpsamplePluginCreator::mFC{};
+std::vector<PluginField> UpsamplePluginCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(UpsamplePluginCreator);
+
+template<typename T>
+void writeToBuffer(char*& buffer, const T& val)
+{
+    *reinterpret_cast<T*>(buffer) = val;
+    buffer += sizeof(T);
+}
+
+// Helper function for deserializing plugin
+template<typename T>
+T readFromBuffer(const char*& buffer)
+{
+    T val = *reinterpret_cast<const T*>(buffer);
+    buffer += sizeof(T);
+    return val;
+}
+
+UpsamplePlugin::UpsamplePlugin(const std::string name, float scale_h, float scale_w)
+    : mLayerName(name)
+    , mScaleFactor_h(scale_h)
+    , mScaleFactor_w(scale_w)
+{
+    mInputShape.c() = -1;
+    mInputShape.h() = -1;
+    mInputShape.w() = -1;
+    mInputVolume = 0;
+}
+
+UpsamplePlugin::UpsamplePlugin(const std::string name, const void* data, size_t length)
+    : mLayerName(name)
+{
+    const char *d = static_cast<const char *>(data);
+    const char *a = d;
+
+    mScaleFactor_h = readFromBuffer<float>(d);
+    mScaleFactor_w = readFromBuffer<float>(d);
+    mInputVolume = readFromBuffer<size_t>(d);
+    mInputShape.c() = readFromBuffer<int>(d);
+    mInputShape.h() = readFromBuffer<int>(d);
+    mInputShape.w() = readFromBuffer<int>(d);
+
+    assert(d == (a + length));
+
+}
+
+const char* UpsamplePlugin::getPluginType() const
+{
+    return UPSAMPLE_PLUGIN_NAME;
+}
+
+const char* UpsamplePlugin::getPluginVersion() const
+{
+    return UPSAMPLE_PLUGIN_VERSION;
+}
+
+int UpsamplePlugin::getNbOutputs() const
+{
+    return 1;
+}
+
+Dims UpsamplePlugin::getOutputDimensions(int index, const Dims* inputs, int nbInputDims)
+{
+    assert(index == 0);
+    assert(nbInputDims == 1);
+    assert(inputs[0].nbDims == 3);
+    return nvinfer1::DimsCHW{inputs[0].d[0],int(inputs[0].d[1]*mScaleFactor_h), int(inputs[0].d[2]*mScaleFactor_w)};
+}
+
+int UpsamplePlugin::initialize()
+{
+    //printf("UpsamplePlugin::initialize\n");
+    return 0;
+}
+
+
+int UpsamplePlugin::enqueue(int batchSize, const void* const* inputs, void** outputs, void*, cudaStream_t stream)
+{
+    //printf("UpsamplePlugin::enqueue\n");
+    int status = -1;
+
+    // Our plugin outputs only one tensor
+    void* output = outputs[0];
+
+    // Launch CUDA kernel wrapper and save its return value
+    status = UpsampleInference(stream, mInputVolume, 
+                                batchSize, mInputShape.c(), mInputShape.h(), mInputShape.w(),
+                                mScaleFactor_h,mScaleFactor_w,
+                                inputs[0], output);
+    return status;
+}
+
+size_t UpsamplePlugin::getSerializationSize() const
+{
+    //printf("UpsamplePlugin::getSerializationSize\n");
+    return sizeof(mScaleFactor_h)  + sizeof(mScaleFactor_w) +
+            sizeof(mInputVolume) + sizeof(mInputShape.c()) + 
+            sizeof(mInputShape.h()) + sizeof(mInputShape.w());
+}
+
+
+void UpsamplePlugin::serialize(void* buffer) const 
+{
+    //printf("UpsamplePlugin::serialize\n");
+    char *d = static_cast<char *>(buffer);
+    const char *a = d;
+
+    writeToBuffer(d, mScaleFactor_h);
+    writeToBuffer(d, mScaleFactor_w);
+    writeToBuffer(d, mInputVolume);
+    writeToBuffer(d, mInputShape.c());
+    writeToBuffer(d, mInputShape.h());
+    writeToBuffer(d, mInputShape.w());
+
+    assert(d == a + getSerializationSize());
+}
+
+void UpsamplePlugin::configureWithFormat(const Dims* inputs, int nbInputs, const Dims* outputs, int nbOutputs, DataType type, PluginFormat format, int)
+{
+    assert(nbOutputs == 1);
+    assert(type == DataType::kFLOAT);
+    assert(format == PluginFormat::kNCHW);
+    assert(inputs[0].nbDims == 3);
+
+    size_t volume = int(inputs[0].d[1]*mScaleFactor_h) * int(inputs[0].d[2]*mScaleFactor_w);
+    mInputVolume = volume;
+    mInputShape.c() = inputs[0].d[0];
+    mInputShape.h() = inputs[0].d[1];
+    mInputShape.w() = inputs[0].d[2];
+}
+
+bool UpsamplePlugin::supportsFormat(DataType type, PluginFormat format) const
+{
+    if (type == DataType::kFLOAT && format == PluginFormat::kNCHW)
+        return true;
+    else
+        return false;
+}
+
+void UpsamplePlugin::terminate() {}
+
+void UpsamplePlugin::destroy() {
+    // This gets called when the network containing plugin is destroyed
+    delete this;
+}
+
+IPluginV2* UpsamplePlugin::clone() const
+{
+    return new UpsamplePlugin(mLayerName, mScaleFactor_h, mScaleFactor_w);
+}
+
+void UpsamplePlugin::setPluginNamespace(const char* libNamespace) 
+{
+    mNamespace = libNamespace;
+}
+
+const char* UpsamplePlugin::getPluginNamespace() const
+{
+    return mNamespace.c_str();
+}
+
+UpsamplePluginCreator::UpsamplePluginCreator()
+{
+    mPluginAttributes.emplace_back(PluginField("scaleFactor", nullptr, PluginFieldType::kFLOAT32, 2));
+
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+const char* UpsamplePluginCreator::getPluginName() const
+{
+    return UPSAMPLE_PLUGIN_NAME;
+}
+
+const char* UpsamplePluginCreator::getPluginVersion() const
+{
+    return UPSAMPLE_PLUGIN_VERSION;
+}
+
+const PluginFieldCollection* UpsamplePluginCreator::getFieldNames()
+{
+    return &mFC;
+}
+
+IPluginV2* UpsamplePluginCreator::createPlugin(const char* name, const PluginFieldCollection* fc)
+{
+    float scaleFactor_h = 0.f;
+    float scaleFactor_w = 0.f;
+    const PluginField* fields = fc->fields;
+
+    assert(fc->nbFields == 1);
+    for (int i = 0; i < fc->nbFields; i++){
+    
+        if (strcmp(fields[i].name, "scaleFactor") == 0) {
+            assert(fields[i].type == PluginFieldType::kFLOAT32);
+            scaleFactor_h = *(static_cast<const float*>(fields[i].data));
+            scaleFactor_w = *(static_cast<const float*>(fields[i].data)+1);
+            //std::cout<<scaleFactor_h<< " , "<<scaleFactor_w<<std::endl;
+        } 
+    }
+    return new UpsamplePlugin(name, scaleFactor_h, scaleFactor_w);
+}
+
+IPluginV2* UpsamplePluginCreator::deserializePlugin(const char* name, const void* serialData, size_t serialLength)
+{
+    return new UpsamplePlugin(name, serialData, serialLength);
+}
+
+void UpsamplePluginCreator::setPluginNamespace(const char* libNamespace) 
+{
+    mNamespace = libNamespace;
+}
+
+const char* UpsamplePluginCreator::getPluginNamespace() const
+{
+    return mNamespace.c_str();
+}
--- a/swin-transformer/semantic-segmentation/UpsamplePlugin.h
+++ b/swin-transformer/semantic-segmentation/UpsamplePlugin.h
@ -0,0 +1,88 @@
+#ifndef UPSAMPLE_PLUGIN_H
+#define UPSAMPLE_PLUGIN_H
+
+#include "NvInferPlugin.h"
+#include <string>
+#include <vector>
+
+
+using namespace nvinfer1;
+
+class UpsamplePlugin : public IPluginV2
+{
+public:
+    UpsamplePlugin(const std::string name, float scale_h,float scale_w);
+
+    UpsamplePlugin(const std::string name, const void* data, size_t length);
+
+    // It doesn't make sense to make UpsamplePlugin without arguments, so we delete default constructor.
+    UpsamplePlugin() = delete;
+
+    int getNbOutputs() const override;
+
+    Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override;
+
+    int initialize() override;
+
+    void terminate() override;
+
+    size_t getWorkspaceSize(int) const override { return 0; };
+
+    int enqueue(int batchSize, const void* const* inputs, void** outputs, void* workspace, cudaStream_t stream) override;
+
+    size_t getSerializationSize() const override;
+
+    void serialize(void* buffer) const override;
+
+    void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, DataType type, PluginFormat format, int maxBatchSize) override;
+
+    bool supportsFormat(DataType type, PluginFormat format) const override;
+
+    const char* getPluginType() const override;
+
+    const char* getPluginVersion() const override;
+
+    void destroy() override;
+
+    nvinfer1::IPluginV2* clone() const override;
+
+    void setPluginNamespace(const char* pluginNamespace) override;
+
+    const char* getPluginNamespace() const override;
+
+private:
+    const std::string mLayerName;
+    bool mAlignCorners;
+    float mScaleFactor_h;
+    float mScaleFactor_w;
+    size_t mInputVolume;
+    DimsCHW mInputShape;
+    std::string mNamespace;
+};
+
+class UpsamplePluginCreator : public IPluginCreator
+{
+public:
+    UpsamplePluginCreator();
+
+    const char* getPluginName() const override;
+
+    const char* getPluginVersion() const override;
+
+    const PluginFieldCollection* getFieldNames() override;
+
+    IPluginV2* createPlugin(const char* name, const PluginFieldCollection* fc) override;
+
+    IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) override;
+    
+    void setPluginNamespace(const char* pluginNamespace) override;
+
+    const char* getPluginNamespace() const override;
+
+private:
+    static PluginFieldCollection mFC;
+    static std::vector<PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+
+#endif
--- a/swin-transformer/semantic-segmentation/UpsmapleKernel.h
+++ b/swin-transformer/semantic-segmentation/UpsmapleKernel.h
@ -0,0 +1,20 @@
+#ifndef UPSAMPLE_KERNEL_H
+#define UPSAMPLE_KERNEL_H
+
+#include <iostream>
+#include "NvInfer.h"
+
+int UpsampleInference(
+    cudaStream_t stream,
+    int n,
+    int input_b,
+    int input_c,
+    int input_h,
+    int input_w,
+    float scale_h,
+    float scale_w,
+    const void* inputs,
+    void* outputs);
+
+
+#endif
--- a/swin-transformer/semantic-segmentation/common.hpp
+++ b/swin-transformer/semantic-segmentation/common.hpp
@ -0,0 +1,888 @@
+#ifndef COMMON_HPP
+#define COMMON_HPP
+
+#include "layerNorm.h"
+#include "NvInfer.h"
+#include "NvInfer.h"
+#include "NvInferPlugin.h"
+#include "cuda_runtime_api.h"
+#include <assert.h>
+#include <map>
+#include <fstream>
+#include<opencv2/core/core.hpp>
+#include<opencv2/imgproc/imgproc.hpp>
+#include<opencv2/imgcodecs/imgcodecs.hpp>
+#include<opencv2/dnn/dnn.hpp>
+
+using namespace nvinfer1;
+#define CHECK(status) \
+    do\
+    {\
+        auto ret = (status);\
+        if (ret != 0)\
+        {\
+            std::cerr << "Cuda failure: " << ret << std::endl;\
+            abort();\
+        }\
+    } while (0)
+
+void mblobFromImages(cv::InputArrayOfArrays images_, cv::OutputArray blob_,
+    cv::Size size, const cv::Scalar& mean_, const cv::Scalar& std_, bool swapRB, bool crop)
+{
+    //CV_TRACE_FUNCTION();
+    std::vector<cv::Mat> images;
+    images_.getMatVector(images);
+    CV_Assert(!images.empty());
+    for (int i = 0; i < images.size(); i++)
+    {
+        cv::Size imgSize = images[i].size();
+        if (size == cv::Size())
+            size = imgSize;
+        if (size != imgSize)
+        {
+            if (crop)
+            {
+                float resizeFactor = std::max(size.width / (float)imgSize.width,
+                    size.height / (float)imgSize.height);
+                resize(images[i], images[i], cv::Size(), resizeFactor, resizeFactor, cv::INTER_LINEAR);
+                cv::Rect crop(cv::Point(0.5 * (images[i].cols - size.width),
+                    0.5 * (images[i].rows - size.height)),
+                    size);
+                images[i] = images[i](crop);
+            }
+            else
+                resize(images[i], images[i], size, 0, 0, cv::INTER_LINEAR);
+        }
+        if (images[i].depth() == CV_8U)
+            images[i].convertTo(images[i], CV_32F);
+        cv::Scalar mean = mean_;
+        cv::Scalar std_num = std_;
+        if (swapRB)
+        {
+            std::swap(mean[0], mean[2]);
+            std::swap(std_num[0], std_num[2]);
+        }
+
+        images[i] -= mean;
+        images[i] /= std_num;
+    }
+
+    size_t i, nimages = images.size();
+    cv::Mat image0 = images[0];
+    int nch = image0.channels();
+    CV_Assert(image0.dims == 2);
+    cv::Mat image;
+    if (nch == 3 || nch == 4)
+    {
+        int sz[] = { (int)nimages, nch, image0.rows, image0.cols };
+        blob_.create(4, sz, CV_32F);
+        cv::Mat blob = blob_.getMat();
+        cv::Mat ch[4];
+
+        for (i = 0; i < nimages; i++)
+        {
+            image = images[i];
+            CV_Assert(image.depth() == CV_32F);
+            nch = image.channels();
+            CV_Assert(image.dims == 2 && (nch == 3 || nch == 4));
+            CV_Assert(image.size() == image0.size());
+
+            for (int j = 0; j < nch; j++)
+                ch[j] = cv::Mat(image.rows, image.cols, CV_32F, blob.ptr((int)i, j));
+            if (swapRB)
+                std::swap(ch[0], ch[2]);
+            split(image, ch);
+        }
+    }
+    else
+    {
+        CV_Assert(nch == 1);
+        int sz[] = { (int)nimages, 1, image0.rows, image0.cols };
+        blob_.create(4, sz, CV_32F);
+        cv::Mat blob = blob_.getMat();
+
+        for (i = 0; i < nimages; i++)
+        {
+            cv::Mat image = images[i];
+            CV_Assert(image.depth() == CV_32F);
+            nch = image.channels();
+            CV_Assert(image.dims == 2 && (nch == 1));
+            CV_Assert(image.size() == image0.size());
+
+            image.copyTo(cv::Mat(image.rows, image.cols, CV_32F, blob.ptr((int)i, 0)));
+        }
+    }
+}
+cv::Mat BlobFromImages(cv::InputArrayOfArrays images, cv::Size size,
+    const cv::Scalar& mean, const cv::Scalar& std_num, bool swapRB, bool crop)
+{
+    //CV_TRACE_FUNCTION();
+    cv::Mat blob;
+    mblobFromImages(images, blob, size, mean, std_num, swapRB, crop);
+    return blob;
+}
+void debug_print(ITensor *input_tensor,std::string head)
+{
+    std::cout << head<< " : ";
+
+       for (int i = 0; i < input_tensor->getDimensions().nbDims; i++)
+       {
+           std::cout << input_tensor->getDimensions().d[i] << " ";
+       }
+       std::cout<<std::endl;
+
+}
+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;
+}
+
+ITensor* m_layerNorm(INetworkDefinition *m_Network,std::map<std::string, Weights> weightMap,ITensor *input, string lname)
+{
+    auto creator = getPluginRegistry()->getPluginCreator("layerNorm_trt","1");
+
+    PluginField pluginMultidata[2];
+
+    const PluginFieldCollection* pluginData = creator->getFieldNames();
+    IPluginV2 *pluginObj = creator->createPlugin(lname.c_str(), pluginData);
+    ITensor* inputTensors[] = {input};
+    auto ln_ms = m_Network->addPluginV2(inputTensors, 1, *pluginObj);
+    auto ln_m = m_Network->addElementWise(*input,*ln_ms->getOutput(0),ElementWiseOperation::kSUB);
+    auto ln = m_Network->addElementWise(*ln_m->getOutput(0),*ln_ms->getOutput(1),ElementWiseOperation::kDIV);
+    Weights W = weightMap[lname + ".weight"];
+    int len = W.count;
+    Dims wb ;
+    wb.nbDims = ln->getOutput(0)->getDimensions().nbDims;
+    for (int i = 0 ; i < wb.nbDims; i++)
+    {
+        if (i != wb.nbDims -1)
+            wb.d[i] = 1;
+        else{
+            wb.d[i] = len;
+        }
+    }
+    auto wgts = m_Network->addConstant(wb,W);
+    auto p_w = m_Network->addElementWise(*ln->getOutput(0),*wgts->getOutput(0),ElementWiseOperation::kPROD);
+    Weights B = weightMap[lname + ".bias"];
+    auto bias = m_Network->addConstant(wb,B);
+    auto sum_bias = m_Network->addElementWise(*p_w->getOutput(0),*bias->getOutput(0),ElementWiseOperation::kSUM);
+    debug_print(sum_bias->getOutput(0),lname);
+    return sum_bias->getOutput(0);
+}
+ITensor* layerNorm(INetworkDefinition *m_Network,std::map<std::string, Weights> weightMap,ITensor *input, string lname)
+{
+    auto mean = m_Network->addReduce(*input, ReduceOperation::kAVG, 2, true);
+    assert(mean);
+
+    auto sub_mean = m_Network->addElementWise(*input, *mean->getOutput(0), ElementWiseOperation::kSUB);
+    assert(sub_mean);
+//    float SCALING_ONE = 1.0;
+//    float SHIFT_ZERO = 0.0;
+//    float POWER_TWO = 2.0;
+//    // implement pow2 with scale
+//    Weights scale{ DataType::kFLOAT, &SCALING_ONE, 1 };
+//    Weights shift{ DataType::kFLOAT, &SHIFT_ZERO, 1 };
+//    Weights power{ DataType::kFLOAT, &POWER_TWO, 1 };
+//    auto pow2 = m_Network->addScaleNd(*sub_mean->getOutput(0), ScaleMode::kUNIFORM, shift, scale, power,0);
+//    assert(pow2);
+    auto pow2 = m_Network->addElementWise(*sub_mean->getOutput(0), *sub_mean->getOutput(0), ElementWiseOperation::kPROD);
+    assert(pow2);
+    debug_print(pow2->getOutput(0),"pow2");
+    auto pow_mean = m_Network->addReduce(*pow2->getOutput(0), ReduceOperation::kAVG, 2, true);
+    assert(pow_mean);
+    debug_print(pow_mean->getOutput(0),"pow_mean");
+    float E = 1e-5;
+    Weights EPS{DataType::kFLOAT,nullptr,1};
+    EPS.values = &E;
+    auto eps = m_Network->addConstant(Dims2{1,1}, EPS);
+    assert(eps);
+
+    auto add_eps = m_Network->addElementWise(*pow_mean->getOutput(0), *eps->getOutput(0), ElementWiseOperation::kSUM);
+    assert(add_eps);
+
+    auto sqrt = m_Network->addUnary(*add_eps->getOutput(0), UnaryOperation::kSQRT);
+    assert(sqrt);
+
+    auto div = m_Network->addElementWise(*sub_mean->getOutput(0), *sqrt->getOutput(0), ElementWiseOperation::kDIV);
+    assert(div);
+    debug_print(div->getOutput(0),"div");
+
+    string weightsFile = lname + ".weight";
+    string biasFile = lname + ".bias";
+
+    int d_model = input->getDimensions().d[input->getDimensions().nbDims - 1];
+    cout<<"d_model = "<<d_model<<endl;
+    float *pval = reinterpret_cast<float*>(malloc(sizeof(float) * d_model));
+    for (int i = 0; i < d_model; i++) {
+        pval[i] = 1.0;
+    }
+    Weights norm1_power{ DataType::kFLOAT, pval, d_model };
+    auto affine = m_Network->addScaleNd(
+        *div->getOutput(0),
+        ScaleMode::kELEMENTWISE,
+        weightMap[biasFile],
+        weightMap[weightsFile],
+        norm1_power,1);
+    assert(affine);
+    return affine->getOutput(0);
+}
+ITensor* conv(INetworkDefinition *m_Network,std::map<std::string, Weights> weightMap,ITensor *input, string lname,
+              int c_out,bool bias = true,int k = 4 , int s = 4, int p = 0)
+{
+    Weights Bias{ DataType::kFLOAT, nullptr, 0 };
+    if(bias)
+        Bias = weightMap[lname + ".bias"];
+    auto out = m_Network->addConvolutionNd(*input,c_out,Dims2{k,k},weightMap[lname + ".weight"],Bias);
+    out->setStrideNd(Dims2{s,s});
+    out->setPaddingNd(Dims2{p,p});
+    out->setNbGroups(1);
+    debug_print(out->getOutput(0),lname);
+    return out->getOutput(0);
+}
+ITensor* shuffle_reshape(INetworkDefinition *m_Network,ITensor *input,Dims reshapeDims)
+{
+    auto out = m_Network->addShuffle(*input);
+    out->setReshapeDimensions(reshapeDims);
+    debug_print(out->getOutput(0),"reshape");
+    return out->getOutput(0);
+}
+ITensor* shuffle_permute(INetworkDefinition *m_Network,ITensor *input,Permutation permutation)
+{
+    auto out = m_Network->addShuffle(*input);
+    out->setFirstTranspose(permutation);
+    debug_print(out->getOutput(0),"permute");
+    return out->getOutput(0);
+}
+ITensor* shuffle_reshapeApermute(INetworkDefinition *m_Network,ITensor *input,Dims reshapeDims,
+                                 Permutation permutation,bool firstReshape)
+{
+    auto out = m_Network->addShuffle(*input);
+    out->setReshapeDimensions(reshapeDims);
+    if(firstReshape)
+        out->setSecondTranspose(permutation);
+    else
+        out->setFirstTranspose(permutation);
+    debug_print(out->getOutput(0),"shuffle");
+    return out->getOutput(0);
+}
+ITensor* trt_transform_imgMask(INetworkDefinition *m_Network,int hw, int window_size, int shift_size)
+{
+    int Hp = hw;
+    int Wp = hw;
+    Weights Mask_param{DataType::kFLOAT,nullptr,Hp*Wp};
+    float *mask_param = new float[Hp*Wp];
+    for(int i = 0; i < Hp ; i++)
+    {
+        for(int j = 0; j < Wp; j++)
+        {
+            if(i<Hp-window_size && j<Wp-window_size)
+                mask_param[i*Wp + j] = 0.0;
+            else if(i<Hp-window_size && j>=Wp-window_size && j < Wp-shift_size)
+                mask_param[i*Wp + j] = 1.0;
+            else if(i<Hp-window_size &&  j >= Wp-shift_size)
+                mask_param[i*Wp + j] = 2.0;
+
+            else if(i >= Hp-window_size && i < Hp-shift_size && j<Wp-window_size)
+                mask_param[i*Wp + j] = 3.0;
+            else if(i >= Hp-window_size && i < Hp-shift_size && j>=Wp-window_size && j < Wp-shift_size)
+                mask_param[i*Wp + j] = 4.0;
+            else if(i >= Hp-window_size && i < Hp-shift_size && j >= Wp-shift_size)
+                mask_param[i*Wp + j] = 5.0;
+
+            else if(i >=  Hp-shift_size && j<Wp-window_size)
+                mask_param[i*Wp + j] = 6.0;
+            else if(i >=  Hp-shift_size && j>=Wp-window_size && j < Wp-shift_size)
+                mask_param[i*Wp + j] = 7.0;
+            else if(i >=  Hp-shift_size && j >= Wp-shift_size)
+                mask_param[i*Wp + j] = 8.0;
+            else{
+                cout<<" i && j not limit"<<endl;
+                return nullptr;
+            }
+        }
+    }
+    Mask_param.values = mask_param;
+    auto img_mask = m_Network->addConstant(Dims4{1,Hp,Wp,1},Mask_param);
+    auto img_mask_shuffle = m_Network->addShuffle(*img_mask->getOutput(0));
+    Dims shuffle1_dims;
+    shuffle1_dims.nbDims = 6;
+    int dims[] = {1,Hp/window_size,window_size,Wp/window_size,window_size,1};
+    for(int i = 0 ; i < 6; i++)
+        shuffle1_dims.d[i] = dims[i];
+    img_mask_shuffle->setReshapeDimensions(shuffle1_dims);
+    img_mask_shuffle->setSecondTranspose(Permutation{0,1,3,2,4,5});
+    auto img_mask_shuffle2 = m_Network->addShuffle(*img_mask_shuffle->getOutput(0));
+    img_mask_shuffle2->setReshapeDimensions(Dims3{-1,1,window_size*window_size});
+    auto img_mask_shuffle3 = m_Network->addShuffle(*img_mask_shuffle->getOutput(0)) ;
+    img_mask_shuffle3->setReshapeDimensions(Dims3{-1,window_size*window_size,1});
+    auto atten_mask = m_Network->addElementWise(*img_mask_shuffle2->getOutput(0),*img_mask_shuffle3->getOutput(0),ElementWiseOperation::kSUB);
+
+    auto creator = getPluginRegistry()->getPluginCreator("fillmaskLayer_TRT", "1");
+    const PluginFieldCollection* pluginData = creator->getFieldNames();
+    IPluginV2 *pluginObj = creator->createPlugin("fillmask", pluginData);
+    ITensor* inputTensors[] = {atten_mask->getOutput(0)};
+    auto fillmask = m_Network->addPluginV2(inputTensors, 1, *pluginObj);
+
+    debug_print(fillmask->getOutput(0),"imgMask");
+    return fillmask->getOutput(0);
+}
+ITensor* trt_transform_pad(INetworkDefinition *m_Network,ITensor *input,int window_size)
+{
+    int h = input->getDimensions().d[0];
+    int w = input->getDimensions().d[1];
+    int c = input->getDimensions().d[2];
+    int pad_h = (window_size - h%window_size)%window_size;
+    int pad_w = (window_size - w%window_size)%window_size;
+
+    ITensor* temp = input;
+    if(pad_h != 0)
+    {
+        Weights pad1{DataType::kFLOAT,nullptr,pad_h*w*c};
+        cout<<pad_h*w*c<<endl;
+        float *p1 = new float[pad_h*w*c];
+        for(int i = 0 ; i < pad_h*w*c; i++)
+            p1[i] = 0.f;
+        pad1.values = p1;
+        auto Pad1 = m_Network->addConstant(Dims3{pad_h,w,c},pad1);
+        ITensor *cat1[2] = {temp,Pad1->getOutput(0)};
+        auto xp1 = m_Network->addConcatenation(cat1,2);
+        xp1->setAxis(0);
+        temp = xp1->getOutput(0);
+    }
+    if(pad_w != 0)
+    {
+        Weights pad2{DataType::kFLOAT,nullptr,pad_w*(h+pad_h)*c};
+        cout<<pad_w*(h+pad_h)*c<<endl;
+        float *p2 = new float[pad_w*(h+pad_h)*c];
+        for(int i = 0 ; i < pad_w*(h+pad_h)*c; i++)
+            p2[i] = 0.0f;
+        pad2.values = p2;
+        auto Pad2 = m_Network->addConstant(Dims3{(h+pad_h),pad_w,c},pad2);
+        ITensor *cat2[] = {temp,Pad2->getOutput(0)};
+        auto xp2 = m_Network->addConcatenation(cat2,2);
+        xp2->setAxis(1);
+        temp = xp2->getOutput(0);
+    }
+    debug_print(temp, "pad");
+    return  temp;
+}
+ITensor* trt_swinRoll(INetworkDefinition *m_Network,ITensor *input,vector<int> shifts, vector<int> dims)
+{
+    int len = shifts.size();
+    Dims input_dim = input->getDimensions();
+    int nbdims = input_dim.nbDims;
+    ITensor *temp = input;
+    for(int i = 0 ; i < len; i++)
+    {
+        Dims start, size,stride;
+        start.nbDims = nbdims;
+        size.nbDims = nbdims;
+        stride.nbDims = nbdims;
+        if(shifts[i] > 0)
+        {
+            for(int j = 0 ; j < nbdims; j++)
+            {
+                if(j != (dims[i] -1 ))
+                {
+                    start.d[j] = 0;
+                    size.d[j] = input_dim.d[j];
+                    stride.d[j] = 1;
+                }
+                else{
+                    start.d[j] = 0;
+                    size.d[j] = input_dim.d[j] - shifts[i];
+                    stride.d[j] = 1;
+                }
+            }
+
+            auto cat1 = m_Network->addSlice(*temp,start,size,stride);
+
+            for(int j = 0 ; j < nbdims; j++)
+            {
+                if(j != (dims[i] - 1))
+                {
+                    start.d[j] = 0;
+                    size.d[j] = input_dim.d[j];
+                    stride.d[j] = 1;
+                }
+                else{
+                    start.d[j] = input_dim.d[j] - shifts[i];
+                    size.d[j] = shifts[i];
+                    stride.d[j] = 1;
+                }
+            }
+            auto cat2 = m_Network->addSlice(*temp,start,size,stride);
+            ITensor *cat[] ={cat2->getOutput(0),cat1->getOutput(0)};
+            auto Cat = m_Network->addConcatenation(cat,2);
+            Cat->setAxis(dims[i] - 1);
+            temp = Cat->getOutput(0);
+        }
+        if(shifts[i] < 0)
+        {
+            for(int j = 0 ; j < nbdims; j++)
+            {
+                if(j != (dims[i] - 1))
+                {
+                    start.d[j] = 0;
+                    size.d[j] = input_dim.d[j];
+                    stride.d[j] = 1;
+                }
+                else{
+                    start.d[j] = 0;
+                    size.d[j] = abs(shifts[i]);
+                    stride.d[j] = 1;
+                }
+            }
+            auto cat1 = m_Network->addSlice(*temp,start,size,stride);
+            debug_print(cat1->getOutput(0), "cat1 dims : ");
+            for(int j = 0 ; j < nbdims; j++)
+            {
+                if(j != (dims[i] - 1))
+                {
+                    start.d[j] = 0;
+                    size.d[j] = input_dim.d[j];
+                    stride.d[j] = 1;
+                }
+                else{
+                    start.d[j] =  abs(shifts[i]);
+                    size.d[j] = input_dim.d[j] - abs(shifts[i]);
+                    stride.d[j] = 1;
+                }
+            }
+            auto cat2 = m_Network->addSlice(*temp,start,size,stride);
+            debug_print(cat2->getOutput(0), "cat2 dims : ");
+            ITensor *cat[] ={cat2->getOutput(0),cat1->getOutput(0)};
+            auto Cat = m_Network->addConcatenation(cat,2);
+            Cat->setAxis(dims[i] - 1);
+            temp = Cat->getOutput(0);
+        }
+    }
+    return temp;
+}
+ITensor* trt_transform_window_partition(INetworkDefinition *m_Network,ITensor *input,int window_size)
+{
+    auto shuffle1 = m_Network->addShuffle(*input);
+    Dims shuffle1_dims;
+    shuffle1_dims.nbDims = 5;
+    int h = input->getDimensions().d[0];
+    int w = input->getDimensions().d[1];
+    int c = input->getDimensions().d[2];
+
+    int dims[] = {h/window_size,window_size,w/window_size,window_size,c};
+    for(int i = 0 ; i < shuffle1_dims.nbDims; i++)
+        shuffle1_dims.d[i] = dims[i];
+    shuffle1->setReshapeDimensions(shuffle1_dims);
+    shuffle1->setSecondTranspose(Permutation{0,2,1,3,4});
+    debug_print(shuffle1->getOutput(0)," shuffle1 dims : ");
+    auto shuffle2 = m_Network->addShuffle(*shuffle1->getOutput(0));
+    shuffle2->setReshapeDimensions(Dims3{-1,window_size*window_size,c});
+
+    debug_print(shuffle2->getOutput(0), "window partition");
+    return shuffle2->getOutput(0);
+}
+ITensor* trt_swinLinear(INetworkDefinition *m_Network,std::map<std::string, Weights> weightMap,
+                        ITensor *input, string lname, bool bias = true)
+{
+    int c = input->getDimensions().d[input->getDimensions().nbDims-1];
+    string fc_wpath = lname + ".weight";
+    Weights fcW = weightMap[fc_wpath];
+    int len_fcw = fcW.count;
+    if(len_fcw == 0)
+    {
+        cout<<"file is not open,please check it's path: "<<fc_wpath<<endl;
+        assert(0);
+    }
+    Dims fcWdims;
+    fcWdims.nbDims = input->getDimensions().nbDims;
+    if(fcWdims.nbDims == 2)
+    {
+        fcWdims.d[0] = len_fcw/c;
+        fcWdims.d[1] = c;
+    }
+    else {
+        fcWdims.d[0] = 1;
+        fcWdims.d[1] = len_fcw/c;
+        fcWdims.d[2] = c;
+    }
+    auto fc_w_constant = m_Network->addConstant(fcWdims,fcW);
+    auto fc_w_mm = m_Network->addMatrixMultiply(*input,MatrixOperation::kNONE,
+                                                *fc_w_constant->getOutput(0),MatrixOperation::kTRANSPOSE);
+
+    string fc_bpath = lname +".bias";
+    Weights fcB = weightMap[fc_bpath];
+    int len_fcb = fcB.count;
+    if(!bias)
+    {
+        cout<<lname<<" bias is Null!"<<endl;
+        debug_print(fc_w_mm->getOutput(0),lname);
+        return fc_w_mm->getOutput(0);
+    }
+    Dims fcBdims;
+    fcBdims.nbDims = input->getDimensions().nbDims;
+    if(fcBdims.nbDims == 2)
+    {
+        fcBdims.d[0] = 1;
+        fcBdims.d[1] = len_fcb;
+    }
+    else {
+        fcBdims.d[0] = 1;
+        fcBdims.d[1] = 1;
+        fcBdims.d[2] = len_fcb;
+    }
+    auto fc_b_constant = m_Network->addConstant(fcBdims,fcB);
+    auto fc = m_Network->addElementWise(*fc_w_mm->getOutput(0),*fc_b_constant->getOutput(0),ElementWiseOperation::kSUM);
+    debug_print(fc->getOutput(0),lname);
+    return fc->getOutput(0);
+}
+ITensor* trt_trainsform_WindowAttention(INetworkDefinition *m_Network,std::map<std::string, Weights> weightMap,ITensor *input,
+                                        ITensor* mask,string lname,int dim, int num_heads,int window_size, int shift_size)
+{
+
+    int b = input->getDimensions().d[0];
+    int n = input->getDimensions().d[1];
+    int c = input->getDimensions().d[2];
+
+    auto qkv = trt_swinLinear(m_Network,weightMap,input,lname+".qkv");
+
+    Dims qkv_dim;
+    qkv_dim.nbDims = 5;
+    int d[5] = {b,n,3,num_heads,c/num_heads};
+    for(int i = 0; i < 5; i++)
+        qkv_dim.d[i] = d[i];
+    Permutation qkv_p;
+    int p[5] = {2, 0, 3, 1, 4};
+    for(int i = 0; i < 5; i++)
+        qkv_p.order[i] = p[i];
+    auto qkv_shuffle = shuffle_reshapeApermute(m_Network,qkv,qkv_dim,qkv_p,true);
+
+    Dims qkvDims = qkv_shuffle->getDimensions();
+    Dims qstart,kstart,vstart,sizes,stride;
+    qstart.nbDims = 5;
+    kstart.nbDims = 5;
+    vstart.nbDims = 5;
+    sizes.nbDims = 5;
+    stride.nbDims = 5;
+    for(int i = 0; i < 5; i++)
+    {
+        if(i == 0)
+        {
+            qstart.d[0] = 0;
+            kstart.d[0] = 1;
+            vstart.d[0] = 2;
+            sizes.d[0] = 1;
+            stride.d[0] =1;
+        }
+        else{
+            qstart.d[i] = 0;
+            kstart.d[i] = 0;
+            vstart.d[i] = 0;
+            sizes.d[i] = qkvDims.d[i];
+            stride.d[i] =1;
+        }
+    }
+    auto q = m_Network->addSlice(*qkv_shuffle,qstart,sizes,stride);
+    auto k = m_Network->addSlice(*qkv_shuffle,kstart,sizes,stride);
+    auto v = m_Network->addSlice(*qkv_shuffle,vstart,sizes,stride);
+
+    // q * s
+    int len = 1;
+    Weights scale_w{DataType::kFLOAT,nullptr,len};
+    float *scale = new float[len];
+    for(int i = 0 ; i < len; i++)
+        scale[i] = 1 / sqrt(dim/num_heads);
+    scale_w.values = scale;
+    Dims scale_dim;
+    scale_dim.nbDims = 5;
+
+    for(int i = 0 ; i < 5; i++)
+        scale_dim.d[i] = 1;
+    auto Scale = m_Network->addConstant(scale_dim,scale_w);
+    auto qs = m_Network->addElementWise(*q->getOutput(0),*Scale->getOutput(0),ElementWiseOperation::kPROD);
+    auto qs_ = m_Network->addShuffle(*qs->getOutput(0));
+    qs_->setReshapeDimensions(Dims4{qkvDims.d[1],qkvDims.d[2],qkvDims.d[3],qkvDims.d[4]});
+    auto k_ = m_Network->addShuffle(*k->getOutput(0));
+    k_->setReshapeDimensions(Dims4{qkvDims.d[1],qkvDims.d[2],qkvDims.d[3],qkvDims.d[4]});
+    auto attn = m_Network->addMatrixMultiply(*qs_->getOutput(0),MatrixOperation::kNONE,
+                                             *k_->getOutput(0),MatrixOperation::kTRANSPOSE);
+    auto relatbias = m_Network->addConstant(Dims2{(2*window_size -1)*(2*window_size -1),num_heads},weightMap[lname + ".relative_position_bias_table"]);
+    Dims r_i_dims;
+    r_i_dims.nbDims = 1;
+    r_i_dims.d[0] = window_size*window_size * window_size*window_size;
+    Weights index{DataType::kINT32,nullptr,r_i_dims.d[0]};
+    int* idx = new int[r_i_dims.d[0]];
+    for (int i = 0; i < r_i_dims.d[0]; i++) {
+        idx[i] =(int)((float*)weightMap[lname+".relative_position_index"].values)[i];
+    }
+    //idx = (int*)weightMap[lname+".relative_position_index"].values;
+    //cout<<"idx = "<<((float*)weightMap[lname+".relative_position_index"].values)[0]<<endl;
+    index.values = idx;
+    auto relatidx = m_Network->addConstant(r_i_dims,index);
+    auto relat = m_Network->addGather(*relatbias->getOutput(0),*relatidx->getOutput(0),0);
+    auto relat_view = shuffle_reshapeApermute(m_Network,relat->getOutput(0),
+                                              Dims4{1,window_size*window_size,window_size*window_size,-1},
+                                              Permutation{0,3,1,2},true);
+    auto attn_rv = m_Network->addElementWise(*attn->getOutput(0),*relat_view,ElementWiseOperation::kSUM);
+    ITensor *Attn_rv = attn_rv->getOutput(0);
+    if (mask != nullptr)
+    {
+        Dims maskdims;
+        maskdims.nbDims = mask->getDimensions().nbDims +1;
+        maskdims.d[0] = mask->getDimensions().d[0];
+        maskdims.d[1] = 1;
+        for(int i = 2; i< maskdims.nbDims; i++)
+        {
+            maskdims.d[i] = mask->getDimensions().d[i-1];
+        }
+        auto maskshuffle = m_Network->addShuffle(*mask);
+        maskshuffle->setReshapeDimensions(maskdims);
+        auto attn_rnM = m_Network->addElementWise(*attn_rv->getOutput(0),*maskshuffle->getOutput(0),ElementWiseOperation::kSUM);
+        Attn_rv = attn_rnM->getOutput(0);
+    }
+    auto attn_rv_s = m_Network->addSoftMax(*Attn_rv);
+    attn_rv_s->setAxes(8);
+    auto v_ = m_Network->addShuffle(*v->getOutput(0));
+    v_->setReshapeDimensions(Dims4{qkvDims.d[1],qkvDims.d[2],qkvDims.d[3],qkvDims.d[4]});
+    auto attn_v = m_Network->addMatrixMultiply(*attn_rv_s->getOutput(0),MatrixOperation::kNONE,
+                                               *v_->getOutput(0),MatrixOperation::kNONE);
+    auto x_reshape = shuffle_reshapeApermute(m_Network,attn_v->getOutput(0),Dims3{b,n,c},Permutation{0,2,1,3},false);
+    auto x_linear = trt_swinLinear(m_Network,weightMap,x_reshape,lname+".proj");
+    return x_linear;
+}
+ITensor* trt_window_reverse(INetworkDefinition *m_Network, ITensor *input, int window_size, int H, int W)
+{
+    Dims viewDims;
+    viewDims.nbDims = 5;
+    int d[5] = {H/window_size,W/window_size,window_size,window_size,-1};
+    for(int i = 0; i < 5; i++)
+        viewDims.d[i] = d[i];
+    auto x_view = shuffle_reshape(m_Network,input,viewDims);
+    auto output = shuffle_reshapeApermute(m_Network,x_view,Dims3{H,W,-1},Permutation{0,2,1,3,4},false);
+    return output;
+}
+ITensor* gelu(INetworkDefinition *m_Network,ITensor *input)
+{
+    auto creator = getPluginRegistry()->getPluginCreator("geluLayer_TRT", "1");
+    const PluginFieldCollection* pluginData = creator->getFieldNames();
+    IPluginV2 *pluginObj = creator->createPlugin("gelu", pluginData);
+    ITensor* inputTensors[] = {input};
+    auto g = m_Network->addPluginV2(inputTensors, 1, *pluginObj);
+    return g->getOutput(0);
+}
+//ITensor* adaptiveAvgPool2d(INetworkDefinition *m_Network,ITensor *input)
+//{
+//    auto creator = getPluginRegistry()->getPluginCreator("adaptiveAvgPooling_TRT", "1");
+//    const PluginFieldCollection* pluginData = creator->getFieldNames();
+//    IPluginV2 *pluginObj = creator->createPlugin("apAvgPool", pluginData);
+//    ITensor* inputTensors[] = {input};
+//    auto g = m_Network->addPluginV2(inputTensors, 1, *pluginObj);
+//    return g->getOutput(0);
+//}
+ITensor* trt_transform_mlp(INetworkDefinition *m_Network,std::map<std::string, Weights> weightMap,ITensor *input,
+                           string lname,int dim,int mlp_ratio = 4)
+{
+//    auto fc1 = m_Network->addFullyConnected(*input,dim * mlp_ratio,
+//                                            weightMap[lname+".fc1.weight"],weightMap[lname+".fc1.bias"]);
+    auto fc1 = trt_swinLinear(m_Network,weightMap,input,lname+".fc1");
+    auto act = gelu(m_Network,fc1);
+//    auto fc2 = m_Network->addFullyConnected(*act,dim ,
+//                                            weightMap[lname+".fc2.weight"],weightMap[lname+".fc2.bias"]);
+    auto fc2 = trt_swinLinear(m_Network,weightMap,act,lname+".fc2");
+    return fc2;
+}
+ITensor* blk(INetworkDefinition *m_Network,std::map<std::string, Weights> weightMap,ITensor *input, ITensor* mask, string lname,
+             int hw,int dim, int num_heads,int window_size,int shift_size,int mlp_ratio = 4)
+{
+    int c = input->getDimensions().d[input->getDimensions().nbDims - 1];
+    auto x = input;
+    auto norm1 = m_layerNorm(m_Network,weightMap,x,lname+".norm1");
+    //auto norm1 = x;
+    auto view1 = shuffle_reshape(m_Network,norm1,Dims3{hw,hw,c});
+    auto pad = trt_transform_pad(m_Network,view1,window_size);
+    int hp = pad->getDimensions().d[0];
+    int wp = pad->getDimensions().d[1];
+    ITensor* shifted_x;
+    ITensor* atten_mask = nullptr;
+    if(shift_size > 0)
+    {
+        shifted_x = trt_swinRoll(m_Network,pad,{-3,-3},{1,2});
+        atten_mask = mask;
+    }
+    else
+    {
+        shifted_x = pad;
+    }
+    auto x_windows = trt_transform_window_partition(m_Network,shifted_x,window_size);
+    auto x_atten_windows = trt_trainsform_WindowAttention(m_Network,weightMap,x_windows,atten_mask,lname+".attn",dim,num_heads,
+                                                          window_size,shift_size);
+    auto x_atten_windows_view = shuffle_reshape(m_Network,x_atten_windows,Dims4{-1,window_size,window_size,c});
+
+    shifted_x = trt_window_reverse(m_Network,x_atten_windows_view,window_size,hp,wp);
+    if(shift_size > 0)
+    {
+        x = trt_swinRoll(m_Network,shifted_x,{3,3},{1,2});
+    }
+    else {
+        x = shifted_x;
+    }
+    if(hw < hp){
+        auto sss = m_Network->addSlice(*x,Dims3{0,0,0},Dims3{hw,hw,c},Dims3{1,1,1});
+        x = sss->getOutput(0);
+    }
+    x = shuffle_reshape(m_Network,x,Dims2{hw*hw,c});
+    x = m_Network->addElementWise(*x,*input,ElementWiseOperation::kSUM)->getOutput(0);
+    auto norm2 = m_layerNorm(m_Network,weightMap,x,lname+".norm2");
+    //auto norm2 = x;
+    auto mlp = trt_transform_mlp(m_Network,weightMap,norm2,lname+".mlp",dim);
+    auto out= m_Network->addElementWise(*x,*mlp,ElementWiseOperation::kSUM)->getOutput(0);
+    debug_print(out, "blk");
+    return out;
+}
+ITensor* downsample(INetworkDefinition* m_Network,std::map<std::string, Weights> weightMap,ITensor *input,
+                    string lname, int hw)
+{
+    int c = input->getDimensions().d[input->getDimensions().nbDims - 1];
+    auto x = shuffle_reshape(m_Network,input,Dims3{hw,hw,c});
+    auto x0 = m_Network->addSlice(*x,Dims3{0,0,0},Dims3{hw/2,hw/2,c},Dims3{2,2,1});
+    auto x1 = m_Network->addSlice(*x,Dims3{1,0,0},Dims3{hw/2,hw/2,c},Dims3{2,2,1});
+    auto x2 = m_Network->addSlice(*x,Dims3{0,1,0},Dims3{hw/2,hw/2,c},Dims3{2,2,1});
+    auto x3 = m_Network->addSlice(*x,Dims3{1,1,0},Dims3{hw/2,hw/2,c},Dims3{2,2,1});
+    ITensor* inputTensors[] = { x0->getOutput(0), x1->getOutput(0), x2->getOutput(0), x3->getOutput(0) };
+    auto cat = m_Network->addConcatenation(inputTensors, 4);
+    cat->setAxis(2);
+    auto cat_view = shuffle_reshape(m_Network,cat->getOutput(0),Dims2{-1,4*c});
+    auto norm = m_layerNorm(m_Network,weightMap,cat_view,lname+".norm");
+    //auto norm = cat_view;
+    auto reduction = trt_swinLinear(m_Network,weightMap,norm,lname+".reduction",false);
+    return reduction;
+}
+ITensor* addBatchNorm2d(
+INetworkDefinition *network,
+std::map<std::string, Weights> weightMap,
+ITensor* input,
+const std::string& lname,
+float eps = 1e-5
+) {
+    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;
+
+    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->getOutput(0);
+}
+ITensor* transform_lateral_conv(INetworkDefinition* m_Network,std::map<std::string, Weights> weightMap,ITensor* input,
+                                string lname, int k = 1, int s = 1,int out_features = 512)
+{
+    Weights empty{DataType::kFLOAT,nullptr,0};
+    auto conv = m_Network->addConvolutionNd(*input,out_features,Dims2{k,k},weightMap[lname+".conv.weight"],empty);
+    conv->setStrideNd(Dims2{s,s});
+    conv->setNbGroups(1);
+    conv->setPaddingNd(Dims2{k/2,k/2});
+    ITensor* bn = addBatchNorm2d(m_Network,weightMap,conv->getOutput(0),lname+".bn");
+    auto act = m_Network->addActivation(*bn,ActivationType::kRELU);
+    return act->getOutput(0);
+}
+ITensor* resize(INetworkDefinition* m_Network, ITensor* input, int grid)
+{
+    float scale_h = 2.0f;
+    float scale_w = 2.0f;
+
+    scale_h = 1.0*grid / input->getDimensions().d[1];
+    scale_w = 1.0*grid / input->getDimensions().d[2];
+
+    auto creator = getPluginRegistry()->getPluginCreator("UpsamplePlugin", "1");
+    PluginField pField[1];
+    float *s = new float[2];
+    s[0] = scale_h;
+    s[1] = scale_w;
+    pField[0].data = s;
+    pField[0].length = 2;
+    pField[0].type = PluginFieldType::kFLOAT32;
+    pField[0].name = "scaleFactor";
+
+    PluginFieldCollection pluginData;
+    pluginData.nbFields = 1;
+    pluginData.fields = pField;
+    IPluginV2 *pluginObj = creator->createPlugin("upSample", &pluginData);
+    ITensor* inputTensors[] = {input};
+    auto upS = m_Network->addPluginV2(inputTensors, 1, *pluginObj);
+    return upS->getOutput(0);
+}
+ITensor* transform_psp(INetworkDefinition* m_Network,std::map<std::string, Weights> weightMap,ITensor* input,
+                       string lname, int output_Avg_Size, int out_features = 512)
+{
+    int inH = input->getDimensions().d[1];
+    int inW = input->getDimensions().d[2];
+    int kH = inH / output_Avg_Size;
+    int kW = inW / output_Avg_Size;
+    auto avgPool = m_Network->addPoolingNd(*input,PoolingType::kAVERAGE,Dims2{kH,kW});
+    avgPool->setStrideNd(Dims2{kH,kW});
+    auto cba = transform_lateral_conv(m_Network,weightMap,avgPool->getOutput(0),lname,1,1,out_features);
+    auto out = resize(m_Network,cba,inH);
+    return out;
+}
+ITensor* up_Add(INetworkDefinition* m_Network,ITensor* input1,ITensor* input2)
+{
+    auto in1 = resize(m_Network,input1,input2->getDimensions().d[1]);
+    auto out = m_Network->addElementWise(*in1,*input2,ElementWiseOperation::kSUM);
+    return out->getOutput(0);
+}
+
+
+#endif // COMMON_HPP
--- a/swin-transformer/semantic-segmentation/fillmask.cu
+++ b/swin-transformer/semantic-segmentation/fillmask.cu
@ -0,0 +1,182 @@
+#include "fillmask.h"
+#include <math.h>
+namespace nvinfer1
+{
+    fillmask::fillmask()
+    {
+    }
+
+    fillmask::~fillmask()
+    {
+    }
+    // create the plugin at runtime from a byte stream
+    fillmask::fillmask(const void* data, size_t length)
+    {
+        const char *d = reinterpret_cast<const char *>(data), *a = d;
+        Tn::read(d, mInputSize);
+        assert(d == a + length);
+    }
+
+    void fillmask::serialize(void* buffer) const
+    {
+        char* d = static_cast<char*>(buffer), *a = d;
+        Tn::write(d, mInputSize);
+        assert(d == a + getSerializationSize());
+    }
+
+    size_t fillmask::getSerializationSize() const
+    {
+        return sizeof(mInputSize);
+    }
+
+    int fillmask::initialize()
+    {
+        return 0;
+    }
+
+    Dims fillmask::getOutputDimensions(int index, const Dims* inputs, int nbInputDims)
+    {
+        assert(nbInputDims == 1);
+        Dims outputDims;
+        outputDims.nbDims = inputs[0].nbDims;
+        for (int i = 0; i < inputs[0].nbDims; i++) {
+            outputDims.d[i] = inputs[0].d[i];
+        }
+        return outputDims;
+    }
+
+    // Set plugin namespace
+    void fillmask::setPluginNamespace(const char* pluginNamespace)
+    {
+        mPluginNamespace = pluginNamespace;
+    }
+
+    const char* fillmask::getPluginNamespace() const
+    {
+        return mPluginNamespace;
+    }
+
+    // Return the DataType of the plugin output at the requested index
+    DataType fillmask::getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const
+    {
+        return DataType::kFLOAT;
+    }
+
+    // Return true if output tensor is broadcast across a batch.
+    bool fillmask::isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const
+    {
+        return false;
+    }
+
+    // Return true if plugin can use input that is broadcast across batch without replication.
+    bool fillmask::canBroadcastInputAcrossBatch(int inputIndex) const
+    {
+        return false;
+    }
+
+    void fillmask::configurePlugin(const PluginTensorDesc* in, int nbInput, const PluginTensorDesc* out, int nbOutput)
+    {
+
+        mInputSize = 1;
+        for (int i = 0; i < in[0].dims.nbDims; i++) {
+            mInputSize *= in[0].dims.d[i];
+        }
+    }
+
+    // Attach the plugin object to an execution context and grant the plugin the access to some context resource.
+    void fillmask::attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator)
+    {
+    }
+
+    // Detach the plugin object from its execution context.
+    void fillmask::detachFromContext() {}
+
+    const char* fillmask::getPluginType() const
+    {
+        return "fillmaskLayer_TRT";
+    }
+
+    const char* fillmask::getPluginVersion() const
+    {
+        return "1";
+    }
+
+    void fillmask::destroy()
+    {
+        delete this;
+    }
+
+    // Clone the plugin
+    IPluginV2IOExt* fillmask::clone() const
+    {
+        fillmask *p = new fillmask();
+        p->setPluginNamespace(mPluginNamespace);
+        p->setInputSize(mInputSize);
+        return p;
+    }
+
+    __global__ void fillmaskKer(const float *in, float *out, int size) {
+        int idx = threadIdx.x + blockIdx.x * blockDim.x;
+        if (idx >= size)
+            return;
+        if (in[idx] != 0.0)
+            out[idx] = -100.0;
+        else
+            out[idx] = 0.0;
+    }
+    void fillmask::forwardGpu(const float *const * inputs, float* output, cudaStream_t stream, int batchSize) {
+
+        int numElem = batchSize * mInputSize;
+        fillmaskKer<<<(numElem + mThreadCount - 1) / mThreadCount, mThreadCount>>>
+            (inputs[0], output, numElem);
+    }
+
+    int fillmask::enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream)
+    {
+        forwardGpu((const float *const *)inputs, (float*)outputs[0], stream, batchSize);
+        return 0;
+    }
+
+    PluginFieldCollection fillmaskCreator::mFC{};
+    std::vector<PluginField> fillmaskCreator::mPluginAttributes;
+
+    fillmaskCreator::fillmaskCreator()
+    {
+        mPluginAttributes.clear();
+        mFC.nbFields = mPluginAttributes.size();
+        mFC.fields = mPluginAttributes.data();
+    }
+
+    const char* fillmaskCreator::getPluginName() const
+    {
+            return "fillmaskLayer_TRT";
+    }
+
+    const char* fillmaskCreator::getPluginVersion() const
+    {
+            return "1";
+    }
+
+    const PluginFieldCollection* fillmaskCreator::getFieldNames()
+    {
+            return &mFC;
+    }
+
+    IPluginV2IOExt* fillmaskCreator::createPlugin(const char* name, const PluginFieldCollection* fc)
+    {
+        fillmask* obj = new fillmask();
+        obj->setPluginNamespace(mNamespace.c_str());
+        return obj;
+    }
+
+    IPluginV2IOExt* fillmaskCreator::deserializePlugin(const char* name, const void* serialData, size_t serialLength)
+    {
+        // This object will be deleted when the network is destroyed, which will
+        fillmask* obj = new fillmask(serialData, serialLength);
+        obj->setPluginNamespace(mNamespace.c_str());
+        return obj;
+    }
+
+
+}
+
--- a/swin-transformer/semantic-segmentation/fillmask.h
+++ b/swin-transformer/semantic-segmentation/fillmask.h
@ -0,0 +1,90 @@
+#ifndef FILLMASK_H
+#define FILLMASK_H
+
+
+#include <vector>
+#include <string>
+#include "NvInfer.h"
+#include "myhpp.h"
+#include <assert.h>
+#include "utilsn.h"
+
+namespace nvinfer1
+{
+    class fillmask:public IPluginV2IOExt
+    {
+    public:
+        explicit fillmask();
+        fillmask(const void* data, size_t length);
+        ~fillmask();
+        int getNbOutputs() const override
+        {
+            return 1;
+        }
+
+        Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override;
+        int initialize() override;
+        virtual void terminate() override {};
+        virtual size_t getWorkspaceSize(int maxBatchSize) const override { return 0;}
+        virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override;
+        virtual size_t getSerializationSize() const override;
+        virtual void serialize(void* buffer) const override;
+
+        bool supportsFormatCombination(int pos, const PluginTensorDesc* inOut, int nbInputs, int nbOutputs) const override {
+            return inOut[pos].format == TensorFormat::kLINEAR && inOut[pos].type == DataType::kFLOAT;
+        }
+
+        const char* getPluginType() const override;
+        const char* getPluginVersion() const override;
+        void destroy() override;
+        IPluginV2IOExt* clone() const override;
+        void setPluginNamespace(const char* pluginNamespace) override;
+        const char* getPluginNamespace() const override;
+        DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const override;
+        bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const override;
+        bool canBroadcastInputAcrossBatch(int inputIndex) const override;
+        void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) override;
+        void configurePlugin(const PluginTensorDesc* in, int nbInput, const PluginTensorDesc* out, int nbOutput) override;
+        void detachFromContext() override;
+
+        void setInputSize(int s) {
+            mInputSize = s;
+        }
+
+    private:
+        void forwardGpu(const float *const * inputs,float * output, cudaStream_t stream,int batchSize = 1);
+        int mThreadCount = 256;
+        int mInputSize;
+        const char* mPluginNamespace;
+    };
+
+    class fillmaskCreator : public IPluginCreator
+    {
+        public:
+            fillmaskCreator();
+            ~fillmaskCreator() override = default;
+            const char* getPluginName() const override;
+            const char* getPluginVersion() const override;
+            const PluginFieldCollection* getFieldNames() override;
+            IPluginV2IOExt* createPlugin(const char* name, const PluginFieldCollection* fc) override;
+            IPluginV2IOExt* deserializePlugin(const char* name, const void* serialData, size_t serialLength) override;
+
+            void setPluginNamespace(const char* libNamespace) override
+            {
+                mNamespace = libNamespace;
+            }
+
+            const char* getPluginNamespace() const override
+            {
+                return mNamespace.c_str();
+            }
+
+        private:
+            std::string mNamespace;
+            static PluginFieldCollection mFC;
+            static std::vector<PluginField> mPluginAttributes;
+    };
+    REGISTER_TENSORRT_PLUGIN(fillmaskCreator);
+};
+
+#endif // FILLMASK_H
--- a/swin-transformer/semantic-segmentation/gelu.cu
+++ b/swin-transformer/semantic-segmentation/gelu.cu
@ -0,0 +1,180 @@
+#include "gelu.h"
+#include <math.h>
+namespace nvinfer1
+{
+    gelu::gelu()
+    {
+    }
+
+    gelu::~gelu()
+    {
+    }
+    // create the plugin at runtime from a byte stream
+    gelu::gelu(const void* data, size_t length)
+    {
+        const char *d = reinterpret_cast<const char *>(data), *a = d;
+        Tn::read(d, mInputSize);
+        assert(d == a + length);
+    }
+
+    void gelu::serialize(void* buffer) const
+    {
+        char* d = static_cast<char*>(buffer), *a = d;
+        Tn::write(d, mInputSize);
+        assert(d == a + getSerializationSize());
+    }
+
+    size_t gelu::getSerializationSize() const
+    {
+        return sizeof(mInputSize);
+    }
+
+    int gelu::initialize()
+    {
+        return 0;
+    }
+
+    Dims gelu::getOutputDimensions(int index, const Dims* inputs, int nbInputDims)
+    {
+        assert(nbInputDims == 1);
+        Dims outputDims;
+        outputDims.nbDims = inputs[0].nbDims;
+        for (int i = 0; i < inputs[0].nbDims; i++) {
+            outputDims.d[i] = inputs[0].d[i];
+        }
+        return outputDims;
+    }
+
+    // Set plugin namespace
+    void gelu::setPluginNamespace(const char* pluginNamespace)
+    {
+        mPluginNamespace = pluginNamespace;
+    }
+
+    const char* gelu::getPluginNamespace() const
+    {
+        return mPluginNamespace;
+    }
+
+    // Return the DataType of the plugin output at the requested index
+    DataType gelu::getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const
+    {
+        return DataType::kFLOAT;
+    }
+
+    // Return true if output tensor is broadcast across a batch.
+    bool gelu::isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const
+    {
+        return false;
+    }
+
+    // Return true if plugin can use input that is broadcast across batch without replication.
+    bool gelu::canBroadcastInputAcrossBatch(int inputIndex) const
+    {
+        return false;
+    }
+
+    void gelu::configurePlugin(const PluginTensorDesc* in, int nbInput, const PluginTensorDesc* out, int nbOutput)
+    {
+
+        mInputSize = 1;
+        for (int i = 0; i < in[0].dims.nbDims; i++) {
+            mInputSize *= in[0].dims.d[i];
+        }
+    }
+
+    // Attach the plugin object to an execution context and grant the plugin the access to some context resource.
+    void gelu::attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator)
+    {
+    }
+
+    // Detach the plugin object from its execution context.
+    void gelu::detachFromContext() {}
+
+    const char* gelu::getPluginType() const
+    {
+        return "geluLayer_TRT";
+    }
+
+    const char* gelu::getPluginVersion() const
+    {
+        return "1";
+    }
+
+    void gelu::destroy()
+    {
+        delete this;
+    }
+
+    // Clone the plugin
+    IPluginV2IOExt* gelu::clone() const
+    {
+        gelu *p = new gelu();
+        p->setPluginNamespace(mPluginNamespace);
+        p->setInputSize(mInputSize);
+        return p;
+    }
+
+    __global__ void geluKer(const float *in, float *out, int size) {
+        int idx = threadIdx.x + blockIdx.x * blockDim.x;
+        if (idx >= size)
+            return;
+        //x * 0.5 * (1.0 + torch.erf(x / math.sqrt(2.0)))
+        out[idx] = in[idx] * 0.5 *(1.0 + erf(in[idx]/1.4142135381698608));
+    }
+    void gelu::forwardGpu(const float *const * inputs, float* output, cudaStream_t stream, int batchSize) {
+
+        int numElem = batchSize * mInputSize;
+        geluKer<<<(numElem + mThreadCount - 1) / mThreadCount, mThreadCount>>>
+            (inputs[0], output, numElem);
+    }
+
+    int gelu::enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream)
+    {
+        forwardGpu((const float *const *)inputs, (float*)outputs[0], stream, batchSize);
+        return 0;
+    }
+
+    PluginFieldCollection geluCreator::mFC{};
+    std::vector<PluginField> geluCreator::mPluginAttributes;
+
+    geluCreator::geluCreator()
+    {
+        mPluginAttributes.clear();
+        mFC.nbFields = mPluginAttributes.size();
+        mFC.fields = mPluginAttributes.data();
+    }
+
+    const char* geluCreator::getPluginName() const
+    {
+            return "geluLayer_TRT";
+    }
+
+    const char* geluCreator::getPluginVersion() const
+    {
+            return "1";
+    }
+
+    const PluginFieldCollection* geluCreator::getFieldNames()
+    {
+            return &mFC;
+    }
+
+    IPluginV2IOExt* geluCreator::createPlugin(const char* name, const PluginFieldCollection* fc)
+    {
+        gelu* obj = new gelu();
+        obj->setPluginNamespace(mNamespace.c_str());
+        return obj;
+    }
+
+    IPluginV2IOExt* geluCreator::deserializePlugin(const char* name, const void* serialData, size_t serialLength)
+    {
+        // This object will be deleted when the network is destroyed, which will
+        gelu* obj = new gelu(serialData, serialLength);
+        obj->setPluginNamespace(mNamespace.c_str());
+        return obj;
+    }
+
+
+}
+
--- a/swin-transformer/semantic-segmentation/gelu.h
+++ b/swin-transformer/semantic-segmentation/gelu.h
@ -0,0 +1,88 @@
+#ifndef GELU_H
+#define GELU_H
+
+#include <vector>
+#include <string>
+#include "NvInfer.h"
+#include "myhpp.h"
+#include <assert.h>
+#include "utilsn.h"
+#define M_PI       3.14159265358979323846   // pi
+namespace nvinfer1
+{
+    class gelu:public IPluginV2IOExt
+    {
+    public:
+        explicit gelu();
+        gelu(const void* data, size_t length);
+        ~gelu();
+        int getNbOutputs() const override
+        {
+            return 1;
+        }
+
+        Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override;
+        int initialize() override;
+        virtual void terminate() override {};
+        virtual size_t getWorkspaceSize(int maxBatchSize) const override { return 0;}
+        virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override;
+        virtual size_t getSerializationSize() const override;
+        virtual void serialize(void* buffer) const override;
+
+        bool supportsFormatCombination(int pos, const PluginTensorDesc* inOut, int nbInputs, int nbOutputs) const override {
+            return inOut[pos].format == TensorFormat::kLINEAR && inOut[pos].type == DataType::kFLOAT;
+        }
+
+        const char* getPluginType() const override;
+        const char* getPluginVersion() const override;
+        void destroy() override;
+        IPluginV2IOExt* clone() const override;
+        void setPluginNamespace(const char* pluginNamespace) override;
+        const char* getPluginNamespace() const override;
+        DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const override;
+        bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const override;
+        bool canBroadcastInputAcrossBatch(int inputIndex) const override;
+        void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) override;
+        void configurePlugin(const PluginTensorDesc* in, int nbInput, const PluginTensorDesc* out, int nbOutput) override;
+        void detachFromContext() override;
+
+        void setInputSize(int s) {
+            mInputSize = s;
+        }
+
+    private:
+        void forwardGpu(const float *const * inputs,float * output, cudaStream_t stream,int batchSize = 1);
+        int mThreadCount = 256;
+        int mInputSize;
+        const char* mPluginNamespace;
+    };
+
+    class geluCreator : public IPluginCreator
+    {
+        public:
+            geluCreator();
+            ~geluCreator() override = default;
+            const char* getPluginName() const override;
+            const char* getPluginVersion() const override;
+            const PluginFieldCollection* getFieldNames() override;
+            IPluginV2IOExt* createPlugin(const char* name, const PluginFieldCollection* fc) override;
+            IPluginV2IOExt* deserializePlugin(const char* name, const void* serialData, size_t serialLength) override;
+
+            void setPluginNamespace(const char* libNamespace) override
+            {
+                mNamespace = libNamespace;
+            }
+
+            const char* getPluginNamespace() const override
+            {
+                return mNamespace.c_str();
+            }
+
+        private:
+            std::string mNamespace;
+            static PluginFieldCollection mFC;
+            static std::vector<PluginField> mPluginAttributes;
+    };
+    REGISTER_TENSORRT_PLUGIN(geluCreator);
+};
+#endif // GELU_H
--- a/swin-transformer/semantic-segmentation/gen_wts.py
+++ b/swin-transformer/semantic-segmentation/gen_wts.py
@ -0,0 +1,19 @@
+import torch
+import struct
+import sys
+
+# Initialize
+pt_file = sys.argv[1]
+# Load model
+model = torch.load(pt_file, map_location=torch.device('cpu'))['model'].float()  # load to FP32
+model.to(device).eval()
+
+with open(pt_file.split('.')[0] + '.wts', 'w') as f:
+    f.write('{}\n'.format(len(model.state_dict().keys())))
+    for k, v in model.state_dict().items():
+        vr = v.reshape(-1).cpu().numpy()
+        f.write('{} {} '.format(k, len(vr)))
+        for vv in vr:
+            f.write(' ')
+            f.write(struct.pack('>f',float(vv)).hex())
+        f.write('\n')
--- a/swin-transformer/semantic-segmentation/include/dirent.h
+++ b/swin-transformer/semantic-segmentation/include/dirent.h
--- a/swin-transformer/semantic-segmentation/layerNorm.cu
+++ b/swin-transformer/semantic-segmentation/layerNorm.cu
@ -0,0 +1,195 @@
+#include <assert.h>
+#include "layerNorm.h"
+#include "utilsn.h"
+#include <assert.h>
+#include <vector>
+
+
+
+
+
+namespace nvinfer1
+{
+
+layernorm::layernorm()
+{
+}
+layernorm::~layernorm()
+{
+
+}
+layernorm::layernorm(const void* data, size_t length)
+{
+    const char *d = reinterpret_cast<const char *>(data), *a = d;
+    Tn::read(d, mInputSize);
+    Tn::read(d,Length);
+
+    assert(d == a + length);
+}
+int layernorm::initialize()
+{
+    return 0;
+}
+void layernorm::serialize(void* buffer) const
+{
+    char* d = static_cast<char*>(buffer), *a = d;
+    Tn::write(d, mInputSize);
+    Tn::write(d,Length);
+    assert(d == a + getSerializationSize());
+}
+size_t layernorm::getSerializationSize() const
+{
+    return sizeof(mInputSize) + sizeof(Length);
+}
+Dims layernorm::getOutputDimensions(int index, const Dims* inputs, int nbInputDims)
+{
+//    outputDims.nbDims  = inputs[0].nbDims;
+//    outputDims.d[0] = inputs[0].d[0];
+//    for (int var = 1; var < inputs[0].nbDims; ++var) {
+//        outputDims.d[var] = 1;
+//    }
+    return Dims2{inputs[0].d[0],1};
+}
+void layernorm::setPluginNamespace(const char* pluginNamespace)
+{
+    mPluginNamespace = pluginNamespace;
+}
+const char* layernorm::getPluginNamespace() const
+{
+    return mPluginNamespace;
+}
+const char* layernorm::getPluginType() const
+{
+    return "layerNorm_trt";
+}
+const char* layernorm::getPluginVersion() const
+{
+    return "1";
+}
+DataType layernorm::getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const
+{
+    return inputTypes[0] ;//== nvinfer1::DataType::kFLOAT ? nvinfer1::DataType::kFLOAT : nvinfer1::DataType::kHALF;
+}
+void layernorm::destroy()
+{
+    delete this;
+}
+IPluginV2IOExt* layernorm::clone() const
+{
+    layernorm *ln = new layernorm();
+    ln->setPluginNamespace(mPluginNamespace);
+    ln->setInputSize(mInputSize,Length);
+    return ln;
+}
+bool layernorm::isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const
+{
+    return false;
+}
+bool layernorm::canBroadcastInputAcrossBatch(int inputIndex) const
+{
+    return false;
+}
+void layernorm::attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator)
+{}
+void layernorm::configurePlugin(const PluginTensorDesc* in, int nbInput, const PluginTensorDesc* out, int nbOutput)
+{
+
+    int size = 1;
+    for(int i = 0 ; i < in[0].dims.nbDims ; i++)
+    {
+        size *= in[0].dims.d[i];
+    }
+    mInputSize = size;
+    Length = in[0].dims.d[in[0].dims.nbDims - 1];
+}
+void layernorm::detachFromContext()
+{}
+
+__device__ welford welford_update(welford a, const float *currValue, int length)
+{
+    #pragma unroll
+    for(int i = 0; i < length; i++){
+        a.count += 1;
+        float delta = currValue[i] - a.mean;
+        a.mean += delta / a.count;
+        float delta2 = currValue[i] - a.mean;
+        a.M2 += delta * delta2;
+    }
+    return a;
+}
+__device__ void mean_std(float* mean, float *std, const float *currValue,int l,int count = 0, float m = 0.0, float s = 0.0)
+{
+    #pragma unroll
+    for(int i = 0; i < l; i++){
+        count += 1;
+        float delta = currValue[i] - m;
+        m += delta / count;
+        float delta2 = currValue[i] - m;
+        s += delta * delta2;
+    }
+    *mean = m;
+    *std = sqrt((s / count) + 1e-5);
+}
+__global__ void lnCudaKer(const float *in, float *mean, float *std, int size,int l)
+{
+    int idx = threadIdx.x + blockIdx.x * blockDim.x;
+    if (idx >= size)
+        return;
+    mean_std(&mean[idx],&std[idx],in+idx*l,l);
+    //printf("idx = %d,mean = %f, std = %f\n",idx,mean[idx],std[idx]);
+}
+void layernorm::forwardGpu(const float *const *inputs, float *mean, float *std, cudaStream_t stream, int batchSize)
+{
+    int numElem = batchSize * mInputSize/Length;
+
+    lnCudaKer<<<(numElem + mThreadCount - 1) / mThreadCount, mThreadCount>>>
+        (inputs[0], mean,std, numElem,Length);
+}
+int layernorm::enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream)
+{
+    forwardGpu((const float *const *)inputs, (float*)outputs[0], (float*)outputs[1], stream, batchSize);
+    return 0;
+}
+
+PluginFieldCollection layernormCreator::mFC{};
+std::vector<PluginField> layernormCreator::mPluginAttributes;
+layernormCreator::layernormCreator()
+{
+    mPluginAttributes.clear();
+
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+const char* layernormCreator::getPluginName() const
+{
+    return "layerNorm_trt";
+}
+const char* layernormCreator::getPluginVersion() const
+{
+    return "1";
+}
+const PluginFieldCollection* layernormCreator::getFieldNames()
+{
+    return &mFC;
+}
+IPluginV2IOExt* layernormCreator::createPlugin(const char* name, const PluginFieldCollection* fc)
+{
+    layernorm* obj = new layernorm();
+    obj->setPluginNamespace(mNamespace.c_str());
+
+    return obj;
+}
+IPluginV2IOExt* layernormCreator::deserializePlugin(const char* name, const void* serialData, size_t serialLength)
+{
+    layernorm* obj = new layernorm(serialData, serialLength);
+    obj->setPluginNamespace(mNamespace.c_str());
+    return obj;
+}
+
+
+
+
+
+
+}
--- a/swin-transformer/semantic-segmentation/layerNorm.h
+++ b/swin-transformer/semantic-segmentation/layerNorm.h
@ -0,0 +1,130 @@
+#ifndef LAYERNORM_H
+#define LAYERNORM_H
+
+#include <vector>
+#include <string>
+#include <iostream>
+#include <NvInfer.h>
+#include <memory>
+#include <string.h>
+#include <cstdint>
+#include <stdlib.h>
+
+using namespace std;
+
+struct welford
+{
+    int count = 0;
+    double mean = 0.f;
+    double M2 = 0.f;
+};
+
+namespace nvinfer1{
+class layernorm : public IPluginV2IOExt
+{
+public:
+    layernorm();
+    layernorm(const void* data, size_t length);
+    ~layernorm();
+    int getNbOutputs() const override
+    {
+        return 2;
+    }
+
+    Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override;
+
+    int initialize() override;
+
+    virtual void terminate() override {};
+
+    virtual size_t getWorkspaceSize(int maxBatchSize) const override { return 0; }
+
+    virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override;
+
+    virtual size_t getSerializationSize() const override;
+
+    virtual void serialize(void* buffer) const override;
+
+    bool supportsFormatCombination(int pos, const PluginTensorDesc* inOut, int nbInputs, int nbOutputs) const override {
+        return inOut[pos].format == TensorFormat::kLINEAR && inOut[pos].type == DataType::kFLOAT;
+    }
+
+    void setPluginNamespace(const char* pluginNamespace) override;
+
+    const char* getPluginNamespace() const override;
+
+    const char* getPluginType() const override;
+
+    const char* getPluginVersion() const override;
+
+    void destroy() override;
+
+    IPluginV2IOExt* clone() const override;
+
+
+
+    DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const override;
+
+    bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const override;
+
+    bool canBroadcastInputAcrossBatch(int inputIndex) const override;
+
+    void attachToContext(
+        cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) override;
+
+    void configurePlugin(const PluginTensorDesc* in, int nbInput, const PluginTensorDesc* out, int nbOutput) override;
+
+    void detachFromContext() override;
+
+    void setInputSize(int s, int l) {
+        mInputSize = s;
+        Length = l;
+    }
+
+
+private:
+    void forwardGpu(const float *const * inputs, float *mean, float *std, cudaStream_t stream, int batchSize = 1);
+    int mThreadCount = 256;
+    int mInputSize;
+    int Length;
+    Dims outputDims ;
+    const char* mPluginNamespace;
+};
+class layernormCreator : public IPluginCreator
+{
+    public:
+        layernormCreator();
+
+        ~layernormCreator() override = default;
+
+        const char* getPluginName() const override;
+
+        const char* getPluginVersion() const override;
+
+        const PluginFieldCollection* getFieldNames() override;
+
+        IPluginV2IOExt* createPlugin(const char* name, const PluginFieldCollection* fc) override;
+
+        IPluginV2IOExt* deserializePlugin(const char* name, const void* serialData, size_t serialLength) override;
+
+        void setPluginNamespace(const char* libNamespace) override
+        {
+            mNamespace = libNamespace;
+        }
+
+        const char* getPluginNamespace() const override
+        {
+            return mNamespace.c_str();
+        }
+
+    private:
+        std::string mNamespace;
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+
+};
+
+REGISTER_TENSORRT_PLUGIN(layernormCreator);
+};
+
+#endif // LAYERNORM_H
--- a/swin-transformer/semantic-segmentation/logging.h
+++ b/swin-transformer/semantic-segmentation/logging.h
@ -0,0 +1,503 @@
+/*
+ * Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifndef TENSORRT_LOGGING_H
+#define TENSORRT_LOGGING_H
+
+#include "NvInferRuntimeCommon.h"
+#include <cassert>
+#include <ctime>
+#include <iomanip>
+#include <iostream>
+#include <ostream>
+#include <sstream>
+#include <string>
+
+using Severity = nvinfer1::ILogger::Severity;
+
+class LogStreamConsumerBuffer : public std::stringbuf
+{
+public:
+    LogStreamConsumerBuffer(std::ostream& stream, const std::string& prefix, bool shouldLog)
+        : mOutput(stream)
+        , mPrefix(prefix)
+        , mShouldLog(shouldLog)
+    {
+    }
+
+    LogStreamConsumerBuffer(LogStreamConsumerBuffer&& other)
+        : mOutput(other.mOutput)
+    {
+    }
+
+    ~LogStreamConsumerBuffer()
+    {
+        // std::streambuf::pbase() gives a pointer to the beginning of the buffered part of the output sequence
+        // std::streambuf::pptr() gives a pointer to the current position of the output sequence
+        // if the pointer to the beginning is not equal to the pointer to the current position,
+        // call putOutput() to log the output to the stream
+        if (pbase() != pptr())
+        {
+            putOutput();
+        }
+    }
+
+    // synchronizes the stream buffer and returns 0 on success
+    // synchronizing the stream buffer consists of inserting the buffer contents into the stream,
+    // resetting the buffer and flushing the stream
+    virtual int sync()
+    {
+        putOutput();
+        return 0;
+    }
+
+    void putOutput()
+    {
+        if (mShouldLog)
+        {
+            // prepend timestamp
+            std::time_t timestamp = std::time(nullptr);
+            tm* tm_local = std::localtime(&timestamp);
+            std::cout << "[";
+            std::cout << std::setw(2) << std::setfill('0') << 1 + tm_local->tm_mon << "/";
+            std::cout << std::setw(2) << std::setfill('0') << tm_local->tm_mday << "/";
+            std::cout << std::setw(4) << std::setfill('0') << 1900 + tm_local->tm_year << "-";
+            std::cout << std::setw(2) << std::setfill('0') << tm_local->tm_hour << ":";
+            std::cout << std::setw(2) << std::setfill('0') << tm_local->tm_min << ":";
+            std::cout << std::setw(2) << std::setfill('0') << tm_local->tm_sec << "] ";
+            // std::stringbuf::str() gets the string contents of the buffer
+            // insert the buffer contents pre-appended by the appropriate prefix into the stream
+            mOutput << mPrefix << str();
+            // set the buffer to empty
+            str("");
+            // flush the stream
+            mOutput.flush();
+        }
+    }
+
+    void setShouldLog(bool shouldLog)
+    {
+        mShouldLog = shouldLog;
+    }
+
+private:
+    std::ostream& mOutput;
+    std::string mPrefix;
+    bool mShouldLog;
+};
+
+//!
+//! \class LogStreamConsumerBase
+//! \brief Convenience object used to initialize LogStreamConsumerBuffer before std::ostream in LogStreamConsumer
+//!
+class LogStreamConsumerBase
+{
+public:
+    LogStreamConsumerBase(std::ostream& stream, const std::string& prefix, bool shouldLog)
+        : mBuffer(stream, prefix, shouldLog)
+    {
+    }
+
+protected:
+    LogStreamConsumerBuffer mBuffer;
+};
+
+//!
+//! \class LogStreamConsumer
+//! \brief Convenience object used to facilitate use of C++ stream syntax when logging messages.
+//!  Order of base classes is LogStreamConsumerBase and then std::ostream.
+//!  This is because the LogStreamConsumerBase class is used to initialize the LogStreamConsumerBuffer member field
+//!  in LogStreamConsumer and then the address of the buffer is passed to std::ostream.
+//!  This is necessary to prevent the address of an uninitialized buffer from being passed to std::ostream.
+//!  Please do not change the order of the parent classes.
+//!
+class LogStreamConsumer : protected LogStreamConsumerBase, public std::ostream
+{
+public:
+    //! \brief Creates a LogStreamConsumer which logs messages with level severity.
+    //!  Reportable severity determines if the messages are severe enough to be logged.
+    LogStreamConsumer(Severity reportableSeverity, Severity severity)
+        : LogStreamConsumerBase(severityOstream(severity), severityPrefix(severity), severity <= reportableSeverity)
+        , std::ostream(&mBuffer) // links the stream buffer with the stream
+        , mShouldLog(severity <= reportableSeverity)
+        , mSeverity(severity)
+    {
+    }
+
+    LogStreamConsumer(LogStreamConsumer&& other)
+        : LogStreamConsumerBase(severityOstream(other.mSeverity), severityPrefix(other.mSeverity), other.mShouldLog)
+        , std::ostream(&mBuffer) // links the stream buffer with the stream
+        , mShouldLog(other.mShouldLog)
+        , mSeverity(other.mSeverity)
+    {
+    }
+
+    void setReportableSeverity(Severity reportableSeverity)
+    {
+        mShouldLog = mSeverity <= reportableSeverity;
+        mBuffer.setShouldLog(mShouldLog);
+    }
+
+private:
+    static std::ostream& severityOstream(Severity severity)
+    {
+        return severity >= Severity::kINFO ? std::cout : std::cerr;
+    }
+
+    static std::string severityPrefix(Severity severity)
+    {
+        switch (severity)
+        {
+        case Severity::kINTERNAL_ERROR: return "[F] ";
+        case Severity::kERROR: return "[E] ";
+        case Severity::kWARNING: return "[W] ";
+        case Severity::kINFO: return "[I] ";
+        case Severity::kVERBOSE: return "[V] ";
+        default: assert(0); return "";
+        }
+    }
+
+    bool mShouldLog;
+    Severity mSeverity;
+};
+
+//! \class Logger
+//!
+//! \brief Class which manages logging of TensorRT tools and samples
+//!
+//! \details This class provides a common interface for TensorRT tools and samples to log information to the console,
+//! and supports logging two types of messages:
+//!
+//! - Debugging messages with an associated severity (info, warning, error, or internal error/fatal)
+//! - Test pass/fail messages
+//!
+//! The advantage of having all samples use this class for logging as opposed to emitting directly to stdout/stderr is
+//! that the logic for controlling the verbosity and formatting of sample output is centralized in one location.
+//!
+//! In the future, this class could be extended to support dumping test results to a file in some standard format
+//! (for example, JUnit XML), and providing additional metadata (e.g. timing the duration of a test run).
+//!
+//! TODO: For backwards compatibility with existing samples, this class inherits directly from the nvinfer1::ILogger
+//! interface, which is problematic since there isn't a clean separation between messages coming from the TensorRT
+//! library and messages coming from the sample.
+//!
+//! In the future (once all samples are updated to use Logger::getTRTLogger() to access the ILogger) we can refactor the
+//! class to eliminate the inheritance and instead make the nvinfer1::ILogger implementation a member of the Logger
+//! object.
+
+class Logger : public nvinfer1::ILogger
+{
+public:
+    Logger(Severity severity = Severity::kWARNING)
+        : mReportableSeverity(severity)
+    {
+    }
+
+    //!
+    //! \enum TestResult
+    //! \brief Represents the state of a given test
+    //!
+    enum class TestResult
+    {
+        kRUNNING, //!< The test is running
+        kPASSED,  //!< The test passed
+        kFAILED,  //!< The test failed
+        kWAIVED   //!< The test was waived
+    };
+
+    //!
+    //! \brief Forward-compatible method for retrieving the nvinfer::ILogger associated with this Logger
+    //! \return The nvinfer1::ILogger associated with this Logger
+    //!
+    //! TODO Once all samples are updated to use this method to register the logger with TensorRT,
+    //! we can eliminate the inheritance of Logger from ILogger
+    //!
+    nvinfer1::ILogger& getTRTLogger()
+    {
+        return *this;
+    }
+
+    //!
+    //! \brief Implementation of the nvinfer1::ILogger::log() virtual method
+    //!
+    //! Note samples should not be calling this function directly; it will eventually go away once we eliminate the
+    //! inheritance from nvinfer1::ILogger
+    //!
+    void log(Severity severity, const char* msg) override
+    {
+        LogStreamConsumer(mReportableSeverity, severity) << "[TRT] " << std::string(msg) << std::endl;
+    }
+
+    //!
+    //! \brief Method for controlling the verbosity of logging output
+    //!
+    //! \param severity The logger will only emit messages that have severity of this level or higher.
+    //!
+    void setReportableSeverity(Severity severity)
+    {
+        mReportableSeverity = severity;
+    }
+
+    //!
+    //! \brief Opaque handle that holds logging information for a particular test
+    //!
+    //! This object is an opaque handle to information used by the Logger to print test results.
+    //! The sample must call Logger::defineTest() in order to obtain a TestAtom that can be used
+    //! with Logger::reportTest{Start,End}().
+    //!
+    class TestAtom
+    {
+    public:
+        TestAtom(TestAtom&&) = default;
+
+    private:
+        friend class Logger;
+
+        TestAtom(bool started, const std::string& name, const std::string& cmdline)
+            : mStarted(started)
+            , mName(name)
+            , mCmdline(cmdline)
+        {
+        }
+
+        bool mStarted;
+        std::string mName;
+        std::string mCmdline;
+    };
+
+    //!
+    //! \brief Define a test for logging
+    //!
+    //! \param[in] name The name of the test.  This should be a string starting with
+    //!                  "TensorRT" and containing dot-separated strings containing
+    //!                  the characters [A-Za-z0-9_].
+    //!                  For example, "TensorRT.sample_googlenet"
+    //! \param[in] cmdline The command line used to reproduce the test
+    //
+    //! \return a TestAtom that can be used in Logger::reportTest{Start,End}().
+    //!
+    static TestAtom defineTest(const std::string& name, const std::string& cmdline)
+    {
+        return TestAtom(false, name, cmdline);
+    }
+
+    //!
+    //! \brief A convenience overloaded version of defineTest() that accepts an array of command-line arguments
+    //!        as input
+    //!
+    //! \param[in] name The name of the test
+    //! \param[in] argc The number of command-line arguments
+    //! \param[in] argv The array of command-line arguments (given as C strings)
+    //!
+    //! \return a TestAtom that can be used in Logger::reportTest{Start,End}().
+    static TestAtom defineTest(const std::string& name, int argc, char const* const* argv)
+    {
+        auto cmdline = genCmdlineString(argc, argv);
+        return defineTest(name, cmdline);
+    }
+
+    //!
+    //! \brief Report that a test has started.
+    //!
+    //! \pre reportTestStart() has not been called yet for the given testAtom
+    //!
+    //! \param[in] testAtom The handle to the test that has started
+    //!
+    static void reportTestStart(TestAtom& testAtom)
+    {
+        reportTestResult(testAtom, TestResult::kRUNNING);
+        assert(!testAtom.mStarted);
+        testAtom.mStarted = true;
+    }
+
+    //!
+    //! \brief Report that a test has ended.
+    //!
+    //! \pre reportTestStart() has been called for the given testAtom
+    //!
+    //! \param[in] testAtom The handle to the test that has ended
+    //! \param[in] result The result of the test. Should be one of TestResult::kPASSED,
+    //!                   TestResult::kFAILED, TestResult::kWAIVED
+    //!
+    static void reportTestEnd(const TestAtom& testAtom, TestResult result)
+    {
+        assert(result != TestResult::kRUNNING);
+        assert(testAtom.mStarted);
+        reportTestResult(testAtom, result);
+    }
+
+    static int reportPass(const TestAtom& testAtom)
+    {
+        reportTestEnd(testAtom, TestResult::kPASSED);
+        return EXIT_SUCCESS;
+    }
+
+    static int reportFail(const TestAtom& testAtom)
+    {
+        reportTestEnd(testAtom, TestResult::kFAILED);
+        return EXIT_FAILURE;
+    }
+
+    static int reportWaive(const TestAtom& testAtom)
+    {
+        reportTestEnd(testAtom, TestResult::kWAIVED);
+        return EXIT_SUCCESS;
+    }
+
+    static int reportTest(const TestAtom& testAtom, bool pass)
+    {
+        return pass ? reportPass(testAtom) : reportFail(testAtom);
+    }
+
+    Severity getReportableSeverity() const
+    {
+        return mReportableSeverity;
+    }
+
+private:
+    //!
+    //! \brief returns an appropriate string for prefixing a log message with the given severity
+    //!
+    static const char* severityPrefix(Severity severity)
+    {
+        switch (severity)
+        {
+        case Severity::kINTERNAL_ERROR: return "[F] ";
+        case Severity::kERROR: return "[E] ";
+        case Severity::kWARNING: return "[W] ";
+        case Severity::kINFO: return "[I] ";
+        case Severity::kVERBOSE: return "[V] ";
+        default: assert(0); return "";
+        }
+    }
+
+    //!
+    //! \brief returns an appropriate string for prefixing a test result message with the given result
+    //!
+    static const char* testResultString(TestResult result)
+    {
+        switch (result)
+        {
+        case TestResult::kRUNNING: return "RUNNING";
+        case TestResult::kPASSED: return "PASSED";
+        case TestResult::kFAILED: return "FAILED";
+        case TestResult::kWAIVED: return "WAIVED";
+        default: assert(0); return "";
+        }
+    }
+
+    //!
+    //! \brief returns an appropriate output stream (cout or cerr) to use with the given severity
+    //!
+    static std::ostream& severityOstream(Severity severity)
+    {
+        return severity >= Severity::kINFO ? std::cout : std::cerr;
+    }
+
+    //!
+    //! \brief method that implements logging test results
+    //!
+    static void reportTestResult(const TestAtom& testAtom, TestResult result)
+    {
+        severityOstream(Severity::kINFO) << "&&&& " << testResultString(result) << " " << testAtom.mName << " # "
+                                         << testAtom.mCmdline << std::endl;
+    }
+
+    //!
+    //! \brief generate a command line string from the given (argc, argv) values
+    //!
+    static std::string genCmdlineString(int argc, char const* const* argv)
+    {
+        std::stringstream ss;
+        for (int i = 0; i < argc; i++)
+        {
+            if (i > 0)
+                ss << " ";
+            ss << argv[i];
+        }
+        return ss.str();
+    }
+
+    Severity mReportableSeverity;
+};
+
+namespace
+{
+
+//!
+//! \brief produces a LogStreamConsumer object that can be used to log messages of severity kVERBOSE
+//!
+//! Example usage:
+//!
+//!     LOG_VERBOSE(logger) << "hello world" << std::endl;
+//!
+inline LogStreamConsumer LOG_VERBOSE(const Logger& logger)
+{
+    return LogStreamConsumer(logger.getReportableSeverity(), Severity::kVERBOSE);
+}
+
+//!
+//! \brief produces a LogStreamConsumer object that can be used to log messages of severity kINFO
+//!
+//! Example usage:
+//!
+//!     LOG_INFO(logger) << "hello world" << std::endl;
+//!
+inline LogStreamConsumer LOG_INFO(const Logger& logger)
+{
+    return LogStreamConsumer(logger.getReportableSeverity(), Severity::kINFO);
+}
+
+//!
+//! \brief produces a LogStreamConsumer object that can be used to log messages of severity kWARNING
+//!
+//! Example usage:
+//!
+//!     LOG_WARN(logger) << "hello world" << std::endl;
+//!
+inline LogStreamConsumer LOG_WARN(const Logger& logger)
+{
+    return LogStreamConsumer(logger.getReportableSeverity(), Severity::kWARNING);
+}
+
+//!
+//! \brief produces a LogStreamConsumer object that can be used to log messages of severity kERROR
+//!
+//! Example usage:
+//!
+//!     LOG_ERROR(logger) << "hello world" << std::endl;
+//!
+inline LogStreamConsumer LOG_ERROR(const Logger& logger)
+{
+    return LogStreamConsumer(logger.getReportableSeverity(), Severity::kERROR);
+}
+
+//!
+//! \brief produces a LogStreamConsumer object that can be used to log messages of severity kINTERNAL_ERROR
+//         ("fatal" severity)
+//!
+//! Example usage:
+//!
+//!     LOG_FATAL(logger) << "hello world" << std::endl;
+//!
+inline LogStreamConsumer LOG_FATAL(const Logger& logger)
+{
+    return LogStreamConsumer(logger.getReportableSeverity(), Severity::kINTERNAL_ERROR);
+}
+
+} // anonymous namespace
+
+#endif // TENSORRT_LOGGING_H
--- a/swin-transformer/semantic-segmentation/main.cpp
+++ b/swin-transformer/semantic-segmentation/main.cpp
@ -0,0 +1,9 @@
+#include <iostream>
+
+
+using namespace std;
+
+
+
+
+
--- a/swin-transformer/semantic-segmentation/myhpp.h
+++ b/swin-transformer/semantic-segmentation/myhpp.h
@ -0,0 +1,36 @@
+#ifndef MYHPP_H
+#define MYHPP_H
+
+#include <assert.h>
+#include <iostream>
+#include<vector>
+#include<map>
+#define _USE_MATH_DEFINES
+#include <math.h>
+#include <cmath>
+#include<string>
+#include<fstream>
+#include<streambuf>
+#include<ctime>
+#include<chrono>
+#include<iomanip>
+#include<cuda_runtime.h>
+#include<opencv2/core/core.hpp>
+#include<opencv2/imgproc/imgproc.hpp>
+#include<opencv2/imgcodecs/imgcodecs.hpp>
+#include<opencv2/dnn/dnn.hpp>
+//#include <opencv2/highgui/highgui.hpp>
+#include<stdio.h>
+#include<cuda.h>
+//#include <cudnn.h>
+#include <cublas_v2.h>
+#include<driver_types.h>
+#include<NvInfer.h>
+#include<NvInferPlugin.h>
+#include<NvOnnxParser.h>
+#include<NvOnnxConfig.h>
+#include<cstdint>
+
+
+
+#endif // MYHPP_H
--- a/swin-transformer/semantic-segmentation/trainsform.cpp
+++ b/swin-transformer/semantic-segmentation/trainsform.cpp
@ -0,0 +1,325 @@
+#include "common.hpp"
+#include "logging.h"
+#include <fstream>
+#include <iostream>
+#include <map>
+#include <sstream>
+#include <vector>
+#include <chrono>
+
+#define USE_FP32
+
+static Logger gLogger;
+
+const char *INPUT_BLOB_NAME = "data";
+const char *OUTPUT_BLOB_NAME = "output";
+static const int bs = 1;
+static const int channels = 96;
+static const int ch = 3;
+static const int INPUT_H = 576;
+static const int INPUT_W = 576;
+static const int NUM_CLASSES = 15;
+static const int outputSize = 576 * 576;
+cudaStream_t m_cudaStream;
+vector<void *> m_bindings;
+IExecutionContext *m_context;
+
+ICudaEngine *createEngine(unsigned int maxBatchSize, IBuilder *builder, IBuilderConfig *config, DataType dt,std::string wtsPath)
+{
+    INetworkDefinition *network = builder->createNetworkV2(0U);
+    ITensor *data = network->addInput(INPUT_BLOB_NAME, dt, Dims3{ch, INPUT_H, INPUT_W});
+    assert(data);
+    std::map<std::string, Weights> weightMap = loadWeights(wtsPath);
+    ITensor* conv1 = conv(network, weightMap, data, "backbone.patch_embed.proj", channels);
+    ITensor* shuffle1 = shuffle_reshapeApermute(network, conv1, Dims2{channels, -1}, Permutation{1, 0}, true);
+    ITensor *ln = m_layerNorm(network, weightMap, shuffle1, "backbone.patch_embed.norm");
+    debug_print(ln, "ln");
+    //layer0
+
+    ITensor *mask0 = trt_transform_imgMask(network, 147, 7, 3);
+    ITensor *blk00 = blk(network, weightMap, ln, mask0, "backbone.layers.0.blocks.0", INPUT_H / 4, channels, 3, 7, 0);
+    debug_print(blk00, "blk00");
+    ITensor *blk01 = blk(network, weightMap, blk00, mask0, "backbone.layers.0.blocks.1", INPUT_H / 4, channels, 3, 7, 3);
+    debug_print(blk01, "blk01");
+    ITensor* out0 = m_layerNorm(network, weightMap, blk01, "backbone.norm0");
+    out0 = shuffle_reshapeApermute(network, out0, Dims3{INPUT_H / 4, INPUT_H / 4, channels}, Permutation{2, 0, 1}, true);
+    ITensor *down_layer0 = downsample(network, weightMap, blk01, "backbone.layers.0.downsample", INPUT_H / 4);
+    debug_print(down_layer0, "down_blk1");
+    //layer1
+    ITensor *mask1 = trt_transform_imgMask(network, 77, 7, 3);
+    ITensor *blk10 = blk(network, weightMap, down_layer0, mask1, "backbone.layers.1.blocks.0", INPUT_H / 8, channels * 2, 6, 7, 0);
+    debug_print(blk10, "blk10");
+    ITensor *blk11 = blk(network, weightMap, blk10, mask1, "backbone.layers.1.blocks.1", INPUT_H / 8, channels * 2, 6, 7, 3);
+    debug_print(blk11, "blk11");
+    ITensor* out1 = m_layerNorm(network, weightMap, blk11, "backbone.norm1");
+    out1 = shuffle_reshapeApermute(network, out1, Dims3{INPUT_H / 8, INPUT_H / 8, channels * 2}, Permutation{2, 0, 1}, true);
+    ITensor *down_layer1 = downsample(network, weightMap, blk11, "backbone.layers.1.downsample", INPUT_H / 8);
+    debug_print(down_layer1, "down_layer1");
+    //layer2
+    ITensor *mask2 = trt_transform_imgMask(network, 42, 7, 3);
+    ITensor *blk20 = blk(network, weightMap, down_layer1, mask2, "backbone.layers.2.blocks.0", INPUT_H / 16, channels * 4, 12, 7, 0);
+    debug_print(blk20, "blk20");
+    ITensor *blk21 = blk(network, weightMap, blk20, mask2, "backbone.layers.2.blocks.1", INPUT_H / 16, channels * 4, 12, 7, 3);
+    debug_print(blk21, "blk21");
+    ITensor *blk22 = blk(network, weightMap, blk21, mask2, "backbone.layers.2.blocks.2", INPUT_H / 16,channels * 4, 12, 7, 0);
+    debug_print(blk22, "blk22");
+    ITensor *blk23 = blk(network, weightMap, blk22, mask2, "backbone.layers.2.blocks.3", INPUT_H / 16, channels * 4, 12, 7, 3);
+    debug_print(blk23, "blk23");
+    ITensor *blk24 = blk(network, weightMap, blk23, mask2, "backbone.layers.2.blocks.4", INPUT_H / 16, channels * 4, 12, 7, 0);
+    debug_print(blk24, "blk24");
+    ITensor *blk25 = blk(network, weightMap, blk24, mask2, "backbone.layers.2.blocks.5", INPUT_H / 16, channels * 4, 12, 7, 3);
+    debug_print(blk25, "blk25");
+    ITensor* out2 = m_layerNorm(network, weightMap, blk25, "backbone.norm2");
+    out2 = shuffle_reshapeApermute(network, out2, Dims3{INPUT_H / 16, INPUT_H / 16, channels * 4}, Permutation{2, 0, 1}, true);
+    ITensor *down_layer2 = downsample(network, weightMap, blk25, "backbone.layers.2.downsample", INPUT_H / 16);
+    debug_print(down_layer2, "down_layer2");
+    //layer3
+    ITensor *mask3 = trt_transform_imgMask(network, 21, 7, 3);
+    ITensor *blk30 = blk(network, weightMap, down_layer2, mask3, "backbone.layers.3.blocks.0", INPUT_H / 32, channels * 8, 24, 7, 0);
+    debug_print(blk30, "blk30");
+    ITensor *blk31 = blk(network, weightMap, blk30, mask3, "backbone.layers.3.blocks.1", INPUT_H / 32, channels * 8, 24, 7, 3);
+    debug_print(blk31, "blk31");
+    ITensor* out3 = m_layerNorm(network, weightMap, blk31, "backbone.norm3");
+    out3 = shuffle_reshapeApermute(network, out3, Dims3{INPUT_H / 32, INPUT_H / 32, channels * 8}, Permutation{2, 0, 1}, true);
+    ITensor* out[4] = {out0, out1, out2, out3};
+    out0 = transform_lateral_conv(network, weightMap, out0, "decode_head.lateral_convs.0");  // 512,INPUT_H/4,INPUT_H/4
+    out1 = transform_lateral_conv(network, weightMap, out1, "decode_head.lateral_convs.1");  // 512,INPUT_H/8,INPUT_H/8
+    out2 = transform_lateral_conv(network, weightMap, out2, "decode_head.lateral_convs.2");  // 512,INPUT_H/16,INPUT_H/16
+    auto psp_out_0 = transform_psp(network, weightMap, out3, "decode_head.psp_modules.0.1", 1);
+    auto psp_out_1 = transform_psp(network, weightMap, out3, "decode_head.psp_modules.1.1", 2);
+    auto psp_out_2 = transform_psp(network, weightMap, out3, "decode_head.psp_modules.2.1", 3);
+    auto psp_out_3 = transform_psp(network, weightMap, out3, "decode_head.psp_modules.3.1", 6);
+    ITensor* psp_outs[5] = {out3, psp_out_0, psp_out_1, psp_out_2, psp_out_3};
+    auto PSP_outs = network->addConcatenation(psp_outs, 5);
+    PSP_outs->setAxis(0);
+    debug_print(PSP_outs->getOutput(0), "PSP_outs");
+    out3 = transform_lateral_conv(network, weightMap, PSP_outs->getOutput(0), "decode_head.bottleneck", 3, 1, 512);  // 512,INPUT_H/32,INPUT_H/32
+    debug_print(out3, "out3");
+    auto laterals2 = up_Add(network, out3, out2);
+    auto laterals1 = up_Add(network, laterals2, out1);
+    auto laterals0 = up_Add(network, laterals1, out0);
+    auto fpn0 = transform_lateral_conv(network, weightMap, laterals0, "decode_head.fpn_convs.0", 3, 1, 512);
+    auto fpn1 = transform_lateral_conv(network, weightMap, laterals1, "decode_head.fpn_convs.1", 3, 1, 512);
+    auto fpn2 = transform_lateral_conv(network, weightMap, laterals2, "decode_head.fpn_convs.2", 3, 1, 512);
+    fpn1 = resize(network, fpn1,fpn0->getDimensions().d[1]);
+    fpn2 = resize(network, fpn2,fpn0->getDimensions().d[1]);
+    auto fpn3 = resize(network, out3, fpn0->getDimensions().d[1]);
+    ITensor* fpn_outs[4] = {fpn0, fpn1, fpn2, fpn3};
+    auto FPN_outs = network->addConcatenation(fpn_outs, 4);
+    FPN_outs->setAxis(0);
+    debug_print(FPN_outs->getOutput(0), "FPN_outs");
+    auto fpn_output = transform_lateral_conv(network, weightMap, FPN_outs->getOutput(0), "decode_head.fpn_bottleneck", 3, 1, 512);
+    debug_print(fpn_output, "fpn_output");
+    auto seg = network->addConvolutionNd(*fpn_output, NUM_CLASSES, Dims2{1, 1}, weightMap["decode_head.conv_seg.weight"], weightMap["decode_head.conv_seg.bias"]);
+    seg->setStrideNd(Dims2{1, 1});
+    debug_print(seg->getOutput(0), "seg");
+    auto seg_resize = resize(network, seg->getOutput(0), INPUT_H);
+    debug_print(seg_resize, "seg_resize");
+    auto output = network->addTopK(*seg_resize, TopKOperation::kMAX, 1, 0X01)->getOutput(1);
+    debug_print(output, "output");
+
+    std::cout << "set name out" << std::endl;
+    output->setName(OUTPUT_BLOB_NAME);
+    network->markOutput(*output);
+    builder->setMaxBatchSize(12);
+    config->setMaxWorkspaceSize((1 << 30)); // 1G
+#ifdef USE_FP16
+    std::cout<< "use fp16"<<std::endl;
+    config->setFlag(BuilderFlag::kFP16);
+#endif
+    ICudaEngine *engine = builder->buildEngineWithConfig(*network, *config);
+    std::cout << "build success!" << std::endl;
+    network->destroy();
+
+    return engine;
+}
+
+void APIToModel(unsigned int maxBatchSize, IHostMemory **modelStream,std::string wtsPath)
+{
+    IBuilder *builder = createInferBuilder(gLogger);
+    IBuilderConfig *config = builder->createBuilderConfig();
+    ICudaEngine *engine = createEngine(maxBatchSize, builder, config, DataType::kFLOAT, wtsPath);
+    assert(engine != nullptr);
+    (*modelStream) = engine->serialize();
+    engine->destroy();
+    builder->destroy();
+}
+
+void createEng(std::string wtsPath, std::string engine_name)
+{
+    char *trtModelStream{nullptr};
+    size_t size{0};
+
+    IHostMemory *modelStream{nullptr};
+    APIToModel(bs, &modelStream, wtsPath);
+    assert(modelStream != nullptr);
+    std::ofstream p(engine_name, std::ios::binary);
+    if (!p)
+    {
+        std::cerr << "could not open plan output file" << std::endl;
+        return;
+    }
+    p.write(reinterpret_cast<const char *>(modelStream->data()), modelStream->size());
+    modelStream->destroy();
+    std::ifstream file(engine_name, 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();
+    }
+}
+
+void inference_init(string ENGPath,ICudaEngine *m_engine)
+{
+    ifstream cache(ENGPath, ios::binary);
+    cache.seekg(0, ios::end);
+    const int engSize = cache.tellg();
+    cache.seekg(0, ios::beg);
+    void *modelMem = malloc(engSize);
+    cache.read((char*)modelMem, engSize);
+    cache.close();
+    IRuntime *runtime = nvinfer1::createInferRuntime(gLogger);
+    m_engine = runtime->deserializeCudaEngine(modelMem, engSize);
+    runtime->destroy();
+    free(modelMem);
+    if (!m_engine) {
+        cout << "deserialize eng error!" << endl;
+        return;
+    }
+    m_context = m_engine->createExecutionContext();
+    if (cudaStreamCreate(&m_cudaStream) != 0) return;
+    int bindings = m_engine->getNbBindings();
+    if (bindings < 2)
+    {
+        cout << "Error! the network have one input and one output at least!" << endl;
+        return;
+    }
+    cout << "1111111111111" << endl;
+    m_bindings.resize(bindings, nullptr);
+    CHECK(cudaMalloc(&m_bindings.at(0), bs * ch * INPUT_H * INPUT_W * sizeof(float)));
+    CHECK(cudaMalloc(&m_bindings.at(1), bs * outputSize * 4));
+}
+
+void doInference(const float *input, int *output)
+{
+    cout << "do infer:" << endl;
+    CHECK(cudaMemcpyAsync(m_bindings.at(0), input, bs * ch * INPUT_H * INPUT_W * sizeof(float),
+                          cudaMemcpyHostToDevice, m_cudaStream));
+
+    m_context->enqueue(bs, m_bindings.data(), m_cudaStream, nullptr);
+
+    CHECK(cudaMemcpyAsync(output, m_bindings.at(1), bs * outputSize * 4,
+                          cudaMemcpyDeviceToHost, m_cudaStream));
+
+    cudaStreamSynchronize(m_cudaStream);
+}
+
+
+int main(int argc, char** argv)
+{
+    cout << "begin" << endl;
+    //string wts = "G:/shaj/trainsform/ktn5n6_29.511.21.8.wts";
+    //string eng = "G:/shaj/trainsform/trainsform.eng";
+    if (argv[1] = "-s") {
+        string wts = argv[2];
+        string eng = argv[3];
+        createEng(wts,eng);
+    } else {
+        string eng = argv[2];
+
+        ICudaEngine *m_engine;
+
+        inference_init(eng,m_engine);
+
+        vector<cv::Mat> testVal;
+        map<string,cv::Mat> dataProb;
+        vector<string> imgs;
+        cv::Mat img;
+        //string pattern_dir = "G:/shaj/trainsform";
+        string pattern_dir = argv[3];
+        string pattern = pattern_dir+ "/*.bmp";
+        vector<cv::String> images_names;
+        cv::glob(pattern, images_names, false);
+        int i = 0;
+        cv::Scalar Mean = cv::Scalar(123.675, 116.28, 103.53);
+        cv::Scalar Std = cv::Scalar(58.395, 57.12, 57.375);
+        cv::Size size = {INPUT_H,INPUT_W};
+
+        for (auto image_name: images_names)
+        {
+            if (i < bs)
+            {
+                cv::Mat Img = cv::imread(image_name, 1);
+
+                testVal.push_back(Img);
+                cout << image_name << endl;
+                imgs.push_back(image_name);
+            }
+        }
+        float *data = new float[bs * ch * INPUT_H * INPUT_W];
+        int *output = new int[bs * outputSize];
+
+        cv::Mat Transed_t = BlobFromImages(testVal, cv::Size{INPUT_H, INPUT_W}, Mean, Std, true, false);
+
+        memcpy(data, Transed_t.data, bs * ch * INPUT_H * INPUT_W * sizeof(float));
+
+        //for(int i = 0 ; i< 20; i++){
+
+        auto start_time = std::chrono::system_clock::now();
+        doInference(data, output);
+
+        auto end_time = std::chrono::system_clock::now();
+        float duration;
+        duration = std::chrono::duration_cast<std::chrono::milliseconds>(end_time - start_time).count();
+        cout << "time:" << duration << endl;
+        //}
+        //    for(int i = 0; i < 100; i++)
+        //        cout<<i<<":"<<output[i]<<endl;
+
+        int n = 0;
+        int *out = new int[outputSize];
+        string outPath = pattern_dir + "/output";
+        for (int i = 0; i < testVal.size(); i++)
+        {
+            cv::Mat img = cv::imread(imgs[i], 1);
+            cv::Mat dst;
+            cv::resize(img,dst,cv::Size{INPUT_H, INPUT_W});
+            //string outPath_n = outPath + "/"+to_string(n) + ".jpg";
+            n += 1;
+            out = output + i * outputSize;
+            for (int i = 0; i < outputSize; i++)
+            {
+                if (out[i] != 0)
+                {
+                    int w = i % (INPUT_H);
+                    int h = i / (INPUT_W);
+                    dst.at<cv::Vec3b>(h, w)[0] =  out[i] * 10;
+                    dst.at<cv::Vec3b>(h, w)[1] =  out[i] * 30;
+                    dst.at<cv::Vec3b>(h, w)[2] =  out[i] * 40;
+                }
+            }
+            //cout<<outPath_n<<endl;
+            string outPath_result = imgs[i].replace(0, pattern_dir.size(), outPath);
+            cout << outPath_result << endl;
+            cv::imwrite(outPath_result, dst);
+        }
+        testVal.clear();
+        imgs.clear();
+    }
+
+    m_context->destroy();
+    m_engine->destroy();
+    for (auto bindings: m_bindings) {
+        cudaFree(bindings);
+    }
+    cudaFree(m_cudaStream);
+
+    cout << "swin_transform" << endl;
+    return 0;
+}
+
--- a/swin-transformer/semantic-segmentation/utilsn.h
+++ b/swin-transformer/semantic-segmentation/utilsn.h
@ -0,0 +1,90 @@
+#ifndef UTILSN_H
+#define UTILSN_H
+
+#include <iostream>
+#include <vector>
+#include <algorithm>
+#include <cudnn.h>
+#include <NvInfer.h>
+#include "myhpp.h"
+using namespace std;
+
+#ifndef CUDA_CHECK
+
+#define CUDA_CHECK(callstr)                                                                    \
+    {                                                                                          \
+        cudaError_t error_code = callstr;                                                      \
+        if (error_code != cudaSuccess) {                                                       \
+            std::cerr << "CUDA error " << error_code << " at " << __FILE__ << ":" << __LINE__; \
+            assert(0);                                                                         \
+        }                                                                                      \
+    }
+
+#endif
+
+namespace Tn
+{
+    class Profiler : public nvinfer1::IProfiler
+    {
+    public:
+        void printLayerTimes(int itrationsTimes)
+        {
+            float totalTime = 0;
+            for (size_t i = 0; i < mProfile.size(); i++)
+            {
+                printf("%-40.40s %4.3fms\n", mProfile[i].first.c_str(), mProfile[i].second / itrationsTimes);
+                totalTime += mProfile[i].second;
+            }
+            printf("Time over all layers: %4.3f\n", totalTime / itrationsTimes);
+        }
+    private:
+        typedef std::pair<std::string, float> Record;
+        std::vector<Record> mProfile;
+
+        virtual void reportLayerTime(const char* layerName, float ms)
+        {
+            auto record = std::find_if(mProfile.begin(), mProfile.end(), [&](const Record& r){ return r.first == layerName; });
+            if (record == mProfile.end())
+               { mProfile.push_back(std::make_pair(layerName, ms));}
+            else
+                record->second += ms;
+        }
+    };
+
+    template<typename T>
+    void write(char*& buffer, const T& val)
+    {
+        *reinterpret_cast<T*>(buffer) = val;
+        buffer += sizeof(T);
+    }
+
+    template<typename T>
+    void read(const char*& buffer, T& val)
+    {
+        val = *reinterpret_cast<const T*>(buffer);
+        buffer += sizeof(T);
+    }
+
+//    void* copyToDevice(const void* data, size_t count)
+//    {
+//        void* deviceData;
+//        cudaMalloc(&deviceData, count);
+//        cudaMemcpy(deviceData, data, count, cudaMemcpyHostToDevice);
+//        return deviceData;
+//    }
+//    void deserializeToDevice(const char*& hostBuffer, void*& deviceWeights, size_t size)
+//    {
+//        deviceWeights = copyToDevice(hostBuffer, size);
+//        hostBuffer += size;
+//    }
+//    size_t type2size(nvinfer1::DataType type) { return sizeof(float); }
+//    void convertAndCopyToBuffer(char*& buffer, const nvinfer1::Weights& weights)
+//    {
+//        memcpy(buffer, weights.values, weights.count * type2size(weights.type));
+//        buffer += weights.count * type2size(weights.type);
+//    }
+
+
+}
+
+#endif // UTILSN_H