Pooling and big test, ELU seem not to work

CMakeLists.txt

@@ -9,7 +9,7 @@ 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/Dense.cpp src/Activation.cpp src/Conv2d.cpp src/Conv3d.cpp src/Flatten.cpp src/MulAdd.cpp src/Pooling.cpp
     src/Network.cpp src/utils.cpp)
 target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDA_TOOLKIT_ROOT_DIR}/lib/libcudnn.so)
 

include/Layer.h

@@ -209,5 +209,40 @@ protected:
     value_type *dstData, *add_vector;  //where results will be putted
 };
 
+
+
+/**
+    Avaible pooling functions (padding on tkDNN is not supported)
+*/
+typedef enum {
+    POOLING_MAX     = 0,
+    POOLING_AVERAGE = 1,                  // count for average includes padded values
+    POOLING_AVERAGE_EXCLUDE_PADDING = 2   // count for average does not include padded values
+} tkdnnPoolingMode_t;
+
+/**
+    Pooling layer
+    currenty supported only 2d pooing (also on 3d input)
+*/
+class Pooling : public Layer {
+
+public:
+    Pooling(Network *net, dataDim_t input_dim, int winH, int winW, 
+            int strideH, int strideW, tkdnnPoolingMode_t pool_mode); 
+    virtual ~Pooling();
+
+    value_type* infer(dataDim_t &dim, value_type* srcData);
+
+protected:
+
+    cudnnPoolingDescriptor_t poolingDesc;
+
+    int winH, winW;
+    int strideH, strideW;
+    tkdnnPoolingMode_t pool_mode;
+    value_type *dstData, *tmpInputData, *tmpOutputData;  //where results will be putted
+    bool poolOn3d;
+};
+
 }
 #endif //LAYER_H

src/Pooling.cpp

@@ -0,0 +1,111 @@
+#include <iostream>
+
+#include "Layer.h"
+#include "kernels.h"
+
+namespace tkDNN {
+
+Pooling::Pooling(   Network *net, dataDim_t input_dim, 
+                    int winH, int winW, int strideH, int strideW, tkdnnPoolingMode_t pool_mode) : 
+    Layer(net, input_dim) {
+
+
+    if(winH != strideH || winW != strideW)
+        FatalError("stride pooling not yet implemented");
+
+    this->winH = winH;
+    this->winW = winW;
+    this->strideH = strideH;
+    this->strideW = strideW;
+    this->pool_mode = pool_mode;
+    
+    checkCUDNN( cudnnCreatePoolingDescriptor(&poolingDesc) );
+
+    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;
+
+    poolOn3d = false;
+    
+    if(l > 1) { 
+        poolOn3d = true;
+    
+        if(n != 1)
+            FatalError("N value on 3d pool must be 1");
+
+        //use batch as l
+        n = l;
+    }
+
+    checkCUDNN( cudnnSetPooling2dDescriptor(poolingDesc, cudnnPoolingMode_t(pool_mode),
+                winH, winW, 0, 0, strideH, strideW) );
+
+    checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc, 
+                net->tensorFormat, net->dataType, n, c, h, w) );
+
+    //get out dim
+    h = h / winH; w = w / winW;
+    
+    checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc,
+                net->tensorFormat, net->dataType, n, c, h, w) );
+       
+
+    output_dim.n = n;
+    output_dim.c = c;
+    output_dim.h = h;
+    output_dim.w = w;
+    output_dim.l = l;
+
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) );
+
+    //pool on 3d data need transposition at the enter and on the exit
+    //allocate for initial and final transposition
+    if(poolOn3d) {
+        output_dim.n = 1;
+
+        checkCuda( cudaMalloc(&tmpInputData, input_dim.tot()*sizeof(value_type)) );
+        checkCuda( cudaMalloc(&tmpOutputData, output_dim.tot()*sizeof(value_type)) );
+    }
+
+}
+
+Pooling::~Pooling() {
+
+    if(poolOn3d) {
+        checkCuda( cudaFree(tmpInputData) );
+        checkCuda( cudaFree(tmpOutputData) );
+    }
+
+    checkCUDNN( cudnnDestroyPoolingDescriptor(poolingDesc) );
+    checkCuda( cudaFree(dstData) );
+}
+
+value_type* Pooling::infer(dataDim_t &dim, value_type* srcData) {
+
+    value_type *poolSrc = srcData;
+    value_type *poolDst = dstData;
+
+    if(poolOn3d) {
+        matrixTranspose(net->cublasHandle, srcData, tmpInputData, dim.h*dim.w*dim.c, dim.l);
+        poolSrc = tmpInputData;
+        poolDst = tmpOutputData;
+    }
+
+    value_type alpha = value_type(1);
+    value_type beta = value_type(0);
+    checkCUDNN( cudnnPoolingForward(net->cudnnHandle, poolingDesc,
+                &alpha, srcTensorDesc, poolSrc,
+                &beta, dstTensorDesc, poolDst) );
+
+    //update dim
+    dim = output_dim;
+
+    if(poolOn3d)
+        matrixTranspose(net->cublasHandle, tmpOutputData, dstData, dim.l, dim.h*dim.w*dim.c);
+
+    return dstData;
+}
+
+}

tests/simple_dense.py

@@ -2,9 +2,9 @@ import keras
 import numpy as np
 import pickle
 from keras.models import Sequential
-from keras.layers import Input, Dense, Activation, Flatten, Dropout, ELU, Reshape
+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
+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
@@ -12,14 +12,24 @@ from weights_exporter import *
 
 def dense_model():
     model = Sequential()
-
+    model.add(Reshape((100, 100, 4, 1), input_shape=(100, 100, 4)))
-    model.add(Reshape((10, 10, 4, 1), input_shape=(10, 10, 4)))
+    model.add(Lambda(lambda x: 2*x - 1.,
-    model.add(Convolution3D(2, (4, 4, 2), subsample=(2, 2, 1), activation="relu",
+            batch_input_shape=(1, 100, 100, 4),         # 100by100by2
-                           bias_initializer='random_uniform'))
+            output_shape=(100, 100, 4, 1)))                         # 100by100by2
-    model.add(Convolution3D(4, (2, 2, 2), subsample=(1, 1, 1), 
+    model.add(Convolution3D(16, kernel_size=(8, 8, 2), subsample=(4, 4, 1), border_mode="valid", 
-                           bias_initializer='random_uniform'))
+                bias_initializer="random_uniform", activation="relu"))
-    model.add(ELU())
+    #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", activation="relu"))
+    model.add(Convolution3D(24, kernel_size=(3, 3, 2), subsample=(1, 1, 1), border_mode="valid", 
+                bias_initializer="random_uniform", activation="relu"))
+    #model.add(ELU())
     model.add(Flatten())
+    model.add(Dense(256, activation="relu", bias_initializer="random_uniform"))
+    model.add(Dense(32, activation="relu", bias_initializer="random_uniform"))
+    #model.add(ELU())
+    model.add(Dense(2, bias_initializer="random_uniform"))
 
     sgd = keras.optimizers.Adam(lr=1e-4, decay=1e-8)
     model.compile(optimizer=sgd, loss="mse")
@@ -31,10 +41,14 @@ if __name__ == '__main__':
 
     model = dense_model()
     wg = model.get_weights()
-    export_conv3d("conv0", wg[0], wg[1])
+    export_conv3d("conv0",  wg[0], wg[1])
-    export_conv3d("conv1", wg[2], wg[3])
+    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(10,10,4)
+    grid = np.random.rand(100, 100,4)
     X = grid[None,:,:]
     i = np.array(grid.flatten(), dtype=np.float32)
     print i

tests/test.cpp

@@ -6,20 +6,35 @@ 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 *d2_bin      = "../tests/dense2.bin";
+const char *c2_bin      = "../tests/conv2.bin";
-const char *d2_bias_bin = "../tests/dense2.bias.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, 10, 10, 4);
+    tkDNN::dataDim_t dim(1, 1, 100, 100, 4);
-    tkDNN::Conv3d     c0 (&net, dim,           2, 4, 4, 2, 2, 2, 1, c0_bin, c0_bias_bin);
+    tkDNN::MulAdd     m0 (&net, dim, 2, -1);
+    tkDNN::Conv3d     c0 (&net, m0.output_dim, 16, 8, 8, 2, 4, 4, 1,    c0_bin, c0_bias_bin);
     tkDNN::Activation a0 (&net, c0.output_dim, tkDNN::ACTIVATION_RELU);
-    tkDNN::Conv3d     c1 (&net, a0.output_dim, 4, 2, 2, 2, 1, 1, 1, c1_bin, c1_bias_bin);
+    tkDNN::Pooling    p0 (&net, a0.output_dim, 2, 2, 2, 2, tkDNN::POOLING_AVERAGE);
-    tkDNN::Activation a1 (&net, c1.output_dim, tkDNN::ACTIVATION_ELU);
+    tkDNN::Conv3d     c1 (&net, p0.output_dim, 16, 4, 4, 2, 2, 2, 1,    c1_bin, c1_bias_bin);
-    tkDNN::Flatten    f1 (&net, a1.output_dim);
+    tkDNN::Activation a1 (&net, c1.output_dim, tkDNN::ACTIVATION_RELU);
-    tkDNN::MulAdd     m1 (&net, f1.output_dim, 2, 1);
+    tkDNN::Conv3d     c2 (&net, a1.output_dim, 24, 3, 3, 2, 1, 1, 1,    c2_bin, c2_bias_bin);
+    tkDNN::Activation a2 (&net, c2.output_dim, tkDNN::ACTIVATION_RELU);
+    tkDNN::Flatten    f2 (&net, a2.output_dim);
+    tkDNN::Dense      d3 (&net, f2.output_dim, 256,                     d3_bin, d3_bias_bin);
+    tkDNN::Activation a3 (&net, d3.output_dim, tkDNN::ACTIVATION_RELU);
+    tkDNN::Dense      d4 (&net, a3.output_dim, 32,                      d4_bin, d4_bias_bin);
+    tkDNN::Activation a4 (&net, d4.output_dim, tkDNN::ACTIVATION_RELU);
+    tkDNN::Dense      d5 (&net, a4.output_dim, 2,                       d5_bin, d5_bias_bin);
+
 
     // Load input
     value_type *data;
@@ -31,14 +46,23 @@ int main() {
     TIMER_START
 
     // Inference
+    data = m0.infer(dim, data); dim.print();
     data = c0.infer(dim, data); dim.print();
     data = a0.infer(dim, data); dim.print();
+    data = p0.infer(dim, data); dim.print();
     data = c1.infer(dim, data); dim.print();
     data = a1.infer(dim, data); dim.print();
-    data = f1.infer(dim, data); dim.print();
+    data = c2.infer(dim, data); dim.print();
-    data = m1.infer(dim, data); dim.print();
+    data = a2.infer(dim, data); dim.print();
-
+    data = f2.infer(dim, data); dim.print();
+    data = d3.infer(dim, data); dim.print();
+    data = a3.infer(dim, data); dim.print();
+    data = d4.infer(dim, data); dim.print();
+    data = a4.infer(dim, data); dim.print();
+    data = d5.infer(dim, data); dim.print();
+    
     TIMER_STOP
+
     // Print result
     printDeviceVector(dim.tot(), data);
     return 0;