From e7ec22d8b8613de3d81a477f462b5caed793f896 Mon Sep 17 00:00:00 2001 From: Wang Xinyu Date: Tue, 15 Mar 2022 12:10:49 +0800 Subject: [PATCH] add swin-transformer (#934) * add swin-transformer * change dir * reformat * fix gen wts and polish readme Co-authored-by: wdhao <873665182@qq.com> --- .../semantic-segmentation/CMakeLists.txt | 75 ++ .../semantic-segmentation/README.md | 36 + .../semantic-segmentation/UpsampleKernel.cu | 146 +++ .../semantic-segmentation/UpsamplePlugin.cpp | 235 ++++ .../semantic-segmentation/UpsamplePlugin.h | 88 ++ .../semantic-segmentation/UpsmapleKernel.h | 20 + .../semantic-segmentation/common.hpp | 888 +++++++++++++ .../semantic-segmentation/fillmask.cu | 182 +++ .../semantic-segmentation/fillmask.h | 90 ++ .../semantic-segmentation/gelu.cu | 180 +++ swin-transformer/semantic-segmentation/gelu.h | 88 ++ .../semantic-segmentation/gen_wts.py | 19 + .../semantic-segmentation/include/dirent.h | 1160 +++++++++++++++++ .../semantic-segmentation/layerNorm.cu | 195 +++ .../semantic-segmentation/layerNorm.h | 130 ++ .../semantic-segmentation/logging.h | 503 +++++++ .../semantic-segmentation/main.cpp | 9 + .../semantic-segmentation/myhpp.h | 36 + .../semantic-segmentation/trainsform.cpp | 325 +++++ .../semantic-segmentation/utilsn.h | 90 ++ 20 files changed, 4495 insertions(+) create mode 100644 swin-transformer/semantic-segmentation/CMakeLists.txt create mode 100644 swin-transformer/semantic-segmentation/README.md create mode 100644 swin-transformer/semantic-segmentation/UpsampleKernel.cu create mode 100644 swin-transformer/semantic-segmentation/UpsamplePlugin.cpp create mode 100644 swin-transformer/semantic-segmentation/UpsamplePlugin.h create mode 100644 swin-transformer/semantic-segmentation/UpsmapleKernel.h create mode 100644 swin-transformer/semantic-segmentation/common.hpp create mode 100644 swin-transformer/semantic-segmentation/fillmask.cu create mode 100644 swin-transformer/semantic-segmentation/fillmask.h create mode 100644 swin-transformer/semantic-segmentation/gelu.cu create mode 100644 swin-transformer/semantic-segmentation/gelu.h create mode 100644 swin-transformer/semantic-segmentation/gen_wts.py create mode 100644 swin-transformer/semantic-segmentation/include/dirent.h create mode 100644 swin-transformer/semantic-segmentation/layerNorm.cu create mode 100644 swin-transformer/semantic-segmentation/layerNorm.h create mode 100644 swin-transformer/semantic-segmentation/logging.h create mode 100644 swin-transformer/semantic-segmentation/main.cpp create mode 100644 swin-transformer/semantic-segmentation/myhpp.h create mode 100644 swin-transformer/semantic-segmentation/trainsform.cpp create mode 100644 swin-transformer/semantic-segmentation/utilsn.h diff --git a/swin-transformer/semantic-segmentation/CMakeLists.txt b/swin-transformer/semantic-segmentation/CMakeLists.txt new file mode 100644 index 0000000..0dac728 --- /dev/null +++ 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) \ No newline at end of file diff --git a/swin-transformer/semantic-segmentation/README.md b/swin-transformer/semantic-segmentation/README.md new file mode 100644 index 0000000..56e88cf --- /dev/null +++ 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) + diff --git a/swin-transformer/semantic-segmentation/UpsampleKernel.cu b/swin-transformer/semantic-segmentation/UpsampleKernel.cu new file mode 100644 index 0000000..9b09d46 --- /dev/null +++ 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 +__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<<< 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(inputs), + static_cast(outputs)); + return 0; +} diff --git a/swin-transformer/semantic-segmentation/UpsamplePlugin.cpp b/swin-transformer/semantic-segmentation/UpsamplePlugin.cpp new file mode 100644 index 0000000..637b680 --- /dev/null +++ b/swin-transformer/semantic-segmentation/UpsamplePlugin.cpp @@ -0,0 +1,235 @@ +#include +#include "UpsmapleKernel.h" +#include "UpsamplePlugin.h" + +#include +#include + +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 UpsamplePluginCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(UpsamplePluginCreator); + +template +void writeToBuffer(char*& buffer, const T& val) +{ + *reinterpret_cast(buffer) = val; + buffer += sizeof(T); +} + +// Helper function for deserializing plugin +template +T readFromBuffer(const char*& buffer) +{ + T val = *reinterpret_cast(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(data); + const char *a = d; + + mScaleFactor_h = readFromBuffer(d); + mScaleFactor_w = readFromBuffer(d); + mInputVolume = readFromBuffer(d); + mInputShape.c() = readFromBuffer(d); + mInputShape.h() = readFromBuffer(d); + mInputShape.w() = readFromBuffer(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(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(fields[i].data)); + scaleFactor_w = *(static_cast(fields[i].data)+1); + //std::cout< +#include + + +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 mPluginAttributes; + std::string mNamespace; +}; + +#endif diff --git a/swin-transformer/semantic-segmentation/UpsmapleKernel.h b/swin-transformer/semantic-segmentation/UpsmapleKernel.h new file mode 100644 index 0000000..7c32666 --- /dev/null +++ b/swin-transformer/semantic-segmentation/UpsmapleKernel.h @@ -0,0 +1,20 @@ +#ifndef UPSAMPLE_KERNEL_H +#define UPSAMPLE_KERNEL_H + +#include +#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 diff --git a/swin-transformer/semantic-segmentation/common.hpp b/swin-transformer/semantic-segmentation/common.hpp new file mode 100644 index 0000000..c5b16a6 --- /dev/null +++ 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 +#include +#include +#include +#include +#include +#include + +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 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< loadWeights(const std::string file) { + std::cout << "Loading weights: " << file << std::endl; + std::map 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(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 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 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 = "<(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 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=Wp-window_size && j < Wp-shift_size) + mask_param[i*Wp + j] = 1.0; + else if(i= Wp-shift_size) + mask_param[i*Wp + j] = 2.0; + + else if(i >= Hp-window_size && i < Hp-shift_size && j= 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= 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"<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<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<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 shifts, vector 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 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: "<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<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 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]<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 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 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 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 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(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(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(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 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 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 diff --git a/swin-transformer/semantic-segmentation/fillmask.cu b/swin-transformer/semantic-segmentation/fillmask.cu new file mode 100644 index 0000000..0f3a6d5 --- /dev/null +++ b/swin-transformer/semantic-segmentation/fillmask.cu @@ -0,0 +1,182 @@ +#include "fillmask.h" +#include +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(data), *a = d; + Tn::read(d, mInputSize); + assert(d == a + length); + } + + void fillmask::serialize(void* buffer) const + { + char* d = static_cast(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 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; + } + + +} + diff --git a/swin-transformer/semantic-segmentation/fillmask.h b/swin-transformer/semantic-segmentation/fillmask.h new file mode 100644 index 0000000..1459ab1 --- /dev/null +++ b/swin-transformer/semantic-segmentation/fillmask.h @@ -0,0 +1,90 @@ +#ifndef FILLMASK_H +#define FILLMASK_H + + +#include +#include +#include "NvInfer.h" +#include "myhpp.h" +#include +#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 mPluginAttributes; + }; + REGISTER_TENSORRT_PLUGIN(fillmaskCreator); +}; + +#endif // FILLMASK_H diff --git a/swin-transformer/semantic-segmentation/gelu.cu b/swin-transformer/semantic-segmentation/gelu.cu new file mode 100644 index 0000000..5a92c54 --- /dev/null +++ b/swin-transformer/semantic-segmentation/gelu.cu @@ -0,0 +1,180 @@ +#include "gelu.h" +#include +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(data), *a = d; + Tn::read(d, mInputSize); + assert(d == a + length); + } + + void gelu::serialize(void* buffer) const + { + char* d = static_cast(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 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; + } + + +} + diff --git a/swin-transformer/semantic-segmentation/gelu.h b/swin-transformer/semantic-segmentation/gelu.h new file mode 100644 index 0000000..5007992 --- /dev/null +++ b/swin-transformer/semantic-segmentation/gelu.h @@ -0,0 +1,88 @@ +#ifndef GELU_H +#define GELU_H + +#include +#include +#include "NvInfer.h" +#include "myhpp.h" +#include +#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 mPluginAttributes; + }; + REGISTER_TENSORRT_PLUGIN(geluCreator); +}; +#endif // GELU_H diff --git a/swin-transformer/semantic-segmentation/gen_wts.py b/swin-transformer/semantic-segmentation/gen_wts.py new file mode 100644 index 0000000..9bb1b2c --- /dev/null +++ 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') diff --git a/swin-transformer/semantic-segmentation/include/dirent.h b/swin-transformer/semantic-segmentation/include/dirent.h new file mode 100644 index 0000000..c885b39 --- /dev/null +++ b/swin-transformer/semantic-segmentation/include/dirent.h @@ -0,0 +1,1160 @@ +/* + * Dirent interface for Microsoft Visual Studio + * + * Copyright (C) 1998-2019 Toni Ronkko + * This file is part of dirent. Dirent may be freely distributed + * under the MIT license. For all details and documentation, see + * https://github.com/tronkko/dirent + */ +#ifndef DIRENT_H +#define DIRENT_H + +/* Hide warnings about unreferenced local functions */ +#if defined(__clang__) +# pragma clang diagnostic ignored "-Wunused-function" +#elif defined(_MSC_VER) +# pragma warning(disable:4505) +#elif defined(__GNUC__) +# pragma GCC diagnostic ignored "-Wunused-function" +#endif + +/* + * Include windows.h without Windows Sockets 1.1 to prevent conflicts with + * Windows Sockets 2.0. + */ +#ifndef WIN32_LEAN_AND_MEAN +# define WIN32_LEAN_AND_MEAN +#endif +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +/* Indicates that d_type field is available in dirent structure */ +#define _DIRENT_HAVE_D_TYPE + +/* Indicates that d_namlen field is available in dirent structure */ +#define _DIRENT_HAVE_D_NAMLEN + +/* Entries missing from MSVC 6.0 */ +#if !defined(FILE_ATTRIBUTE_DEVICE) +# define FILE_ATTRIBUTE_DEVICE 0x40 +#endif + +/* File type and permission flags for stat(), general mask */ +#if !defined(S_IFMT) +# define S_IFMT _S_IFMT +#endif + +/* Directory bit */ +#if !defined(S_IFDIR) +# define S_IFDIR _S_IFDIR +#endif + +/* Character device bit */ +#if !defined(S_IFCHR) +# define S_IFCHR _S_IFCHR +#endif + +/* Pipe bit */ +#if !defined(S_IFFIFO) +# define S_IFFIFO _S_IFFIFO +#endif + +/* Regular file bit */ +#if !defined(S_IFREG) +# define S_IFREG _S_IFREG +#endif + +/* Read permission */ +#if !defined(S_IREAD) +# define S_IREAD _S_IREAD +#endif + +/* Write permission */ +#if !defined(S_IWRITE) +# define S_IWRITE _S_IWRITE +#endif + +/* Execute permission */ +#if !defined(S_IEXEC) +# define S_IEXEC _S_IEXEC +#endif + +/* Pipe */ +#if !defined(S_IFIFO) +# define S_IFIFO _S_IFIFO +#endif + +/* Block device */ +#if !defined(S_IFBLK) +# define S_IFBLK 0 +#endif + +/* Link */ +#if !defined(S_IFLNK) +# define S_IFLNK 0 +#endif + +/* Socket */ +#if !defined(S_IFSOCK) +# define S_IFSOCK 0 +#endif + +/* Read user permission */ +#if !defined(S_IRUSR) +# define S_IRUSR S_IREAD +#endif + +/* Write user permission */ +#if !defined(S_IWUSR) +# define S_IWUSR S_IWRITE +#endif + +/* Execute user permission */ +#if !defined(S_IXUSR) +# define S_IXUSR 0 +#endif + +/* Read group permission */ +#if !defined(S_IRGRP) +# define S_IRGRP 0 +#endif + +/* Write group permission */ +#if !defined(S_IWGRP) +# define S_IWGRP 0 +#endif + +/* Execute group permission */ +#if !defined(S_IXGRP) +# define S_IXGRP 0 +#endif + +/* Read others permission */ +#if !defined(S_IROTH) +# define S_IROTH 0 +#endif + +/* Write others permission */ +#if !defined(S_IWOTH) +# define S_IWOTH 0 +#endif + +/* Execute others permission */ +#if !defined(S_IXOTH) +# define S_IXOTH 0 +#endif + +/* Maximum length of file name */ +#if !defined(PATH_MAX) +# define PATH_MAX MAX_PATH +#endif +#if !defined(FILENAME_MAX) +# define FILENAME_MAX MAX_PATH +#endif +#if !defined(NAME_MAX) +# define NAME_MAX FILENAME_MAX +#endif + +/* File type flags for d_type */ +#define DT_UNKNOWN 0 +#define DT_REG S_IFREG +#define DT_DIR S_IFDIR +#define DT_FIFO S_IFIFO +#define DT_SOCK S_IFSOCK +#define DT_CHR S_IFCHR +#define DT_BLK S_IFBLK +#define DT_LNK S_IFLNK + +/* Macros for converting between st_mode and d_type */ +#define IFTODT(mode) ((mode) & S_IFMT) +#define DTTOIF(type) (type) + +/* + * File type macros. Note that block devices, sockets and links cannot be + * distinguished on Windows and the macros S_ISBLK, S_ISSOCK and S_ISLNK are + * only defined for compatibility. These macros should always return false + * on Windows. + */ +#if !defined(S_ISFIFO) +# define S_ISFIFO(mode) (((mode) & S_IFMT) == S_IFIFO) +#endif +#if !defined(S_ISDIR) +# define S_ISDIR(mode) (((mode) & S_IFMT) == S_IFDIR) +#endif +#if !defined(S_ISREG) +# define S_ISREG(mode) (((mode) & S_IFMT) == S_IFREG) +#endif +#if !defined(S_ISLNK) +# define S_ISLNK(mode) (((mode) & S_IFMT) == S_IFLNK) +#endif +#if !defined(S_ISSOCK) +# define S_ISSOCK(mode) (((mode) & S_IFMT) == S_IFSOCK) +#endif +#if !defined(S_ISCHR) +# define S_ISCHR(mode) (((mode) & S_IFMT) == S_IFCHR) +#endif +#if !defined(S_ISBLK) +# define S_ISBLK(mode) (((mode) & S_IFMT) == S_IFBLK) +#endif + +/* Return the exact length of the file name without zero terminator */ +#define _D_EXACT_NAMLEN(p) ((p)->d_namlen) + +/* Return the maximum size of a file name */ +#define _D_ALLOC_NAMLEN(p) ((PATH_MAX)+1) + + +#ifdef __cplusplus +extern "C" { +#endif + + +/* Wide-character version */ +struct _wdirent { + /* Always zero */ + long d_ino; + + /* File position within stream */ + long d_off; + + /* Structure size */ + unsigned short d_reclen; + + /* Length of name without \0 */ + size_t d_namlen; + + /* File type */ + int d_type; + + /* File name */ + wchar_t d_name[PATH_MAX+1]; +}; +typedef struct _wdirent _wdirent; + +struct _WDIR { + /* Current directory entry */ + struct _wdirent ent; + + /* Private file data */ + WIN32_FIND_DATAW data; + + /* True if data is valid */ + int cached; + + /* Win32 search handle */ + HANDLE handle; + + /* Initial directory name */ + wchar_t *patt; +}; +typedef struct _WDIR _WDIR; + +/* Multi-byte character version */ +struct dirent { + /* Always zero */ + long d_ino; + + /* File position within stream */ + long d_off; + + /* Structure size */ + unsigned short d_reclen; + + /* Length of name without \0 */ + size_t d_namlen; + + /* File type */ + int d_type; + + /* File name */ + char d_name[PATH_MAX+1]; +}; +typedef struct dirent dirent; + +struct DIR { + struct dirent ent; + struct _WDIR *wdirp; +}; +typedef struct DIR DIR; + + +/* Dirent functions */ +static DIR *opendir (const char *dirname); +static _WDIR *_wopendir (const wchar_t *dirname); + +static struct dirent *readdir (DIR *dirp); +static struct _wdirent *_wreaddir (_WDIR *dirp); + +static int readdir_r( + DIR *dirp, struct dirent *entry, struct dirent **result); +static int _wreaddir_r( + _WDIR *dirp, struct _wdirent *entry, struct _wdirent **result); + +static int closedir (DIR *dirp); +static int _wclosedir (_WDIR *dirp); + +static void rewinddir (DIR* dirp); +static void _wrewinddir (_WDIR* dirp); + +static int scandir (const char *dirname, struct dirent ***namelist, + int (*filter)(const struct dirent*), + int (*compare)(const struct dirent**, const struct dirent**)); + +static int alphasort (const struct dirent **a, const struct dirent **b); + +static int versionsort (const struct dirent **a, const struct dirent **b); + + +/* For compatibility with Symbian */ +#define wdirent _wdirent +#define WDIR _WDIR +#define wopendir _wopendir +#define wreaddir _wreaddir +#define wclosedir _wclosedir +#define wrewinddir _wrewinddir + + +/* Internal utility functions */ +static WIN32_FIND_DATAW *dirent_first (_WDIR *dirp); +static WIN32_FIND_DATAW *dirent_next (_WDIR *dirp); + +static int dirent_mbstowcs_s( + size_t *pReturnValue, + wchar_t *wcstr, + size_t sizeInWords, + const char *mbstr, + size_t count); + +static int dirent_wcstombs_s( + size_t *pReturnValue, + char *mbstr, + size_t sizeInBytes, + const wchar_t *wcstr, + size_t count); + +static void dirent_set_errno (int error); + + +/* + * Open directory stream DIRNAME for read and return a pointer to the + * internal working area that is used to retrieve individual directory + * entries. + */ +static _WDIR* +_wopendir( + const wchar_t *dirname) +{ + _WDIR *dirp; + DWORD n; + wchar_t *p; + + /* Must have directory name */ + if (dirname == NULL || dirname[0] == '\0') { + dirent_set_errno (ENOENT); + return NULL; + } + + /* Allocate new _WDIR structure */ + dirp = (_WDIR*) malloc (sizeof (struct _WDIR)); + if (!dirp) { + return NULL; + } + + /* Reset _WDIR structure */ + dirp->handle = INVALID_HANDLE_VALUE; + dirp->patt = NULL; + dirp->cached = 0; + + /* + * Compute the length of full path plus zero terminator + * + * Note that on WinRT there's no way to convert relative paths + * into absolute paths, so just assume it is an absolute path. + */ +#if WINAPI_FAMILY_PARTITION(WINAPI_PARTITION_DESKTOP) + /* Desktop */ + n = GetFullPathNameW (dirname, 0, NULL, NULL); +#else + /* WinRT */ + n = wcslen (dirname); +#endif + + /* Allocate room for absolute directory name and search pattern */ + dirp->patt = (wchar_t*) malloc (sizeof (wchar_t) * n + 16); + if (dirp->patt == NULL) { + goto exit_closedir; + } + + /* + * Convert relative directory name to an absolute one. This + * allows rewinddir() to function correctly even when current + * working directory is changed between opendir() and rewinddir(). + * + * Note that on WinRT there's no way to convert relative paths + * into absolute paths, so just assume it is an absolute path. + */ +#if WINAPI_FAMILY_PARTITION(WINAPI_PARTITION_DESKTOP) + /* Desktop */ + n = GetFullPathNameW (dirname, n, dirp->patt, NULL); + if (n <= 0) { + goto exit_closedir; + } +#else + /* WinRT */ + wcsncpy_s (dirp->patt, n+1, dirname, n); +#endif + + /* Append search pattern \* to the directory name */ + p = dirp->patt + n; + switch (p[-1]) { + case '\\': + case '/': + case ':': + /* Directory ends in path separator, e.g. c:\temp\ */ + /*NOP*/; + break; + + default: + /* Directory name doesn't end in path separator */ + *p++ = '\\'; + } + *p++ = '*'; + *p = '\0'; + + /* Open directory stream and retrieve the first entry */ + if (!dirent_first (dirp)) { + goto exit_closedir; + } + + /* Success */ + return dirp; + + /* Failure */ +exit_closedir: + _wclosedir (dirp); + return NULL; +} + +/* + * Read next directory entry. + * + * Returns pointer to static directory entry which may be overwritten by + * subsequent calls to _wreaddir(). + */ +static struct _wdirent* +_wreaddir( + _WDIR *dirp) +{ + struct _wdirent *entry; + + /* + * Read directory entry to buffer. We can safely ignore the return value + * as entry will be set to NULL in case of error. + */ + (void) _wreaddir_r (dirp, &dirp->ent, &entry); + + /* Return pointer to statically allocated directory entry */ + return entry; +} + +/* + * Read next directory entry. + * + * Returns zero on success. If end of directory stream is reached, then sets + * result to NULL and returns zero. + */ +static int +_wreaddir_r( + _WDIR *dirp, + struct _wdirent *entry, + struct _wdirent **result) +{ + WIN32_FIND_DATAW *datap; + + /* Read next directory entry */ + datap = dirent_next (dirp); + if (datap) { + size_t n; + DWORD attr; + + /* + * Copy file name as wide-character string. If the file name is too + * long to fit in to the destination buffer, then truncate file name + * to PATH_MAX characters and zero-terminate the buffer. + */ + n = 0; + while (n < PATH_MAX && datap->cFileName[n] != 0) { + entry->d_name[n] = datap->cFileName[n]; + n++; + } + entry->d_name[n] = 0; + + /* Length of file name excluding zero terminator */ + entry->d_namlen = n; + + /* File type */ + attr = datap->dwFileAttributes; + if ((attr & FILE_ATTRIBUTE_DEVICE) != 0) { + entry->d_type = DT_CHR; + } else if ((attr & FILE_ATTRIBUTE_DIRECTORY) != 0) { + entry->d_type = DT_DIR; + } else { + entry->d_type = DT_REG; + } + + /* Reset dummy fields */ + entry->d_ino = 0; + entry->d_off = 0; + entry->d_reclen = sizeof (struct _wdirent); + + /* Set result address */ + *result = entry; + + } else { + + /* Return NULL to indicate end of directory */ + *result = NULL; + + } + + return /*OK*/0; +} + +/* + * Close directory stream opened by opendir() function. This invalidates the + * DIR structure as well as any directory entry read previously by + * _wreaddir(). + */ +static int +_wclosedir( + _WDIR *dirp) +{ + int ok; + if (dirp) { + + /* Release search handle */ + if (dirp->handle != INVALID_HANDLE_VALUE) { + FindClose (dirp->handle); + } + + /* Release search pattern */ + free (dirp->patt); + + /* Release directory structure */ + free (dirp); + ok = /*success*/0; + + } else { + + /* Invalid directory stream */ + dirent_set_errno (EBADF); + ok = /*failure*/-1; + + } + return ok; +} + +/* + * Rewind directory stream such that _wreaddir() returns the very first + * file name again. + */ +static void +_wrewinddir( + _WDIR* dirp) +{ + if (dirp) { + /* Release existing search handle */ + if (dirp->handle != INVALID_HANDLE_VALUE) { + FindClose (dirp->handle); + } + + /* Open new search handle */ + dirent_first (dirp); + } +} + +/* Get first directory entry (internal) */ +static WIN32_FIND_DATAW* +dirent_first( + _WDIR *dirp) +{ + WIN32_FIND_DATAW *datap; + DWORD error; + + /* Open directory and retrieve the first entry */ + dirp->handle = FindFirstFileExW( + dirp->patt, FindExInfoStandard, &dirp->data, + FindExSearchNameMatch, NULL, 0); + if (dirp->handle != INVALID_HANDLE_VALUE) { + + /* a directory entry is now waiting in memory */ + datap = &dirp->data; + dirp->cached = 1; + + } else { + + /* Failed to open directory: no directory entry in memory */ + dirp->cached = 0; + datap = NULL; + + /* Set error code */ + error = GetLastError (); + switch (error) { + case ERROR_ACCESS_DENIED: + /* No read access to directory */ + dirent_set_errno (EACCES); + break; + + case ERROR_DIRECTORY: + /* Directory name is invalid */ + dirent_set_errno (ENOTDIR); + break; + + case ERROR_PATH_NOT_FOUND: + default: + /* Cannot find the file */ + dirent_set_errno (ENOENT); + } + + } + return datap; +} + +/* + * Get next directory entry (internal). + * + * Returns + */ +static WIN32_FIND_DATAW* +dirent_next( + _WDIR *dirp) +{ + WIN32_FIND_DATAW *p; + + /* Get next directory entry */ + if (dirp->cached != 0) { + + /* A valid directory entry already in memory */ + p = &dirp->data; + dirp->cached = 0; + + } else if (dirp->handle != INVALID_HANDLE_VALUE) { + + /* Get the next directory entry from stream */ + if (FindNextFileW (dirp->handle, &dirp->data) != FALSE) { + /* Got a file */ + p = &dirp->data; + } else { + /* The very last entry has been processed or an error occurred */ + FindClose (dirp->handle); + dirp->handle = INVALID_HANDLE_VALUE; + p = NULL; + } + + } else { + + /* End of directory stream reached */ + p = NULL; + + } + + return p; +} + +/* + * Open directory stream using plain old C-string. + */ +static DIR* +opendir( + const char *dirname) +{ + struct DIR *dirp; + + /* Must have directory name */ + if (dirname == NULL || dirname[0] == '\0') { + dirent_set_errno (ENOENT); + return NULL; + } + + /* Allocate memory for DIR structure */ + dirp = (DIR*) malloc (sizeof (struct DIR)); + if (!dirp) { + return NULL; + } + { + int error; + wchar_t wname[PATH_MAX + 1]; + size_t n; + + /* Convert directory name to wide-character string */ + error = dirent_mbstowcs_s( + &n, wname, PATH_MAX + 1, dirname, PATH_MAX + 1); + if (error) { + /* + * Cannot convert file name to wide-character string. This + * occurs if the string contains invalid multi-byte sequences or + * the output buffer is too small to contain the resulting + * string. + */ + goto exit_free; + } + + + /* Open directory stream using wide-character name */ + dirp->wdirp = _wopendir (wname); + if (!dirp->wdirp) { + goto exit_free; + } + + } + + /* Success */ + return dirp; + + /* Failure */ +exit_free: + free (dirp); + return NULL; +} + +/* + * Read next directory entry. + */ +static struct dirent* +readdir( + DIR *dirp) +{ + struct dirent *entry; + + /* + * Read directory entry to buffer. We can safely ignore the return value + * as entry will be set to NULL in case of error. + */ + (void) readdir_r (dirp, &dirp->ent, &entry); + + /* Return pointer to statically allocated directory entry */ + return entry; +} + +/* + * Read next directory entry into called-allocated buffer. + * + * Returns zero on success. If the end of directory stream is reached, then + * sets result to NULL and returns zero. + */ +static int +readdir_r( + DIR *dirp, + struct dirent *entry, + struct dirent **result) +{ + WIN32_FIND_DATAW *datap; + + /* Read next directory entry */ + datap = dirent_next (dirp->wdirp); + if (datap) { + size_t n; + int error; + + /* Attempt to convert file name to multi-byte string */ + error = dirent_wcstombs_s( + &n, entry->d_name, PATH_MAX + 1, datap->cFileName, PATH_MAX + 1); + + /* + * If the file name cannot be represented by a multi-byte string, + * then attempt to use old 8+3 file name. This allows traditional + * Unix-code to access some file names despite of unicode + * characters, although file names may seem unfamiliar to the user. + * + * Be ware that the code below cannot come up with a short file + * name unless the file system provides one. At least + * VirtualBox shared folders fail to do this. + */ + if (error && datap->cAlternateFileName[0] != '\0') { + error = dirent_wcstombs_s( + &n, entry->d_name, PATH_MAX + 1, + datap->cAlternateFileName, PATH_MAX + 1); + } + + if (!error) { + DWORD attr; + + /* Length of file name excluding zero terminator */ + entry->d_namlen = n - 1; + + /* File attributes */ + attr = datap->dwFileAttributes; + if ((attr & FILE_ATTRIBUTE_DEVICE) != 0) { + entry->d_type = DT_CHR; + } else if ((attr & FILE_ATTRIBUTE_DIRECTORY) != 0) { + entry->d_type = DT_DIR; + } else { + entry->d_type = DT_REG; + } + + /* Reset dummy fields */ + entry->d_ino = 0; + entry->d_off = 0; + entry->d_reclen = sizeof (struct dirent); + + } else { + + /* + * Cannot convert file name to multi-byte string so construct + * an erroneous directory entry and return that. Note that + * we cannot return NULL as that would stop the processing + * of directory entries completely. + */ + entry->d_name[0] = '?'; + entry->d_name[1] = '\0'; + entry->d_namlen = 1; + entry->d_type = DT_UNKNOWN; + entry->d_ino = 0; + entry->d_off = -1; + entry->d_reclen = 0; + + } + + /* Return pointer to directory entry */ + *result = entry; + + } else { + + /* No more directory entries */ + *result = NULL; + + } + + return /*OK*/0; +} + +/* + * Close directory stream. + */ +static int +closedir( + DIR *dirp) +{ + int ok; + if (dirp) { + + /* Close wide-character directory stream */ + ok = _wclosedir (dirp->wdirp); + dirp->wdirp = NULL; + + /* Release multi-byte character version */ + free (dirp); + + } else { + + /* Invalid directory stream */ + dirent_set_errno (EBADF); + ok = /*failure*/-1; + + } + return ok; +} + +/* + * Rewind directory stream to beginning. + */ +static void +rewinddir( + DIR* dirp) +{ + /* Rewind wide-character string directory stream */ + _wrewinddir (dirp->wdirp); +} + +/* + * Scan directory for entries. + */ +static int +scandir( + const char *dirname, + struct dirent ***namelist, + int (*filter)(const struct dirent*), + int (*compare)(const struct dirent**, const struct dirent**)) +{ + struct dirent **files = NULL; + size_t size = 0; + size_t allocated = 0; + const size_t init_size = 1; + DIR *dir = NULL; + struct dirent *entry; + struct dirent *tmp = NULL; + size_t i; + int result = 0; + + /* Open directory stream */ + dir = opendir (dirname); + if (dir) { + + /* Read directory entries to memory */ + while (1) { + + /* Enlarge pointer table to make room for another pointer */ + if (size >= allocated) { + void *p; + size_t num_entries; + + /* Compute number of entries in the enlarged pointer table */ + if (size < init_size) { + /* Allocate initial pointer table */ + num_entries = init_size; + } else { + /* Double the size */ + num_entries = size * 2; + } + + /* Allocate first pointer table or enlarge existing table */ + p = realloc (files, sizeof (void*) * num_entries); + if (p != NULL) { + /* Got the memory */ + files = (dirent**) p; + allocated = num_entries; + } else { + /* Out of memory */ + result = -1; + break; + } + + } + + /* Allocate room for temporary directory entry */ + if (tmp == NULL) { + tmp = (struct dirent*) malloc (sizeof (struct dirent)); + if (tmp == NULL) { + /* Cannot allocate temporary directory entry */ + result = -1; + break; + } + } + + /* Read directory entry to temporary area */ + if (readdir_r (dir, tmp, &entry) == /*OK*/0) { + + /* Did we get an entry? */ + if (entry != NULL) { + int pass; + + /* Determine whether to include the entry in result */ + if (filter) { + /* Let the filter function decide */ + pass = filter (tmp); + } else { + /* No filter function, include everything */ + pass = 1; + } + + if (pass) { + /* Store the temporary entry to pointer table */ + files[size++] = tmp; + tmp = NULL; + + /* Keep up with the number of files */ + result++; + } + + } else { + + /* + * End of directory stream reached => sort entries and + * exit. + */ + qsort (files, size, sizeof (void*), + (int (*) (const void*, const void*)) compare); + break; + + } + + } else { + /* Error reading directory entry */ + result = /*Error*/ -1; + break; + } + + } + + } else { + /* Cannot open directory */ + result = /*Error*/ -1; + } + + /* Release temporary directory entry */ + free (tmp); + + /* Release allocated memory on error */ + if (result < 0) { + for (i = 0; i < size; i++) { + free (files[i]); + } + free (files); + files = NULL; + } + + /* Close directory stream */ + if (dir) { + closedir (dir); + } + + /* Pass pointer table to caller */ + if (namelist) { + *namelist = files; + } + return result; +} + +/* Alphabetical sorting */ +static int +alphasort( + const struct dirent **a, const struct dirent **b) +{ + return strcoll ((*a)->d_name, (*b)->d_name); +} + +/* Sort versions */ +static int +versionsort( + const struct dirent **a, const struct dirent **b) +{ + /* FIXME: implement strverscmp and use that */ + return alphasort (a, b); +} + +/* Convert multi-byte string to wide character string */ +static int +dirent_mbstowcs_s( + size_t *pReturnValue, + wchar_t *wcstr, + size_t sizeInWords, + const char *mbstr, + size_t count) +{ + int error; + +#if defined(_MSC_VER) && _MSC_VER >= 1400 + + /* Microsoft Visual Studio 2005 or later */ + error = mbstowcs_s (pReturnValue, wcstr, sizeInWords, mbstr, count); + +#else + + /* Older Visual Studio or non-Microsoft compiler */ + size_t n; + + /* Convert to wide-character string (or count characters) */ + n = mbstowcs (wcstr, mbstr, sizeInWords); + if (!wcstr || n < count) { + + /* Zero-terminate output buffer */ + if (wcstr && sizeInWords) { + if (n >= sizeInWords) { + n = sizeInWords - 1; + } + wcstr[n] = 0; + } + + /* Length of resulting multi-byte string WITH zero terminator */ + if (pReturnValue) { + *pReturnValue = n + 1; + } + + /* Success */ + error = 0; + + } else { + + /* Could not convert string */ + error = 1; + + } + +#endif + return error; +} + +/* Convert wide-character string to multi-byte string */ +static int +dirent_wcstombs_s( + size_t *pReturnValue, + char *mbstr, + size_t sizeInBytes, /* max size of mbstr */ + const wchar_t *wcstr, + size_t count) +{ + int error; + +#if defined(_MSC_VER) && _MSC_VER >= 1400 + + /* Microsoft Visual Studio 2005 or later */ + error = wcstombs_s (pReturnValue, mbstr, sizeInBytes, wcstr, count); + +#else + + /* Older Visual Studio or non-Microsoft compiler */ + size_t n; + + /* Convert to multi-byte string (or count the number of bytes needed) */ + n = wcstombs (mbstr, wcstr, sizeInBytes); + if (!mbstr || n < count) { + + /* Zero-terminate output buffer */ + if (mbstr && sizeInBytes) { + if (n >= sizeInBytes) { + n = sizeInBytes - 1; + } + mbstr[n] = '\0'; + } + + /* Length of resulting multi-bytes string WITH zero-terminator */ + if (pReturnValue) { + *pReturnValue = n + 1; + } + + /* Success */ + error = 0; + + } else { + + /* Cannot convert string */ + error = 1; + + } + +#endif + return error; +} + +/* Set errno variable */ +static void +dirent_set_errno( + int error) +{ +#if defined(_MSC_VER) && _MSC_VER >= 1400 + + /* Microsoft Visual Studio 2005 and later */ + _set_errno (error); + +#else + + /* Non-Microsoft compiler or older Microsoft compiler */ + errno = error; + +#endif +} + + +#ifdef __cplusplus +} +#endif +#endif /*DIRENT_H*/ diff --git a/swin-transformer/semantic-segmentation/layerNorm.cu b/swin-transformer/semantic-segmentation/layerNorm.cu new file mode 100644 index 0000000..2f1a200 --- /dev/null +++ b/swin-transformer/semantic-segmentation/layerNorm.cu @@ -0,0 +1,195 @@ +#include +#include "layerNorm.h" +#include "utilsn.h" +#include +#include + + + + + +namespace nvinfer1 +{ + +layernorm::layernorm() +{ +} +layernorm::~layernorm() +{ + +} +layernorm::layernorm(const void* data, size_t length) +{ + const char *d = reinterpret_cast(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(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 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; +} + + + + + + +} diff --git a/swin-transformer/semantic-segmentation/layerNorm.h b/swin-transformer/semantic-segmentation/layerNorm.h new file mode 100644 index 0000000..ff5b13b --- /dev/null +++ b/swin-transformer/semantic-segmentation/layerNorm.h @@ -0,0 +1,130 @@ +#ifndef LAYERNORM_H +#define LAYERNORM_H + +#include +#include +#include +#include +#include +#include +#include +#include + +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 mPluginAttributes; + +}; + +REGISTER_TENSORRT_PLUGIN(layernormCreator); +}; + +#endif // LAYERNORM_H diff --git a/swin-transformer/semantic-segmentation/logging.h b/swin-transformer/semantic-segmentation/logging.h new file mode 100644 index 0000000..66d6c1f --- /dev/null +++ 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 +#include +#include +#include +#include +#include +#include + +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(×tamp); + 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 diff --git a/swin-transformer/semantic-segmentation/main.cpp b/swin-transformer/semantic-segmentation/main.cpp new file mode 100644 index 0000000..0cb6b87 --- /dev/null +++ b/swin-transformer/semantic-segmentation/main.cpp @@ -0,0 +1,9 @@ +#include + + +using namespace std; + + + + + diff --git a/swin-transformer/semantic-segmentation/myhpp.h b/swin-transformer/semantic-segmentation/myhpp.h new file mode 100644 index 0000000..c26750b --- /dev/null +++ b/swin-transformer/semantic-segmentation/myhpp.h @@ -0,0 +1,36 @@ +#ifndef MYHPP_H +#define MYHPP_H + +#include +#include +#include +#include +#define _USE_MATH_DEFINES +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +//#include +#include +#include +//#include +#include +#include +#include +#include +#include +#include +#include + + + +#endif // MYHPP_H diff --git a/swin-transformer/semantic-segmentation/trainsform.cpp b/swin-transformer/semantic-segmentation/trainsform.cpp new file mode 100644 index 0000000..61eb183 --- /dev/null +++ b/swin-transformer/semantic-segmentation/trainsform.cpp @@ -0,0 +1,325 @@ +#include "common.hpp" +#include "logging.h" +#include +#include +#include +#include +#include +#include + +#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 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 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"<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(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 testVal; + map dataProb; + vector imgs; + cv::Mat img; + //string pattern_dir = "G:/shaj/trainsform"; + string pattern_dir = argv[3]; + string pattern = pattern_dir+ "/*.bmp"; + vector 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(end_time - start_time).count(); + cout << "time:" << duration << endl; + //} + // for(int i = 0; i < 100; i++) + // cout<(h, w)[0] = out[i] * 10; + dst.at(h, w)[1] = out[i] * 30; + dst.at(h, w)[2] = out[i] * 40; + } + } + //cout<destroy(); + m_engine->destroy(); + for (auto bindings: m_bindings) { + cudaFree(bindings); + } + cudaFree(m_cudaStream); + + cout << "swin_transform" << endl; + return 0; +} + diff --git a/swin-transformer/semantic-segmentation/utilsn.h b/swin-transformer/semantic-segmentation/utilsn.h new file mode 100644 index 0000000..336a81b --- /dev/null +++ b/swin-transformer/semantic-segmentation/utilsn.h @@ -0,0 +1,90 @@ +#ifndef UTILSN_H +#define UTILSN_H + +#include +#include +#include +#include +#include +#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 Record; + std::vector 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 + void write(char*& buffer, const T& val) + { + *reinterpret_cast(buffer) = val; + buffer += sizeof(T); + } + + template + void read(const char*& buffer, T& val) + { + val = *reinterpret_cast(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