resnet merge

CMakeLists.txt

@@ -89,8 +89,24 @@ target_link_libraries(test_yolo3_coco4 tkDNN)
 add_executable(test_yolo3_berkeley tests/yolo3_berkeley/yolo3_berkeley.cpp)
 target_link_libraries(test_yolo3_berkeley tkDNN)
 
+<<<<<<< HEAD
 add_executable(test_yolo3_flir tests/yolo3_flir/yolo3_flir.cpp)
 target_link_libraries(test_yolo3_flir tkDNN)
+=======
+add_executable(test_yolo3_tetrapack tests/yolo3_tetrapack/yolo3_tetrapack.cpp)
+target_link_libraries(test_yolo3_tetrapack tkDNN)
+
+add_executable(test_yolo3_tetrapack_resize tests/yolo3_tetrapack_resize/yolo3_tetrapack_resize.cpp)
+target_link_libraries(test_yolo3_tetrapack_resize tkDNN)
+
+add_executable(test_yolo3_BCDS6 tests/yolo3_BCDS6/yolo3_BCDS6.cpp)
+target_link_libraries(test_yolo3_BCDS6 tkDNN)
+
+add_executable(test_resnet101 tests/resnet101/resnet101.cpp)
+target_link_libraries(test_resnet101 tkDNN)
+
+
+>>>>>>> e1d0fa3... resnet first try
 ################################################################################
 
 

include/tkDNN/Layer.h

@@ -227,7 +227,9 @@ public:
     int paddingH, paddingW;
 
     Pooling(Network *net, int winH, int winW, 
-            int strideH, int strideW, tkdnnPoolingMode_t pool_mode); 
+            int strideH, int strideW, 
+            int paddingH = 0, int paddingW = 0,
+            tkdnnPoolingMode_t pool_mode = POOLING_MAX); 
     virtual ~Pooling();
     virtual layerType_t getLayerType() { return LAYER_POOLING; };
 

src/Pooling.cpp

@@ -6,6 +6,7 @@
 namespace tk { namespace dnn {
 
 Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW,
+                  int paddingH, int paddingW,
                   tkdnnPoolingMode_t pool_mode) : 
     Layer(net) {
 
@@ -14,8 +15,8 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW,
     this->strideH = strideH;
     this->strideW = strideW;
     this->pool_mode = pool_mode;
-    this->paddingH = 0;
-    this->paddingW = 0;
+    this->paddingH = paddingH;
+    this->paddingW = paddingW;
 
     checkCUDNN( cudnnCreatePoolingDescriptor(&poolingDesc) );
 
@@ -38,7 +39,7 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW,
     }
 
     checkCUDNN( cudnnSetPooling2dDescriptor(poolingDesc, cudnnPoolingMode_t(pool_mode),
-                CUDNN_NOT_PROPAGATE_NAN, winH, winW, 0, 0, strideH, strideW) );
+                CUDNN_NOT_PROPAGATE_NAN, winH, winW, paddingH, paddingW, strideH, strideW) );
 
     checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc, 
                 net->tensorFormat, net->dataType, n, c, h, w) );

tests/resnet101/resnet101.cpp

@@ -0,0 +1,398 @@
+#include <iostream>
+#include "tkdnn.h"
+
+const char *input_bin = "../tests/resnet101/debug/input.bin";
+const char *conv1_bin = "../tests/resnet101/layers/conv1.bin";
+
+//layer1
+const char *layer1_0_conv1_bin = "../tests/resnet101/layers/layer1-0-conv1.bin";
+const char *layer1_0_conv2_bin = "../tests/resnet101/layers/layer1-0-conv2.bin";
+const char *layer1_0_conv3_bin = "../tests/resnet101/layers/layer1-0-conv3.bin";
+const char *layer1_0_downsample_0_bin = "../tests/resnet101/layers/layer1-0-downsample-0.bin";
+
+const char *layer1_1_conv1_bin = "../tests/resnet101/layers/layer1-1-conv1.bin";
+const char *layer1_1_conv2_bin = "../tests/resnet101/layers/layer1-1-conv2.bin";
+const char *layer1_1_conv3_bin = "../tests/resnet101/layers/layer1-1-conv3.bin";
+
+const char *layer1_2_conv1_bin = "../tests/resnet101/layers/layer1-2-conv1.bin";
+const char *layer1_2_conv2_bin = "../tests/resnet101/layers/layer1-2-conv2.bin";
+const char *layer1_2_conv3_bin = "../tests/resnet101/layers/layer1-2-conv3.bin";
+
+//layer2
+const char *layer2_0_conv1_bin = "../tests/resnet101/layers/layer2-0-conv1.bin";
+const char *layer2_0_conv2_bin = "../tests/resnet101/layers/layer2-0-conv2.bin";
+const char *layer2_0_conv3_bin = "../tests/resnet101/layers/layer2-0-conv3.bin";
+const char *layer2_0_downsample_0_bin = "../tests/resnet101/layers/layer2-0-downsample-0.bin";
+
+const char *layer2_1_conv1_bin = "../tests/resnet101/layers/layer2-1-conv1.bin";
+const char *layer2_1_conv2_bin = "../tests/resnet101/layers/layer2-1-conv2.bin";
+const char *layer2_1_conv3_bin = "../tests/resnet101/layers/layer2-1-conv3.bin";
+
+const char *layer2_2_conv1_bin = "../tests/resnet101/layers/layer2-2-conv1.bin";
+const char *layer2_2_conv2_bin = "../tests/resnet101/layers/layer2-2-conv2.bin";
+const char *layer2_2_conv3_bin = "../tests/resnet101/layers/layer2-2-conv3.bin";
+
+const char *layer2_3_conv1_bin = "../tests/resnet101/layers/layer2-3-conv1.bin";
+const char *layer2_3_conv2_bin = "../tests/resnet101/layers/layer2-3-conv2.bin";
+const char *layer2_3_conv3_bin = "../tests/resnet101/layers/layer2-3-conv3.bin";
+
+//layer3
+const char *layer3_0_conv1_bin = "../tests/resnet101/layers/layer3-0-conv1.bin";
+const char *layer3_0_conv2_bin = "../tests/resnet101/layers/layer3-0-conv2.bin";
+const char *layer3_0_conv3_bin = "../tests/resnet101/layers/layer3-0-conv3.bin";
+const char *layer3_0_downsample_0_bin = "../tests/resnet101/layers/layer3-0-downsample-0.bin";
+
+const char *layer3_1_conv1_bin = "../tests/resnet101/layers/layer3-1-conv1.bin";
+const char *layer3_1_conv2_bin = "../tests/resnet101/layers/layer3-1-conv2.bin";
+const char *layer3_1_conv3_bin = "../tests/resnet101/layers/layer3-1-conv3.bin";
+
+const char *layer3_2_conv1_bin = "../tests/resnet101/layers/layer3-2-conv1.bin";
+const char *layer3_2_conv2_bin = "../tests/resnet101/layers/layer3-2-conv2.bin";
+const char *layer3_2_conv3_bin = "../tests/resnet101/layers/layer3-2-conv3.bin";
+
+const char *layer3_3_conv1_bin = "../tests/resnet101/layers/layer3-3-conv1.bin";
+const char *layer3_3_conv2_bin = "../tests/resnet101/layers/layer3-3-conv2.bin";
+const char *layer3_3_conv3_bin = "../tests/resnet101/layers/layer3-3-conv3.bin";
+
+const char *layer3_4_conv1_bin = "../tests/resnet101/layers/layer3-4-conv1.bin";
+const char *layer3_4_conv2_bin = "../tests/resnet101/layers/layer3-4-conv2.bin";
+const char *layer3_4_conv3_bin = "../tests/resnet101/layers/layer3-4-conv3.bin";
+
+const char *layer3_5_conv1_bin = "../tests/resnet101/layers/layer3-5-conv1.bin";
+const char *layer3_5_conv2_bin = "../tests/resnet101/layers/layer3-5-conv2.bin";
+const char *layer3_5_conv3_bin = "../tests/resnet101/layers/layer3-5-conv3.bin";
+
+const char *layer3_6_conv1_bin = "../tests/resnet101/layers/layer3-6-conv1.bin";
+const char *layer3_6_conv2_bin = "../tests/resnet101/layers/layer3-6-conv2.bin";
+const char *layer3_6_conv3_bin = "../tests/resnet101/layers/layer3-6-conv3.bin";
+
+const char *layer3_7_conv1_bin = "../tests/resnet101/layers/layer3-7-conv1.bin";
+const char *layer3_7_conv2_bin = "../tests/resnet101/layers/layer3-7-conv2.bin";
+const char *layer3_7_conv3_bin = "../tests/resnet101/layers/layer3-7-conv3.bin";
+
+const char *layer3_8_conv1_bin = "../tests/resnet101/layers/layer3-8-conv1.bin";
+const char *layer3_8_conv2_bin = "../tests/resnet101/layers/layer3-8-conv2.bin";
+const char *layer3_8_conv3_bin = "../tests/resnet101/layers/layer3-8-conv3.bin";
+
+const char *layer3_9_conv1_bin = "../tests/resnet101/layers/layer3-9-conv1.bin";
+const char *layer3_9_conv2_bin = "../tests/resnet101/layers/layer3-9-conv2.bin";
+const char *layer3_9_conv3_bin = "../tests/resnet101/layers/layer3-9-conv3.bin";
+
+const char *layer3_10_conv1_bin = "../tests/resnet101/layers/layer3-10-conv1.bin";
+const char *layer3_10_conv2_bin = "../tests/resnet101/layers/layer3-10-conv2.bin";
+const char *layer3_10_conv3_bin = "../tests/resnet101/layers/layer3-10-conv3.bin";
+
+const char *layer3_11_conv1_bin = "../tests/resnet101/layers/layer3-11-conv1.bin";
+const char *layer3_11_conv2_bin = "../tests/resnet101/layers/layer3-11-conv2.bin";
+const char *layer3_11_conv3_bin = "../tests/resnet101/layers/layer3-11-conv3.bin";
+
+const char *layer3_12_conv1_bin = "../tests/resnet101/layers/layer3-12-conv1.bin";
+const char *layer3_12_conv2_bin = "../tests/resnet101/layers/layer3-12-conv2.bin";
+const char *layer3_12_conv3_bin = "../tests/resnet101/layers/layer3-12-conv3.bin";
+
+const char *layer3_13_conv1_bin = "../tests/resnet101/layers/layer3-13-conv1.bin";
+const char *layer3_13_conv2_bin = "../tests/resnet101/layers/layer3-13-conv2.bin";
+const char *layer3_13_conv3_bin = "../tests/resnet101/layers/layer3-13-conv3.bin";
+
+const char *layer3_14_conv1_bin = "../tests/resnet101/layers/layer3-14-conv1.bin";
+const char *layer3_14_conv2_bin = "../tests/resnet101/layers/layer3-14-conv2.bin";
+const char *layer3_14_conv3_bin = "../tests/resnet101/layers/layer3-14-conv3.bin";
+
+const char *layer3_15_conv1_bin = "../tests/resnet101/layers/layer3-15-conv1.bin";
+const char *layer3_15_conv2_bin = "../tests/resnet101/layers/layer3-15-conv2.bin";
+const char *layer3_15_conv3_bin = "../tests/resnet101/layers/layer3-15-conv3.bin";
+
+const char *layer3_16_conv1_bin = "../tests/resnet101/layers/layer3-16-conv1.bin";
+const char *layer3_16_conv2_bin = "../tests/resnet101/layers/layer3-16-conv2.bin";
+const char *layer3_16_conv3_bin = "../tests/resnet101/layers/layer3-16-conv3.bin";
+
+const char *layer3_17_conv1_bin = "../tests/resnet101/layers/layer3-17-conv1.bin";
+const char *layer3_17_conv2_bin = "../tests/resnet101/layers/layer3-17-conv2.bin";
+const char *layer3_17_conv3_bin = "../tests/resnet101/layers/layer3-17-conv3.bin";
+
+const char *layer3_18_conv1_bin = "../tests/resnet101/layers/layer3-18-conv1.bin";
+const char *layer3_18_conv2_bin = "../tests/resnet101/layers/layer3-18-conv2.bin";
+const char *layer3_18_conv3_bin = "../tests/resnet101/layers/layer3-18-conv3.bin";
+
+const char *layer3_19_conv1_bin = "../tests/resnet101/layers/layer3-19-conv1.bin";
+const char *layer3_19_conv2_bin = "../tests/resnet101/layers/layer3-19-conv2.bin";
+const char *layer3_19_conv3_bin = "../tests/resnet101/layers/layer3-19-conv3.bin";
+
+const char *layer3_20_conv1_bin = "../tests/resnet101/layers/layer3-20-conv1.bin";
+const char *layer3_20_conv2_bin = "../tests/resnet101/layers/layer3-20-conv2.bin";
+const char *layer3_20_conv3_bin = "../tests/resnet101/layers/layer3-20-conv3.bin";
+
+const char *layer3_21_conv1_bin = "../tests/resnet101/layers/layer3-21-conv1.bin";
+const char *layer3_21_conv2_bin = "../tests/resnet101/layers/layer3-21-conv2.bin";
+const char *layer3_21_conv3_bin = "../tests/resnet101/layers/layer3-21-conv3.bin";
+
+const char *layer3_22_conv1_bin = "../tests/resnet101/layers/layer3-22-conv1.bin";
+const char *layer3_22_conv2_bin = "../tests/resnet101/layers/layer3-22-conv2.bin";
+const char *layer3_22_conv3_bin = "../tests/resnet101/layers/layer3-22-conv3.bin";
+
+
+//layer4
+const char *layer4_0_conv1_bin = "../tests/resnet101/layers/layer4-0-conv1.bin";
+const char *layer4_0_conv2_bin = "../tests/resnet101/layers/layer4-0-conv2.bin";
+const char *layer4_0_conv3_bin = "../tests/resnet101/layers/layer4-0-conv3.bin";
+const char *layer4_0_downsample_0_bin = "../tests/resnet101/layers/layer4-0-downsample-0.bin";
+
+const char *layer4_1_conv1_bin = "../tests/resnet101/layers/layer4-1-conv1.bin";
+const char *layer4_1_conv2_bin = "../tests/resnet101/layers/layer4-1-conv2.bin";
+const char *layer4_1_conv3_bin = "../tests/resnet101/layers/layer4-1-conv3.bin";
+
+const char *layer4_2_conv1_bin = "../tests/resnet101/layers/layer4-2-conv1.bin";
+const char *layer4_2_conv2_bin = "../tests/resnet101/layers/layer4-2-conv2.bin";
+const char *layer4_2_conv3_bin = "../tests/resnet101/layers/layer4-2-conv3.bin";
+
+//final
+const char *fc_bin = "../tests/resnet101/layers/fc.bin";
+
+const char *output_bin = "../tests/resnet101/debug/layer1-0-conv3.bin";
+
+int main()
+{
+
+    // Network layout
+    tk::dnn::dataDim_t dim(1, 3, 224, 224, 1);
+    tk::dnn::Network net(dim);
+
+    tk::dnn::Conv2d conv1(&net, 64, 7, 7, 2, 2, 3, 3, conv1_bin, true);
+    tk::dnn::Activation relu3(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Pooling maxpool4(&net, 3, 3, 2, 2, 1, 1, tk::dnn::POOLING_MAX);
+
+
+    //layer 1
+    tk::dnn::Conv2d layer1_0_conv1(&net, 64, 1, 1, 1, 1, 0, 0, layer1_0_conv1_bin, true);
+    tk::dnn::Conv2d layer1_0_conv2(&net, 64, 3, 3, 1, 1, 1, 1, layer1_0_conv2_bin, true);
+    tk::dnn::Conv2d layer1_0_conv3(&net, 256, 1, 1, 1, 1, 0, 0, layer1_0_conv3_bin, true);
+
+/*
+    tk::dnn::Activation layer1_0_relu(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d layer1_0_downsample_0(&net, 256, 1, 1, 1, 1, 1, 1, layer1_0_downsample_0, true);
+
+    tk::dnn::Conv2d layer1_1_conv1(&net, 64, 1, 1, 1, 1, 1, 1, layer1_1_conv1_bin, true);
+    tk::dnn::Conv2d layer1_1_conv2(&net, 64, 3, 3, 1, 1, 1, 1, layer1_1_conv2_bin, true);
+    tk::dnn::Conv2d layer1_1_conv3(&net, 256, 1, 1, 1, 1, 1, 1, layer1_1_conv3_bin, true);
+    tk::dnn::Activation layer1_1_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer1_2_conv1(&net, 64, 1, 1, 1, 1, 1, 1, layer1_2_conv1_bin, true);
+    tk::dnn::Conv2d layer1_2_conv2(&net, 64, 3, 3, 1, 1, 1, 1, layer1_2_conv2_bin, true);
+    tk::dnn::Conv2d layer1_2_conv3(&net, 256, 1, 1, 1, 1, 1, 1, layer1_2_conv3_bin, true);
+    tk::dnn::Activation layer1_2_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    
+    //layer 2
+    tk::dnn::Conv2d layer2_0_conv1(&net, 128, 1, 1, 1, 1, 1, 1, layer2_0_conv1_bin, true);
+    tk::dnn::Conv2d layer2_0_conv2(&net, 128, 3, 3, 2, 2, 1, 1, layer2_0_conv2_bin, true);
+    tk::dnn::Conv2d layer2_0_conv3(&net, 512, 1, 1, 1, 1, 1, 1, layer2_0_conv3_bin, true);
+    tk::dnn::Activation layer2_0_relu(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d layer2_0_downsample_0(&net, 512, 1, 1, 2, 2, 1, 1, layer2_0_downsample_0, true);
+
+    tk::dnn::Conv2d layer2_1_conv1(&net, 128, 1, 1, 1, 1, 1, 1, layer2_1_conv1_bin, true);
+    tk::dnn::Conv2d layer2_1_conv2(&net, 128, 3, 3, 1, 1, 1, 1, layer2_1_conv2_bin, true);
+    tk::dnn::Conv2d layer2_1_conv3(&net, 512, 1, 1, 1, 1, 1, 1, layer2_1_conv3_bin, true);
+    tk::dnn::Activation layer2_1_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer2_2_conv1(&net, 128, 1, 1, 1, 1, 1, 1, layer2_2_conv1_bin, true);
+    tk::dnn::Conv2d layer2_2_conv2(&net, 128, 3, 3, 1, 1, 1, 1, layer2_2_conv2_bin, true);
+    tk::dnn::Conv2d layer2_2_conv3(&net, 512, 1, 1, 1, 1, 1, 1, layer2_2_conv3_bin, true);
+    tk::dnn::Activation layer2_2_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer2_3_conv1(&net, 128, 1, 1, 1, 1, 1, 1, layer2_3_conv1_bin, true);
+    tk::dnn::Conv2d layer2_3_conv2(&net, 128, 3, 3, 1, 1, 1, 1, layer2_3_conv2_bin, true);
+    tk::dnn::Conv2d layer2_3_conv3(&net, 512, 1, 1, 1, 1, 1, 1, layer2_3_conv3_bin, true);
+    tk::dnn::Activation layer2_3_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    //layer 3
+    tk::dnn::Conv2d layer3_0_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_0_conv1_bin, true);
+    tk::dnn::Conv2d layer3_0_conv2(&net, 256, 3, 3, 2, 2, 1, 1, layer3_0_conv2_bin, true);
+    tk::dnn::Conv2d layer3_0_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_0_conv3_bin, true);
+    tk::dnn::Activation layer3_0_relu(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d layer3_0_downsample_0(&net, 1024, 1, 1, 2, 2, 1, 1, layer3_0_downsample_0, true);
+
+    tk::dnn::Conv2d layer3_1_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_1_conv1_bin, true);
+    tk::dnn::Conv2d layer3_1_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_1_conv2_bin, true);
+    tk::dnn::Conv2d layer3_1_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_1_conv3_bin, true);
+    tk::dnn::Activation layer3_1_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_2_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_2_conv1_bin, true);
+    tk::dnn::Conv2d layer3_2_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_2_conv2_bin, true);
+    tk::dnn::Conv2d layer3_2_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_2_conv3_bin, true);
+    tk::dnn::Activation layer3_2_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_3_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_3_conv1_bin, true);
+    tk::dnn::Conv2d layer3_3_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_3_conv2_bin, true);
+    tk::dnn::Conv2d layer3_3_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_3_conv3_bin, true);
+    tk::dnn::Activation layer3_3_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_4_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_4_conv1_bin, true);
+    tk::dnn::Conv2d layer3_4_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_4_conv2_bin, true);
+    tk::dnn::Conv2d layer3_4_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_4_conv3_bin, true);
+    tk::dnn::Activation layer3_4_relu(&net, CUDNN_ACTIVATION_RELU);
+    
+    tk::dnn::Conv2d layer3_5_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_5_conv1_bin, true);
+    tk::dnn::Conv2d layer3_5_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_5_conv2_bin, true);
+    tk::dnn::Conv2d layer3_5_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_5_conv3_bin, true);
+    tk::dnn::Activation layer3_5_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_6_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_6_conv1_bin, true);
+    tk::dnn::Conv2d layer3_6_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_6_conv2_bin, true);
+    tk::dnn::Conv2d layer3_6_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_6_conv3_bin, true);
+    tk::dnn::Activation layer3_6_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_7_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_7_conv1_bin, true);
+    tk::dnn::Conv2d layer3_7_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_7_conv2_bin, true);
+    tk::dnn::Conv2d layer3_7_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_7_conv3_bin, true);
+    tk::dnn::Activation layer3_7_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_8_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_8_conv1_bin, true);
+    tk::dnn::Conv2d layer3_8_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_8_conv2_bin, true);
+    tk::dnn::Conv2d layer3_8_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_8_conv3_bin, true);
+    tk::dnn::Activation layer3_8_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_9_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_9_conv1_bin, true);
+    tk::dnn::Conv2d layer3_9_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_9_conv2_bin, true);
+    tk::dnn::Conv2d layer3_9_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_9_conv3_bin, true);
+    tk::dnn::Activation layer3_9_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_10_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_10_conv1_bin, true);
+    tk::dnn::Conv2d layer3_10_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_10_conv2_bin, true);
+    tk::dnn::Conv2d layer3_10_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_10_conv3_bin, true);
+    tk::dnn::Activation layer3_10_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_11_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_11_conv1_bin, true);
+    tk::dnn::Conv2d layer3_11_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_11_conv2_bin, true);
+    tk::dnn::Conv2d layer3_11_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_11_conv3_bin, true);
+    tk::dnn::Activation layer3_11_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_12_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_12_conv1_bin, true);
+    tk::dnn::Conv2d layer3_12_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_12_conv2_bin, true);
+    tk::dnn::Conv2d layer3_12_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_12_conv3_bin, true);
+    tk::dnn::Activation layer3_12_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_13_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_13_conv1_bin, true);
+    tk::dnn::Conv2d layer3_13_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_13_conv2_bin, true);
+    tk::dnn::Conv2d layer3_13_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_13_conv3_bin, true);
+    tk::dnn::Activation layer3_13_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_14_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_14_conv1_bin, true);
+    tk::dnn::Conv2d layer3_14_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_14_conv2_bin, true);
+    tk::dnn::Conv2d layer3_14_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_14_conv3_bin, true);
+    tk::dnn::Activation layer3_14_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_15_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_15_conv1_bin, true);
+    tk::dnn::Conv2d layer3_15_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_15_conv2_bin, true);
+    tk::dnn::Conv2d layer3_15_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_15_conv3_bin, true);
+    tk::dnn::Activation layer3_15_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_16_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_16_conv1_bin, true);
+    tk::dnn::Conv2d layer3_16_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_16_conv2_bin, true);
+    tk::dnn::Conv2d layer3_16_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_16_conv3_bin, true);
+    tk::dnn::Activation layer3_16_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_17_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_17_conv1_bin, true);
+    tk::dnn::Conv2d layer3_17_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_17_conv2_bin, true);
+    tk::dnn::Conv2d layer3_17_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_17_conv3_bin, true);
+    tk::dnn::Activation layer3_17_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_18_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_18_conv1_bin, true);
+    tk::dnn::Conv2d layer3_18_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_18_conv2_bin, true);
+    tk::dnn::Conv2d layer3_18_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_18_conv3_bin, true);
+    tk::dnn::Activation layer3_18_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_19_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_19_conv1_bin, true);
+    tk::dnn::Conv2d layer3_19_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_19_conv2_bin, true);
+    tk::dnn::Conv2d layer3_19_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_19_conv3_bin, true);
+    tk::dnn::Activation layer3_19_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_20_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_20_conv1_bin, true);
+    tk::dnn::Conv2d layer3_20_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_20_conv2_bin, true);
+    tk::dnn::Conv2d layer3_20_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_20_conv3_bin, true);
+    tk::dnn::Activation layer3_20_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_21_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_21_conv1_bin, true);
+    tk::dnn::Conv2d layer3_21_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_21_conv2_bin, true);
+    tk::dnn::Conv2d layer3_21_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_21_conv3_bin, true);
+    tk::dnn::Activation layer3_21_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer3_22_conv1(&net, 256, 1, 1, 1, 1, 1, 1, layer3_22_conv1_bin, true);
+    tk::dnn::Conv2d layer3_22_conv2(&net, 256, 3, 3, 1, 1, 1, 1, layer3_22_conv2_bin, true);
+    tk::dnn::Conv2d layer3_22_conv3(&net, 1024, 1, 1, 1, 1, 1, 1, layer3_22_conv3_bin, true);
+    tk::dnn::Activation layer3_22_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    //layer 4
+    tk::dnn::Conv2d layer4_0_conv1(&net, 512, 1, 1, 1, 1, 1, 1, layer4_0_conv1_bin, true);
+    tk::dnn::Conv2d layer4_0_conv2(&net, 512, 3, 3, 2, 2, 1, 1, layer4_0_conv2_bin, true);
+    tk::dnn::Conv2d layer4_0_conv3(&net, 2048, 1, 1, 1, 1, 1, 1, layer4_0_conv3_bin, true);
+    tk::dnn::Activation layer4_0_relu(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d layer4_0_downsample_0(&net, 2048, 1, 1, 2, 2, 1, 1, layer4_0_downsample_0, true);
+
+    tk::dnn::Conv2d layer4_1_conv1(&net, 512, 1, 1, 1, 1, 1, 1, layer4_1_conv1_bin, true);
+    tk::dnn::Conv2d layer4_1_conv2(&net, 512, 3, 3, 1, 1, 1, 1, layer4_1_conv2_bin, true);
+    tk::dnn::Conv2d layer4_1_conv3(&net, 2048, 1, 1, 1, 1, 1, 1, layer4_1_conv3_bin, true);
+    tk::dnn::Activation layer4_1_relu(&net, CUDNN_ACTIVATION_RELU);
+
+    tk::dnn::Conv2d layer4_2_conv1(&net, 512, 1, 1, 1, 1, 1, 1, layer4_2_conv1_bin, true);
+    tk::dnn::Conv2d layer4_2_conv2(&net, 512, 3, 3, 1, 1, 1, 1, layer4_2_conv2_bin, true);
+    tk::dnn::Conv2d layer4_2_conv3(&net, 2048, 1, 1, 1, 1, 1, 1, layer4_2_conv3_bin, true);
+    tk::dnn::Activation layer4_2_relu(&net, CUDNN_ACTIVATION_RELU);
+
+
+    //final
+    tk::dnn::Pooling avgpool(&net, 3, 3, 2, 2, tk::dnn::POOLING_AVERAGE);
+    tk::dnn::Dense   fc(&net, 1000, fc_bin);
+*/
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+
+    //print network model
+    net.print();
+/*
+    //convert network to tensorRT
+    tk::dnn::NetworkRT netRT(&net, "resnet101.rt");
+*/
+    
+    tk::dnn::dataDim_t out_dim;
+    out_dim = layer1_0_conv3.output_dim;
+    dnnType *cudnn_out, *rt_out;
+
+    tk::dnn::dataDim_t dim1 = dim; //input dim
+    printCenteredTitle(" CUDNN inference ", '=', 30);
+    {
+        dim1.print();
+        TIMER_START
+        net.infer(dim1, data);
+        TIMER_STOP
+        dim1.print();
+    }
+    cudnn_out = layer1_0_conv3.dstData;
+/*
+    tk::dnn::dataDim_t dim2 = dim;
+    printCenteredTitle(" TENSORRT inference ", '=', 30);
+    {
+        dim2.print();
+        TIMER_START
+        netRT.infer(dim2, data);
+        TIMER_STOP
+        dim2.print();
+    }
+    rt_out = (dnnType *)netRT.buffersRT[1];
+*/
+
+    printCenteredTitle(std::string(" RESNET CHECK RESULTS ").c_str(), '=', 30);
+    dnnType *out, *out_h;
+    int odim = out_dim.tot();
+    readBinaryFile(output_bin, odim, &out_h, &out);
+    std::cout << "CUDNN vs correct";
+    checkResult(odim, cudnn_out, out);
+/*
+    std::cout << "TRT   vs correct";
+    checkResult(odim, rt_out, out);
+    std::cout << "CUDNN vs TRT    ";
+    checkResult(odim, cudnn_out, rt_out);
+*/
+    return 0;
+}

tests/resnet101/resnet101_weightsexporter.py

@@ -0,0 +1,122 @@
+import torch
+import urllib
+from PIL import Image
+from torchvision import transforms
+from torchsummary import summary
+import numpy as np
+
+def hook(module, input, output):
+    setattr(module, "_value_hook", output)
+
+
+
+def print_wb(model, folder):
+    for name, param in model.named_parameters():
+        print ("Layer", name)
+        t = name.split('.')[0:-1]
+        arg = name.split('.')[-1]
+        t = '-'.join(t)
+        print ("    type: ", t)
+
+        if arg == 'weight':
+            w = param.data.numpy()
+            print ("    weights shape:", np.shape(w))
+            w.tofile(folder + "/" + t + ".bin", format="f")
+        elif arg == 'bias':
+            b = param.data.numpy()
+            print ("    bias shape:", np.shape(b))
+            b.tofile(folder + "/" + t + ".bias.bin", format="f")
+        else:
+            print("Ops!")
+
+
+def print_wb_output(model, input_batch):
+    for n, m in model.named_modules():
+        m.register_forward_hook(hook)
+
+    model(input_batch)
+    i = input_batch.data.numpy()
+    i = np.array(i, dtype=np.float32)
+    print(i.shape)
+    i.tofile("debug/input.bin", format="f")
+
+    for n, m in model.named_modules():
+        in_output = m._value_hook
+        print(n, ' ----------------------------------------------------------------')
+        o = in_output.data.numpy()
+        o = np.array(o, dtype=np.float32)
+        t = '-'.join(n.split('.'))
+        o.tofile("debug/" + t + ".bin", format="f")
+
+        # print(m._parameters)
+        print(m.type)
+
+        w = np.array([])
+        b = np.array([])
+
+
+        if 'weight' in m._parameters and m._parameters['weight'] is not None:
+            w = m._parameters['weight'].data.numpy()
+            w = np.array(w, dtype=np.float32)
+            print ("    weights shape:", np.shape(w))
+        
+        if 'bias' in m._parameters and m._parameters['bias'] is not None:
+            b = m._parameters['bias'].data.numpy()
+            b = np.array(b, dtype=np.float32)
+            print ("    bias shape:", np.shape(b))
+        else:
+            b = np.array(o, dtype=np.float32)*0
+            print ("    bias shape:", np.shape(b))
+
+        if 'BatchNorm2d' in str(m.type):
+            s = np.array(o, dtype=np.float32)*0+1
+
+        f = open("layers/" + t + ".bin", mode='wb')
+        
+        if 'BatchNorm2d' in str(m.type):
+            s.tofile(f, format="f")
+
+        w.tofile(f, format="f")
+        b.tofile(f, format="f")
+        f.close()
+
+
+
+
+
+if __name__ == '__main__':
+
+    model = torch.hub.load('pytorch/vision', 'resnet101', pretrained=True)
+    model.eval()
+
+    # Download an example image from the pytorch website
+    url, filename = ("https://github.com/pytorch/hub/raw/master/dog.jpg", "dog.jpg")
+    try: urllib.URLopener().retrieve(url, filename)
+    except: urllib.request.urlretrieve(url, filename)
+
+    # sample execution (requires torchvision)
+    input_image = Image.open(filename)
+    preprocess = transforms.Compose([
+        transforms.Resize(256),
+        transforms.CenterCrop(224),
+        transforms.ToTensor(),
+        transforms.Normalize(mean=[0.485, 0.456, 0.406], std=[0.229, 0.224, 0.225]),
+    ])
+    input_tensor = preprocess(input_image)
+    input_batch = input_tensor.unsqueeze(0) # create a mini-batch as expected by the model
+
+    # move the input and model to GPU for speed if available
+    if torch.cuda.is_available():
+        input_batch = input_batch.to('cuda')
+        model.to('cuda')
+
+    with torch.no_grad():
+        output = model(input_batch)
+
+    # Tensor of shape 1000, with confidence scores over Imagenet's 1000 classes
+    print(output)
+
+
+    print_wb_output(model, input_batch)
+
+    print(list(model.children()))