Merge with master works

Signed-off-by:  Micaela Verucchi <micaelaverucchi@gmail.com>
		Davide Sapienza <sapienza.dav@gmail.com>

.gitignore

@@ -8,4 +8,6 @@ build/
 *.h5
 *.tar.gz
 *.weights
-.idea/
+.idea/
+*.hdf5
+*.pk

CMakeLists.txt

@@ -30,6 +30,9 @@ cuda_add_library(kernels SHARED ${tkdnn_CUSRC})
 #-------------------------------------------------------------------------------
 # External Libraries
 #-------------------------------------------------------------------------------
+find_package(Eigen3 REQUIRED)
+include_directories(${EIGEN3_INCLUDE_DIR})
+
 find_package(OpenCV REQUIRED)
 set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV")
 
@@ -130,6 +133,8 @@ target_link_libraries(test_dla34 tkDNN)
 add_executable(test_dla34_cnet tests/dla34_cnet/dla34_cnet.cpp)
 target_link_libraries(test_dla34_cnet tkDNN)
 
+add_executable(test_imuodom tests/imuodom/imuodom.cpp)
+target_link_libraries(test_imuodom tkDNN)
 ################################################################################
 
 

include/tkDNN/ImuOdom.h

@@ -0,0 +1,143 @@
+#include <iostream>
+#include <signal.h>
+#include <stdlib.h>     /* srand, rand */
+#include <unistd.h>
+#include <mutex>
+#include <Eigen/Dense>
+#include "utils.h"
+#include "tkdnn.h"
+
+namespace tk { namespace dnn {
+
+/**
+ * 
+ * @author Francesco Gatti
+ */
+class ImuOdom {
+        
+    public:
+        tk::dnn::Network *net = nullptr;
+
+        // Network input dim
+        tk::dnn::dataDim_t dim0;
+        tk::dnn::dataDim_t dim1;
+        tk::dnn::dataDim_t dim2;
+
+        // Network output dim
+        tk::dnn::dataDim_t odim0;
+        tk::dnn::dataDim_t odim1;
+
+        // input pointers
+        dnnType *i0_d, *i1_d, *i2_d;
+        // output pointers
+        dnnType *o0_d, *o1_d;
+
+        // output eigen CPU
+        Eigen::MatrixXf deltaP, deltaQ;
+
+        Eigen::MatrixXd odomPOS, odomROT;
+        Eigen::Isometry3f tf = Eigen::Isometry3f::Identity();
+
+        ImuOdom() {}
+
+        virtual ~ImuOdom() {}
+
+        /**
+         * Method used for inizialize the class
+         * 
+         * @return Success of the initialization
+         */
+        bool init(std::string layers_path) {
+
+            dim0 = tk::dnn::dataDim_t(1, 4, 1, 100);
+            dim1 = tk::dnn::dataDim_t(1, 3, 1, 100);
+            dim2 = tk::dnn::dataDim_t(1, 3, 1, 100);
+
+            checkCuda( cudaMalloc(&i0_d, dim0.tot()*sizeof(dnnType)) );
+            checkCuda( cudaMalloc(&i1_d, dim1.tot()*sizeof(dnnType)) );
+            checkCuda( cudaMalloc(&i2_d, dim2.tot()*sizeof(dnnType)) );
+
+            std::string c0_bin = layers_path + "/conv1d_7.bin";
+            std::string c1_bin = layers_path + "/conv1d_8.bin";
+            std::string c2_bin = layers_path + "/conv1d_9.bin";
+            std::string c3_bin = layers_path + "/conv1d_10.bin";
+            std::string c4_bin = layers_path + "/conv1d_11.bin";
+            std::string c5_bin = layers_path + "/conv1d_12.bin";
+            std::string l0_bin = layers_path + "/bidirectional_3.bin";
+            std::string l1_bin = layers_path + "/bidirectional_4.bin";
+            std::string d0_bin = layers_path + "/dense_3.bin";
+            std::string d1_bin = layers_path + "/dense_4.bin";
+
+            net = new tk::dnn::Network(dim0);
+            tk::dnn::Input   *x0   = new tk::dnn::Input  (net, dim0, i0_d);
+            tk::dnn::Conv2d  *x0_0 = new tk::dnn::Conv2d (net, 128, 1, 11, 1, 1, 0, 0, c0_bin);
+            tk::dnn::Conv2d  *x0_1 = new tk::dnn::Conv2d (net, 128, 1, 11, 1, 1, 0, 0, c1_bin);
+            tk::dnn::Pooling *x0_2 = new tk::dnn::Pooling(net, 1, 3, 1, 3, tk::dnn::tkdnnPoolingMode_t::POOLING_MAX);
+
+            tk::dnn::Input   *x1   = new tk::dnn::Input  (net, dim1, i1_d);
+            tk::dnn::Conv2d  *x1_0 = new tk::dnn::Conv2d (net, 128, 1, 11, 1, 1, 0, 0, c2_bin);
+            tk::dnn::Conv2d  *x1_1 = new tk::dnn::Conv2d (net, 128, 1, 11, 1, 1, 0, 0, c3_bin);
+            tk::dnn::Pooling *x1_2 = new tk::dnn::Pooling(net, 1, 3, 1, 3, tk::dnn::tkdnnPoolingMode_t::POOLING_MAX);
+
+            tk::dnn::Input   *x2   = new tk::dnn::Input  (net, dim2, i2_d);
+            tk::dnn::Conv2d  *x2_0 = new tk::dnn::Conv2d (net, 128, 1, 11, 1, 1, 0, 0, c4_bin);
+            tk::dnn::Conv2d  *x2_1 = new tk::dnn::Conv2d (net, 128, 1, 11, 1, 1, 0, 0, c5_bin);
+            tk::dnn::Pooling *x2_2 = new tk::dnn::Pooling(net, 1, 3, 1, 3, tk::dnn::tkdnnPoolingMode_t::POOLING_MAX);
+
+            tk::dnn::Layer *concat_l[3] = { x0_2, x1_2, x2_2 };
+            tk::dnn::Route *concat = new tk::dnn::Route(net, concat_l, 3);
+
+            tk::dnn::LSTM *lstm0 = new tk::dnn::LSTM(net, 128, true, l0_bin);
+            tk::dnn::LSTM *lstm1 = new tk::dnn::LSTM(net, 128, false, l1_bin);
+
+            tk::dnn::Dense *d0 = new tk::dnn::Dense(net, 3, d0_bin);
+
+            tk::dnn::Layer *lstm1_l[1] = { lstm1 };
+            tk::dnn::Route *lstm1_link = new tk::dnn::Route(net, lstm1_l, 1);
+            tk::dnn::Dense *d1 = new tk::dnn::Dense(net, 4, d1_bin);
+            
+            net->print();
+
+            // output data
+            o0_d = d0->dstData;
+            o1_d = d1->dstData;
+            odim0 = d0->output_dim;
+            odim1 = d1->output_dim;
+
+            deltaP.resize(odim0.tot(), 1);
+            deltaQ.resize(odim1.tot(), 1);
+
+            odomPOS = Eigen::MatrixXd::Zero(3, 1); 
+            odomROT = Eigen::MatrixXd::Identity(3, 3);
+        }
+
+        void update(dnnType *x0, dnnType *x1, dnnType *x2) {
+
+            checkCuda( cudaMemcpy(i0_d, x0, dim0.tot()*sizeof(dnnType), cudaMemcpyHostToDevice) );
+            checkCuda( cudaMemcpy(i1_d, x1, dim1.tot()*sizeof(dnnType), cudaMemcpyHostToDevice) );
+            checkCuda( cudaMemcpy(i2_d, x2, dim2.tot()*sizeof(dnnType), cudaMemcpyHostToDevice) );
+
+            // Inference
+            tk::dnn::dataDim_t dim;
+            net->infer(dim, nullptr);
+
+            checkCuda( cudaMemcpy(deltaP.data(), o0_d, odim0.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost) );
+            checkCuda( cudaMemcpy(deltaQ.data(), o1_d, odim1.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost) ); 
+
+            // compute odom
+            Eigen::Quaterniond q;
+            q.w() = deltaQ(0);
+            q.x() = deltaQ(1);
+            q.y() = deltaQ(2);
+            q.z() = deltaQ(3);    
+            odomPOS = odomPOS + odomROT*deltaP.cast<double>();
+            odomROT = odomROT * q.normalized().toRotationMatrix();
+
+            // compose tf
+            tf.matrix().block(0, 0, 3, 3) = odomROT.cast<float>();
+            tf.matrix().block(0, 3, 3, 1) = odomPOS.cast<float>();
+        }
+
+};
+
+}}

include/tkDNN/Layer.h

@@ -9,10 +9,12 @@
 namespace tk { namespace dnn {
 
 enum layerType_t {
+    LAYER_INPUT,
     LAYER_DENSE,
     LAYER_CONV2D,
     LAYER_DECONV2D,
     LAYER_DEFORMCONV2D,
+    LAYER_LSTM,
     LAYER_ACTIVATION,
     LAYER_ACTIVATION_CRELU,
     LAYER_ACTIVATION_LEAKY,
@@ -55,10 +57,12 @@ public:
     std::string getLayerName() {
         layerType_t type = getLayerType();
         switch(type) {
+            case LAYER_INPUT:               return "Input";
             case LAYER_DENSE:               return "Dense";
             case LAYER_CONV2D:              return "Conv2d";
             case LAYER_DECONV2D:            return "DeConv2d";
             case LAYER_DEFORMCONV2D:        return "DeformConv2d";
+            case LAYER_LSTM:                return "LSTM";
             case LAYER_ACTIVATION:          return "Activation";
             case LAYER_ACTIVATION_CRELU:    return "ActivationCReLU";
             case LAYER_ACTIVATION_LEAKY:    return "ActivationLeaky";
@@ -123,6 +127,28 @@ public:
 };
 
 
+/**
+    Input layer (it doesnt need weigths)
+*/
+class Input : public Layer {
+
+public:
+
+    Input(Network *net, dataDim_t &dim, dnnType* srcData) : Layer(net) {
+        input_dim = dim;
+        output_dim = dim;
+        dstData = srcData;
+    }
+    virtual ~Input() {}
+    virtual layerType_t getLayerType() { return LAYER_INPUT; };
+
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData) {
+        dim = output_dim;
+        return dstData;
+    }
+};
+
+
 /**
     Dense (full interconnection) layer
 */
@@ -174,6 +200,14 @@ protected:
 
 /**
     Convolutional 2D layer
+
+    WEIGHTS shape: OUTCH, INCH, KH, KW ...
+    BIAS shape: OUTCH
+
+    with BATCHNORM:
+        scales:   OUTCH
+        means:    OUTCH
+        variance: OUTCH
 */
 class Conv2d : public LayerWgs {
 
@@ -203,6 +237,71 @@ protected:
     size_t ws_sizeInBytes;
 };
 
+/**
+    Bidirectional LSTM layer
+    ONLY BIDIRECTIONAL (TODO: more configurable)
+    currently implemented as 2 inferences: forward and backward (TODO: only 1 cudnn inference)
+
+    implementation info:
+    https://github.com/jiangnanhugo/seq2seq_cuda/blob/e4dbdcfa0517c972bfd4beea9f11a5233954093c/src/rnn.cpp
+    https://github.com/Jeffery-Song/mxnet-test/blob/aab666faad44011f7a67b527b5f6c960367d0422/src/operator/cudnn_rnn-inl.h
+    https://stackoverflow.com/a/38737941
+    https://colah.github.io/posts/2015-08-Understanding-LSTMs/
+
+    PARAMS (numlayers*2):
+        layer0:
+            ( INCH, ? )    ???
+            ( HIDDEN, ? )  ???
+            ( HIDDEN * 8 ) ???
+        layer2:
+            ( INCH, ? )    ???
+            ( HIDDEN, ? )  ???
+            ( HIDDEN * 8 ) ???
+
+    OUTPUT shape:
+        (N, C, 1, W) ---> LSTM(HIDDEN, returnSeq=True)  ---> (N, 2*HIDDEN, 1, W)   # W is seqLength 
+        (N, C, 1, W) ---> LSTM(HIDDEN, returnSeq=False) ---> (N, 2*HIDDEN, 1, 1)
+*/
+class LSTM : public Layer {
+
+public:
+    LSTM(Network *net, int hiddensize, bool returnSeq, std::string fname_weights);
+    virtual ~LSTM();
+    virtual layerType_t getLayerType() { return LAYER_LSTM; };
+
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+
+    const bool bidirectional = true; /**> is the net bidir */
+    bool returnSeq = false;       /**> if false return only the result of last timestep */
+    int stateSize = 0; /**> number of hidden states */
+    int seqLen = 0;    /**> number of timesteps */
+    int numLayers = 1; /**> number of internal layers */
+
+protected:
+    cudnnRNNDescriptor_t rnnDesc;
+    cudnnDropoutDescriptor_t dropoutDesc;
+    dnnType  *dropout_states_, *work_space_;
+
+    size_t workspace_byte_, dropout_byte_;
+    int workspace_size_, dropout_size_;
+    
+    std::vector<cudnnTensorDescriptor_t> x_desc_vec_, y_desc_vec_;
+    cudnnTensorDescriptor_t hx_desc_, cx_desc_;
+    cudnnTensorDescriptor_t hy_desc_, cy_desc_;
+    dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr;
+    int stateDataDim;
+
+    cudnnFilterDescriptor_t w_desc_;
+    dnnType *w_ptr;
+    dnnType *w_h;
+    dnnType *wf_ptr, *wb_ptr; // params pointer forward and backward layer
+
+    // used during inference
+    dataDim_t one_output_dim; // output dim of as single inference
+    dnnType *srcF, *srcB; // input of single inference 
+    dnnType *dstF, *dstB_NR, *dstB; // output of single inference, dstB_NR = dstB not reversed
+};
+
 
 /**
     Convolutional 2D layer
@@ -370,8 +469,8 @@ public:
     virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
 
 public:
-    static const int MAX_INPUT_LAYERS = 16;
+    static const int MAX_LAYERS = 32;
-    Layer *layers[MAX_INPUT_LAYERS];  //ids of layers to be merged
+    Layer *layers[MAX_LAYERS];  //ids of layers to be merged
     int layers_n; //number of layers
 };
 

include/tkDNN/utils.h

@@ -93,11 +93,11 @@
     }                                                                  \
 }
 
-void printCenteredTitle(const char *title, char fill, int dim);
+void printCenteredTitle(const char *title, char fill, int dim = 30);
 bool fileExist(const char *fname);
 void downloadWeightsifDoNotExist(const std::string& input_bin, const std::string& test_folder, const std::string& weights_url);
-void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0, bool skipLoad = false);
+void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0);
-int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true);
+int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true, int limit = 10);
 void printDeviceVector(int size, dnnType* vec_d, bool device = true);
 float getColor(const int c, const int x, const int max);
 void resize(int size, dnnType **data);

scripts/test_all_tests.sh

@@ -1,6 +1,7 @@
 #!/bin/bash
 
 cd build
+# rm *rt
 
 RED='\033[1;31m'
 GREEN='\033[1;32m'
@@ -31,9 +32,9 @@ print_output $res_yolo3 test_yolo3
 res_yolo3_flir=$?
 print_output $res_yolo3_flir test_yolo3_flir
 
-./test_yolo3_512 &>> $out_file
+# ./test_yolo3_512 &>> $out_file
-res_yolo3_512=$?
+# res_yolo3_512=$?
-print_output $res_yolo3_512 test_yolo3_512
+# print_output $res_yolo3_512 test_yolo3_512
 
 ./test_yolo3_tiny &>> $out_file
 res_yolo3_tiny=$?
@@ -67,20 +68,20 @@ print_output $res_mnist test_mnist
 res_yolo=$?
 print_output $res_yolo test_yolo
 
-./test_yolo3_berkeley &>> $out_file
+# ./test_yolo3_berkeley &>> $out_file
-res_yolo3_berkeley=$?
+# res_yolo3_berkeley=$?
-print_output $res_yolo3_berkeley test_yolo3_berkeley
+# print_output $res_yolo3_berkeley test_yolo3_berkeley
 
 ./test_yolo_voc &>> $out_file
 res_yolo_voc=$?
 print_output $res_yolo_voc test_yolo_voc
 
-./test_dla34_cnet &>> $out_file
+# ./test_dla34_cnet &>> $out_file
-res_dla34_cnet=$?
+# res_dla34_cnet=$?
-print_output $res_dla34_cnet test_dla34_cnet
+# print_output $res_dla34_cnet test_dla34_cnet
 
-./test_yolo3_coco4 &>> $out_file
+# ./test_yolo3_coco4 &>> $out_file
-res_yolo3_coco4=$?
+# res_yolo3_coco4=$?
-print_output $res_yolo3_coco4 test_yolo3_coco4
+# print_output $res_yolo3_coco4 test_yolo3_coco4
 
 echo "If errors occured, check logfile $out_file" 

src/LSTM.cpp

@@ -0,0 +1,327 @@
+#include <iostream>
+
+#include "Layer.h"
+
+namespace tk { namespace dnn {
+
+LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weights) :
+    Layer(net) {
+
+    this->returnSeq = returnSeq;
+    int batchSize = input_dim.n;
+    int inputSize = input_dim.c;
+    seqLen        = input_dim.w;
+    stateSize     = hiddensize;
+
+    // init Tensor Descriptors
+    std::vector<cudnnTensorDescriptor_t> x_vec(seqLen);
+    std::vector<cudnnTensorDescriptor_t> y_vec(seqLen);
+
+    int dimA[3];
+    int strideA[3];
+    for (int i = 0; i < seqLen; i++) {
+        checkCUDNN(cudnnCreateTensorDescriptor(&x_vec[i]));
+        checkCUDNN(cudnnCreateTensorDescriptor(&y_vec[i]));
+
+        dimA[0] = batchSize;
+        dimA[1] = inputSize;
+        dimA[2] = 1;
+        dimA[0] = batchSize;
+        dimA[1] = inputSize;
+        strideA[0] = dimA[2] * dimA[1];
+        strideA[1] = dimA[2];
+        strideA[2] = 1;
+        checkCUDNN(cudnnSetTensorNdDescriptor(x_vec[i],
+            net->dataType, 3, dimA, strideA));
+
+        dimA[0] = batchSize;
+        dimA[1] = stateSize;
+        dimA[2] = 1;
+        strideA[0] = dimA[2] * dimA[1];
+        strideA[1] = dimA[2];
+        strideA[2] = 1;
+        checkCUDNN(cudnnSetTensorNdDescriptor(y_vec[i],
+            net->dataType, 3, dimA, strideA));
+    }
+    // apply tensordesc
+    x_desc_vec_ = x_vec;
+    y_desc_vec_ = y_vec;
+
+
+    // set the state tensors
+    dimA[0] = numLayers;
+    dimA[1] = batchSize;
+    dimA[2] = stateSize;
+    strideA[0] = dimA[2] * dimA[1];
+    strideA[1] = dimA[2];
+    strideA[2] = 1;
+    checkCUDNN(cudnnCreateTensorDescriptor(&hx_desc_));
+    checkCUDNN(cudnnCreateTensorDescriptor(&cx_desc_));
+    checkCUDNN(cudnnCreateTensorDescriptor(&hy_desc_));
+    checkCUDNN(cudnnCreateTensorDescriptor(&cy_desc_));
+    checkCUDNN(cudnnSetTensorNdDescriptor(hx_desc_, net->dataType, 3, dimA, strideA));
+    checkCUDNN(cudnnSetTensorNdDescriptor(cx_desc_, net->dataType, 3, dimA, strideA));
+    checkCUDNN(cudnnSetTensorNdDescriptor(hy_desc_, net->dataType, 3, dimA, strideA));
+    checkCUDNN(cudnnSetTensorNdDescriptor(cy_desc_, net->dataType, 3, dimA, strideA));
+    // allocate     dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr;
+    stateDataDim = dimA[0]*dimA[1]*dimA[2];
+    checkCuda( cudaMalloc(&hx_ptr, stateDataDim*sizeof(dnnType)) );
+    checkCuda( cudaMalloc(&cx_ptr, stateDataDim*sizeof(dnnType)) );
+    checkCuda( cudaMalloc(&hy_ptr, stateDataDim*sizeof(dnnType)) );
+    checkCuda( cudaMalloc(&cy_ptr, stateDataDim*sizeof(dnnType)) );
+    
+
+
+    // Create Dropout descriptors // TODO: ??? IS IT NECESSARY ???
+    float dropoutprob = 0.1f; // random val ????
+    checkCUDNN(cudnnCreateDropoutDescriptor(&dropoutDesc));
+    checkCUDNN(cudnnDropoutGetStatesSize(net->cudnnHandle, &dropout_byte_));
+    dropout_size_ = dropout_byte_ / sizeof(dnnType);
+    checkCuda( cudaMalloc(&dropout_states_, dropout_byte_) );
+    uint64_t seed_ = 17 + rand() % 4096;  // NOLINT(runtime/threadsafe_fn)
+    checkCUDNN(cudnnSetDropoutDescriptor(dropoutDesc,
+        net->cudnnHandle, dropoutprob, dropout_states_, dropout_byte_, seed_));
+
+
+    // RNN descriptors
+    checkCUDNN(cudnnCreateRNNDescriptor(&rnnDesc));
+
+    checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle, 
+        rnnDesc, stateSize, numLayers, dropoutDesc,
+        cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT,
+        //(bidirectional ? cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL : cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL),
+        cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL,
+        cudnnRNNMode_t::CUDNN_LSTM,
+        cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD,
+        net->dataType));
+
+
+    // Get temp space sizes
+    checkCUDNN(cudnnGetRNNWorkspaceSize(net->cudnnHandle,
+        rnnDesc, seqLen, x_desc_vec_.data(), &workspace_byte_));
+    workspace_size_ = workspace_byte_ / sizeof(dnnType);
+    checkCuda( cudaMalloc(&work_space_, workspace_byte_) );    
+    
+
+    // Check that number of params are correct
+    size_t cudnn_param_size;
+    checkCUDNN(cudnnGetRNNParamsSize(net->cudnnHandle,
+        rnnDesc,x_desc_vec_[0], &cudnn_param_size, net->dataType));
+    int cudnn_params = cudnn_param_size/sizeof(dnnType);
+    //std::cout<<"LSTM params size: "<<cudnn_params << ", bytes: "<<cudnn_param_size<<"\n";
+
+    // Set param descriptors
+    checkCUDNN(cudnnCreateFilterDescriptor(&w_desc_));
+    int dim_w[3] = {1, 1, 1};
+    dim_w[0] = cudnn_params;
+    checkCUDNN(cudnnSetFilterNdDescriptor(w_desc_,
+        net->dataType, net->tensorFormat, 3, dim_w));
+
+    // load params
+    std::cout<<"Reading weights: PARAMS="<<cudnn_params*2<<"\n";
+    readBinaryFile(fname_weights, cudnn_params*2, &w_h, &w_ptr);
+    // set forward and backward params
+    wf_ptr = w_ptr;
+    wb_ptr = w_ptr + cudnn_params;
+    //std::cout<<"wf: "<<wf_ptr<<" wb "<<wb_ptr<<"\n"; 
+
+    // set output dim 
+    output_dim = input_dim;
+    output_dim.c = stateSize*(bidirectional ? 2 : 1);
+
+    // if retunseq is disabled only the last timestep is returned
+    if(!returnSeq) {
+        output_dim.h = 1;
+        output_dim.w = 1;
+    }
+
+    //allocate data for infer result
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
+
+    // used during inference
+    one_output_dim = input_dim;
+    one_output_dim.c = stateSize;
+    checkCuda( cudaMalloc(&srcF, input_dim.tot()*sizeof(dnnType)) );
+    checkCuda( cudaMalloc(&srcB, input_dim.tot()*sizeof(dnnType)) );
+    checkCuda( cudaMalloc(&dstF, one_output_dim.tot()*sizeof(dnnType)) );
+    checkCuda( cudaMalloc(&dstB_NR, one_output_dim.tot()*sizeof(dnnType)) );
+    checkCuda( cudaMalloc(&dstB, one_output_dim.tot()*sizeof(dnnType)) );
+
+
+    /*
+    // Query weight layout
+    cudnnFilterDescriptor_t m_desc;
+    checkCUDNN(cudnnCreateFilterDescriptor(&m_desc));
+    dnnType *p;
+    int n = 8; // lstm layers
+    
+    printCenteredTitle("WEIGHTS", '=', 20);
+    for (int i = 0; i < numLayers; ++i) {
+        for (int j = 0; j < n; ++j) {
+            
+            checkCUDNN(cudnnGetRNNLinLayerMatrixParams(net->cudnnHandle, rnnDesc,
+                i, x_desc_vec_[0], w_desc_, 0, j, m_desc, (void**)&p));
+            
+            std::cout << "ptr: " << ((int64_t)(p - NULL))/sizeof(dnnType)<<"\n";
+            
+            cudnnDataType_t t;
+            cudnnTensorFormat_t f;
+            int ndim = 5;
+            int dims[5] = {0, 0, 0, 0, 0};
+            checkCUDNN(cudnnGetFilterNdDescriptor(m_desc, ndim, &t, &f, &ndim, &dims[0]));
+            std::cout << "(layer, linlayer): " <<  i << " " << j << "\n";
+
+            int tot = 1;
+            for (int i = 0; i < ndim; ++i) {
+                std::cout << dims[i] << " ";
+                tot *= dims[i];
+            }
+            std::cout<<"\t-> "<<tot<<"\n\n";
+        }
+    }
+
+    printCenteredTitle("BIAS", '=', 20);
+    for (int i = 0; i < numLayers; ++i) {
+        for (int j = 0; j < n; ++j) {
+            checkCUDNN(cudnnGetRNNLinLayerBiasParams(net->cudnnHandle, rnnDesc, 
+                i, x_desc_vec_[0], w_desc_, 0, j, m_desc, (void**)&p));
+
+            std::cout << "ptr: " << ((int64_t)(p - NULL))/sizeof(dnnType)<<"\n";
+
+            cudnnDataType_t t;
+            cudnnTensorFormat_t f;
+            int ndim = 5;
+            int dims[5] = {0, 0, 0, 0, 0};
+            checkCUDNN(cudnnGetFilterNdDescriptor(m_desc, ndim, &t, &f, &ndim, &dims[0]));
+            std::cout << "(layer, linlayer): " <<  i << " " << j << "\n";
+
+            int tot = 1;
+            for (int i = 0; i < ndim; ++i) {
+                std::cout << dims[i] << " ";
+                tot *= dims[i];
+            }
+            std::cout<<"\t-> "<<tot<<"\n\n";
+        }
+    }
+    
+    checkCUDNN(cudnnDestroyFilterDescriptor(m_desc));
+    */
+}
+
+LSTM::~LSTM() {
+    checkCuda(cudaFree(hx_ptr));
+    checkCuda(cudaFree(cx_ptr));
+    checkCuda(cudaFree(hy_ptr));
+    checkCuda(cudaFree(cy_ptr));
+    checkCuda(cudaFree(w_ptr ));
+
+    checkCuda(cudaFree(work_space_    ));
+    checkCuda(cudaFree(dropout_states_));
+
+    checkCuda(cudaFree(srcF));
+    checkCuda(cudaFree(srcB));
+    checkCuda(cudaFree(dstF));
+    checkCuda(cudaFree(dstB_NR));
+    checkCuda(cudaFree(dstB));
+    checkCuda(cudaFree(dstData));
+}
+
+dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) {
+
+    // transpose input 
+    matrixTranspose(net->cublasHandle, srcData, srcF, dim.c, dim.h*dim.w*dim.l);
+    
+    // build srcB as reversed srcF
+    for(int i=0; i<input_dim.w; i++) {
+        int off_0 = i*(input_dim.c);
+        int off_1 = (i+1)*(input_dim.c);
+        checkCuda( cudaMemcpy(srcB + dim.tot() - off_1, srcF + off_0, 
+            input_dim.c*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+    }
+
+    // forward
+    { 
+        // reset states
+        checkCuda( cudaMemset(hx_ptr, 0, stateDataDim*sizeof(float)) );	
+        checkCuda( cudaMemset(cx_ptr, 0, stateDataDim*sizeof(float)) );
+
+
+        checkCUDNN(cudnnRNNForwardInference(net->cudnnHandle,
+            rnnDesc,
+            seqLen,                     // number of time steps (nT)
+            x_desc_vec_.data(),         // input array of desc (nT*nC_in)
+            srcF,                    // input pointer
+            hx_desc_,                   // initial hidden state desc     
+            hx_ptr,                     // initial hidden state pointer 
+            cx_desc_,                   // initial cell state desc      
+            cx_ptr,                     // initial cell state pointer   
+            w_desc_,                    // weights desc
+            wf_ptr,                     // weights pointer
+            y_desc_vec_.data(),         // output desc     (nT*nC_out)
+            dstF,                       // output pointer
+            hy_desc_,                   // final hidden state desc        
+            hy_ptr,                     // final hidden state pointer 
+            cy_desc_,                   // final cell state desc          
+            cy_ptr,                     // final cell state pointer   
+            work_space_,                // workspace pointer
+            workspace_byte_));          // workspace size    
+    }
+
+    // backward
+    { 
+        // reset states
+        checkCuda( cudaMemset(hx_ptr, 0, stateDataDim*sizeof(float)) );	
+        checkCuda( cudaMemset(cx_ptr, 0, stateDataDim*sizeof(float)) );
+
+        checkCUDNN(cudnnRNNForwardInference(net->cudnnHandle,
+            rnnDesc,
+            seqLen,                     // number of time steps (nT)
+            x_desc_vec_.data(),         // input array of desc (nT*nC_in)
+            srcB,                    // input pointer
+            hx_desc_,                   // initial hidden state desc     
+            hx_ptr,                     // initial hidden state pointer 
+            cx_desc_,                   // initial cell state desc      
+            cx_ptr,                     // initial cell state pointer   
+            w_desc_,                    // weights desc
+            wb_ptr,                     // weights pointer
+            y_desc_vec_.data(),         // output desc     (nT*nC_out)
+            dstB_NR,                       // output pointer
+            hy_desc_,                   // final hidden state desc        
+            hy_ptr,                     // final hidden state pointer 
+            cy_desc_,                   // final cell state desc          
+            cy_ptr,                     // final cell state pointer   
+            work_space_,                // workspace pointer
+            workspace_byte_));          // workspace size    
+    }
+
+
+    // reverse order of dstB
+    for(int i=0; i<one_output_dim.w; i++) {
+        int off_0 = i*(one_output_dim.c);
+        int off_1 = (i+1)*(one_output_dim.c);
+        checkCuda( cudaMemcpy(dstB + one_output_dim.tot() - off_1, dstB_NR + off_0, 
+            one_output_dim.c*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+    }    
+
+    // if retunseq is disabled only the last timestep is returned
+    if(returnSeq) {
+        // forward transpose
+        matrixTranspose(net->cublasHandle, dstF, dstData, 
+                         one_output_dim.h* one_output_dim.w*one_output_dim.l, one_output_dim.c);
+        // backward transpose
+        matrixTranspose(net->cublasHandle, dstB, dstData + one_output_dim.tot(), 
+                         one_output_dim.h* one_output_dim.w*one_output_dim.l, one_output_dim.c);    
+    } else {
+        // copy last of forward 
+        checkCuda( cudaMemcpy(dstData, dstF + one_output_dim.tot() - one_output_dim.c, 
+            one_output_dim.c*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+        // copy first of backward
+        checkCuda( cudaMemcpy(dstData + one_output_dim.c, dstB,  
+            one_output_dim.c*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+    }
+
+    dim = output_dim;
+    return dstData;
+}
+
+}}

src/LayerWgs.cpp

@@ -17,24 +17,24 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs,
 
     std::cout<<"Reading weights: I="<<inputs<<" O="<<outputs<<" KERNEL="<<kh<<"x"<<kw<<"x"<<kl<<"\n";
     int seek = 0;
-    readBinaryFile(weights_path.c_str(), inputs*outputs*kh*kw*kl, &data_h, &data_d, seek, net->dontLoadWeights);
+    readBinaryFile(weights_path.c_str(), inputs*outputs*kh*kw*kl, &data_h, &data_d, seek);
     seek += inputs*outputs*kh*kw*kl;
     this->additional_bias = additional_bias;
     if(additional_bias) {
-        readBinaryFile(weights_path.c_str(), outputs, &bias2_h, &bias2_d, seek, net->dontLoadWeights);
+        readBinaryFile(weights_path.c_str(), outputs, &bias2_h, &bias2_d, seek);
         seek += outputs;   
     }
     
-    readBinaryFile(weights_path.c_str(), outputs, &bias_h, &bias_d, seek, net->dontLoadWeights);
+    readBinaryFile(weights_path.c_str(), outputs, &bias_h, &bias_d, seek);
 
     this->batchnorm = batchnorm;
     if(batchnorm) {
         seek += outputs;
-        readBinaryFile(weights_path.c_str(), outputs, &scales_h, &scales_d, seek, net->dontLoadWeights);
+        readBinaryFile(weights_path.c_str(), outputs, &scales_h, &scales_d, seek);
         seek += outputs;
-        readBinaryFile(weights_path.c_str(), outputs, &mean_h, &mean_d, seek, net->dontLoadWeights);
+        readBinaryFile(weights_path.c_str(), outputs, &mean_h, &mean_d, seek);
         seek += outputs;
-        readBinaryFile(weights_path.c_str(), outputs, &variance_h, &variance_d, seek, net->dontLoadWeights);
+        readBinaryFile(weights_path.c_str(), outputs, &variance_h, &variance_d, seek);
 
         float eps = TKDNN_BN_MIN_EPSILON;
 

src/Route.cpp

@@ -7,11 +7,14 @@ namespace tk { namespace dnn {
 
 Route::Route(Network *net, Layer **layers, int layers_n, bool final) : Layer(net, final) {
 
-    this->layers_n = layers_n;
+    // copy input layers
-    if(layers_n > MAX_INPUT_LAYERS)
+    if(layers_n > MAX_LAYERS) {
-        FatalError("Route: MAX INPUT LAYERS overload");
+        FatalError("ROUTE: reached max number of input layers");
-    for(int i=0; i<layers_n; i++)
+    }
+    for(int i=0; i<layers_n; i++) {
         this->layers[i] = layers[i];
+    }
+    this->layers_n = layers_n;
     
     //get dims
     output_dim.l = 1;

src/utils.cpp

@@ -32,43 +32,41 @@ void downloadWeightsifDoNotExist(const std::string& input_bin, const std::string
 }
 
 
-void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek, bool skipLoad){
+void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek)
+{
+    std::ifstream dataFile (fname, std::ios::in | std::ios::binary);
+    std::stringstream error_s;
+    if (!dataFile)
+    {
+        error_s << "Error opening file " << fname; 
+        FatalError(error_s.str());
+    }
+
+    if(seek != 0) {
+        dataFile.seekg(seek*sizeof(dnnType), dataFile.cur);
+    }
+
     int size_b = size*sizeof(dnnType);
     *data_h = new dnnType[size];
-
+    if (!dataFile.read ((char*) *data_h, size_b)) 
-    if(!skipLoad) {
+    {
-        std::ifstream dataFile(fname, std::ios::in | std::ios::binary);
+        error_s << "Error reading file " << fname << " with n of float: "<<size;
-        std::stringstream error_s;
+        error_s << " seek: "<<seek << " size: "<<size_b<<"\n";
-        if (!dataFile) {
+        FatalError(error_s.str());
-            error_s << "Error opening file " << fname;
-            FatalError(error_s.str());
-        }
-
-        if (seek != 0) {
-            dataFile.seekg(seek * sizeof(dnnType), dataFile.cur);
-        }
-
-        // printf("data_h %d size_b %d\n", *data_h,size_b);
-        if (!dataFile.read((char *) *data_h, size_b)) {
-            
-            error_s << "Error reading file " << fname;
-            FatalError(error_s.str());
-        }
-    } else {
-        std::cout<<COL_RED<<"WARNING: skipping data load, this should only used in debug\n"<<COL_END;
     }
     
     checkCuda( cudaMalloc(data_d, size_b) );
     checkCuda( cudaMemcpy(*data_d, *data_h, size_b, cudaMemcpyHostToDevice) );
 }
 
+
 void printDeviceVector(int size, dnnType* vec_d, bool device){
     dnnType *vec;
     if(device) {
         vec = new dnnType[size];
-        cudaDeviceSynchronize();
+        checkCuda(cudaDeviceSynchronize());
-        cudaMemcpy(vec, vec_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost);
+        checkCuda(cudaMemcpy(vec, vec_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost));
-        cudaDeviceSynchronize();
+        checkCuda(cudaDeviceSynchronize());
     } else {
         vec = vec_d;
     }
@@ -82,7 +80,7 @@ void printDeviceVector(int size, dnnType* vec_d, bool device){
         delete [] vec;
 }
 
-int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) {
+int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device, int limit) {
 
     dnnType *data_h, *correct_h;
     const float eps = 0.02f;
@@ -90,10 +88,11 @@ int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) {
     if(device) {
         data_h = new dnnType[size];
         correct_h = new dnnType[size];
-        cudaDeviceSynchronize();
+        checkCuda(cudaDeviceSynchronize());
-        cudaMemcpy(data_h, data_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost);
+        checkCuda(cudaMemcpy(data_h, data_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost));
-        cudaMemcpy(correct_h, correct_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost);
+        checkCuda(cudaMemcpy(correct_h, correct_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost));
-        cudaDeviceSynchronize();
+        checkCuda(cudaDeviceSynchronize());
+
     } else {
         data_h = data_d;
         correct_h = correct_d;
@@ -105,7 +104,7 @@ int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) {
             diffs += 1;
             if(diffs == 1)
                 std::cout<<"\n";
-            if(diffs < 10)
+            if(diffs < limit)
                 std::cout<<" | [ "<<i<<" ]: "<<data_h[i]<<" "<<correct_h[i]<<"\n";
         }
     }

tests/imuodom/imuodom.cpp

@@ -0,0 +1,77 @@
+#include<iostream>
+#include "tkDNN/ImuOdom.h"
+
+const char *i0_bin   = "../tests/imuodom/layers/input0.bin";
+const char *i1_bin   = "../tests/imuodom/layers/input1.bin";
+const char *i2_bin   = "../tests/imuodom/layers/input2.bin";
+const char *o0_bin   = "../tests/imuodom/layers/output0.bin";
+const char *o1_bin   = "../tests/imuodom/layers/output1.bin";
+
+const char *c0_bin = "../tests/imuodom/layers/conv1d_7.bin";
+const char *c1_bin = "../tests/imuodom/layers/conv1d_8.bin";
+const char *c2_bin = "../tests/imuodom/layers/conv1d_9.bin";
+const char *c3_bin = "../tests/imuodom/layers/conv1d_10.bin";
+const char *c4_bin = "../tests/imuodom/layers/conv1d_11.bin";
+const char *c5_bin = "../tests/imuodom/layers/conv1d_12.bin";
+const char *l0_bin = "../tests/imuodom/layers/bidirectional_3.bin";
+const char *l1_bin = "../tests/imuodom/layers/bidirectional_4.bin";
+const char *d0_bin = "../tests/imuodom/layers/dense_3.bin";
+const char *d1_bin = "../tests/imuodom/layers/dense_4.bin";
+
+
+int main() {
+
+    tk::dnn::ImuOdom ImuNet;
+    ImuNet.init("../tests/imuodom/layers/");
+
+    const int N = 19513;
+
+    // Network layout
+    tk::dnn::dataDim_t dim0(1, 4, 1, 100);
+    tk::dnn::dataDim_t dim1(1, 3, 1, 100);
+    tk::dnn::dataDim_t dim2(1, 3, 1, 100);
+
+    // Load input
+    dnnType *i0_d, *i1_d, *i2_d;
+    dnnType *i0_h, *i1_h, *i2_h;
+    readBinaryFile(i0_bin, dim0.tot()*N, &i0_h, &i0_d);
+    readBinaryFile(i1_bin, dim1.tot()*N, &i1_h, &i1_d);
+    readBinaryFile(i2_bin, dim2.tot()*N, &i2_h, &i2_d);
+
+    dnnType *data;
+    tk::dnn::dataDim_t dim;
+
+    dnnType *out0, *out1;
+    dnnType *out0_h, *out1_h;
+    readBinaryFile(o0_bin, ImuNet.odim0.tot()*N, &out0_h, &out0);
+    readBinaryFile(o1_bin, ImuNet.odim1.tot()*N, &out1_h, &out1);
+
+    
+    std::ofstream path("path.txt");
+
+    for(int i=0; i<N; i++) {
+        std::cout<<"i: "<<i<<"\n";
+        //TIMER_START
+        // Inference
+        ImuNet.update(i0_h, i1_h, i2_h);
+        //TIMER_STOP
+
+        // log path
+        path<<ImuNet.odomPOS(0)<<" "<<ImuNet.odomPOS(1)<<" "<< ImuNet.odomPOS(2)<<"\n";
+        path.flush();
+
+        // Print real test
+        //printCenteredTitle( (std::string(" CHECK RESULT ") + std::to_string(i) + " ").c_str() , '=');
+        //ImuNet.odim0.print();
+        //checkResult(ImuNet.odim0.tot(), out0, ImuNet.o0_d);
+        //ImuNet.odim1.print();
+        //checkResult(ImuNet.odim0.tot(), out1, ImuNet.o1_d);
+
+        i0_h += ImuNet.dim0.tot();
+        i1_h += ImuNet.dim1.tot();
+        i2_h += ImuNet.dim2.tot();
+        out0 += ImuNet.odim0.tot();
+        out1 += ImuNet.odim1.tot();
+    }
+    return 0;
+}

tests/imuodom/infer.py

@@ -0,0 +1,87 @@
+import keras
+from keras.models import load_model
+import keras.backend.tensorflow_backend as KTF
+import numpy as np
+import argparse
+import tensorflow as tf
+import os
+import random
+import struct
+from keras.models import Sequential, Model
+import pickle
+
+def bin_write(f, data):
+    data = data.flatten()
+    fmt = 'f'*len(data)
+    bin = struct.pack(fmt, *data)
+    f.write(bin)
+
+
+if __name__ == '__main__':
+
+
+    print("DATA FORMAT: ", keras.backend.image_data_format())
+
+    print("Load model: ", "ferrariS1.hdf5")
+    model = load_model("ferrariS1.hdf5")
+    model.summary()
+    weights = model.get_weights()
+
+    indata = pickle.load(open("input.pk", 'rb'))
+    outdata = pickle.load(open("output.pk", 'rb'))
+
+    x_angle = indata[0]
+    x_gyro = indata[1]
+    x_acc = indata[2]
+    
+    [yhat_delta_p, yhat_delta_q] = model.predict(indata, batch_size=1, verbose=1)
+    predictdata = [yhat_delta_p, yhat_delta_q]
+
+    error = outdata[0] - predictdata[0]
+    print("error delta_p: ", error.sum())
+    error = outdata[1] - predictdata[1]
+    print("error delta_q: ", error.sum())
+
+
+    #layer_name = 'dense_4'
+    #intermediate_layer_model = Model(inputs=model.input,
+    #                                 outputs=model.get_layer(layer_name).output)
+    #intermediate_output = intermediate_layer_model.predict([x_angle, x_gyro, x_acc])
+
+
+    x_angle = np.array([x_angle])
+    x_gyro = np.array([x_gyro])
+    x_acc = np.array([x_acc])
+    #intermediate_output = np.array([intermediate_output])
+
+    x_angle =  x_angle.transpose(1, 3, 0, 2)
+    x_gyro =  x_gyro.transpose(1, 3, 0, 2)
+    x_acc =  x_acc.transpose(1, 3, 0, 2)
+    #intermediate_output = intermediate_output.transpose(0, 3, 1, 2)
+    #print("Aggregate:")
+    #print(intermediate_output.tolist())
+ 
+    print("x0: ", np.shape(x_angle))
+    #print("out: ",np.shape(intermediate_output))
+
+    x_angle = np.array(x_angle.flatten(), dtype=np.float32)
+    x_gyro = np.array(x_gyro.flatten(), dtype=np.float32)
+    x_acc = np.array(x_acc.flatten(), dtype=np.float32)
+    yhat_delta_p = np.array(yhat_delta_p.flatten(), dtype=np.float32)
+    yhat_delta_q = np.array(yhat_delta_q.flatten(), dtype=np.float32)
+    #intermediate_output = np.array(intermediate_output.flatten(), dtype=np.float32)
+
+
+    f = open("layers/input0.bin", mode='wb')
+    bin_write(f, x_angle)
+    f = open("layers/input1.bin", mode='wb')
+    bin_write(f, x_gyro)
+    f = open("layers/input2.bin", mode='wb')
+    bin_write(f, x_acc)
+    f = open("layers/output0.bin", mode='wb')
+    bin_write(f, yhat_delta_p)
+    f = open("layers/output1.bin", mode='wb')
+    bin_write(f, yhat_delta_q)
+    #f = open("layers/output.bin", mode='wb')
+    #bin_write(f, intermediate_output)
+

tests/mobilenetv2ssd/mobilenetv2ssd.cpp

@@ -134,7 +134,7 @@ const char *regression_header5 = "../tests/mobilenetv2ssd/layers/regression_head
 int main()
 {
 
-    downloadWeightsifDoNotExist(input_bin, "../tests/mobilenetv2ssd", "https://cloud.hipert.unimore.it/s/B6mj33k7beECXsY/download");
+    downloadWeightsifDoNotExist(input_bin, "../tests/mobilenetv2ssd", "https://cloud.hipert.unimore.it/s/x4ZfxBKN23zAJQp/download");
 
     int classes = 21;
 

tests/simple/test_model.py

@@ -1,43 +1,54 @@
 import keras
 import numpy as np
 from keras.models import Sequential
-from keras.layers import Input, Dense, Activation, Flatten, Dropout, ELU, Reshape, Lambda
+from keras.layers import Input, Dense, Activation, Flatten, Dropout, ELU, Reshape, Lambda, Conv1D
+from keras.layers import Bidirectional, CuDNNLSTM
 from keras.layers.convolutional import Convolution2D, Convolution3D
 from keras.layers.pooling import MaxPooling2D, MaxPooling3D, AveragePooling3D
 from keras.models import Sequential, Model
 from keras.layers import Cropping2D
 import keras.backend.tensorflow_backend as KTF
+import struct
+from keras.models import Sequential, Model
 
-def dense_model():
+def bin_write(f, data):
-    model = Sequential()
+    data = data.flatten()
+    fmt = 'f'*len(data)
+    bin = struct.pack(fmt, *data)
+    f.write(bin)
+
+def create_model():
+    x1 = Input((3, 8), name='x1')
+    conv = Conv1D(4, 2)(x1)
+    lstm = Bidirectional(CuDNNLSTM(5, return_sequences=True))(conv)    
+    lstm2 = Bidirectional(CuDNNLSTM(5, return_sequences=False))(lstm)    
+    model = Model([x1], [lstm2])
+    model.summary()
 
-    model.add(Reshape((10, 10, 1), input_shape=(10, 10)))
-    model.add(Convolution2D(2, (4, 4), subsample=(2, 2), 
-                           bias_initializer='random_uniform', activation="relu"))
-    model.add(Convolution2D(4, (2, 2), subsample=(1, 1), 
-                           bias_initializer='random_uniform', activation="relu"))
-    model.add(Flatten())
-    model.add(Dense(4, bias_initializer='random_uniform', activation="relu"))    
-    sgd = keras.optimizers.Adam(lr=1e-4, decay=1e-8)
-    model.compile(optimizer=sgd, loss="mse")
     return model
 
-
 if __name__ == '__main__':
-    print "DATA FORMAT: ", keras.backend.image_data_format()
+    print ("DATA FORMAT: ", keras.backend.image_data_format())
 
-    model = dense_model()
+    model = create_model()
-    model.save("net.h5")
+    model.save("net.hdf5")
 
-    grid = np.random.rand(10,10)
+    np.random.seed(2)
-    X = grid[None,:,:]
+    x = np.random.rand(1,1,3,8)
-    i = np.array(grid.flatten(), dtype=np.float32)
+    r = model.predict( x[0], batch_size=1)
-    print i
+
-    i.tofile("input.bin", format="f")
+    r = np.array([r])
-    print "Input: ", X    
+    x =  x.transpose(0, 3, 1, 2)
+    #r =  r.transpose(0, 3, 1, 2)
+    print("in: ", np.shape(x))
+    print("out: ", np.shape(r))
+    print("output: ", r.tolist())
+
+    x = np.array(x.flatten(), dtype=np.float32)
+    f = open("input.bin", mode='wb')
+    bin_write(f, x)
+
+    r = np.array(r.flatten(), dtype=np.float32)
+    f = open("output.bin", mode='wb')
+    bin_write(f, r)
 
-    r = model.predict( X, batch_size=1)
-    print np.shape(r)  
-    print "Result: ", r
-    print "Result shape: ", np.shape(r) 
-    r.tofile("output.bin", format="f")

tests/simple/test_simple.cpp

@@ -2,22 +2,21 @@
 #include "tkdnn.h"
 
 const char *input_bin   = "../tests/simple/input.bin";
-const char *c0_bin      = "../tests/simple/layers/c0.bin";
+const char *c0_bin      = "../tests/simple/layers/conv1d_1.bin";
-const char *c1_bin      = "../tests/simple/layers/c1.bin";
+const char *l1_bin      = "../tests/simple/layers/bidirectional_1.bin";
-const char *d2_bin      = "../tests/simple/layers/d2.bin";
+const char *l2_bin      = "../tests/simple/layers/bidirectional_2.bin";
 const char *output_bin  = "../tests/simple/output.bin";
 
 int main() {
 
     // Network layout
-    tk::dnn::dataDim_t dim(1, 1, 10, 10, 1);
+    tk::dnn::dataDim_t dim(1, 8, 1, 3);
     tk::dnn::Network net(dim);
-    tk::dnn::Conv2d     l0(&net, 2, 4, 4, 2, 2, 0, 0, c0_bin);
+    tk::dnn::Conv2d     l0(&net, 4, 1, 2, 1, 1, 0, 0, c0_bin);
-    tk::dnn::Activation l1(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::LSTM       l1(&net, 5, true, l1_bin);
-    tk::dnn::Conv2d     l2(&net, 4, 2, 2, 1, 1, 0, 0, c1_bin);
+    tk::dnn::LSTM       l2(&net, 5, false, l2_bin);
-    tk::dnn::Activation l3(&net, CUDNN_ACTIVATION_RELU);
+
-    tk::dnn::Dense      l5(&net, 4, d2_bin);
+    net.print();
-    tk::dnn::Activation l6(&net, CUDNN_ACTIVATION_RELU);
 
     net.print();
 

tests/weights_exporter.py

@@ -5,96 +5,89 @@ import numpy as np
 import argparse
 import tensorflow as tf
 import os
-import msgpack
-import lmdb
 import random
+import struct
+from keras.models import Sequential, Model
 
-def export_dense(name, weights, bias):
+def bin_write(f, data):
-    print "########    EXPORT", name, "LAYER    ########"
+    data = data.flatten()
-    print "Original weighs:"
+    fmt = 'f'*len(data)
-    print weights
+    bin = struct.pack(fmt, *data)
-    print bias, "\n"
+    f.write(bin)
 
-    #input, filters
+def export_layer(name, weights, bias):
-    I, C = np.shape(weights) 
+    print ("########    EXPORT", name, "LAYER    ########")
-    B = np.shape(bias)
-    print "w shape: ", I, C
-    print "b shape: ", B
 
-    wgs = [ [ j[i] for j in weights ]  for i in xrange(C) ]
+    print("wgs pretranpose: ", np.shape(weights))
-    wgs = np.array(wgs, dtype=np.float32)
+    # convert NHWC to NCHW
-    
+    if(weights.ndim == 4):
-    print "REPOSITIONED WEIGHTS:"
+        weights = weights.transpose(3,2,0,1)
-    print wgs
+    elif(weights.ndim == 3):
+        weights = weights.transpose(2,1,0)
+    elif(weights.ndim == 2):
+        weights = weights.transpose(1,0)
+    else:
+        print("Ndim", weights.ndim)
+        raise("not implemented with dim" )
 
+    print("weights: ", np.shape(weights))
+    print("bias: ", np.shape(bias))
+
+    weights = np.array(weights.flatten(), dtype=np.float32)
     bias = np.array(bias, dtype=np.float32)
-    wgs.tofile(name + ".bin", format="f")
+    print(len(weights) + len(bias))
-    bias.tofile(name + ".bias.bin", format="f")
-    print "WEIGHTS saved\n"
 
-def export_conv2d(name, weights, bias):
+    f = open(name + ".bin", mode='wb')
-    print "########    EXPORT", name, "LAYER    ########"
+    bin_write(f, weights)
-    print "Original weighs:"
+    bin_write(f, bias)
-    print weights
+    print ("WEIGHTS saved\n")
-    print bias, "\n"
-  
-    # height, width, input, filters
-    H, W, N, C = np.shape(weights) 
-    B = np.shape(bias)
-    print "w shape: ", N, C, H, W
-    print "b shape: ", B
 
-    wgs = weights.transpose()
+def export_bidir(name, params, paramsb):
-    wgs = wgs.transpose(0, 1, 3, 2)
+    print ("########    EXPORT", name, "LAYER    ########")
-    print "Final shape:", np.shape(wgs)    
+ 
-    wgs = np.array(wgs.flatten(), dtype=np.float32)
+    f = open(name + ".bin", mode='wb')
-    
-    print "REPOSITIONED WEIGHTS:"
-    print wgs
 
-    bias = np.array(bias, dtype=np.float32)
+    print("FORWARD")
-
+    ker = params[0]
-    wgs.tofile(name + ".bin", format="f")
+    rec_ker = params[1]
-    bias.tofile(name + ".bias.bin", format="f")
+    bias = params[2]
-    print "WEIGHTS saved\n"
+    print ("export kernels: ", np.shape(ker))
-
+    units = np.shape(ker)[1] // 4
-def export_conv3d(name, weights, bias):
+    bin_write(f, ker[:,:units])
-    print "########    EXPORT", name, "LAYER    ########"
+    bin_write(f, ker[:,units:units*2])
-    print "Original weighs:"
+    bin_write(f, ker[:,units*2:units*3])
-    print weights
+    bin_write(f, ker[:,units*3:])
-    print bias, "\n"
+    print ("export recurrent kernels: ", np.shape(rec_ker))
-  
+    bin_write(f, rec_ker[:,:units])
-    print np.shape(weights)
+    bin_write(f, rec_ker[:,units:units*2])
-    # height, width, input, thickness, filters
+    bin_write(f, rec_ker[:,units*2:units*3])
-    H, W, T, N, C = np.shape(weights) 
+    bin_write(f, rec_ker[:,units*3:])
-    B = np.shape(bias)
+    print ("export kernels: ", np.shape(ker))
-    print "w shape: ", T, C, H, W   #thickness is number of images for cudnn
+    bin_write(f, bias)
-    print "b shape: ", B
+    print("WEIGHTS saved\n")
-
-    wgs = weights.transpose()
-    wgs = wgs.transpose(0, 1, 4, 3, 2)
-    print "Final shape:", np.shape(wgs)    
-    wgs = np.array(wgs.flatten(), dtype=np.float32)
-    
-    print "REPOSITIONED WEIGHTS:"
-    print wgs
-
-    bias = np.array(bias, dtype=np.float32)
-
-    wgs.tofile(name + ".bin", format="f")
-    bias.tofile(name + ".bias.bin", format="f")
-    print "WEIGHTS saved\n"
-
-
-def get_session(gpu_fraction=0.5):
-    gpu_options = tf.GPUOptions(allow_growth=True) 
-        #per_process_gpu_memory_fraction=gpu_fraction)
-    return tf.Session(config=tf.ConfigProto(gpu_options=gpu_options))
 
+    print("BACKWARD")
+    ker = paramsb[0]
+    rec_ker = paramsb[1]
+    bias = paramsb[2]
+    print ("export kernels: ", np.shape(ker))
+    units = np.shape(ker)[1] // 4
+    bin_write(f, ker[:,:units])
+    bin_write(f, ker[:,units:units*2])
+    bin_write(f, ker[:,units*2:units*3])
+    bin_write(f, ker[:,units*3:])
+    print ("export recurrent kernels: ", np.shape(rec_ker))
+    bin_write(f, rec_ker[:,:units])
+    bin_write(f, rec_ker[:,units:units*2])
+    bin_write(f, rec_ker[:,units*2:units*3])
+    bin_write(f, rec_ker[:,units*3:])
+    print ("export kernels: ", np.shape(ker))
+    bin_write(f, bias)
+    print("WEIGHTS saved\n")
 
 #https://github.com/fchollet/keras/wiki/Converting-convolution-kernels-from-Theano-to-TensorFlow-and-vice-versa
 if __name__ == '__main__':
-    KTF.set_session(get_session())
+    print("DATA FORMAT: ", keras.backend.image_data_format())
 
     parser = argparse.ArgumentParser(description='KERAS WEIGHTS EXPORTER TO CUDNN')
     parser.add_argument('model',type=str,
@@ -103,31 +96,43 @@ if __name__ == '__main__':
 
     args = parser.parse_args()
     
-    print "DATA FORMAT: ", keras.backend.image_data_format()
+    print("DATA FORMAT: ", keras.backend.image_data_format())
     
-    print "Load model: ", args.model
+    print("Load model: ", args.model)
     model = load_model(args.model)
+    model.summary()
+
+
     weights = model.get_weights()
 
     ws = np.shape(weights)
-    print "Weights shape:", ws
+    print("Weights shape:", ws)
 
     if not os.path.exists(args.output):
         os.makedirs(args.output)
 
-    num = 0
+
     name_num = 0
     for l in model.layers:
-        name = l.name        
+        print("\n\nNAME: ", l.name)
-        if name.startswith("conv3d"):
+        print("input: ", l.input_shape, " output: ", l.output_shape)
-            export_conv3d(args.output + "/conv" + str(name_num), weights[num], weights[num+1])
+        wgs = l.get_weights()
-        elif name.startswith("conv2d"):
+        print("wgs num: ", len(wgs))
-            export_conv2d(args.output + "/conv" + str(name_num), weights[num], weights[num+1])
+
-        elif name.startswith("dense"):
+        name = l.name 
-            export_dense(args.output + "/dense" + str(name_num), weights[num], weights[num+1])
+        if name.startswith("conv3d"):
-        else:
+            export_layer(args.output + "/" + name, wgs[0], wgs[1])
-            print "skip:", name, "has no weights"
+        elif name.startswith("conv2d"):
-            continue
+            export_layer(args.output + "/" + name, wgs[0], wgs[1])
-        name_num += 1
+        elif name.startswith("conv1d"):
-        num += 2
+            export_layer(args.output + "/" + name, wgs[0], wgs[1])     
+        elif name.startswith("dense"):
+            export_layer(args.output + "/" + name, wgs[0], wgs[1])
+        elif name.startswith("bidirectional"):
+            wgs = l.forward_layer.get_weights()
+            export_bidir(args.output + "/" + name, l.forward_layer.get_weights(), l.backward_layer.get_weights())
+        else:
+            print ("skip:", name, "has no weights")
+            continue
+