conv2d implementation

CMakeLists.txt

@@ -8,7 +8,9 @@ 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)
+add_library(tkDNN SHARED src/Layer.cpp src/LayerWgs.cpp 
+    src/Dense.cpp src/Activation.cpp src/Conv2d.cpp
+    src/Network.cpp src/utils.cpp)
 target_link_libraries(tkDNN kernels)
 
 add_executable(tkDNNtest tests/test.cpp)

include/Layer.h

@@ -9,6 +9,11 @@ namespace tkDNN {
 
 /**
     Data rapresentation beetween layers
+    n = batch size
+    c = channels
+    h = heigth (lines)
+    w = width  (rows)
+    l = lenght (3rd dimension)
 */
 struct dataDim_t {
 
@@ -28,6 +33,7 @@ struct dataDim_t {
     }
 };
 
+
 /** 
     Simple layer Father class
 */
@@ -49,6 +55,7 @@ protected:
     cudnnTensorDescriptor_t srcTensorDesc, dstTensorDesc;
 };
 
+
 /**
     Father class of all layer that need to load trained weights
 */
@@ -68,6 +75,7 @@ protected:
     value_type *bias_h, *bias_d;
 };
 
+
 /**
     Dense (full interconnection) layer
 */
@@ -85,7 +93,7 @@ protected:
 };
 
 /**
-    Activation layer (it doesnt need weigths)
+    Avaible activation functions
 */
 typedef enum {
     ACTIVATION_SIGMOID = 0,
@@ -94,6 +102,9 @@ typedef enum {
     ACTIVATION_ELU     = 100
 } tkdnnActivationMode_t;
 
+/**
+    Activation layer (it doesnt need weigths)
+*/
 class Activation : public Layer {
 
 public:
@@ -107,5 +118,32 @@ protected:
     value_type *dstData;  //where results will be putted
 };
 
+
+/**
+    Convolutional 2D layer
+*/
+class Conv2d : public LayerWgs {
+
+public:
+    Conv2d(Network *net, dataDim_t in_dim, int out_ch,
+            int kernelH, int kernelW, int strideH, int strideW,
+            const char* fname_weights, const char* fname_bias); 
+    virtual ~Conv2d();
+
+    value_type* infer(dataDim_t &dim, value_type* srcData);
+
+protected:
+    value_type *dstData;  //where results will be putted
+    int kernelH, kernelW, strideH, strideW;
+
+    cudnnTensorDescriptor_t biasTensorDesc;
+    cudnnFilterDescriptor_t filterDesc;
+    cudnnConvolutionDescriptor_t convDesc;
+    cudnnConvolutionFwdAlgo_t algo;
+
+    void*  workSpace;
+    size_t ws_sizeInBytes;
+};
+
 }
 #endif //LAYER_H

src/Conv2d.cpp

@@ -0,0 +1,112 @@
+#include <iostream>
+
+#include "Layer.h"
+
+namespace tkDNN {
+
+Conv2d::Conv2d( Network *net, dataDim_t in_dim, int out_ch, 
+                int kernelH, int kernelW, int strideH, int strideW,
+                const char* fname_weights, const char* fname_bias) : 
+    
+    LayerWgs(net, in_dim, in_dim.c, out_ch, kernelH, kernelW, 1, 
+             fname_weights, fname_bias) {
+
+    this->kernelH = kernelH;
+    this->kernelW = kernelW;
+    this->strideH = strideH;
+    this->strideW = strideW;
+
+    checkCUDNN( cudnnCreateTensorDescriptor(&biasTensorDesc) );
+    checkCUDNN( cudnnCreateFilterDescriptor(&filterDesc) );
+    checkCUDNN( cudnnCreateConvolutionDescriptor(&convDesc) );
+
+    int n = input_dim.n;
+    int c = input_dim.c;
+    int h = input_dim.h;
+    int w = input_dim.w;
+
+    checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc,
+                net->tensorFormat, net->dataType, n, c, h, w) );
+
+    checkCUDNN( cudnnSetFilter4dDescriptor(filterDesc,
+                net->dataType, out_ch, input_dim.c, 
+                kernelH, kernelW) );
+
+    checkCUDNN( cudnnSetConvolution2dDescriptor(convDesc,
+                0,0, // padding
+                strideH, strideW, // stride
+                1,1, // upscale
+                CUDNN_CROSS_CORRELATION) );
+
+    // find dimension of convolution output
+    checkCUDNN( cudnnGetConvolution2dForwardOutputDim(
+                convDesc, srcTensorDesc, filterDesc,
+                &n, &c, &h, &w) );
+
+    checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc,
+                net->tensorFormat, net->dataType, n, c, h, w) );
+                
+    checkCUDNN( cudnnGetConvolutionForwardAlgorithm(net->cudnnHandle,
+                srcTensorDesc, filterDesc, convDesc, dstTensorDesc,
+                CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo) );
+ 
+    workSpace = NULL;
+    ws_sizeInBytes = 0;
+
+    checkCUDNN( cudnnGetConvolutionForwardWorkspaceSize(net->cudnnHandle,
+                srcTensorDesc, filterDesc, convDesc, dstTensorDesc,
+                algo, &ws_sizeInBytes) );
+
+    if (ws_sizeInBytes!=0) {
+        checkCuda( cudaMalloc(&workSpace, ws_sizeInBytes) );
+    }
+
+
+    checkCUDNN( cudnnSetTensor4dDescriptor(biasTensorDesc,
+                net->tensorFormat, net->dataType,
+                1, out_ch, 1, 1) );
+
+
+    output_dim.n = n;
+    output_dim.c = c;
+    output_dim.h = h;
+    output_dim.w = w;
+    output_dim.l = 1;
+
+    //allocate data for infer result
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) );
+}
+
+Conv2d::~Conv2d() {
+    
+    if (ws_sizeInBytes!=0)
+        checkCuda( cudaFree(workSpace) );
+
+    checkCuda( cudaFree(dstData) );
+}
+
+value_type* Conv2d::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, dstTensorDesc, dstData) );
+
+    //update data dimensions    
+    dim = output_dim;
+
+    return dstData;
+}
+
+}

src/kernels/activation_elu.cu

@@ -1,5 +1,11 @@
 #include "kernels.h"
 
+/**
+    Exponential Linear Unit compute kernel
+    it does the following operation for each x input element:
+        x < 0 :    y = e^(x) -1
+        x > 0 :    y = x
+*/
 __global__
 void activation_elu(value_type *input, value_type *output, int size) {
 
@@ -7,9 +13,14 @@ void activation_elu(value_type *input, value_type *output, int size) {
     
     if(i<size) 
         output[i] =  (input[i]>0)*input[i] + (input[i]<0)*(expf(input[i]) -1);
+    
+    // the if x > or < is condensed in one operation for better threads flow 
  }
 
 
+/**
+    ELU activation function
+*/
 void activationELUForward(value_type* srcData, value_type* dstData, int size)
 {
     activation_elu<<<(size+255)/256, 256>>>(srcData, dstData, size);

tests/simple_dense.py

@@ -2,7 +2,7 @@ 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 import Input, Dense, Activation, Flatten, Dropout, ELU, Reshape
 from keras.layers.convolutional import Convolution2D, Convolution3D
 from keras.layers.pooling import MaxPooling2D, MaxPooling3D
 from keras.models import Sequential, Model
@@ -12,33 +12,34 @@ from weights_exporter import *
 
 def dense_model():
     model = Sequential()
-    model.add(Dense(256, input_shape=(1, 512)))
+
+    model.add(Reshape((10, 10, 1), input_shape=(10, 10)))
+    model.add(Convolution2D(2, (4, 4), subsample=(2, 2), 
+                           bias_initializer='random_uniform'))
     model.add(ELU())
-    model.add(Dense(32))
+    model.add(Convolution2D(4, (2, 2), subsample=(1, 1), 
-    model.add(ELU())
+                           bias_initializer='random_uniform', activation="relu"))
-    model.add(Dense(2))
-    
     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_conv2d("conv0", wg[0], wg[1])
+    export_conv2d("conv1", wg[2], wg[3])
 
-    export_dense("dense0", wg[0], wg[1])
+    grid = np.random.rand(10,10)
-    export_dense("dense1", wg[2], wg[3])
+    X = grid[None,:,:]
-    export_dense("dense2", wg[4], wg[5])
+    i = np.array(grid.flatten(), dtype=np.float32)
-
+    print i
-    model.set_weights(wg)
-   
-    X = np.random.rand(1, 512)
-    i = np.array(X, dtype=np.float32)
     i.tofile("input.bin", format="f")
+    print "Input: ", X    
 
-    print "Input: ", i
+    r = model.predict( X, batch_size=1)
-    r = model.predict( X[None, :], batch_size=1)
+    print np.shape(r)  
-
     print "Result: ", r
     print "Result shape: ", np.shape(r) 

tests/test.cpp

@@ -2,10 +2,10 @@
 #include "Layer.h"
 
 const char *input_bin   = "../tests/input.bin";
-const char *d0_bin      = "../tests/dense0.bin";
+const char *c0_bin      = "../tests/conv0.bin";
-const char *d0_bias_bin = "../tests/dense0.bias.bin";
+const char *c0_bias_bin = "../tests/conv0.bias.bin";
-const char *d1_bin      = "../tests/dense1.bin";
+const char *c1_bin      = "../tests/conv1.bin";
-const char *d1_bias_bin = "../tests/dense1.bias.bin";
+const char *c1_bias_bin = "../tests/conv1.bias.bin";
 const char *d2_bin      = "../tests/dense2.bin";
 const char *d2_bias_bin = "../tests/dense2.bias.bin";
 
@@ -13,13 +13,12 @@ int main() {
 
     // Network layout
     tkDNN::Network net;
-    tkDNN::dataDim_t dim(1, 512, 1, 1);
+    tkDNN::dataDim_t dim(1, 1, 10, 10);
-    tkDNN::Dense      d0 (&net, dim,           256, d0_bin, d0_bias_bin);
+    tkDNN::Conv2d     c0 (&net, dim,           2, 4, 4, 2, 2, c0_bin, c0_bias_bin);
-    tkDNN::Activation a0 (&net, d0.output_dim, tkDNN::ACTIVATION_ELU);
+    tkDNN::Activation a0 (&net, c0.output_dim, tkDNN::ACTIVATION_ELU);
-    tkDNN::Dense      d1 (&net, a0.output_dim, 32,  d1_bin, d1_bias_bin);
+    tkDNN::Conv2d     c1 (&net, a0.output_dim, 4, 2, 2, 1, 1, c1_bin, c1_bias_bin);
-    tkDNN::Activation a1 (&net, d1.output_dim, tkDNN::ACTIVATION_ELU);
+    tkDNN::Activation a1 (&net, c1.output_dim, tkDNN::ACTIVATION_RELU);
-    tkDNN::Dense      d2 (&net, a1.output_dim, 2,   d2_bin, d2_bias_bin);
+
-    
     // Load input
     value_type *data;
     value_type *input_h;
@@ -27,13 +26,15 @@ int main() {
 
     dim.print(); //print initial dimension
     
+    TIMER_START
+
     // Inference
-    data = d0.infer(dim, data); dim.print();
+    data = c0.infer(dim, data); dim.print();
     data = a0.infer(dim, data); dim.print();
-    data = d1.infer(dim, data); dim.print();
+    data = c1.infer(dim, data); dim.print();
     data = a1.infer(dim, data); dim.print();
-    data = d2.infer(dim, data); dim.print();
+
-  
+    TIMER_STOP
     // Print result
     printDeviceVector(dim.tot(), data);
     return 0;