better tests

CMakeLists.txt

@@ -9,9 +9,12 @@ cuda_add_library(kernels SHARED src/kernels/activation_elu.cu)
 
 include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS})
 add_library(tkDNN SHARED src/Layer.cpp src/LayerWgs.cpp
-    src/Dense.cpp src/Activation.cpp src/Conv2d.cpp src/Conv3d.cpp src/Flatten.cpp src/MulAdd.cpp src/Pooling.cpp
+    src/Dense.cpp src/Activation.cpp src/Conv2d.cpp src/Flatten.cpp src/MulAdd.cpp src/Pooling.cpp src/Softmax.cpp
     src/Network.cpp src/utils.cpp)
-target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDA_TOOLKIT_ROOT_DIR}/lib/libcudnn.so)
+target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn)
 
-add_executable(tkDNNtest tests/test.cpp)
+add_executable(test_simple tests/test/test.cpp)
-target_link_libraries(tkDNNtest tkDNN)
+target_link_libraries(test_simple tkDNN)
+
+add_executable(test_mnist tests/mnist/test.cpp)
+target_link_libraries(test_mnist tkDNN)

include/Layer.h

@@ -92,15 +92,6 @@ protected:
     value_type *dstData;  //where results will be putted
 };
 
-/**
-    Avaible activation functions
-*/
-typedef enum {
-    ACTIVATION_SIGMOID = 0,
-    ACTIVATION_RELU    = 1,
-    ACTIVATION_TANH    = 2,
-    ACTIVATION_ELU     = 100
-} tkdnnActivationMode_t;
 
 /**
     Activation layer (it doesnt need weigths)
@@ -108,13 +99,14 @@ typedef enum {
 class Activation : public Layer {
 
 public:
-    Activation(Network *net, dataDim_t input_dim, tkdnnActivationMode_t act_mode); 
+    Activation(Network *net, dataDim_t input_dim, cudnnActivationMode_t act_mode); 
     virtual ~Activation();
 
     virtual value_type* infer(dataDim_t &dim, value_type* srcData);
 
 protected:
-    tkdnnActivationMode_t act_mode;
+    cudnnActivationMode_t act_mode;
+    cudnnActivationDescriptor_t activDesc;
     value_type *dstData;  //where results will be putted
 };
 
@@ -145,35 +137,6 @@ protected:
     size_t ws_sizeInBytes;
 };
 
-/**
-    Convolutional 3D layer
-*/
-class Conv3d : public LayerWgs {
-
-public:
-    Conv3d(Network *net, dataDim_t in_dim, int out_ch,
-            int kernelH, int kernelW, int kernelL, 
-            int strideH, int strideW, int strideL,
-            const char* fname_weights, const char* fname_bias); 
-    virtual ~Conv3d();
-
-    virtual value_type* infer(dataDim_t &dim, value_type* srcData);
-
-protected:
-    value_type *dstData;  //where results will be putted
-    int kernelH, kernelW, kernelL; 
-    int strideH, strideW, strideL;
-
-    cudnnFilterDescriptor_t filterDesc;
-    cudnnConvolutionDescriptor_t convDesc;
-    cudnnConvolutionFwdAlgo_t algo;
-    cudnnTensorDescriptor_t biasTensorDesc;
-    cudnnTensorDescriptor_t biasDstTensorDesc;
-
-    void*  workSpace;
-    size_t ws_sizeInBytes;
-};
-
 
 /**
     Flatten layer
@@ -244,5 +207,20 @@ protected:
     bool poolOn3d;
 };
 
+/**
+    Softmax layer
+*/
+class Softmax : public Layer {
+
+public:
+    Softmax(Network *net, dataDim_t input_dim); 
+    virtual ~Softmax();
+
+    virtual value_type* infer(dataDim_t &dim, value_type* srcData);
+
+protected:
+    value_type *dstData;  //where results will be putted
+};
+
 }
 #endif //LAYER_H

src/Activation.cpp

@@ -5,7 +5,7 @@
 
 namespace tkDNN {
 
-Activation::Activation(Network *net, dataDim_t input_dim, tkdnnActivationMode_t act_mode) : 
+Activation::Activation(Network *net, dataDim_t input_dim, cudnnActivationMode_t act_mode) : 
     Layer(net, input_dim) {
 
     this->act_mode = act_mode;
@@ -23,30 +23,34 @@ Activation::Activation(Network *net, dataDim_t input_dim, tkdnnActivationMode_t
                                             input_dim.n*input_dim.l, 
                                             input_dim.c,
                                             input_dim.h, input_dim.w) );
+
+
+    checkCUDNN( cudnnCreateActivationDescriptor(&activDesc) );
+    checkCUDNN( cudnnSetActivationDescriptor(activDesc,
+                                             act_mode,
+                                             CUDNN_PROPAGATE_NAN,
+                                             0.0) );
 }
 
 Activation::~Activation() {
 
     checkCuda( cudaFree(dstData) );
+
+    checkCUDNN( cudnnDestroyActivationDescriptor(activDesc) );
 }
 
 value_type* Activation::infer(dataDim_t &dim, value_type* srcData) {
 
-    if(act_mode == ACTIVATION_ELU) {
+    value_type alpha = value_type(1);
-        activationELUForward(srcData, dstData, dim.tot());
+    value_type beta  = value_type(0);
-
+    checkCUDNN( cudnnActivationForward(net->cudnnHandle,
-    } else {
+                                        activDesc,
-        value_type alpha = value_type(1);
+                                        &alpha,
-        value_type beta  = value_type(0);
+                                        srcTensorDesc,
-        checkCUDNN( cudnnActivationForward(net->cudnnHandle,
+                                        srcData,
-                                            cudnnActivationMode_t(act_mode),
+                                        &beta,
-                                            &alpha,
+                                        dstTensorDesc,
-                                            srcTensorDesc,
+                                        dstData) );    
-                                            srcData,
-                                            &beta,
-                                            dstTensorDesc,
-                                            dstData) );    
-    }
     return dstData;
 }
 

src/Conv2d.cpp

@@ -29,7 +29,7 @@ Conv2d::Conv2d( Network *net, dataDim_t in_dim, int out_ch,
                 net->tensorFormat, net->dataType, n, c, h, w) );
 
     checkCUDNN( cudnnSetFilter4dDescriptor(filterDesc,
-                net->dataType, out_ch, input_dim.c, 
+                net->dataType, net->tensorFormat, out_ch, input_dim.c, 
                 kernelH, kernelW) );
 
     checkCUDNN( cudnnSetConvolution2dDescriptor(convDesc,
@@ -103,7 +103,7 @@ value_type* Conv2d::infer(dataDim_t &dim, value_type* srcData) {
     // bias
     alpha = value_type(1);
     beta  = value_type(1);
-    checkCUDNN( cudnnAddTensor(net->cudnnHandle, CUDNN_ADD_SAME_C,
+    checkCUDNN( cudnnAddTensor(net->cudnnHandle,
                 &alpha, biasTensorDesc, bias_d,
                 &beta, dstTensorDesc, dstData) );
 

src/Conv3d.cpp

@@ -1,157 +0,0 @@
-#include <iostream>
-
-#include "Layer.h"
-
-namespace tkDNN {
-
-Conv3d::Conv3d( Network *net, dataDim_t in_dim, int out_ch,
-                int kernelH, int kernelW, int kernelL, 
-                int strideH, int strideW, int strideL,
-                const char* fname_weights, const char* fname_bias) : 
-    
-    LayerWgs(net, in_dim, in_dim.c, out_ch, kernelH, kernelW, kernelL, 
-             fname_weights, fname_bias) {
-
-    this->kernelH = kernelH;
-    this->kernelW = kernelW;
-    this->kernelL = kernelL;
-    this->strideH = strideH;
-    this->strideW = strideW;
-    this->strideL = strideL;
-
-    checkCUDNN( cudnnCreateTensorDescriptor(&biasDstTensorDesc) );
-    checkCUDNN( cudnnCreateFilterDescriptor(&filterDesc) );
-    checkCUDNN( cudnnCreateConvolutionDescriptor(&convDesc) );
-    checkCUDNN( cudnnCreateTensorDescriptor(&biasTensorDesc) );
-
-    int n = input_dim.n;
-    int c = input_dim.c;
-    int h = input_dim.h;
-    int w = input_dim.w;
-    int l = input_dim.l;
-
-    //create a tensor Nd descriptor with N = 4  
-    int dimA[5];
-    dimA[0] = n; dimA[1] = c; dimA[2] = h; dimA[3] = w; dimA[4] = l;
-    int strideA[5];
-    strideA[0] = c*h*w*l; 
-    strideA[1] = h*w*l; 
-    strideA[2] = w*l; 
-    strideA[3] = l; 
-    strideA[4] = 1;
-    checkCUDNN( cudnnSetTensorNdDescriptor(srcTensorDesc, 
-                net->dataType, 5, dimA, strideA));
-
-    //filter descriptor
-    int filterDim[5];
-    filterDim[0] = out_ch;     
-    filterDim[1] = in_dim.c; 
-    filterDim[2] = kernelH; 
-    filterDim[3] = kernelW;
-    filterDim[4] = kernelL;
-    checkCUDNN( cudnnSetFilterNdDescriptor(filterDesc, 
-                net->dataType, 5, filterDim));
-
-    //convolutional descriptor
-    int padA[3] = {0, 0, 0};
-    int filterStride[3] = {strideH, strideW, strideL};
-    int upscale[3] = {1, 1, 1};
-    checkCUDNN( cudnnSetConvolutionNdDescriptor(convDesc, 3, 
-                padA, filterStride, upscale, CUDNN_CROSS_CORRELATION));
-
-
-    //get output dimension
-    int outputDim[5];
-    checkCUDNN(cudnnGetConvolutionNdForwardOutputDim(convDesc, srcTensorDesc, filterDesc, 5, outputDim));
-    n = outputDim[0];
-    c = outputDim[1];
-    h = outputDim[2];
-    w = outputDim[3];
-    l = outputDim[4];
-
-    //destination sensor
-    int outputStride[5];
-    outputStride[0] = c*h*w*l;
-    outputStride[1] = h*w*l;
-    outputStride[2] = w*l;
-    outputStride[3] = l;
-    outputStride[4] = 1;
-    checkCUDNN( cudnnSetTensorNdDescriptor(dstTensorDesc, net->dataType,
-                5, outputDim, outputStride));
-
-
-    //conv algo
-    checkCUDNN( cudnnGetConvolutionForwardAlgorithm(net->cudnnHandle,
-                srcTensorDesc, filterDesc, convDesc, dstTensorDesc,
-                CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo) );
-   
-    checkCUDNN( cudnnGetConvolutionForwardWorkspaceSize(net->cudnnHandle,
-                                            srcTensorDesc,
-                                            filterDesc,
-                                            convDesc,
-                                            dstTensorDesc,
-                                            algo,
-                                            &ws_sizeInBytes) );
-    if (ws_sizeInBytes!=0) 
-        checkCuda( cudaMalloc(&workSpace, ws_sizeInBytes) );
-    
-
-    // bias on N dimensional is not SUPPORTED so i have to use 2d method
-    //the trick is to upscale the 2d matrix width by the factor of 3d thickness 
-    checkCUDNN( cudnnSetTensor4dDescriptor(biasDstTensorDesc,
-                net->tensorFormat, net->dataType, n, c, h*l, w) );
-
-
-    checkCUDNN( cudnnSetTensor4dDescriptor(biasTensorDesc,
-                net->tensorFormat, net->dataType,
-                1, c, 1, 1) );
-
-    output_dim.n = n;
-    output_dim.c = c;
-    output_dim.h = h;
-    output_dim.w = w;
-    output_dim.l = l;
-
-    //allocate data for infer result
-    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) );
-}
-
-Conv3d::~Conv3d() {
-    
-    checkCUDNN( cudnnDestroyFilterDescriptor(filterDesc) );
-    checkCUDNN( cudnnDestroyConvolutionDescriptor(convDesc) );
-    checkCUDNN( cudnnDestroyTensorDescriptor(biasTensorDesc) );
-    checkCUDNN( cudnnDestroyTensorDescriptor(biasDstTensorDesc) );
-
-    if (ws_sizeInBytes!=0)
-        checkCuda( cudaFree(workSpace) );
-
-    checkCuda( cudaFree(dstData) );
-}
-
-value_type* Conv3d::infer(dataDim_t &dim, value_type* srcData) {
-
-
-    // convolution
-    value_type alpha = value_type(1);
-    value_type beta  = value_type(0);
-    checkCUDNN( cudnnConvolutionForward(net->cudnnHandle,
-                &alpha, srcTensorDesc, srcData, filterDesc,
-                data_d, convDesc, algo, workSpace, ws_sizeInBytes,
-                &beta, dstTensorDesc, dstData) );
-
-
-    // bias
-    alpha = value_type(1);
-    beta  = value_type(1);
-    checkCUDNN( cudnnAddTensor(net->cudnnHandle, CUDNN_ADD_SAME_C,
-                &alpha, biasTensorDesc, bias_d,
-                &beta, biasDstTensorDesc, dstData) );
-
-    //update data dimensions    
-    dim = output_dim;
-
-    return dstData;
-}
-
-}

src/Pooling.cpp

@@ -40,7 +40,7 @@ Pooling::Pooling(   Network *net, dataDim_t input_dim,
     }
 
     checkCUDNN( cudnnSetPooling2dDescriptor(poolingDesc, cudnnPoolingMode_t(pool_mode),
-                winH, winW, 0, 0, strideH, strideW) );
+                CUDNN_NOT_PROPAGATE_NAN, winH, winW, 0, 0, strideH, strideW) );
 
     checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc, 
                 net->tensorFormat, net->dataType, n, c, h, w) );

src/Softmax.cpp

@@ -0,0 +1,48 @@
+#include <iostream>
+
+#include "Layer.h"
+#include "kernels.h"
+
+namespace tkDNN {
+
+Softmax::Softmax(Network *net, dataDim_t input_dim) : 
+    Layer(net, input_dim) {
+
+    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) );
+
+    checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc,
+                                        net->tensorFormat,
+                                        net->dataType,
+                                        input_dim.n*input_dim.l, 
+                                        input_dim.c,
+                                        input_dim.h, input_dim.w) );
+    checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc,
+                                            net->tensorFormat,
+                                            net->dataType,
+                                            input_dim.n*input_dim.l, 
+                                            input_dim.c,
+                                            input_dim.h, input_dim.w) );
+}
+
+Softmax::~Softmax() {
+
+    checkCuda( cudaFree(dstData) );
+}
+
+value_type* Softmax::infer(dataDim_t &dim, value_type* srcData) {
+
+    value_type alpha = value_type(1);
+    value_type beta  = value_type(0);
+    checkCUDNN( cudnnSoftmaxForward(net->cudnnHandle,
+                                    CUDNN_SOFTMAX_ACCURATE ,
+                                    CUDNN_SOFTMAX_MODE_CHANNEL,
+                                    &alpha,
+                                    srcTensorDesc,
+                                    srcData,
+                                    &beta,
+                                    dstTensorDesc,
+                                    dstData) );
+    return dstData;
+}
+
+}

tests/build_models.sh

@@ -0,0 +1,11 @@
+#!/bin/bash
+echo "build test Model"
+cd test
+python test_model.py
+cd ..
+cd mnist
+python mnist_model.py
+cd ..
+echo "export weights"
+python weights_exporter.py test/net.h5 --output test/layers
+python caffe_weights_exporter.py mnist/lenet.prototxt mnist/lenet.caffemodel --output mnist/layers

tests/caffe_weights_exporter.py

@@ -0,0 +1,38 @@
+import argparse
+import os
+import msgpack
+import lmdb
+import random
+import caffe
+import numpy as np
+
+if __name__ == '__main__':
+    parser = argparse.ArgumentParser(description='CAFFE WEIGHTS EXPORTER TO CUDNN')
+    parser.add_argument('model',type=str,
+        help='Path to prototxt network model')
+    parser.add_argument('weights',type=str,
+        help='Path to caffemodel file')
+        
+    parser.add_argument('--output', type=str, help="output directory", default="layers")
+
+    args = parser.parse_args()
+
+    if not os.path.exists(args.output):
+        os.makedirs(args.output)
+
+    print "\n\n ====== NET LOADED ====== "
+    net = caffe.Net(args.model, args.weights, caffe.TEST)
+    n_lay = len(net.params)
+    print "Number of layers: ", n_lay
+    for i in xrange(n_lay):
+        key = net.params.keys()[i]
+        print "Layer", key
+        t = net.layer_dict[key].type 
+        print "    type: ", t
+        w = net.params[key][0].data
+        b = net.params[key][1].data 
+        print "    weights shape:", np.shape(w)
+        print "    bias shape:", np.shape(b)
+        
+        w.tofile(args.output + "/" + t + str(i) + ".bin", format="f")
+        b.tofile(args.output + "/" + t + str(i) + ".bias.bin", format="f")

tests/mnist/mnist_model.py

@@ -0,0 +1,32 @@
+#!/usr/bin/env python
+# mail:    admin@9crk.com
+# author:  9crk.from China.ShenZhen
+# time:    2017-03-22
+
+import caffe
+import numpy as np
+import cv2
+import sys
+import Image
+import matplotlib.pyplot as plt
+
+model = 'lenet.prototxt';
+weights = 'lenet.caffemodel';
+net = caffe.Net(model,weights,caffe.TEST);
+caffe.set_mode_gpu()
+img = np.array(np.random.rand(28,28), dtype=np.float32)
+#revert the image,and normalize it to 0-1 range
+
+print "INPUT: ", img
+img.tofile("input.bin", format="f")
+print "SHAPE: ", np.shape(img)
+out = net.forward_all(data=np.asarray([img]))
+
+out = out[out.keys()[0]]
+print out
+print np.shape(out)
+out.tofile("output.bin", format="f")
+#print out['prob'][0]
+#print out['prob'][0].argmax()
+
+

tests/mnist/test.cpp

@@ -0,0 +1,58 @@
+#include<iostream>
+#include "tkdnn.h"
+
+const char *input_bin   = "../tests/mnist/input.bin";
+const char *c0_bin      = "../tests/mnist/layers/Convolution0.bin";
+const char *c0_bias_bin = "../tests/mnist/layers/Convolution0.bias.bin";
+const char *c1_bin      = "../tests/mnist/layers/Convolution1.bin";
+const char *c1_bias_bin = "../tests/mnist/layers/Convolution1.bias.bin";
+const char *d2_bin      = "../tests/mnist/layers/InnerProduct2.bin";
+const char *d2_bias_bin = "../tests/mnist/layers/InnerProduct2.bias.bin";
+const char *d3_bin      = "../tests/mnist/layers/InnerProduct3.bin";
+const char *d3_bias_bin = "../tests/mnist/layers/InnerProduct3.bias.bin";
+const char *output_bin   = "../tests/mnist/output.bin";
+
+int main() {
+
+    // Network layout
+    tkDNN::Network net;
+    tkDNN::dataDim_t dim(1, 1, 28, 28, 1);
+    tkDNN::Layer *l;
+    l = new tkDNN::Conv2d     (&net, dim, 20, 5, 5, 1, 1, c0_bin, c0_bias_bin);
+    l = new tkDNN::Pooling    (&net, l->output_dim, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    l = new tkDNN::Conv2d     (&net, l->output_dim, 50, 5, 5, 1, 1, c1_bin, c1_bias_bin);
+    l = new tkDNN::Pooling    (&net, l->output_dim, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    l = new tkDNN::Dense      (&net, l->output_dim, 500, d2_bin, d2_bias_bin);
+    l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU);
+    l = new tkDNN::Dense      (&net, l->output_dim, 10, d3_bin, d3_bias_bin);
+    l = new tkDNN::Softmax    (&net, l->output_dim);
+ 
+    // Load input
+    value_type *data;
+    value_type *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+
+    printDeviceVector(dim.tot(), data);
+    dim.print(); //print initial dimension
+    
+    TIMER_START
+
+    // Inference
+    data = net.infer(dim, data);
+    
+    TIMER_STOP
+    dim.print();   
+
+    // Print result
+    std::cout<<"\n======= RESULT =======\n";
+    printDeviceVector(dim.tot(), data);
+ 
+    // Print real test
+    std::cout<<"\n==== CHECK RESULT ====\n";
+    value_type *out;
+    value_type *out_h;
+    readBinaryFile(output_bin, dim.tot(), &out_h, &out);
+    printDeviceVector(dim.tot(), out);
+ 
+    return 0;
+}

tests/test.cpp

@@ -1,57 +0,0 @@
-#include<iostream>
-#include "tkdnn.h"
-
-const char *input_bin   = "../tests/input.bin";
-const char *c0_bin      = "../tests/conv0.bin";
-const char *c0_bias_bin = "../tests/conv0.bias.bin";
-const char *c1_bin      = "../tests/conv1.bin";
-const char *c1_bias_bin = "../tests/conv1.bias.bin";
-const char *c2_bin      = "../tests/conv2.bin";
-const char *c2_bias_bin = "../tests/conv2.bias.bin";
-const char *d3_bin      = "../tests/dense3.bin";
-const char *d3_bias_bin = "../tests/dense3.bias.bin";
-const char *d4_bin      = "../tests/dense4.bin";
-const char *d4_bias_bin = "../tests/dense4.bias.bin";
-const char *d5_bin      = "../tests/dense5.bin";
-const char *d5_bias_bin = "../tests/dense5.bias.bin";
-
-int main() {
-
-    // Network layout
-    tkDNN::Network net;
-    tkDNN::dataDim_t dim(1, 1, 100, 100, 4);
-    tkDNN::Layer *l;
-    l = new tkDNN::MulAdd     (&net, dim, 2, -1);
-    l = new tkDNN::Conv3d     (&net, l->output_dim, 16, 8, 8, 2, 4, 4, 1,    c0_bin, c0_bias_bin);
-    l = new tkDNN::Activation (&net, l->output_dim, tkDNN::ACTIVATION_ELU);
-    l = new tkDNN::Pooling    (&net, l->output_dim, 2, 2, 2, 2, tkDNN::POOLING_AVERAGE);
-    l = new tkDNN::Conv3d     (&net, l->output_dim, 16, 4, 4, 2, 2, 2, 1,    c1_bin, c1_bias_bin);
-    l = new tkDNN::Activation (&net, l->output_dim, tkDNN::ACTIVATION_ELU);
-    l = new tkDNN::Conv3d     (&net, l->output_dim, 24, 3, 3, 2, 1, 1, 1,    c2_bin, c2_bias_bin);
-    l = new tkDNN::Activation (&net, l->output_dim, tkDNN::ACTIVATION_ELU);
-    l = new tkDNN::Flatten    (&net, l->output_dim);
-    l = new tkDNN::Dense      (&net, l->output_dim, 256,                     d3_bin, d3_bias_bin);
-    l = new tkDNN::Activation (&net, l->output_dim, tkDNN::ACTIVATION_ELU);
-    l = new tkDNN::Dense      (&net, l->output_dim, 32,                      d4_bin, d4_bias_bin);
-    l = new tkDNN::Activation (&net, l->output_dim, tkDNN::ACTIVATION_RELU);
-    l = new tkDNN::Dense      (&net, l->output_dim, 2,                       d5_bin, d5_bias_bin);
-
-
-    // Load input
-    value_type *data;
-    value_type *input_h;
-    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
-
-    dim.print(); //print initial dimension
-    
-    TIMER_START
-
-    // Inference
-    data = net.infer(dim, data); dim.print();
-
-    TIMER_STOP
-
-    // Print result
-    printDeviceVector(dim.tot(), data);
-    return 0;
-}

tests/test/test.cpp

@@ -0,0 +1,54 @@
+#include<iostream>
+#include "tkdnn.h"
+
+const char *input_bin   = "../tests/test/input.bin";
+const char *c0_bin      = "../tests/test/layers/conv0.bin";
+const char *c0_bias_bin = "../tests/test/layers/conv0.bias.bin";
+const char *c1_bin      = "../tests/test/layers/conv1.bin";
+const char *c1_bias_bin = "../tests/test/layers/conv1.bias.bin";
+const char *d2_bin      = "../tests/test/layers/dense2.bin";
+const char *d2_bias_bin = "../tests/test/layers/dense2.bias.bin";
+const char *output_bin   = "../tests/test/output.bin";
+
+int main() {
+
+    // Network layout
+    tkDNN::Network net;
+    tkDNN::dataDim_t dim(1, 1, 10, 10, 1);
+    tkDNN::Layer *l;
+    l = new tkDNN::Conv2d     (&net, dim, 2, 4, 4, 2, 2, c0_bin, c0_bias_bin);
+    l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU);
+    l = new tkDNN::Conv2d     (&net, l->output_dim, 4, 2, 2, 1, 1, c1_bin, c1_bias_bin);
+    l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU);
+    l = new tkDNN::Flatten      (&net, l->output_dim);
+    l = new tkDNN::Dense      (&net, l->output_dim, 4, d2_bin, d2_bias_bin);
+    l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU);
+
+    // Load input
+    value_type *data;
+    value_type *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+
+    printDeviceVector(dim.tot(), data);
+    dim.print(); //print initial dimension
+    
+    TIMER_START
+
+    // Inference
+    data = net.infer(dim, data); dim.print();
+    
+
+    TIMER_STOP
+
+    // Print result
+    std::cout<<"\n======= RESULT =======\n";
+    printDeviceVector(dim.tot(), data);
+
+    // Print real test
+    std::cout<<"\n==== CHECK RESULT ====\n";
+    value_type *out;
+    value_type *out_h;
+    readBinaryFile(output_bin, dim.tot(), &out_h, &out);
+    printDeviceVector(dim.tot(), out);
+    return 0;
+}

tests/test/test_model.py

@@ -0,0 +1,43 @@
+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.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
+
+def dense_model():
+    model = Sequential()
+
+    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()
+
+    model = dense_model()
+    model.save("net.h5")
+
+    grid = np.random.rand(10,10)
+    X = grid[None,:,:]
+    i = np.array(grid.flatten(), dtype=np.float32)
+    print i
+    i.tofile("input.bin", format="f")
+    print "Input: ", X    
+
+    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/test_model.py

@@ -1,61 +0,0 @@
-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.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
-from weights_exporter import *
-
-def dense_model():
-    model = Sequential()
-    model.add(Reshape((100, 100, 4, 1), input_shape=(100, 100, 4)))
-    model.add(Lambda(lambda x: 2*x - 1.,
-            batch_input_shape=(1, 100, 100, 4),         # 100by100by2
-            output_shape=(100, 100, 4, 1)))                         # 100by100by2
-    model.add(Convolution3D(16, kernel_size=(8, 8, 2), subsample=(4, 4, 1), border_mode="valid", 
-                bias_initializer="random_uniform"))
-    model.add(ELU())
-    model.add(AveragePooling3D(pool_size=(2, 2, 1)))
-    model.add(Convolution3D(16, kernel_size=(4, 4, 2), subsample=(2, 2, 1), border_mode="valid", 
-                bias_initializer="random_uniform"))
-    model.add(ELU())
-    model.add(Convolution3D(24, kernel_size=(3, 3, 2), subsample=(1, 1, 1), border_mode="valid", 
-                bias_initializer="random_uniform"))
-    model.add(ELU())
-    model.add(Flatten())
-    model.add(Dense(256, bias_initializer="random_uniform"))
-    model.add(ELU())
-    model.add(Dense(32, activation="relu", bias_initializer="random_uniform"))
-    model.add(Dense(2, bias_initializer="random_uniform"))
-
-    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()
-
-    model = dense_model()
-    wg = model.get_weights()
-    export_conv3d("conv0",  wg[0], wg[1])
-    export_conv3d("conv1",  wg[2], wg[3])
-    export_conv3d("conv2",  wg[4], wg[5])
-    export_dense ("dense3", wg[6], wg[7])
-    export_dense ("dense4", wg[8], wg[9])
-    export_dense ("dense5", wg[10], wg[11])
-
-    grid = np.random.rand(100, 100,4)
-    X = grid[None,:,:]
-    i = np.array(grid.flatten(), dtype=np.float32)
-    print i
-    i.tofile("input.bin", format="f")
-    print "Input: ", X    
-
-    r = model.predict( X, batch_size=1)
-    print np.shape(r)  
-    print "Result: ", r
-    print "Result shape: ", np.shape(r) 

tests/weights_exporter.py

@@ -99,9 +99,7 @@ if __name__ == '__main__':
     parser = argparse.ArgumentParser(description='KERAS WEIGHTS EXPORTER TO CUDNN')
     parser.add_argument('model',type=str,
         help='Path to model h5 file. Model should be on the same path.')
-    parser.add_argument('layers', type=str, help="layers list [ dense, conv2d ]", nargs='+')
     parser.add_argument('--output', type=str, help="output directory", default="layers")
-    parser.add_argument('--test_db', type=str, help="input db to test", default=None)
 
     args = parser.parse_args()
     
@@ -119,34 +117,17 @@ if __name__ == '__main__':
 
     num = 0
     name_num = 0
-    for i in args.layers:
+    for l in model.layers:
-        if i == "conv3d":
+        name = l.name        
+        if name.startswith("conv3d"):
             export_conv3d(args.output + "/conv" + str(name_num), weights[num], weights[num+1])
-        elif i == "conv2d":
+        elif name.startswith("conv2d"):
             export_conv2d(args.output + "/conv" + str(name_num), weights[num], weights[num+1])
-        elif i == "dense":
+        elif name.startswith("dense"):
             export_dense(args.output + "/dense" + str(name_num), weights[num], weights[num+1])
         else:
-            print "error: ", i, "is not a layer type"
+            print "skip:", name, "has no weights"
-            break
+            continue
         name_num += 1
         num += 2
 
-    if args.test_db != None:
-        print "Test on db: ", args.test_db
-        db = lmdb.open(args.test_db, subdir=False, readonly=True, lock=False) 
-        txn = db.begin()
-
-        s = random.randint(0, txn.stat()["entries"]-1)
-        print "camp number: ", s
-        s = txn.get(str(s))
-        c = msgpack.unpackb(s)
-    
-        print "Steer, throttle: ", c["actuators"]
-        print "Speed (m/s): ", c["speed"]
-        grid = np.asarray(c["bitmap"], np.float32)
-        i = np.array(grid.flatten(), dtype=np.float32)
-        i.tofile(args.output + "input.bin", format="f")
-        X = grid[None, :, :]
-    
-        print "Prediction: ", model.predict(X)