From a9c0db0bf6abc3f21611b1c13f55aaba383159e5 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 13 Feb 2020 19:27:18 +0100 Subject: [PATCH 01/11] LSTM cudnn test --- .gitignore | 3 +- CMakeLists.txt | 3 + include/tkDNN/Layer.h | 76 +++++++++++++++++ src/LSTM.cpp | 142 ++++++++++++++++++++++++++++++++ src/utils.cpp | 3 +- tests/imuodom/imuodom.cpp | 70 ++++++++++++++++ tests/imuodom/infer.py | 74 +++++++++++++++++ tests/simple/test_model.py | 60 +++++++------- tests/simple/test_simple.cpp | 14 +--- tests/weights_exporter.py | 153 ++++++++++++++--------------------- 10 files changed, 464 insertions(+), 134 deletions(-) create mode 100644 src/LSTM.cpp create mode 100644 tests/imuodom/imuodom.cpp create mode 100644 tests/imuodom/infer.py diff --git a/.gitignore b/.gitignore index 02f2a8e..de78410 100644 --- a/.gitignore +++ b/.gitignore @@ -8,4 +8,5 @@ build/ *.h5 *.tar.gz *.weights -.idea/ \ No newline at end of file +.idea/ +*.hdf5 \ No newline at end of file diff --git a/CMakeLists.txt b/CMakeLists.txt index b4b3c38..e3d8607 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -85,6 +85,9 @@ target_link_libraries(test_yolo3_berkeley tkDNN) add_executable(test_yolo3_flir tests/yolo3_flir/yolo3_flir.cpp) target_link_libraries(test_yolo3_flir tkDNN) + +add_executable(test_imuodom tests/imuodom/imuodom.cpp) +target_link_libraries(test_imuodom tkDNN) ################################################################################ diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index ee73109..f15c781 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -9,8 +9,10 @@ namespace tk { namespace dnn { enum layerType_t { + LAYER_INPUT, LAYER_DENSE, LAYER_CONV2D, + LAYER_LSTM, LAYER_ACTIVATION, LAYER_FLATTEN, LAYER_MULADD, @@ -47,8 +49,10 @@ public: std::string getLayerName() { layerType_t type = getLayerType(); switch(type) { + case LAYER_INPUT: return "Input"; case LAYER_DENSE: return "Dense"; case LAYER_CONV2D: return "Conv2d"; + case LAYER_LSTM: return "LSTM"; case LAYER_ACTIVATION: return "Activation"; case LAYER_FLATTEN: return "Flatten"; case LAYER_MULADD: return "MulAdd"; @@ -105,6 +109,27 @@ public: }; +/** + Input layer (it doesnt need weigths) +*/ +class Input : public Layer { + +public: + + Input(Network *net, dataDim_t &dim, dnnType* srcData) : Layer(net) { + input_dim = dim; + output_dim = dim; + dstData = srcData; + } + virtual ~Input() {} + virtual layerType_t getLayerType() { return LAYER_INPUT; }; + + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData) { + return dstData; + } +}; + + /** Dense (full interconnection) layer */ @@ -148,6 +173,14 @@ protected: /** Convolutional 2D layer + + WEIGHTS shape: OUTCH, INCH, KH, KW ... + BIAS shape: OUTCH + + with BATCHNORM: + scales: OUTCH + means: OUTCH + variance: OUTCH */ class Conv2d : public LayerWgs { @@ -172,6 +205,49 @@ protected: size_t ws_sizeInBytes; }; +/** + Bidirectional LSTM layer + https://github.com/jiangnanhugo/seq2seq_cuda/blob/e4dbdcfa0517c972bfd4beea9f11a5233954093c/src/rnn.cpp + + numlayers = 1 # hardcoded as 1 + + PARAMS (numlayers*2): + layer0: + ( INCH, ? ) ??? + ( HIDDEN, ? ) ??? + ( HIDDEN * 8 ) ??? + layer2: + ( INCH, ? ) ??? + ( HIDDEN, ? ) ??? + ( HIDDEN * 8 ) ??? + + output shape: ( 2*HIDDEN, INH, INW ) +*/ +class LSTM : public Layer { + +public: + LSTM(Network *net, int hiddensize, std::string fname_weights); + virtual ~LSTM(); + virtual layerType_t getLayerType() { return LAYER_LSTM; }; + + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); + + int kernelH, kernelW, strideH, strideW, paddingH, paddingW; + +protected: + cudnnFilterDescriptor_t paramDesc; + cudnnTensorDescriptor_t hiddenStateTensorDesc, cellStateTensorDesc; + cudnnRNNDescriptor_t rnnDesc; + cudnnRNNDataDescriptor_t rnnDataDesc; + cudnnDropoutDescriptor_t dropDesc; + cudnnRNNAlgo_t algo; + + dnnType *hiddenStateData, *cellStateData; + dnnType *paramsSpace; + void* workSpace; + size_t ws_sizeInBytes; +}; + /** Flatten layer diff --git a/src/LSTM.cpp b/src/LSTM.cpp new file mode 100644 index 0000000..3176819 --- /dev/null +++ b/src/LSTM.cpp @@ -0,0 +1,142 @@ +#include + +#include "Layer.h" + +namespace tk { namespace dnn { + +LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : + Layer(net) { + + checkCUDNN( cudnnCreateFilterDescriptor(¶mDesc)); + checkCUDNN( cudnnCreateRNNDescriptor(&rnnDesc) ); + checkCUDNN( cudnnCreateRNNDataDescriptor(&rnnDataDesc) ); + checkCUDNN( cudnnCreateDropoutDescriptor(&dropDesc)); + + 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, 1, h, w) ); + + int numlayers = 1; + checkCUDNN( cudnnSetRNNDescriptor(net->cudnnHandle, rnnDesc, hiddensize, numlayers, dropDesc, + cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT, + cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL, cudnnRNNMode_t::CUDNN_LSTM, + cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD, net->dataType) ); + + // find dimension of params + size_t params_size = 0; + checkCUDNN( cudnnGetRNNParamsSize(net->cudnnHandle, rnnDesc, srcTensorDesc, ¶ms_size, net->dataType) ); + std::cout<<"Params size bytes: "<dataType, net->tensorFormat, 3, dimW)); + checkCuda( cudaMalloc(¶msSpace, params_size) ); + + + int numlinearlayers = 8; + + for(int i=0; icudnnHandle, rnnDesc, + i, srcTensorDesc, paramDesc, paramsSpace, + j, linLayerMatDesc, (void **)&linLayerMat)); + + if(linLayerMat == nullptr) { + FatalError("LSTM No weights in hidden layer"); + } + + cudnnDataType_t dataType; + cudnnTensorFormat_t format; + int nbDims; + int filterDimA[3]; + checkCUDNN(cudnnGetFilterNdDescriptor(linLayerMatDesc, 3, &dataType, + &format, &nbDims, filterDimA)); + std::cout<<"Wgs Dims: "<cudnnHandle, rnnDesc, + i, srcTensorDesc, paramDesc, paramsSpace, + j, linLayerBiasDesc, (void **)&linLayerBias)); + + if(linLayerMat == nullptr) { + FatalError("LSTM No bias in hidden layer"); + } + + checkCUDNN(cudnnGetFilterNdDescriptor(linLayerBiasDesc, 3, &dataType, + &format, &nbDims, filterDimA)); + std::cout<<"bias Dims: "<tensorFormat, net->dataType, 2*n, c, h, w) ); + checkCuda( cudaMalloc(&hiddenStateData, 2*input_dim.tot()*sizeof(dnnType)) ); + checkCUDNN( cudnnCreateTensorDescriptor(&cellStateTensorDesc)); + checkCUDNN( cudnnSetTensor4dDescriptor(cellStateTensorDesc, + net->tensorFormat, net->dataType, 2*n, c, h, w) ); + checkCuda( cudaMalloc(&cellStateData, 2*input_dim.tot()*sizeof(dnnType)) ); + + + output_dim = input_dim; + output_dim.c = hiddensize*2; + checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc, + net->tensorFormat, net->dataType, output_dim.n, output_dim.c, output_dim.h, output_dim.w) ); + + + + + //allocate data for infer result + checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); +} + +LSTM::~LSTM() { + + checkCuda( cudaFree(dstData) ); +} + +dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) { + + checkCUDNN(cudnnRNNForwardInference( + net->cudnnHandle, rnnDesc, 1, + &srcTensorDesc, srcData, + hiddenStateTensorDesc, hiddenStateData, + cellStateTensorDesc, cellStateData, + paramDesc, paramsSpace, + &dstTensorDesc, dstData, + hiddenStateTensorDesc, hiddenStateData, + cellStateTensorDesc, cellStateData, + workSpace, ws_sizeInBytes + )); + + return dstData; +} + +}} diff --git a/src/utils.cpp b/src/utils.cpp index e6f71f8..444318d 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -39,7 +39,8 @@ void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** dat *data_h = new dnnType[size]; if (!dataFile.read ((char*) *data_h, size_b)) { - error_s << "Error reading file " << fname; + error_s << "Error reading file " << fname << " with n of float: "< +#include "tkdnn.h" + +const char *i0_bin = "../tests/imuodom/layers/input0.bin"; +const char *i1_bin = "../tests/imuodom/layers/input1.bin"; +const char *i2_bin = "../tests/imuodom/layers/input2.bin"; +const char *o0_bin = "../tests/imuodom/layers/output0.bin"; +const char *o1_bin = "../tests/imuodom/layers/output1.bin"; +const char *output_bin = "../tests/imuodom/layers/output.bin"; + +const char *c0_bin = "../tests/imuodom/layers/conv1d_7.bin"; +const char *c1_bin = "../tests/imuodom/layers/conv1d_8.bin"; +const char *c2_bin = "../tests/imuodom/layers/conv1d_9.bin"; +const char *c3_bin = "../tests/imuodom/layers/conv1d_10.bin"; +const char *c4_bin = "../tests/imuodom/layers/conv1d_11.bin"; +const char *c5_bin = "../tests/imuodom/layers/conv1d_12.bin"; + +int main() { + + // Network layout + tk::dnn::dataDim_t dim0(1, 4, 1, 100); + tk::dnn::dataDim_t dim1(1, 3, 1, 100); + tk::dnn::dataDim_t dim2(1, 3, 1, 100); + + // Load input + dnnType *i0_d, *i1_d, *i2_d; + dnnType *i0_h, *i1_h, *i2_h; + readBinaryFile(i0_bin, dim0.tot(), &i0_h, &i0_d); + readBinaryFile(i1_bin, dim1.tot(), &i1_h, &i1_d); + readBinaryFile(i2_bin, dim2.tot(), &i2_h, &i2_d); + + tk::dnn::Network net(dim0); + tk::dnn::Input x0 (&net, dim0, i0_d); + tk::dnn::Conv2d x0_0(&net, 128, 1, 11, 1, 1, 0, 0, c0_bin); + tk::dnn::Conv2d x0_1(&net, 128, 1, 11, 1, 1, 0, 0, c1_bin); + tk::dnn::Pooling x0_2(&net, 1, 3, 1, 3, tk::dnn::tkdnnPoolingMode_t::POOLING_MAX); + + tk::dnn::Input x1 (&net, dim1, i1_d); + tk::dnn::Conv2d x1_0(&net, 128, 1, 11, 1, 1, 0, 0, c2_bin); + tk::dnn::Conv2d x1_1(&net, 128, 1, 11, 1, 1, 0, 0, c3_bin); + tk::dnn::Pooling x1_2(&net, 1, 3, 1, 3, tk::dnn::tkdnnPoolingMode_t::POOLING_MAX); + + tk::dnn::Input x2 (&net, dim2, i2_d); + tk::dnn::Conv2d x2_0(&net, 128, 1, 11, 1, 1, 0, 0, c4_bin); + tk::dnn::Conv2d x2_1(&net, 128, 1, 11, 1, 1, 0, 0, c5_bin); + tk::dnn::Pooling x2_2(&net, 1, 3, 1, 3, tk::dnn::tkdnnPoolingMode_t::POOLING_MAX); + + tk::dnn::Layer *concat_l[3] = { &x0_2, &x1_2, &x2_2 }; + tk::dnn::Route concat (&net, concat_l, 3); + + tk::dnn::LSTM lstm0(&net, 128, "ciao"); + + net.print(); + + dnnType *data; + tk::dnn::dataDim_t dim; + + TIMER_START + // Inference + data = net.infer(dim, data); dim.print(); + TIMER_STOP + + // Print real test + std::cout<<"\n==== CHECK RESULT ====\n"; + dnnType *out; + dnnType *out_h; + readBinaryFile(output_bin, dim.tot(), &out_h, &out); + checkResult(dim.tot(), data, out); + return 0; +} diff --git a/tests/imuodom/infer.py b/tests/imuodom/infer.py new file mode 100644 index 0000000..9d9929d --- /dev/null +++ b/tests/imuodom/infer.py @@ -0,0 +1,74 @@ +import keras +from keras.models import load_model +import keras.backend.tensorflow_backend as KTF +import numpy as np +import argparse +import tensorflow as tf +import os +import random +import struct +from keras.models import Sequential, Model + +def bin_write(f, data): + data = data.flatten() + fmt = 'f'*len(data) + bin = struct.pack(fmt, *data) + f.write(bin) + + +if __name__ == '__main__': + + + print("DATA FORMAT: ", keras.backend.image_data_format()) + + print("Load model: ", "ferrariS1.hdf5") + model = load_model("ferrariS1.hdf5") + model.summary() + + weights = model.get_weights() + + x_angle = np.random.rand(1,100,4) + x_gyro = np.random.rand(1,100,3) + x_acc = np.random.rand(1,100,3) + + [yhat_delta_p, yhat_delta_q] = model.predict([x_angle, x_gyro, x_acc], batch_size=1, verbose=1) + + layer_name = 'bidirectional_3' + intermediate_layer_model = Model(inputs=model.input, + outputs=model.get_layer(layer_name).output) + intermediate_output = intermediate_layer_model.predict([x_angle, x_gyro, x_acc]) + + x_angle = np.array([x_angle]) + x_gyro = np.array([x_gyro]) + x_acc = np.array([x_acc]) + intermediate_output = np.array([intermediate_output]) + + x_angle = x_angle.transpose(0, 3, 1, 2) + x_gyro = x_gyro.transpose(0, 3, 1, 2) + x_acc = x_acc.transpose(0, 3, 1, 2) + intermediate_output = intermediate_output.transpose(0, 3, 1, 2) + + print("x0: ", np.shape(x_angle)) + print("out: ",np.shape(intermediate_output)) + + x_angle = np.array(x_angle.flatten(), dtype=np.float32) + x_gyro = np.array(x_gyro.flatten(), dtype=np.float32) + x_acc = np.array(x_acc.flatten(), dtype=np.float32) + yhat_delta_p = np.array(yhat_delta_p.flatten(), dtype=np.float32) + yhat_delta_q = np.array(yhat_delta_q.flatten(), dtype=np.float32) + intermediate_output = np.array(intermediate_output.flatten(), dtype=np.float32) + + + f = open("layers/input0.bin", mode='wb') + bin_write(f, x_angle) + f = open("layers/input1.bin", mode='wb') + bin_write(f, x_gyro) + f = open("layers/input2.bin", mode='wb') + bin_write(f, x_acc) + f = open("layers/output0.bin", mode='wb') + bin_write(f, yhat_delta_p) + f = open("layers/output1.bin", mode='wb') + bin_write(f, yhat_delta_q) + f = open("layers/output.bin", mode='wb') + bin_write(f, intermediate_output) + diff --git a/tests/simple/test_model.py b/tests/simple/test_model.py index 60f17d8..3b1d053 100644 --- a/tests/simple/test_model.py +++ b/tests/simple/test_model.py @@ -1,43 +1,49 @@ import keras import numpy as np from keras.models import Sequential -from keras.layers import Input, Dense, Activation, Flatten, Dropout, ELU, Reshape, Lambda +from keras.layers import Input, Dense, Activation, Flatten, Dropout, ELU, Reshape, Lambda, Conv1D from keras.layers.convolutional import Convolution2D, Convolution3D 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 +import struct +from keras.models import Sequential, Model -def dense_model(): - model = Sequential() +def bin_write(f, data): + data = data.flatten() + fmt = 'f'*len(data) + bin = struct.pack(fmt, *data) + f.write(bin) + +def create_model(): + x1 = Input((6, 16), name='x1') + conv = Conv1D(4, 2)(x1) + model = Model([x1], [conv]) + model.summary() - 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")) - model.add(Flatten()) - model.add(Dense(4, 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() + print ("DATA FORMAT: ", keras.backend.image_data_format()) - model = dense_model() - model.save("net.h5") + model = create_model() + model.save("net.hdf5") - 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 + x = np.random.rand(1,1,6,16) + r = model.predict( x[0], batch_size=1) + r = np.array([r]) + + x = x.transpose(0, 3, 1, 2) + r = r.transpose(0, 3, 1, 2) + print("in: ", np.shape(x)) + print("out: ", np.shape(r)) + + x = np.array(x.flatten(), dtype=np.float32) + f = open("input.bin", mode='wb') + bin_write(f, x) + + r = np.array(r.flatten(), dtype=np.float32) + f = open("output.bin", mode='wb') + bin_write(f, r) - r = model.predict( X, batch_size=1) - print np.shape(r) - print "Result: ", r - print "Result shape: ", np.shape(r) - r.tofile("output.bin", format="f") diff --git a/tests/simple/test_simple.cpp b/tests/simple/test_simple.cpp index a7628ea..3427d6d 100644 --- a/tests/simple/test_simple.cpp +++ b/tests/simple/test_simple.cpp @@ -2,23 +2,15 @@ #include "tkdnn.h" const char *input_bin = "../tests/simple/input.bin"; -const char *c0_bin = "../tests/simple/layers/c0.bin"; -const char *c1_bin = "../tests/simple/layers/c1.bin"; -const char *d2_bin = "../tests/simple/layers/d2.bin"; +const char *c0_bin = "../tests/simple/layers/conv1d_1.bin"; const char *output_bin = "../tests/simple/output.bin"; int main() { // Network layout - tk::dnn::dataDim_t dim(1, 1, 10, 10, 1); + tk::dnn::dataDim_t dim(1, 16, 1, 6); tk::dnn::Network net(dim); - tk::dnn::Conv2d l0(&net, 2, 4, 4, 2, 2, 0, 0, c0_bin); - tk::dnn::Activation l1(&net, CUDNN_ACTIVATION_RELU); - tk::dnn::Conv2d l2(&net, 4, 2, 2, 1, 1, 0, 0, c1_bin); - tk::dnn::Activation l3(&net, CUDNN_ACTIVATION_RELU); - tk::dnn::Flatten l4(&net); - tk::dnn::Dense l5(&net, 4, d2_bin); - tk::dnn::Activation l6(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d l0(&net, 4, 1, 2, 1, 1, 0, 0, c0_bin); // Load input dnnType *data; diff --git a/tests/weights_exporter.py b/tests/weights_exporter.py index df8bab3..912ef82 100644 --- a/tests/weights_exporter.py +++ b/tests/weights_exporter.py @@ -5,96 +5,51 @@ import numpy as np import argparse import tensorflow as tf import os -import msgpack -import lmdb import random +import struct +from keras.models import Sequential, Model -def export_dense(name, weights, bias): - print "######## EXPORT", name, "LAYER ########" - print "Original weighs:" - print weights - print bias, "\n" +def bin_write(f, data): + data = data.flatten() + fmt = 'f'*len(data) + bin = struct.pack(fmt, *data) + f.write(bin) - #input, filters - I, C = np.shape(weights) - B = np.shape(bias) - print "w shape: ", I, C - print "b shape: ", B +def export_layer(name, weights, bias): + print ("######## EXPORT", name, "LAYER ########") - wgs = [ [ j[i] for j in weights ] for i in xrange(C) ] - wgs = np.array(wgs, dtype=np.float32) - - print "REPOSITIONED WEIGHTS:" - print wgs + print("wgs pretranpose: ", np.shape(weights)) + # convert NHWC to NCHW + if(weights.ndim == 4): + weights = weights.transpose(3,2,0,1) + elif(weights.ndim == 3): + weights = weights.transpose(2,1,0) + else: + print("Ndim", weights.ndim) + raise("not implemented with dim" ) + print("weights: ", np.shape(weights)) + print("bias: ", np.shape(bias)) + + weights = np.array(weights.flatten(), dtype=np.float32) bias = np.array(bias, dtype=np.float32) - wgs.tofile(name + ".bin", format="f") - bias.tofile(name + ".bias.bin", format="f") - print "WEIGHTS saved\n" + print(len(weights) + len(bias)) -def export_conv2d(name, weights, bias): - print "######## EXPORT", name, "LAYER ########" - print "Original weighs:" - print weights - print bias, "\n" - - # height, width, input, filters - H, W, N, C = np.shape(weights) - B = np.shape(bias) - print "w shape: ", N, C, H, W - print "b shape: ", B + f = open(name + ".bin", mode='wb') + bin_write(f, weights) + bin_write(f, bias) + print ("WEIGHTS saved\n") - wgs = weights.transpose() - wgs = wgs.transpose(0, 1, 3, 2) - print "Final shape:", np.shape(wgs) - wgs = np.array(wgs.flatten(), dtype=np.float32) - - print "REPOSITIONED WEIGHTS:" - print wgs - - bias = np.array(bias, dtype=np.float32) - - wgs.tofile(name + ".bin", format="f") - bias.tofile(name + ".bias.bin", format="f") - print "WEIGHTS saved\n" - -def export_conv3d(name, weights, bias): - print "######## EXPORT", name, "LAYER ########" - print "Original weighs:" - print weights - print bias, "\n" - - print np.shape(weights) - # height, width, input, thickness, filters - H, W, T, N, C = np.shape(weights) - B = np.shape(bias) - print "w shape: ", T, C, H, W #thickness is number of images for cudnn - print "b shape: ", B - - wgs = weights.transpose() - wgs = wgs.transpose(0, 1, 4, 3, 2) - print "Final shape:", np.shape(wgs) - wgs = np.array(wgs.flatten(), dtype=np.float32) - - print "REPOSITIONED WEIGHTS:" - print wgs - - bias = np.array(bias, dtype=np.float32) - - wgs.tofile(name + ".bin", format="f") - bias.tofile(name + ".bias.bin", format="f") - print "WEIGHTS saved\n" - - -def get_session(gpu_fraction=0.5): - gpu_options = tf.GPUOptions(allow_growth=True) - #per_process_gpu_memory_fraction=gpu_fraction) - return tf.Session(config=tf.ConfigProto(gpu_options=gpu_options)) +def export_bidir(name, weights): + print ("######## EXPORT", name, "LAYER ########") + + for w in weights: + print(np.shape(w)) #https://github.com/fchollet/keras/wiki/Converting-convolution-kernels-from-Theano-to-TensorFlow-and-vice-versa if __name__ == '__main__': - KTF.set_session(get_session()) + print("DATA FORMAT: ", keras.backend.image_data_format()) parser = argparse.ArgumentParser(description='KERAS WEIGHTS EXPORTER TO CUDNN') parser.add_argument('model',type=str, @@ -103,31 +58,41 @@ if __name__ == '__main__': args = parser.parse_args() - print "DATA FORMAT: ", keras.backend.image_data_format() + print("DATA FORMAT: ", keras.backend.image_data_format()) - print "Load model: ", args.model + print("Load model: ", args.model) model = load_model(args.model) + model.summary() + weights = model.get_weights() ws = np.shape(weights) - print "Weights shape:", ws + print("Weights shape:", ws) if not os.path.exists(args.output): os.makedirs(args.output) - num = 0 + name_num = 0 for l in model.layers: - name = l.name - if name.startswith("conv3d"): - export_conv3d(args.output + "/conv" + str(name_num), weights[num], weights[num+1]) - elif name.startswith("conv2d"): - export_conv2d(args.output + "/conv" + str(name_num), weights[num], weights[num+1]) - elif name.startswith("dense"): - export_dense(args.output + "/dense" + str(name_num), weights[num], weights[num+1]) - else: - print "skip:", name, "has no weights" - continue - name_num += 1 - num += 2 + print("\n\nNAME: ", l.name) + print("input: ", l.input_shape, " output: ", l.output_shape) + wgs = l.get_weights() + print("wgs num: ", len(wgs)) + + name = l.name + if name.startswith("conv3d"): + export_layer(args.output + "/" + name, wgs[0], wgs[1]) + elif name.startswith("conv2d"): + export_layer(args.output + "/" + name, wgs[0], wgs[1]) + elif name.startswith("conv1d"): + export_layer(args.output + "/" + name, wgs[0], wgs[1]) + elif name.startswith("dense"): + export_layer(args.output + "/" + name, wgs[0], wgs[1]) + elif name.startswith("bidirectional"): + export_bidir(args.output + "/" + name, wgs) + else: + print ("skip:", name, "has no weights") + continue + From 03d39d991c5932aa4c94f7b86fc8c64e9d48238b Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 13 Feb 2020 23:04:29 +0100 Subject: [PATCH 02/11] LSTM to be tested --- include/tkDNN/Layer.h | 40 ++++--- src/LSTM.cpp | 246 ++++++++++++++++++++++++------------------ 2 files changed, 170 insertions(+), 116 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index f15c781..c538fc9 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -207,9 +207,11 @@ protected: /** Bidirectional LSTM layer + + implementation info: https://github.com/jiangnanhugo/seq2seq_cuda/blob/e4dbdcfa0517c972bfd4beea9f11a5233954093c/src/rnn.cpp - - numlayers = 1 # hardcoded as 1 + https://github.com/Jeffery-Song/mxnet-test/blob/aab666faad44011f7a67b527b5f6c960367d0422/src/operator/cudnn_rnn-inl.h + https://stackoverflow.com/a/38737941 PARAMS (numlayers*2): layer0: @@ -221,7 +223,9 @@ protected: ( HIDDEN, ? ) ??? ( HIDDEN * 8 ) ??? - output shape: ( 2*HIDDEN, INH, INW ) + OUTPUT shape: + (N, C, 1, W) ---> LSTM(HIDDEN, returnSeq=True) ---> (N, 2*HIDDEN, 1, W) # W is seqLength + (N, C, 1, W) ---> LSTM(HIDDEN, returnSeq=False) ---> (N, 2*HIDDEN, 1, 1) */ class LSTM : public Layer { @@ -232,20 +236,28 @@ public: virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); - int kernelH, kernelW, strideH, strideW, paddingH, paddingW; + const bool bidirectional = 1; /**> is the net bidir */ + int stateSize = 0; /**> number of hidden states */ + int seqLen = 0; /**> number of timesteps */ + int numLayers = 1; /**> number of internal layers */ protected: - cudnnFilterDescriptor_t paramDesc; - cudnnTensorDescriptor_t hiddenStateTensorDesc, cellStateTensorDesc; - cudnnRNNDescriptor_t rnnDesc; - cudnnRNNDataDescriptor_t rnnDataDesc; - cudnnDropoutDescriptor_t dropDesc; - cudnnRNNAlgo_t algo; + cudnnRNNDescriptor_t rnnDesc; + cudnnDropoutDescriptor_t dropoutDesc; + dnnType *dropout_states_, *work_space_; - dnnType *hiddenStateData, *cellStateData; - dnnType *paramsSpace; - void* workSpace; - size_t ws_sizeInBytes; + size_t workspace_byte_, reserve_space_byte_, dropout_byte_; + int workspace_size_, dropout_size_; + + std::vector x_desc_vec_, y_desc_vec_, dx_desc_vec_, dy_desc_vec_; + cudnnTensorDescriptor_t hx_desc_, cx_desc_; + cudnnTensorDescriptor_t hy_desc_, cy_desc_; + cudnnTensorDescriptor_t dhx_desc_, dcx_desc_; + cudnnTensorDescriptor_t dhy_desc_, dcy_desc_; + dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; + + cudnnFilterDescriptor_t w_desc_, dw_desc_; + dnnType *w_ptr, *dw_ptr; }; diff --git a/src/LSTM.cpp b/src/LSTM.cpp index 3176819..be12cec 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -7,111 +7,143 @@ namespace tk { namespace dnn { LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : Layer(net) { - checkCUDNN( cudnnCreateFilterDescriptor(¶mDesc)); - checkCUDNN( cudnnCreateRNNDescriptor(&rnnDesc) ); - checkCUDNN( cudnnCreateRNNDataDescriptor(&rnnDataDesc) ); - checkCUDNN( cudnnCreateDropoutDescriptor(&dropDesc)); + int batchSize = input_dim.n; + int inputSize = input_dim.c; + seqLen = input_dim.w; + stateSize = hiddensize; - 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, 1, h, w) ); + std::cout<<"LSTM seqLen: "<cudnnHandle, rnnDesc, hiddensize, numlayers, dropDesc, - cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT, - cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL, cudnnRNNMode_t::CUDNN_LSTM, - cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD, net->dataType) ); + // init Tensor Descriptors + std::vector x_vec(seqLen); + std::vector y_vec(seqLen); + std::vector dx_vec(seqLen); + std::vector dy_vec(seqLen); - // find dimension of params - size_t params_size = 0; - checkCUDNN( cudnnGetRNNParamsSize(net->cudnnHandle, rnnDesc, srcTensorDesc, ¶ms_size, net->dataType) ); - std::cout<<"Params size bytes: "<dataType, net->tensorFormat, 3, dimW)); - checkCuda( cudaMalloc(¶msSpace, params_size) ); + checkCUDNN(cudnnSetTensorNdDescriptor(x_vec[i], + net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dx_vec[i], + net->dataType, 3, dimA, strideA)); + dimA[0] = batchSize; + dimA[1] = bidirectional ? stateSize*2 : stateSize; + dimA[2] = 1; + strideA[0] = dimA[2] * dimA[1]; + strideA[1] = dimA[2]; + strideA[2] = 1; - - int numlinearlayers = 8; - - for(int i=0; icudnnHandle, rnnDesc, - i, srcTensorDesc, paramDesc, paramsSpace, - j, linLayerMatDesc, (void **)&linLayerMat)); - - if(linLayerMat == nullptr) { - FatalError("LSTM No weights in hidden layer"); - } - - cudnnDataType_t dataType; - cudnnTensorFormat_t format; - int nbDims; - int filterDimA[3]; - checkCUDNN(cudnnGetFilterNdDescriptor(linLayerMatDesc, 3, &dataType, - &format, &nbDims, filterDimA)); - std::cout<<"Wgs Dims: "<cudnnHandle, rnnDesc, - i, srcTensorDesc, paramDesc, paramsSpace, - j, linLayerBiasDesc, (void **)&linLayerBias)); - - if(linLayerMat == nullptr) { - FatalError("LSTM No bias in hidden layer"); - } - - checkCUDNN(cudnnGetFilterNdDescriptor(linLayerBiasDesc, 3, &dataType, - &format, &nbDims, filterDimA)); - std::cout<<"bias Dims: "<dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dy_vec[i], + net->dataType, 3, dimA, strideA)); } + // apply tensordesc + x_desc_vec_ = x_vec; + y_desc_vec_ = y_vec; + dx_desc_vec_ = dx_vec; + dy_desc_vec_ = dy_vec; + // set the state tensors + dimA[0] = numLayers * (bidirectional ? 2 : 1); + dimA[1] = batchSize; + dimA[2] = stateSize; + strideA[0] = dimA[2] * dimA[1]; + strideA[1] = dimA[2]; + strideA[2] = 1; + checkCUDNN(cudnnCreateTensorDescriptor(&hx_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&cx_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&hy_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&cy_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&dhx_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&dcx_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&dhy_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&dcy_desc_)); + checkCUDNN(cudnnSetTensorNdDescriptor(hx_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(cx_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(hy_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(cy_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dhx_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dcx_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dhy_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dcy_desc_, net->dataType, 3, dimA, strideA)); + // allocate dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; + checkCuda( cudaMalloc(&hx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&cx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&hy_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&cy_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); - checkCUDNN( cudnnCreateTensorDescriptor(&hiddenStateTensorDesc)); - checkCUDNN( cudnnSetTensor4dDescriptor(hiddenStateTensorDesc, - net->tensorFormat, net->dataType, 2*n, c, h, w) ); - checkCuda( cudaMalloc(&hiddenStateData, 2*input_dim.tot()*sizeof(dnnType)) ); - checkCUDNN( cudnnCreateTensorDescriptor(&cellStateTensorDesc)); - checkCUDNN( cudnnSetTensor4dDescriptor(cellStateTensorDesc, - net->tensorFormat, net->dataType, 2*n, c, h, w) ); - checkCuda( cudaMalloc(&cellStateData, 2*input_dim.tot()*sizeof(dnnType)) ); + + // Create Dropout descriptors // TODO: ??? IS IT NECESSARY ??? + float dropoutprob = 0.1f; // random val ???? + checkCUDNN(cudnnCreateDropoutDescriptor(&dropoutDesc)); + checkCUDNN(cudnnDropoutGetStatesSize(net->cudnnHandle, &dropout_byte_)); + dropout_size_ = dropout_byte_ / sizeof(dnnType); + checkCuda( cudaMalloc(&dropout_states_, dropout_byte_) ); + uint64_t seed_ = 17 + rand() % 4096; // NOLINT(runtime/threadsafe_fn) + checkCUDNN(cudnnSetDropoutDescriptor(dropoutDesc, + net->cudnnHandle, dropoutprob, dropout_states_, dropout_byte_, seed_)); + + + // RNN descriptors + checkCUDNN(cudnnCreateRNNDescriptor(&rnnDesc)); + + checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle, + rnnDesc, stateSize, numLayers, dropoutDesc, + cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT, + cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL, + cudnnRNNMode_t::CUDNN_LSTM, + cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD, + net->dataType)); + + + // Get temp space sizes + checkCUDNN(cudnnGetRNNWorkspaceSize(net->cudnnHandle, + rnnDesc, seqLen, x_desc_vec_.data(), &workspace_byte_)); + workspace_size_ = workspace_byte_ / sizeof(dnnType); + checkCuda( cudaMalloc(&work_space_, workspace_byte_) ); + + + // Check that number of params are correct + size_t cudnn_param_size; + checkCUDNN(cudnnGetRNNParamsSize(net->cudnnHandle, + rnnDesc,x_desc_vec_[0], &cudnn_param_size, net->dataType)); + int cudnn_params = cudnn_param_size/sizeof(dnnType); + std::cout<<"LSTM params size: "<dataType, net->tensorFormat, 3, dim_w)); + checkCUDNN(cudnnSetFilterNdDescriptor(dw_desc_, + net->dataType, net->tensorFormat, 3, dim_w)); + // allocate params dnnType *w_ptr, *dw_ptr; + checkCuda( cudaMalloc(&w_ptr, cudnn_params*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&dw_ptr, cudnn_params*sizeof(dnnType)) ); output_dim = input_dim; - output_dim.c = hiddensize*2; - checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc, - net->tensorFormat, net->dataType, output_dim.n, output_dim.c, output_dim.h, output_dim.w) ); - - - + output_dim.c = stateSize*2; //allocate data for infer result checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); @@ -123,19 +155,29 @@ LSTM::~LSTM() { } dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) { + std::cout<<"LSTM infer\n"; - checkCUDNN(cudnnRNNForwardInference( - net->cudnnHandle, rnnDesc, 1, - &srcTensorDesc, srcData, - hiddenStateTensorDesc, hiddenStateData, - cellStateTensorDesc, cellStateData, - paramDesc, paramsSpace, - &dstTensorDesc, dstData, - hiddenStateTensorDesc, hiddenStateData, - cellStateTensorDesc, cellStateData, - workSpace, ws_sizeInBytes - )); + checkCUDNN(cudnnRNNForwardInference(net->cudnnHandle, + rnnDesc, + seqLen, + x_desc_vec_.data(), // input array of desc + srcData, // input pointer + hx_desc_, // initial hidden state desc + hx_ptr, // initial hidden state pointer + cx_desc_, // initial cell state desc + cx_ptr, // initial cell state pointer + w_desc_, // weights desc + w_ptr, // weights pointer + y_desc_vec_.data(), // output desc + dstData, // output pointer + hy_desc_, // final hidden state desc + hy_ptr, // final hidden state pointer + cy_desc_, // final cell state desc + cy_ptr, // final cell state pointer + work_space_, // workspace pointer + workspace_byte_)); // workspace size + dim = output_dim; return dstData; } From c1c2173e4d9ea4f59d95ee069fa9e0fc1f58f31b Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 13 Feb 2020 23:10:48 +0100 Subject: [PATCH 03/11] removed unused var --- include/tkDNN/Layer.h | 10 ++++------ src/LSTM.cpp | 42 +++++++++++++----------------------------- 2 files changed, 17 insertions(+), 35 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index c538fc9..596fbec 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -246,18 +246,16 @@ protected: cudnnDropoutDescriptor_t dropoutDesc; dnnType *dropout_states_, *work_space_; - size_t workspace_byte_, reserve_space_byte_, dropout_byte_; + size_t workspace_byte_, dropout_byte_; int workspace_size_, dropout_size_; - std::vector x_desc_vec_, y_desc_vec_, dx_desc_vec_, dy_desc_vec_; + std::vector x_desc_vec_, y_desc_vec_; cudnnTensorDescriptor_t hx_desc_, cx_desc_; cudnnTensorDescriptor_t hy_desc_, cy_desc_; - cudnnTensorDescriptor_t dhx_desc_, dcx_desc_; - cudnnTensorDescriptor_t dhy_desc_, dcy_desc_; dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; - cudnnFilterDescriptor_t w_desc_, dw_desc_; - dnnType *w_ptr, *dw_ptr; + cudnnFilterDescriptor_t w_desc_; + dnnType *w_ptr; }; diff --git a/src/LSTM.cpp b/src/LSTM.cpp index be12cec..46a0593 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -17,16 +17,12 @@ LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : // init Tensor Descriptors std::vector x_vec(seqLen); std::vector y_vec(seqLen); - std::vector dx_vec(seqLen); - std::vector dy_vec(seqLen); int dimA[3]; int strideA[3]; for (int i = 0; i < seqLen; i++) { checkCUDNN(cudnnCreateTensorDescriptor(&x_vec[i])); checkCUDNN(cudnnCreateTensorDescriptor(&y_vec[i])); - checkCUDNN(cudnnCreateTensorDescriptor(&dx_vec[i])); - checkCUDNN(cudnnCreateTensorDescriptor(&dy_vec[i])); dimA[0] = batchSize; dimA[1] = inputSize; @@ -36,29 +32,21 @@ LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : strideA[0] = dimA[2] * dimA[1]; strideA[1] = dimA[2]; strideA[2] = 1; - checkCUDNN(cudnnSetTensorNdDescriptor(x_vec[i], net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dx_vec[i], - net->dataType, 3, dimA, strideA)); + dimA[0] = batchSize; dimA[1] = bidirectional ? stateSize*2 : stateSize; dimA[2] = 1; strideA[0] = dimA[2] * dimA[1]; strideA[1] = dimA[2]; strideA[2] = 1; - checkCUDNN(cudnnSetTensorNdDescriptor(y_vec[i], net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dy_vec[i], - net->dataType, 3, dimA, strideA)); } - // apply tensordesc x_desc_vec_ = x_vec; y_desc_vec_ = y_vec; - dx_desc_vec_ = dx_vec; - dy_desc_vec_ = dy_vec; // set the state tensors @@ -72,18 +60,10 @@ LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : checkCUDNN(cudnnCreateTensorDescriptor(&cx_desc_)); checkCUDNN(cudnnCreateTensorDescriptor(&hy_desc_)); checkCUDNN(cudnnCreateTensorDescriptor(&cy_desc_)); - checkCUDNN(cudnnCreateTensorDescriptor(&dhx_desc_)); - checkCUDNN(cudnnCreateTensorDescriptor(&dcx_desc_)); - checkCUDNN(cudnnCreateTensorDescriptor(&dhy_desc_)); - checkCUDNN(cudnnCreateTensorDescriptor(&dcy_desc_)); checkCUDNN(cudnnSetTensorNdDescriptor(hx_desc_, net->dataType, 3, dimA, strideA)); checkCUDNN(cudnnSetTensorNdDescriptor(cx_desc_, net->dataType, 3, dimA, strideA)); checkCUDNN(cudnnSetTensorNdDescriptor(hy_desc_, net->dataType, 3, dimA, strideA)); checkCUDNN(cudnnSetTensorNdDescriptor(cy_desc_, net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dhx_desc_, net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dcx_desc_, net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dhy_desc_, net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dcy_desc_, net->dataType, 3, dimA, strideA)); // allocate dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; checkCuda( cudaMalloc(&hx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); checkCuda( cudaMalloc(&cx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); @@ -130,28 +110,32 @@ LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : // Set param descriptors checkCUDNN(cudnnCreateFilterDescriptor(&w_desc_)); - checkCUDNN(cudnnCreateFilterDescriptor(&dw_desc_)); int dim_w[3] = {1, 1, 1}; dim_w[0] = cudnn_params; checkCUDNN(cudnnSetFilterNdDescriptor(w_desc_, net->dataType, net->tensorFormat, 3, dim_w)); - checkCUDNN(cudnnSetFilterNdDescriptor(dw_desc_, - net->dataType, net->tensorFormat, 3, dim_w)); - // allocate params dnnType *w_ptr, *dw_ptr; + // allocate params dnnType *w_ptr; checkCuda( cudaMalloc(&w_ptr, cudnn_params*sizeof(dnnType)) ); - checkCuda( cudaMalloc(&dw_ptr, cudnn_params*sizeof(dnnType)) ); - + // set output dim output_dim = input_dim; output_dim.c = stateSize*2; - + //allocate data for infer result checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); } LSTM::~LSTM() { + checkCuda(cudaFree(hx_ptr)); + checkCuda(cudaFree(cx_ptr)); + checkCuda(cudaFree(hy_ptr)); + checkCuda(cudaFree(cy_ptr)); + checkCuda(cudaFree(w_ptr )); - checkCuda( cudaFree(dstData) ); + checkCuda(cudaFree(work_space_ )); + checkCuda(cudaFree(dropout_states_)); + + checkCuda(cudaFree(dstData)); } dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) { From 4fa5d2c231900e8ee9bfbab908b2ac97add974db Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 13 Feb 2020 23:21:28 +0100 Subject: [PATCH 04/11] lstm return seq --- include/tkDNN/Layer.h | 3 ++- src/LSTM.cpp | 16 ++++++++++++---- tests/imuodom/imuodom.cpp | 3 ++- 3 files changed, 16 insertions(+), 6 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 596fbec..240d9a2 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -230,13 +230,14 @@ protected: class LSTM : public Layer { public: - LSTM(Network *net, int hiddensize, std::string fname_weights); + LSTM(Network *net, int hiddensize, bool returnSeq, std::string fname_weights); virtual ~LSTM(); virtual layerType_t getLayerType() { return LAYER_LSTM; }; virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); const bool bidirectional = 1; /**> is the net bidir */ + bool returnSeq = false; /**> if false return only the result of last timestep */ int stateSize = 0; /**> number of hidden states */ int seqLen = 0; /**> number of timesteps */ int numLayers = 1; /**> number of internal layers */ diff --git a/src/LSTM.cpp b/src/LSTM.cpp index 46a0593..c31bccf 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -4,9 +4,10 @@ namespace tk { namespace dnn { -LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : +LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weights) : Layer(net) { + this->returnSeq = returnSeq; int batchSize = input_dim.n; int inputSize = input_dim.c; seqLen = input_dim.w; @@ -117,12 +118,19 @@ LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : // allocate params dnnType *w_ptr; checkCuda( cudaMalloc(&w_ptr, cudnn_params*sizeof(dnnType)) ); + + + //allocate data for infer result + int dstDim = input_dim.n * stateSize*2 * input_dim.h * input_dim.w; + checkCuda( cudaMalloc(&dstData, dstDim*sizeof(dnnType)) ); + // set output dim output_dim = input_dim; output_dim.c = stateSize*2; - - //allocate data for infer result - checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); + if(!returnSeq) { + output_dim.h = 1; + output_dim.w = 1; + } } LSTM::~LSTM() { diff --git a/tests/imuodom/imuodom.cpp b/tests/imuodom/imuodom.cpp index 14146b2..6fc3fba 100644 --- a/tests/imuodom/imuodom.cpp +++ b/tests/imuodom/imuodom.cpp @@ -48,7 +48,8 @@ int main() { tk::dnn::Layer *concat_l[3] = { &x0_2, &x1_2, &x2_2 }; tk::dnn::Route concat (&net, concat_l, 3); - tk::dnn::LSTM lstm0(&net, 128, "ciao"); + tk::dnn::LSTM lstm0(&net, 128, true, "ciao"); + tk::dnn::LSTM lstm1(&net, 128, false, "ciao"); net.print(); From 4746121d438c72d0287a6ff7edb193a2da68029d Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Sat, 15 Feb 2020 20:37:08 +0100 Subject: [PATCH 05/11] LSTM params --- include/tkDNN/Layer.h | 5 ++- src/LSTM.cpp | 94 +++++++++++++++++++++++++++++++++------ tests/imuodom/imuodom.cpp | 20 ++++++--- tests/imuodom/infer.py | 1 + tests/weights_exporter.py | 16 +++++-- 5 files changed, 111 insertions(+), 25 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 240d9a2..319df6a 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -212,6 +212,7 @@ protected: https://github.com/jiangnanhugo/seq2seq_cuda/blob/e4dbdcfa0517c972bfd4beea9f11a5233954093c/src/rnn.cpp https://github.com/Jeffery-Song/mxnet-test/blob/aab666faad44011f7a67b527b5f6c960367d0422/src/operator/cudnn_rnn-inl.h https://stackoverflow.com/a/38737941 + https://colah.github.io/posts/2015-08-Understanding-LSTMs/ PARAMS (numlayers*2): layer0: @@ -236,7 +237,7 @@ public: virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); - const bool bidirectional = 1; /**> is the net bidir */ + const bool bidirectional = false; /**> is the net bidir */ bool returnSeq = false; /**> if false return only the result of last timestep */ int stateSize = 0; /**> number of hidden states */ int seqLen = 0; /**> number of timesteps */ @@ -254,9 +255,11 @@ protected: cudnnTensorDescriptor_t hx_desc_, cx_desc_; cudnnTensorDescriptor_t hy_desc_, cy_desc_; dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; + int stateDataDim; cudnnFilterDescriptor_t w_desc_; dnnType *w_ptr; + dnnType *w_h; }; diff --git a/src/LSTM.cpp b/src/LSTM.cpp index c31bccf..ce2a9f5 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -66,10 +66,12 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig checkCUDNN(cudnnSetTensorNdDescriptor(hy_desc_, net->dataType, 3, dimA, strideA)); checkCUDNN(cudnnSetTensorNdDescriptor(cy_desc_, net->dataType, 3, dimA, strideA)); // allocate dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; - checkCuda( cudaMalloc(&hx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); - checkCuda( cudaMalloc(&cx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); - checkCuda( cudaMalloc(&hy_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); - checkCuda( cudaMalloc(&cy_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); + stateDataDim = dimA[0]*dimA[1]*dimA[2]; + checkCuda( cudaMalloc(&hx_ptr, stateDataDim*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&cx_ptr, stateDataDim*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&hy_ptr, stateDataDim*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&cy_ptr, stateDataDim*sizeof(dnnType)) ); + // Create Dropout descriptors // TODO: ??? IS IT NECESSARY ??? @@ -89,7 +91,7 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle, rnnDesc, stateSize, numLayers, dropoutDesc, cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT, - cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL, + (bidirectional ? cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL : cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL), cudnnRNNMode_t::CUDNN_LSTM, cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD, net->dataType)); @@ -115,22 +117,81 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig dim_w[0] = cudnn_params; checkCUDNN(cudnnSetFilterNdDescriptor(w_desc_, net->dataType, net->tensorFormat, 3, dim_w)); - // allocate params dnnType *w_ptr; - checkCuda( cudaMalloc(&w_ptr, cudnn_params*sizeof(dnnType)) ); - + // load params + readBinaryFile(fname_weights, cudnn_params, &w_h, &w_ptr); //allocate data for infer result - int dstDim = input_dim.n * stateSize*2 * input_dim.h * input_dim.w; + int dstDim = input_dim.n * stateSize*(bidirectional ? 2 : 1) * input_dim.h * input_dim.w; checkCuda( cudaMalloc(&dstData, dstDim*sizeof(dnnType)) ); // set output dim output_dim = input_dim; - output_dim.c = stateSize*2; + output_dim.c = stateSize*(bidirectional ? 2 : 1); if(!returnSeq) { output_dim.h = 1; output_dim.w = 1; } + + + + + // Query weight layout + cudnnFilterDescriptor_t m_desc; + checkCUDNN(cudnnCreateFilterDescriptor(&m_desc)); + dnnType *p; + int n = 8; // lstm layers + + printCenteredTitle("WEIGHTS", '=', 20); + for (int i = 0; i < numLayers*(bidirectional?2:1); ++i) { + for (int j = 0; j < n; ++j) { + + checkCUDNN(cudnnGetRNNLinLayerMatrixParams(net->cudnnHandle, rnnDesc, + i, x_desc_vec_[0], w_desc_, 0, j, m_desc, (void**)&p)); + + std::cout << "ptr: " << ((int64_t)(p - NULL))/sizeof(dnnType)<<"\n"; + + cudnnDataType_t t; + cudnnTensorFormat_t f; + int ndim = 5; + int dims[5] = {0, 0, 0, 0, 0}; + checkCUDNN(cudnnGetFilterNdDescriptor(m_desc, ndim, &t, &f, &ndim, &dims[0])); + std::cout << "(layer, linlayer): " << i << " " << j << "\n"; + + int tot = 1; + for (int i = 0; i < ndim; ++i) { + std::cout << dims[i] << " "; + tot *= dims[i]; + } + std::cout<<"\t-> "< Date: Sun, 16 Feb 2020 16:28:39 +0100 Subject: [PATCH 06/11] works but it need cleaning --- include/tkDNN/Layer.h | 3 +- include/tkDNN/utils.h | 2 +- src/LSTM.cpp | 163 ++++++++++++++++++++++++++++------- src/utils.cpp | 4 +- tests/imuodom/imuodom.cpp | 26 +++--- tests/imuodom/infer.py | 25 +++--- tests/simple/test_model.py | 15 ++-- tests/simple/test_simple.cpp | 8 +- tests/weights_exporter.py | 44 ++++++++-- 9 files changed, 219 insertions(+), 71 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 319df6a..f12027f 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -237,7 +237,7 @@ public: virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); - const bool bidirectional = false; /**> is the net bidir */ + const bool bidirectional = true; /**> is the net bidir */ bool returnSeq = false; /**> if false return only the result of last timestep */ int stateSize = 0; /**> number of hidden states */ int seqLen = 0; /**> number of timesteps */ @@ -260,6 +260,7 @@ protected: cudnnFilterDescriptor_t w_desc_; dnnType *w_ptr; dnnType *w_h; + dnnType *wf_ptr, *wb_ptr; // params pointer forward and backward layer }; diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h index dc34a31..3fa9d34 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -91,7 +91,7 @@ void printCenteredTitle(const char *title, char fill, int dim); bool fileExist(const char *fname); void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0); -int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true); +int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true, int limit = 10); void printDeviceVector(int size, dnnType* vec_d, bool device = true); void resize(int size, dnnType **data); diff --git a/src/LSTM.cpp b/src/LSTM.cpp index ce2a9f5..00d8f76 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -37,7 +37,7 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig net->dataType, 3, dimA, strideA)); dimA[0] = batchSize; - dimA[1] = bidirectional ? stateSize*2 : stateSize; + dimA[1] = stateSize; dimA[2] = 1; strideA[0] = dimA[2] * dimA[1]; strideA[1] = dimA[2]; @@ -51,7 +51,7 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig // set the state tensors - dimA[0] = numLayers * (bidirectional ? 2 : 1); + dimA[0] = numLayers; dimA[1] = batchSize; dimA[2] = stateSize; strideA[0] = dimA[2] * dimA[1]; @@ -91,7 +91,8 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle, rnnDesc, stateSize, numLayers, dropoutDesc, cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT, - (bidirectional ? cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL : cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL), + //(bidirectional ? cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL : cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL), + cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL, cudnnRNNMode_t::CUDNN_LSTM, cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD, net->dataType)); @@ -119,23 +120,26 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig net->dataType, net->tensorFormat, 3, dim_w)); // load params - readBinaryFile(fname_weights, cudnn_params, &w_h, &w_ptr); - - //allocate data for infer result - int dstDim = input_dim.n * stateSize*(bidirectional ? 2 : 1) * input_dim.h * input_dim.w; - checkCuda( cudaMalloc(&dstData, dstDim*sizeof(dnnType)) ); + readBinaryFile(fname_weights, cudnn_params*2, &w_h, &w_ptr); + // set forward and backward params + wf_ptr = w_ptr; + wb_ptr = w_ptr + cudnn_params; + std::cout<<"wf: "<cublasHandle, srcData, trans, dim.c, dim.h*dim.w*dim.l); + srcData = trans; + + // reposition in invered order + dnnType *srcBack; + checkCuda( cudaMalloc(&srcBack, dim.tot()*sizeof(dnnType))); + for(int i=0; icudnnHandle, + rnnDesc, + seqLen, // number of time steps (nT) + x_desc_vec_.data(), // input array of desc (nT*nC_in) + srcData, // input pointer + hx_desc_, // initial hidden state desc + hx_ptr, // initial hidden state pointer + cx_desc_, // initial cell state desc + cx_ptr, // initial cell state pointer + w_desc_, // weights desc + wf_ptr, // weights pointer + y_desc_vec_.data(), // output desc (nT*nC_out) + dstF, // output pointer + hy_desc_, // final hidden state desc + hy_ptr, // final hidden state pointer + cy_desc_, // final cell state desc + cy_ptr, // final cell state pointer + work_space_, // workspace pointer + workspace_byte_)); // workspace size + } + std::cout<<"OUTPUT F:\n"; + printDeviceVector(singleOutput.tot(), dstF); + + std::cout<<"INPUT:\n"; + printDeviceVector(input_dim.tot(), srcBack); + + // backward + { + // reset states + checkCuda( cudaMemset(hx_ptr, 0, stateDataDim*sizeof(float)) ); + checkCuda( cudaMemset(cx_ptr, 0, stateDataDim*sizeof(float)) ); + + checkCUDNN(cudnnRNNForwardInference(net->cudnnHandle, + rnnDesc, + seqLen, // number of time steps (nT) + x_desc_vec_.data(), // input array of desc (nT*nC_in) + srcBack, // input pointer + hx_desc_, // initial hidden state desc + hx_ptr, // initial hidden state pointer + cx_desc_, // initial cell state desc + cx_ptr, // initial cell state pointer + w_desc_, // weights desc + wb_ptr, // weights pointer + y_desc_vec_.data(), // output desc (nT*nC_out) + dstB, // output pointer + hy_desc_, // final hidden state desc + hy_ptr, // final hidden state pointer + cy_desc_, // final cell state desc + cy_ptr, // final cell state pointer + work_space_, // workspace pointer + workspace_byte_)); // workspace size + } + + + // reposition in invered order + dnnType *dstBack; + checkCuda( cudaMalloc(&dstBack, singleOutput.tot()*sizeof(dnnType))); + for(int i=0; icublasHandle, dstF, trans, + singleOutput.h*singleOutput.w*singleOutput.l, singleOutput.c); + // backward transpose + matrixTranspose(net->cublasHandle, dstB, trans + singleOutput.tot(), + singleOutput.h*singleOutput.w*singleOutput.l, singleOutput.c); + dstData = trans; + } else { + // copy last of forward + checkCuda( cudaMemcpy(trans, dstF + singleOutput.tot() - singleOutput.c, singleOutput.c*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + // copy first of backward + checkCuda( cudaMemcpy(trans + singleOutput.c, dstB, singleOutput.c*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + dstData = trans; + } - checkCUDNN(cudnnRNNForwardInference(net->cudnnHandle, - rnnDesc, - seqLen, // number of time steps (nT) - x_desc_vec_.data(), // input array of desc (nT*nC_in) - srcData, // input pointer - hx_desc_, // initial hidden state desc - hx_ptr, // initial hidden state pointer - cx_desc_, // initial cell state desc - cx_ptr, // initial cell state pointer - w_desc_, // weights desc - w_ptr, // weights pointer - y_desc_vec_.data(), // output desc (nT*nC_out) - dstData, // output pointer - hy_desc_, // final hidden state desc - hy_ptr, // final hidden state pointer - cy_desc_, // final cell state desc - cy_ptr, // final cell state pointer - work_space_, // workspace pointer - workspace_byte_)); // workspace size dim = output_dim; return dstData; diff --git a/src/utils.cpp b/src/utils.cpp index 444318d..6789e8c 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -68,7 +68,7 @@ void printDeviceVector(int size, dnnType* vec_d, bool device) delete [] vec; } -int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) { +int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device, int limit) { dnnType *data_h, *correct_h; const float eps = 0.02f; @@ -92,7 +92,7 @@ int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) { diffs += 1; if(diffs == 1) std::cout<<"\n"; - if(diffs < 10) + if(diffs < limit) std::cout<<" | [ "< Date: Sun, 16 Feb 2020 17:08:19 +0100 Subject: [PATCH 07/11] structure ok, result wrong --- include/tkDNN/Layer.h | 9 +++- src/LSTM.cpp | 103 ++++++++++++++++---------------------- tests/imuodom/imuodom.cpp | 4 +- 3 files changed, 55 insertions(+), 61 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index f12027f..7e827d8 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -207,7 +207,9 @@ protected: /** Bidirectional LSTM layer - + ONLY BIDIRECTIONAL (TODO: more configurable) + currently implemented as 2 inferences: forward and backward (TODO: only 1 cudnn inference) + implementation info: https://github.com/jiangnanhugo/seq2seq_cuda/blob/e4dbdcfa0517c972bfd4beea9f11a5233954093c/src/rnn.cpp https://github.com/Jeffery-Song/mxnet-test/blob/aab666faad44011f7a67b527b5f6c960367d0422/src/operator/cudnn_rnn-inl.h @@ -261,6 +263,11 @@ protected: dnnType *w_ptr; dnnType *w_h; dnnType *wf_ptr, *wb_ptr; // params pointer forward and backward layer + + // used during inference + dataDim_t one_output_dim; // output dim of as single inference + dnnType *srcF, *srcB; // input of single inference + dnnType *dstF, *dstB_NR, *dstB; // output of single inference, dstB_NR = dstB not reversed }; diff --git a/src/LSTM.cpp b/src/LSTM.cpp index 00d8f76..44502fa 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -13,8 +13,6 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig seqLen = input_dim.w; stateSize = hiddensize; - std::cout<<"LSTM seqLen: "< x_vec(seqLen); std::vector y_vec(seqLen); @@ -110,7 +108,7 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig checkCUDNN(cudnnGetRNNParamsSize(net->cudnnHandle, rnnDesc,x_desc_vec_[0], &cudnn_param_size, net->dataType)); int cudnn_params = cudnn_param_size/sizeof(dnnType); - std::cout<<"LSTM params size: "<cudnnHandle, rnnDesc, i, x_desc_vec_[0], w_desc_, 0, j, m_desc, (void**)&p)); @@ -209,37 +217,27 @@ LSTM::~LSTM() { checkCuda(cudaFree(work_space_ )); checkCuda(cudaFree(dropout_states_)); + checkCuda(cudaFree(srcF)); + checkCuda(cudaFree(srcB)); + checkCuda(cudaFree(dstF)); + checkCuda(cudaFree(dstB_NR)); + checkCuda(cudaFree(dstB)); checkCuda(cudaFree(dstData)); } dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) { - std::cout<<"LSTM infer\n"; + // transpose input + matrixTranspose(net->cublasHandle, srcData, srcF, dim.c, dim.h*dim.w*dim.l); - dnnType *trans; - checkCuda( cudaMalloc(&trans, dim.tot()*sizeof(dnnType))); - matrixTranspose(net->cublasHandle, srcData, trans, dim.c, dim.h*dim.w*dim.l); - srcData = trans; - - // reposition in invered order - dnnType *srcBack; - checkCuda( cudaMalloc(&srcBack, dim.tot()*sizeof(dnnType))); + // build srcB as reversed srcF for(int i=0; i