mmr/tkDNN
Archived
1
0
Fork 0
Code Issues 46 Pull Requests 12 Releases 4 Activity

cudnn5 branch for tx2

Browse Source
This commit is contained in:
Francesco Gatti
2017-07-14 11:11:51 +02:00
parent 33513468e6
commit e4df86a07c
9 changed files with 49 additions and 276 deletions
Download Patch File Download Diff File Expand all files Collapse all files
CMakeLists.txt
+3 -3
View File
@@ -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)
target_link_libraries(tkDNNtest tkDNN)
include/Layer.h
+4 -41
View File
@@ -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
#endif //LAYER_H
src/Activation.cpp
+19 -17
View File
@@ -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;
}
}
}
src/Conv2d.cpp
+3 -3
View File
@@ -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;
}
}
}
src/Conv3d.cpp
-157
View File
@@ -1,157 +0,0 @@
#include <iostream>
#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;
}
}
src/Dense.cpp
+1 -1
View File
@@ -56,4 +56,4 @@ value_type* Dense::infer(dataDim_t &dim, value_type* srcData) {
return dstData;
}
}
}
src/Pooling.cpp
+2 -2
View File
@@ -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;
}
}
}
tests/test.cpp
+8 -25
View File
@@ -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;
}
}
tests/test_model.py
+9 -27
View File
@@ -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)
print "Result shape: ", np.shape(r)