update readme

Author: Wish

.vscode/tasks.json

@@ -7,6 +7,9 @@
             "label": "build",
             "type": "shell",
             "command": "make pro -j25"
+
+            // for cmake
+            //"command": "cd build && make pro -j25"
         }
     ]
 }

README.md

@@ -76,31 +76,7 @@ auto box = engine->commit(image).get();
 ## 效果图
 ![](workspace/yq.jpg)
 
-## YoloV5-ONNX推理支持-第一种，使用提供的onnx
-- 这个yolov5m.onnx模型使用官方最新版本直接导出得到
-- CMake
-    - 在CMakeLists.txt中配置依赖路径tensorRT、cuda、cudnn、protobuf
-    ```bash
-    git clone [email protected]:shouxieai/tensorRT_cpp.git
-    cd tensorRT_cpp
-
-    mkdir build
-    cd build
-    cmake ..
-    make yolo -j32
-
-    # 或者make alphapose -j32
-    ```
-
-- Makefile
-    - 在Makefile中配置好依赖的tensorRT、cuda、cudnn、protobuf
-    ```bash
-    git clone [email protected]:shouxieai/tensorRT_cpp.git
-    cd tensorRT_cpp
-    make yolo -j32
-    ```
-
-## YoloV5-ONNX推理支持-第二种，自行从官方导出onnx
+## YoloV5支持
 - yolov5的onnx，你的pytorch版本>=1.7时，导出的onnx模型可以直接被当前框架所使用
 - 你的pytorch版本低于1.7时，或者对于yolov5其他版本（2.0、3.0、4.0），可以对opset进行简单改动后直接被框架所支持
 - 如果你想实现低版本pytorch的tensorRT推理、动态batchsize等更多更高级的问题，请打开我们[博客地址](http://zifuture.com:8090)后找到二维码进群交流
@@ -137,7 +113,7 @@ torch.onnx.export(dynamic_axes={'images': {0: 'batch'},  # shape(1,3,640,640)
 3. 导出onnx模型
 ```bash
 cd yolov5
-python export.py --weights=yolov5s.pt --dynamic --opset=11
+python export.py --weights=yolov5s.pt --dynamic --include=onnx --opset=11
 ```
 4. 复制模型并执行
 ```bash

src/application/app_lesson.cpp

@@ -3,9 +3,64 @@
 #include <infer/trt_infer.hpp>
 #include <builder/trt_builder.hpp>
 #include "app_yolo/yolo.hpp"
+#include <cuda_runtime.h>
 
 using namespace std;
 
+// static void test_tensor1(){
+
+//     size_t cpu_bytes = 1024;
+//     size_t gpu_bytes = 2048;
+
+//     ///////////////////////////////////////////////////////////////////
+//     // 封装效果，自动分配和释放
+//     TRT::MixMemory memory;
+//     void* host_ptr   = memory.cpu(cpu_bytes);
+//     void* device_ptr = memory.gpu(gpu_bytes);
+
+//     ///////////////////////////////////////////////////////////////////
+//     // 不封装效果
+//     void* host_ptr   = nullptr;
+//     void* device_ptr = nullptr;
+//     cudaMallocHost(&host_ptr, cpu_bytes);
+//     cudaMalloc(&device_ptr, gpu_bytes);
+
+//     cudaFreeHost(&host_ptr);
+//     cudaFree(&device_ptr);
+//     ///////////////////////////////////////////////////////////////////
+// }
+
+static void test_tensor2(){
+
+    ///////////////////////////////////////////////////////////////////
+    // 内存的自动复制，依靠head属性标记数据最新的位置
+    // 若访问的数据不是最新的，则会自动发生复制操作
+    TRT::Tensor tensor({1, 3, 5, 5}, TRT::DataType::Float);
+    INFO("tensor.head = %s", TRT::data_head_string(tensor.head()));   // 输出 Init，内存没有分配
+
+    tensor.cpu<float>()[0] = 512;               // 访问cpu时，分配cpu内存
+    INFO("tensor.head = %s", TRT::data_head_string(tensor.head()));   // 输出 Host
+
+    float* device_ptr = tensor.gpu<float>();    // 访问gpu时，最新数据在Host，发生复制动作并标记最新数据在Device
+    INFO("tensor.head = %s", TRT::data_head_string(tensor.head()));   // 输出 Device
+    INFO("device_ptr[0] = %f", device_ptr[0]);                        // 输出 512.00000
+}
+
+static void test_tensor3(){
+
+    ///////////////////////////////////////////////////////////////////
+    // 计算维度的偏移量
+    TRT::Tensor tensor({1, 3, 5, 5, 2, 5}, TRT::DataType::Float);
+    auto ptr_origin   = tensor.cpu<float>();
+    auto ptr_channel2 = tensor.cpu<float>(0, 2, 3, 2, 1, 3);
+
+    INFO("Offset = %d", ptr_channel2 - ptr_origin);                          // 输出678
+    INFO("Offset = %d", tensor.offset(0, 2, 3, 2, 1, 3));                    // 输出678
+
+    int offset_compute = ((((0 * 3 + 2) * 5 + 3) * 5 + 2) * 2 + 1) * 5 + 3;  
+    INFO("Compute = %d", offset_compute);                                    // 输出678
+}
+
 static void lesson1(){
 
     /** 模型编译，onnx到trtmodel **/
@@ -180,19 +235,11 @@ void lesson_cache1frame(){
 int app_lesson(){
 
     iLogger::set_log_level(iLogger::LogLevel::Verbose);
-    lesson1();
+    test_tensor3();
+    // lesson1();
     // lesson2();
     // lesson3();
     // lesson_cache1frame();
     return 0;
 }
 
-
-
-
-
-
-
-
-
- 

src/application/app_yolo.cpp

@@ -169,11 +169,47 @@ static void test(Yolo::Type type, TRT::Mode mode, const string& model){
     forward_engine(model_file, type);
 }
 
+void my_test(){
+
+    TRT::compile(
+        TRT::Mode::FP32,
+        5,
+        "/data/sxai/temp/yolov5-5.0/yolov5s.onnx",
+        "my-yolov5-5.0s.trtmodel"
+    );
+    INFO("Done");
+
+    auto yolo = Yolo::create_infer(
+        "my-yolov5-5.0s.trtmodel", 
+        Yolo::Type::V5,
+        0, 0.25f, 0.5f
+    );
+
+    auto image = cv::imread("/data/sxai/tensorRT/workspace/inference/car.jpg");
+    auto bboxes = yolo->commits({image, image})[1].get();
+
+    for(auto& box : bboxes){
+
+        uint8_t r, g, b;
+        tie(r, g, b) = iLogger::random_color(box.class_label);
+
+        cv::rectangle(
+            image, 
+            cv::Point(box.left, box.top),
+            cv::Point(box.right, box.bottom),
+            cv::Scalar(b, g, r),
+            3
+        );
+    }
+    cv::imwrite("my-yolov5s-car.jpg", image);
+}
+
 int app_yolo(){
 
+    // my_test();
     //iLogger::set_log_level(iLogger::LogLevel::Info);
     test(Yolo::Type::X, TRT::Mode::FP32, "yolox_m");
-    //test(Yolo::Type::V5, TRT::Mode::FP32, "yolov5s");
+    // test(Yolo::Type::V5, TRT::Mode::FP32, "yolov5s");
     // test(Yolo::Type::X, TRT::Mode::FP16, "yolox_s");
     // test(Yolo::Type::V5, TRT::Mode::FP16, "yolov5s");
     // test_int8(Yolo::Type::X, "yolox_s");

src/tensorRT/common/trt_tensor.hpp

@@ -79,7 +79,7 @@ namespace TRT {
         virtual ~Tensor();
 
         int numel() const;
-        int ndims() const{return shape_.size();}
+        inline int ndims() const{return shape_.size();}
         inline int size(int index)  const{return shape_[index];}
         inline int shape(int index) const{return shape_[index];}
 

src/tensorRT/onnxplugin/plugins/DCNv2.cu

@@ -229,8 +229,8 @@ __global__ void DCNIm2colKernel(
 	}
 }
 
-template<typename _T>
-static __global__ void biasKernel(_T* data_input, const _T* bias, const int f_area, int edge) {
+template<typename DataType>
+static __global__ void biasKernel(DataType* data_input, const DataType* bias, const int f_area, int edge) {
 
 	KernelPositionBlock;
 	int bias_index = position / f_area;
@@ -276,7 +276,7 @@ inline void segemm_native(cublasHandle_t handle,
 	cublasCheck(cublasGemmEx(handle, transa, transb, m, n, k, &halpha, A, CUDA_R_16F, lda, B, CUDA_R_16F, ldb, &hbeta, C, CUDA_R_16F, ldc, CUDA_R_16F, CUBLAS_GEMM_DFALT));
 }
 
-template<typename _T>
+template<typename DataType>
 static void enqueue_native(cublasHandle_t handle, const std::vector<GTensor>& inputs, std::vector<GTensor>& outputs, const std::vector<GTensor>& weights, void* workspace, cudaStream_t stream) {
 	auto& data = inputs[0];
 	auto& om = inputs[1];
@@ -295,16 +295,16 @@ static void enqueue_native(cublasHandle_t handle, const std::vector<GTensor>& in
 
 	cublasCheck(cublasSetStream(handle, stream));
 	for (int ibatch = 0; ibatch < data.batch(); ++ibatch) {
-		_T* maskWorkspacePtr = (_T*)workspace + (maskSize + im2colSize) * ibatch;
-		_T* im2colWorkspacePtr = (_T*)workspace + (maskSize + im2colSize) * ibatch + maskSize;
+		DataType* maskWorkspacePtr = (DataType*)workspace + (maskSize + im2colSize) * ibatch;
+		DataType* im2colWorkspacePtr = (DataType*)workspace + (maskSize + im2colSize) * ibatch + maskSize;
 
-		_T* inputMask = om.ptr<_T>(ibatch, om.channel() / 3 * 2);
+		DataType* inputMask = om.ptr<DataType>(ibatch, om.channel() / 3 * 2);
 		checkCudaKernel(
 			sigmoidKernel<<<CUDATools::grid_dims(maskSize), CUDATools::block_dims(maskSize), 0, stream>>>(inputMask, maskWorkspacePtr, maskSize);
 		);
 
-		_T* datainput = data.ptr<_T>(ibatch);
-		_T* offset = om.ptr<_T>(ibatch);
+		DataType* datainput = data.ptr<DataType>(ibatch);
+		DataType* offset = om.ptr<DataType>(ibatch);
 
 		auto jobs = (size_t)data.channel() * out.height() * out.width();
 		checkCudaKernel(
@@ -314,17 +314,17 @@ static void enqueue_native(cublasHandle_t handle, const std::vector<GTensor>& in
 			);
 		);
 
-		_T* weightKernel = weights[0].ptr<_T>();
-		segemm_native(handle, CUBLAS_OP_N, CUBLAS_OP_N, n, m, k, alpha, im2colWorkspacePtr, n, weightKernel, k, beta, out.ptr<_T>(ibatch), n);
+		DataType* weightKernel = weights[0].ptr<DataType>();
+		segemm_native(handle, CUBLAS_OP_N, CUBLAS_OP_N, n, m, k, alpha, im2colWorkspacePtr, n, weightKernel, k, beta, out.ptr<DataType>(ibatch), n);
 
 		if (weights.size() > 1) {
-			_T* weightBias = weights[1].ptr<_T>();
+			DataType* weightBias = weights[1].ptr<DataType>();
 			size_t edge = out.count(1);
 			size_t area = out.count(2);
 
 			checkCudaKernel(
 				biasKernel<<<CUDATools::grid_dims(edge), CUDATools::block_dims(edge), 0, stream>>>(
-					out.ptr<_T>(ibatch), weightBias, area, edge
+					out.ptr<DataType>(ibatch), weightBias, area, edge
 				);
 			);
 		}