From ba215c35c033ce267ebd29fada666a4ce899e9ab Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 29 Jun 2017 11:37:29 +0000 Subject: [PATCH] Pooling and big test, ELU seem not to work --- CMakeLists.txt | 2 +- include/Layer.h | 35 +++++++++++++ src/Pooling.cpp | 111 ++++++++++++++++++++++++++++++++++++++++++ tests/simple_dense.py | 38 ++++++++++----- tests/test.cpp | 46 ++++++++++++----- 5 files changed, 208 insertions(+), 24 deletions(-) create mode 100644 src/Pooling.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index c79fc8c..240f650 100644 --- a/CMakeLists.txt +++ b/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) diff --git a/include/Layer.h b/include/Layer.h index 927154e..9ec37d4 100644 --- a/include/Layer.h +++ b/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 \ No newline at end of file diff --git a/src/Pooling.cpp b/src/Pooling.cpp new file mode 100644 index 0000000..2626445 --- /dev/null +++ b/src/Pooling.cpp @@ -0,0 +1,111 @@ +#include + +#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; +} + +} \ No newline at end of file diff --git a/tests/simple_dense.py b/tests/simple_dense.py index 208bff9..112b5f4 100644 --- a/tests/simple_dense.py +++ b/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((10, 10, 4, 1), input_shape=(10, 10, 4))) - model.add(Convolution3D(2, (4, 4, 2), subsample=(2, 2, 1), activation="relu", - bias_initializer='random_uniform')) - model.add(Convolution3D(4, (2, 2, 2), subsample=(1, 1, 1), - bias_initializer='random_uniform')) - model.add(ELU()) + 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", activation="relu")) + #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("conv1", wg[2], wg[3]) + 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(10,10,4) + grid = np.random.rand(100, 100,4) X = grid[None,:,:] i = np.array(grid.flatten(), dtype=np.float32) print i diff --git a/tests/test.cpp b/tests/test.cpp index 9b54887..09b4c33 100644 --- a/tests/test.cpp +++ b/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 *d2_bias_bin = "../tests/dense2.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, 10, 10, 4); - tkDNN::Conv3d c0 (&net, dim, 2, 4, 4, 2, 2, 2, 1, c0_bin, c0_bias_bin); + tkDNN::dataDim_t dim(1, 1, 100, 100, 4); + 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::Activation a1 (&net, c1.output_dim, tkDNN::ACTIVATION_ELU); - tkDNN::Flatten f1 (&net, a1.output_dim); - tkDNN::MulAdd m1 (&net, f1.output_dim, 2, 1); + tkDNN::Pooling p0 (&net, a0.output_dim, 2, 2, 2, 2, tkDNN::POOLING_AVERAGE); + tkDNN::Conv3d c1 (&net, p0.output_dim, 16, 4, 4, 2, 2, 2, 1, c1_bin, c1_bias_bin); + tkDNN::Activation a1 (&net, c1.output_dim, tkDNN::ACTIVATION_RELU); + 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 = m1.infer(dim, data); dim.print(); - + data = c2.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;