diff --git a/.gitignore b/.gitignore index 3b47b1f..011a55f 100644 --- a/.gitignore +++ b/.gitignore @@ -1,3 +1,4 @@ *~ build/ .vscode/ +*.bin diff --git a/CMakeLists.txt b/CMakeLists.txt index d5fef62..fd31073 100644 --- a/CMakeLists.txt +++ b/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) \ No newline at end of file +target_link_libraries(tkDNNtest tkDNN + ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDA_TOOLKIT_ROOT_DIR}/lib/libcudnn.so) \ No newline at end of file diff --git a/include/Layer.h b/include/Layer.h index 9eab1d9..ca0f27e 100644 --- a/include/Layer.h +++ b/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 }; diff --git a/include/kernels.h b/include/kernels.h new file mode 100644 index 0000000..4085ef7 --- /dev/null +++ b/include/kernels.h @@ -0,0 +1,3 @@ +#include "utils.h" + +void activationELUForward(value_type* srcData, value_type* dstData, int size); \ No newline at end of file diff --git a/src/Activation.cpp b/src/Activation.cpp index dd8ced2..cf09b3d 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -1,10 +1,11 @@ #include #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; } diff --git a/src/Dense.cpp b/src/Dense.cpp index 7edf888..e9c6124 100644 --- a/src/Dense.cpp +++ b/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); diff --git a/src/Layer.cpp b/src/Layer.cpp index 37999e6..1629812 100644 --- a/src/Layer.cpp +++ b/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) ); } diff --git a/src/kernels/activation_elu.cu b/src/kernels/activation_elu.cu new file mode 100644 index 0000000..7c1a378 --- /dev/null +++ b/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(i0)*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() ); +} \ No newline at end of file diff --git a/tests/simple_dense.py b/tests/simple_dense.py new file mode 100644 index 0000000..0c5dbc8 --- /dev/null +++ b/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) diff --git a/tests/test.cpp b/tests/test.cpp index 9c5ebfe..4e246a2 100644 --- a/tests/test.cpp +++ b/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; } \ No newline at end of file