ELU implemented

.gitignore

@@ -1,3 +1,4 @@
 *~
 build/
 .vscode/
+*.bin

CMakeLists.txt

@@ -4,10 +4,14 @@ project (tkDNN)
 
 find_package(CUDA QUIET REQUIRED)
 
+cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS})
+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/Network.cpp src/utils.cpp)
+target_link_libraries(tkDNN kernels)
 
 add_executable(tkDNNtest tests/test.cpp)
 message(${CUDA_LIBRARIES})
-target_link_libraries(tkDNNtest tkDNN ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} /usr/local/cuda-8.0/cudnn/libcudnn.so)
+target_link_libraries(tkDNNtest tkDNN 
+    ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDA_TOOLKIT_ROOT_DIR}/lib/libcudnn.so)

include/Layer.h

@@ -42,9 +42,10 @@ public:
         return NULL;
     }
 
+    dataDim_t input_dim, output_dim;
+
 protected:
     Network *net;
-    dataDim_t input_dim;
     cudnnTensorDescriptor_t srcTensorDesc, dstTensorDesc;
 };
 
@@ -81,22 +82,28 @@ public:
 
 protected:
     value_type *dstData;  //where results will be putted
-    int out_ch;
 };
 
 /**
     Activation layer (it doesnt need weigths)
 */
+typedef enum {
+    ACTIVATION_SIGMOID = 0,
+    ACTIVATION_RELU    = 1,
+    ACTIVATION_TANH    = 2,
+    ACTIVATION_ELU     = 100
+} tkdnnActivationMode_t;
+
 class Activation : public Layer {
 
 public:
-    Activation(Network *net, dataDim_t input_dim, cudnnActivationMode_t act_mode); 
+    Activation(Network *net, dataDim_t input_dim, tkdnnActivationMode_t act_mode); 
     virtual ~Activation();
 
     value_type* infer(dataDim_t &dim, value_type* srcData);
 
 protected:
-    cudnnActivationMode_t act_mode;
+    tkdnnActivationMode_t act_mode;
     value_type *dstData;  //where results will be putted
 };
 

include/kernels.h

@@ -0,0 +1,3 @@
+#include "utils.h"
+
+void activationELUForward(value_type* srcData, value_type* dstData, int size);

src/Activation.cpp

@@ -1,10 +1,11 @@
 #include <iostream>
 
 #include "Layer.h"
+#include "kernels.h"
 
 namespace tkDNN {
 
-Activation::Activation(Network *net, dataDim_t input_dim, cudnnActivationMode_t act_mode) : 
+Activation::Activation(Network *net, dataDim_t input_dim, tkdnnActivationMode_t act_mode) : 
     Layer(net, input_dim) {
 
     this->act_mode = act_mode;
@@ -29,17 +30,21 @@ Activation::~Activation() {
 
 value_type* Activation::infer(dataDim_t &dim, value_type* srcData) {
 
-    value_type alpha = value_type(1);
-    value_type beta  = value_type(0);
-    checkCUDNN( cudnnActivationForward(net->cudnnHandle,
-                                        act_mode,
-                                        &alpha,
-                                        srcTensorDesc,
-                                        srcData,
-                                        &beta,
-                                        dstTensorDesc,
-                                        dstData) );    
+    if(act_mode == ACTIVATION_ELU) {
+        activationELUForward(srcData, dstData, dim.tot());
 
+    } else {
+        value_type alpha = value_type(1);
+        value_type beta  = value_type(0);
+        checkCUDNN( cudnnActivationForward(net->cudnnHandle,
+                                            cudnnActivationMode_t(act_mode),
+                                            &alpha,
+                                            srcTensorDesc,
+                                            srcData,
+                                            &beta,
+                                            dstTensorDesc,
+                                            dstData) );    
+    }
     return dstData;
 }
 

src/Dense.cpp

@@ -8,9 +8,14 @@ Dense::Dense(Network *net, dataDim_t in_dim,
     int out_ch, const char* fname_weights, const char* fname_bias) : 
     LayerWgs(net, in_dim, in_dim.tot(), out_ch, 1, 1, 1, fname_weights, fname_bias) {
 
-    this->out_ch = out_ch;
+    output_dim.n = 1;
+    output_dim.c = out_ch;
+    output_dim.h = 1;
+    output_dim.w = 1;
+    output_dim.l = 1;
+
     //allocate data for infer result
-    checkCuda( cudaMalloc(&dstData, outputs*sizeof(value_type)) );
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) );
 }
 
 Dense::~Dense() {
@@ -24,9 +29,9 @@ value_type* Dense::infer(dataDim_t &dim, value_type* srcData) {
         FatalError("Not Implemented"); 
     
     int dim_x = dim.tot();
-    int dim_y = outputs;
+    int dim_y = output_dim.tot();
 
-    if (dim_x != inputs)
+    if (dim_x != input_dim.tot())
         FatalError("Input mismatch");
 
     value_type alpha = value_type(1), beta = value_type(1);

src/Layer.cpp

@@ -8,7 +8,8 @@ Layer::Layer(Network *net, dataDim_t in_dim) {
 
     this->net = net;
     this->input_dim = in_dim;
-
+    this->output_dim = in_dim;
+    
     checkCUDNN( cudnnCreateTensorDescriptor(&srcTensorDesc) );
     checkCUDNN( cudnnCreateTensorDescriptor(&dstTensorDesc) );
 }

src/kernels/activation_elu.cu

@@ -0,0 +1,17 @@
+#include "kernels.h"
+
+__global__
+void activation_elu(value_type *input, value_type *output, int size) {
+
+    int i = threadIdx.x*(blockIdx.x +1);
+    
+    if(i<size) 
+        output[i] =  (input[i]>0)*input[i] + (input[i]<0)*(expf(input[i]) -1);
+ }
+
+
+void activationELUForward(value_type* srcData, value_type* dstData, int size)
+{
+    activation_elu<<<(size+255)/256, 256>>>(srcData, dstData, size);
+    checkCuda( cudaDeviceSynchronize() );
+}

tests/simple_dense.py

@@ -0,0 +1,44 @@
+import keras
+import numpy as np
+import pickle
+from keras.models import Sequential
+from keras.layers import Input, Dense, Activation, Flatten, Dropout, ELU
+from keras.layers.convolutional import Convolution2D, Convolution3D
+from keras.layers.pooling import MaxPooling2D, MaxPooling3D
+from keras.models import Sequential, Model
+from keras.layers import Cropping2D
+import keras.backend.tensorflow_backend as KTF
+
+
+def dense_model(inp, out):
+    model = Sequential()
+    model.add(Dense(out, input_shape=(1, inp)))
+    model.add(ELU())
+
+    sgd = keras.optimizers.Adam(lr=1e-4, decay=1e-8)
+    model.compile(optimizer=sgd, loss="mse")
+
+    return model
+
+
+if __name__ == '__main__':
+
+    model = dense_model(8, 2)
+    wg = model.get_weights()
+    w = np.squeeze(wg[0])
+    w = np.array([ i[0] for i in w ] + [ i[1] for i in w ], dtype=np.float32) 
+    b = np.squeeze(wg[1])
+    
+    print "weigths: ", w
+    print "bias: ", b
+    w.tofile("dense.bin", format="f")
+    b.tofile("dense.bias.bin", format="f")
+   
+    X = np.array([[[0,1,2,3,4,5,6,7]]], dtype=np.float32)
+    i = np.squeeze(X[0][0])
+    print "input: ", i    
+    i.tofile("input.bin", format="f")
+    
+    r = model.predict( X, batch_size=1)
+    print "Result: ", r
+    print "Result shape: ", np.shape(r) 

tests/test.cpp

@@ -3,10 +3,22 @@
 
 int main() {
 
-    tkDNN::dataDim_t dim(1, 1, 10, 10);
-    dim.print();
     tkDNN::Network net;
-    tkDNN::Dense d(&net, dim, 2, "ci", "lol");
+    tkDNN::dataDim_t dim(1, 8, 1, 1);
+    tkDNN::Dense      d(&net, dim, 2, "../tests/dense.bin", "../tests/dense.bias.bin");
+    tkDNN::Activation a(&net, d.output_dim, tkDNN::ACTIVATION_ELU);
+
+    value_type *data;
+    value_type *input_h;
+    readBinaryFile("../tests/input.bin", 8, &input_h, &data);
+
+    dim.print();
+    data = d.infer(dim, data);
+    dim.print();
+    data = a.infer(dim, data);
+    dim.print();
+
+    printDeviceVector(dim.tot(), data);
 
     return 0;
 }