diff --git a/CMakeLists.txt b/CMakeLists.txt index fd31073..cc2153d 100644 --- a/CMakeLists.txt +++ b/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) diff --git a/include/Layer.h b/include/Layer.h index ca0f27e..ba1cf14 100644 --- a/include/Layer.h +++ b/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 \ No newline at end of file diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp new file mode 100644 index 0000000..0853788 --- /dev/null +++ b/src/Conv2d.cpp @@ -0,0 +1,112 @@ +#include + +#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; +} + +} \ No newline at end of file diff --git a/src/kernels/activation_elu.cu b/src/kernels/activation_elu.cu index 7c1a378..85df29f 100644 --- a/src/kernels/activation_elu.cu +++ b/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(i0)*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); diff --git a/tests/simple_dense.py b/tests/simple_dense.py index cd40ced..14af613 100644 --- a/tests/simple_dense.py +++ b/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(ELU()) - model.add(Dense(2)) - + model.add(Convolution2D(4, (2, 2), subsample=(1, 1), + 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() 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]) - export_dense("dense1", wg[2], wg[3]) - export_dense("dense2", wg[4], wg[5]) - - model.set_weights(wg) - - X = np.random.rand(1, 512) - i = np.array(X, dtype=np.float32) + 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 - print "Input: ", i - r = model.predict( X[None, :], batch_size=1) - + r = model.predict( X, batch_size=1) + print np.shape(r) print "Result: ", r print "Result shape: ", np.shape(r) \ No newline at end of file diff --git a/tests/test.cpp b/tests/test.cpp index 048ae00..3c24101 100644 --- a/tests/test.cpp +++ b/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 *d0_bias_bin = "../tests/dense0.bias.bin"; -const char *d1_bin = "../tests/dense1.bin"; -const char *d1_bias_bin = "../tests/dense1.bias.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 *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::Dense d0 (&net, dim, 256, d0_bin, d0_bias_bin); - tkDNN::Activation a0 (&net, d0.output_dim, tkDNN::ACTIVATION_ELU); - tkDNN::Dense d1 (&net, a0.output_dim, 32, d1_bin, d1_bias_bin); - tkDNN::Activation a1 (&net, d1.output_dim, tkDNN::ACTIVATION_ELU); - tkDNN::Dense d2 (&net, a1.output_dim, 2, d2_bin, d2_bias_bin); - + tkDNN::dataDim_t dim(1, 1, 10, 10); + tkDNN::Conv2d c0 (&net, dim, 2, 4, 4, 2, 2, c0_bin, c0_bias_bin); + tkDNN::Activation a0 (&net, c0.output_dim, tkDNN::ACTIVATION_ELU); + tkDNN::Conv2d c1 (&net, a0.output_dim, 4, 2, 2, 1, 1, c1_bin, c1_bias_bin); + tkDNN::Activation a1 (&net, c1.output_dim, tkDNN::ACTIVATION_RELU); + // 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;