diff --git a/CMakeLists.txt b/CMakeLists.txt index f0db05e..d843668 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -9,9 +9,9 @@ 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/Pooling.cpp + src/Dense.cpp src/Activation.cpp src/Conv2d.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) +target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn) add_executable(tkDNNtest tests/test.cpp) -target_link_libraries(tkDNNtest tkDNN) \ No newline at end of file +target_link_libraries(tkDNNtest tkDNN) diff --git a/include/Layer.h b/include/Layer.h index 7059695..3ee566a 100644 --- a/include/Layer.h +++ b/include/Layer.h @@ -92,15 +92,6 @@ protected: value_type *dstData; //where results will be putted }; -/** - Avaible activation functions -*/ -typedef enum { - ACTIVATION_SIGMOID = 0, - ACTIVATION_RELU = 1, - ACTIVATION_TANH = 2, - ACTIVATION_ELU = 100 -} tkdnnActivationMode_t; /** Activation layer (it doesnt need weigths) @@ -108,13 +99,14 @@ typedef enum { class Activation : public Layer { public: - Activation(Network *net, dataDim_t input_dim, tkdnnActivationMode_t act_mode); + Activation(Network *net, dataDim_t input_dim, cudnnActivationMode_t act_mode); virtual ~Activation(); virtual value_type* infer(dataDim_t &dim, value_type* srcData); protected: - tkdnnActivationMode_t act_mode; + cudnnActivationMode_t act_mode; + cudnnActivationDescriptor_t activDesc; value_type *dstData; //where results will be putted }; @@ -145,35 +137,6 @@ protected: size_t ws_sizeInBytes; }; -/** - Convolutional 3D layer -*/ -class Conv3d : public LayerWgs { - -public: - Conv3d(Network *net, dataDim_t in_dim, int out_ch, - int kernelH, int kernelW, int kernelL, - int strideH, int strideW, int strideL, - const char* fname_weights, const char* fname_bias); - virtual ~Conv3d(); - - virtual value_type* infer(dataDim_t &dim, value_type* srcData); - -protected: - value_type *dstData; //where results will be putted - int kernelH, kernelW, kernelL; - int strideH, strideW, strideL; - - cudnnFilterDescriptor_t filterDesc; - cudnnConvolutionDescriptor_t convDesc; - cudnnConvolutionFwdAlgo_t algo; - cudnnTensorDescriptor_t biasTensorDesc; - cudnnTensorDescriptor_t biasDstTensorDesc; - - void* workSpace; - size_t ws_sizeInBytes; -}; - /** Flatten layer @@ -245,4 +208,4 @@ protected: }; } -#endif //LAYER_H \ No newline at end of file +#endif //LAYER_H diff --git a/src/Activation.cpp b/src/Activation.cpp index 2ee971e..caf8aff 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -5,7 +5,7 @@ namespace tkDNN { -Activation::Activation(Network *net, dataDim_t input_dim, tkdnnActivationMode_t act_mode) : +Activation::Activation(Network *net, dataDim_t input_dim, cudnnActivationMode_t act_mode) : Layer(net, input_dim) { this->act_mode = act_mode; @@ -23,6 +23,13 @@ Activation::Activation(Network *net, dataDim_t input_dim, tkdnnActivationMode_t input_dim.n*input_dim.l, input_dim.c, input_dim.h, input_dim.w) ); + + + checkCUDNN( cudnnCreateActivationDescriptor(&activDesc) ); + checkCUDNN( cudnnSetActivationDescriptor(activDesc, + act_mode, + CUDNN_PROPAGATE_NAN, + 0.0) ); } Activation::~Activation() { @@ -32,22 +39,17 @@ Activation::~Activation() { value_type* Activation::infer(dataDim_t &dim, value_type* srcData) { - 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) ); - } + value_type alpha = value_type(1); + value_type beta = value_type(0); + checkCUDNN( cudnnActivationForward(net->cudnnHandle, + activDesc, + &alpha, + srcTensorDesc, + srcData, + &beta, + dstTensorDesc, + dstData) ); return dstData; } -} \ No newline at end of file +} diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index 677c395..3b0e816 100644 --- a/src/Conv2d.cpp +++ b/src/Conv2d.cpp @@ -29,7 +29,7 @@ Conv2d::Conv2d( Network *net, dataDim_t in_dim, int out_ch, net->tensorFormat, net->dataType, n, c, h, w) ); checkCUDNN( cudnnSetFilter4dDescriptor(filterDesc, - net->dataType, out_ch, input_dim.c, + net->dataType, net->tensorFormat, out_ch, input_dim.c, kernelH, kernelW) ); checkCUDNN( cudnnSetConvolution2dDescriptor(convDesc, @@ -103,7 +103,7 @@ value_type* Conv2d::infer(dataDim_t &dim, value_type* srcData) { // bias alpha = value_type(1); beta = value_type(1); - checkCUDNN( cudnnAddTensor(net->cudnnHandle, CUDNN_ADD_SAME_C, + checkCUDNN( cudnnAddTensor(net->cudnnHandle, &alpha, biasTensorDesc, bias_d, &beta, dstTensorDesc, dstData) ); @@ -113,4 +113,4 @@ value_type* Conv2d::infer(dataDim_t &dim, value_type* srcData) { return dstData; } -} \ No newline at end of file +} diff --git a/src/Conv3d.cpp b/src/Conv3d.cpp deleted file mode 100644 index fe6c204..0000000 --- a/src/Conv3d.cpp +++ /dev/null @@ -1,157 +0,0 @@ -#include - -#include "Layer.h" - -namespace tkDNN { - -Conv3d::Conv3d( Network *net, dataDim_t in_dim, int out_ch, - int kernelH, int kernelW, int kernelL, - int strideH, int strideW, int strideL, - const char* fname_weights, const char* fname_bias) : - - LayerWgs(net, in_dim, in_dim.c, out_ch, kernelH, kernelW, kernelL, - fname_weights, fname_bias) { - - this->kernelH = kernelH; - this->kernelW = kernelW; - this->kernelL = kernelL; - this->strideH = strideH; - this->strideW = strideW; - this->strideL = strideL; - - checkCUDNN( cudnnCreateTensorDescriptor(&biasDstTensorDesc) ); - checkCUDNN( cudnnCreateFilterDescriptor(&filterDesc) ); - checkCUDNN( cudnnCreateConvolutionDescriptor(&convDesc) ); - checkCUDNN( cudnnCreateTensorDescriptor(&biasTensorDesc) ); - - 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; - - //create a tensor Nd descriptor with N = 4 - int dimA[5]; - dimA[0] = n; dimA[1] = c; dimA[2] = h; dimA[3] = w; dimA[4] = l; - int strideA[5]; - strideA[0] = c*h*w*l; - strideA[1] = h*w*l; - strideA[2] = w*l; - strideA[3] = l; - strideA[4] = 1; - checkCUDNN( cudnnSetTensorNdDescriptor(srcTensorDesc, - net->dataType, 5, dimA, strideA)); - - //filter descriptor - int filterDim[5]; - filterDim[0] = out_ch; - filterDim[1] = in_dim.c; - filterDim[2] = kernelH; - filterDim[3] = kernelW; - filterDim[4] = kernelL; - checkCUDNN( cudnnSetFilterNdDescriptor(filterDesc, - net->dataType, 5, filterDim)); - - //convolutional descriptor - int padA[3] = {0, 0, 0}; - int filterStride[3] = {strideH, strideW, strideL}; - int upscale[3] = {1, 1, 1}; - checkCUDNN( cudnnSetConvolutionNdDescriptor(convDesc, 3, - padA, filterStride, upscale, CUDNN_CROSS_CORRELATION)); - - - //get output dimension - int outputDim[5]; - checkCUDNN(cudnnGetConvolutionNdForwardOutputDim(convDesc, srcTensorDesc, filterDesc, 5, outputDim)); - n = outputDim[0]; - c = outputDim[1]; - h = outputDim[2]; - w = outputDim[3]; - l = outputDim[4]; - - //destination sensor - int outputStride[5]; - outputStride[0] = c*h*w*l; - outputStride[1] = h*w*l; - outputStride[2] = w*l; - outputStride[3] = l; - outputStride[4] = 1; - checkCUDNN( cudnnSetTensorNdDescriptor(dstTensorDesc, net->dataType, - 5, outputDim, outputStride)); - - - //conv algo - checkCUDNN( cudnnGetConvolutionForwardAlgorithm(net->cudnnHandle, - srcTensorDesc, filterDesc, convDesc, dstTensorDesc, - CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo) ); - - checkCUDNN( cudnnGetConvolutionForwardWorkspaceSize(net->cudnnHandle, - srcTensorDesc, - filterDesc, - convDesc, - dstTensorDesc, - algo, - &ws_sizeInBytes) ); - if (ws_sizeInBytes!=0) - checkCuda( cudaMalloc(&workSpace, ws_sizeInBytes) ); - - - // bias on N dimensional is not SUPPORTED so i have to use 2d method - //the trick is to upscale the 2d matrix width by the factor of 3d thickness - checkCUDNN( cudnnSetTensor4dDescriptor(biasDstTensorDesc, - net->tensorFormat, net->dataType, n, c, h*l, w) ); - - - checkCUDNN( cudnnSetTensor4dDescriptor(biasTensorDesc, - net->tensorFormat, net->dataType, - 1, c, 1, 1) ); - - output_dim.n = n; - output_dim.c = c; - output_dim.h = h; - output_dim.w = w; - output_dim.l = l; - - //allocate data for infer result - checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) ); -} - -Conv3d::~Conv3d() { - - checkCUDNN( cudnnDestroyFilterDescriptor(filterDesc) ); - checkCUDNN( cudnnDestroyConvolutionDescriptor(convDesc) ); - checkCUDNN( cudnnDestroyTensorDescriptor(biasTensorDesc) ); - checkCUDNN( cudnnDestroyTensorDescriptor(biasDstTensorDesc) ); - - if (ws_sizeInBytes!=0) - checkCuda( cudaFree(workSpace) ); - - checkCuda( cudaFree(dstData) ); -} - -value_type* Conv3d::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, biasDstTensorDesc, dstData) ); - - //update data dimensions - dim = output_dim; - - return dstData; -} - -} \ No newline at end of file diff --git a/src/Dense.cpp b/src/Dense.cpp index e9c6124..1ad8965 100644 --- a/src/Dense.cpp +++ b/src/Dense.cpp @@ -56,4 +56,4 @@ value_type* Dense::infer(dataDim_t &dim, value_type* srcData) { return dstData; } -} \ No newline at end of file +} diff --git a/src/Pooling.cpp b/src/Pooling.cpp index 2626445..2d964fc 100644 --- a/src/Pooling.cpp +++ b/src/Pooling.cpp @@ -40,7 +40,7 @@ Pooling::Pooling( Network *net, dataDim_t input_dim, } checkCUDNN( cudnnSetPooling2dDescriptor(poolingDesc, cudnnPoolingMode_t(pool_mode), - winH, winW, 0, 0, strideH, strideW) ); + CUDNN_NOT_PROPAGATE_NAN, winH, winW, 0, 0, strideH, strideW) ); checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc, net->tensorFormat, net->dataType, n, c, h, w) ); @@ -108,4 +108,4 @@ value_type* Pooling::infer(dataDim_t &dim, value_type* srcData) { return dstData; } -} \ No newline at end of file +} diff --git a/tests/test.cpp b/tests/test.cpp index c91da28..4f189d2 100644 --- a/tests/test.cpp +++ b/tests/test.cpp @@ -6,52 +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 *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, 100, 100, 4); + tkDNN::dataDim_t dim(1, 1, 10, 10, 1); tkDNN::Layer *l; - l = new tkDNN::MulAdd (&net, dim, 2, -1); - l = new tkDNN::Conv3d (&net, l->output_dim, 16, 8, 8, 2, 4, 4, 1, c0_bin, c0_bias_bin); - l = new tkDNN::Activation (&net, l->output_dim, tkDNN::ACTIVATION_ELU); - l = new tkDNN::Pooling (&net, l->output_dim, 2, 2, 2, 2, tkDNN::POOLING_AVERAGE); - l = new tkDNN::Conv3d (&net, l->output_dim, 16, 4, 4, 2, 2, 2, 1, c1_bin, c1_bias_bin); - l = new tkDNN::Activation (&net, l->output_dim, tkDNN::ACTIVATION_ELU); - l = new tkDNN::Conv3d (&net, l->output_dim, 24, 3, 3, 2, 1, 1, 1, c2_bin, c2_bias_bin); - l = new tkDNN::Activation (&net, l->output_dim, tkDNN::ACTIVATION_ELU); - l = new tkDNN::Flatten (&net, l->output_dim); - l = new tkDNN::Dense (&net, l->output_dim, 256, d3_bin, d3_bias_bin); - l = new tkDNN::Activation (&net, l->output_dim, tkDNN::ACTIVATION_ELU); - l = new tkDNN::Dense (&net, l->output_dim, 32, d4_bin, d4_bias_bin); - l = new tkDNN::Activation (&net, l->output_dim, tkDNN::ACTIVATION_RELU); - l = new tkDNN::Dense (&net, l->output_dim, 2, d5_bin, d5_bias_bin); - + l = new tkDNN::Conv2d (&net, dim, 2, 4, 4, 2, 2, c0_bin, c0_bias_bin); + l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU); + l = new tkDNN::Conv2d (&net, l->output_dim, 4, 2, 2, 1, 1, c1_bin, c1_bias_bin); + l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU); // Load input value_type *data; value_type *input_h; readBinaryFile(input_bin, dim.tot(), &input_h, &data); + printDeviceVector(dim.tot(), data); dim.print(); //print initial dimension TIMER_START // Inference data = net.infer(dim, data); dim.print(); + TIMER_STOP // Print result printDeviceVector(dim.tot(), data); return 0; -} \ No newline at end of file +} diff --git a/tests/test_model.py b/tests/test_model.py index ec93533..c2a69af 100644 --- a/tests/test_model.py +++ b/tests/test_model.py @@ -11,26 +11,12 @@ from weights_exporter import * def dense_model(): model = Sequential() - 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")) - 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")) - model.add(ELU()) - model.add(Convolution3D(24, kernel_size=(3, 3, 2), subsample=(1, 1, 1), border_mode="valid", - bias_initializer="random_uniform")) - model.add(ELU()) - model.add(Flatten()) - model.add(Dense(256, bias_initializer="random_uniform")) - model.add(ELU()) - model.add(Dense(32, activation="relu", bias_initializer="random_uniform")) - model.add(Dense(2, bias_initializer="random_uniform")) + model.add(Reshape((10, 10, 1), input_shape=(10, 10))) + model.add(Convolution2D(2, (4, 4), subsample=(2, 2), + bias_initializer='random_uniform', activation="relu")) + 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 @@ -41,14 +27,10 @@ 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("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]) + export_conv2d("conv0", wg[0], wg[1]) + export_conv2d("conv1", wg[2], wg[3]) - grid = np.random.rand(100, 100,4) + grid = np.random.rand(10,10) X = grid[None,:,:] i = np.array(grid.flatten(), dtype=np.float32) print i @@ -58,4 +40,4 @@ if __name__ == '__main__': 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 + print "Result shape: ", np.shape(r)