merge
This commit is contained in:
Francesco Gatti
+4
-2
@@ -14,12 +14,14 @@ Activation::Activation(Network *net, dataDim_t input_dim, tkdnnActivationMode_t
|
||||
checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc,
|
||||
net->tensorFormat,
|
||||
net->dataType,
|
||||
input_dim.n, input_dim.c,
|
||||
input_dim.n*input_dim.l,
|
||||
input_dim.c,
|
||||
input_dim.h, input_dim.w) );
|
||||
checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc,
|
||||
net->tensorFormat,
|
||||
net->dataType,
|
||||
input_dim.n, input_dim.c,
|
||||
input_dim.n*input_dim.l,
|
||||
input_dim.c,
|
||||
input_dim.h, input_dim.w) );
|
||||
}
|
||||
|
||||
|
||||
+157
@@ -0,0 +1,157 @@
|
||||
#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;
|
||||
}
|
||||
|
||||
}
|
||||
@@ -0,0 +1,37 @@
|
||||
#include <iostream>
|
||||
|
||||
#include "Layer.h"
|
||||
#include "kernels.h"
|
||||
|
||||
namespace tkDNN {
|
||||
|
||||
Flatten::Flatten(Network *net, dataDim_t input_dim) :
|
||||
Layer(net, input_dim) {
|
||||
|
||||
checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) );
|
||||
|
||||
output_dim.n = 1;
|
||||
output_dim.c = input_dim.tot();
|
||||
output_dim.h = 1;
|
||||
output_dim.w = 1;
|
||||
output_dim.l = 1;
|
||||
|
||||
}
|
||||
|
||||
Flatten::~Flatten() {
|
||||
|
||||
checkCuda( cudaFree(dstData) );
|
||||
}
|
||||
|
||||
value_type* Flatten::infer(dataDim_t &dim, value_type* srcData) {
|
||||
|
||||
//transpose per channel
|
||||
matrixTranspose(net->cublasHandle, srcData, dstData, dim.c, dim.h*dim.w*dim.l);
|
||||
|
||||
//update data dimensions
|
||||
dim = output_dim;
|
||||
|
||||
return dstData;
|
||||
}
|
||||
|
||||
}
|
||||
@@ -0,0 +1,45 @@
|
||||
#include <iostream>
|
||||
|
||||
#include "Layer.h"
|
||||
#include "kernels.h"
|
||||
|
||||
namespace tkDNN {
|
||||
|
||||
MulAdd::MulAdd(Network *net, dataDim_t input_dim, value_type mul, value_type add) :
|
||||
Layer(net, input_dim) {
|
||||
|
||||
this->mul = mul;
|
||||
this->add = add;
|
||||
|
||||
int size = input_dim.tot();
|
||||
|
||||
// create a vector with all value setted to add
|
||||
value_type *add_vector_h = new value_type[size];
|
||||
for(int i=0; i<size; i++)
|
||||
add_vector_h[i] = add;
|
||||
|
||||
checkCuda( cudaMalloc(&add_vector, size*sizeof(value_type)));
|
||||
checkCuda( cudaMemcpy(add_vector, add_vector_h, size*sizeof(value_type), cudaMemcpyHostToDevice));
|
||||
delete [] add_vector_h;
|
||||
|
||||
|
||||
checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) );
|
||||
}
|
||||
|
||||
MulAdd::~MulAdd() {
|
||||
|
||||
checkCuda( cudaFree(add_vector) );
|
||||
checkCuda( cudaFree(dstData) );
|
||||
}
|
||||
|
||||
value_type* MulAdd::infer(dataDim_t &dim, value_type* srcData) {
|
||||
|
||||
matrixMulAdd(net->cublasHandle, srcData, dstData, add_vector, input_dim.tot(), mul);
|
||||
|
||||
//update data dimensions
|
||||
dim = output_dim;
|
||||
|
||||
return dstData;
|
||||
}
|
||||
|
||||
}
|
||||
+5
-1
@@ -1,5 +1,6 @@
|
||||
#include <iostream>
|
||||
|
||||
#include "tkdnn.h"
|
||||
#include "Network.h"
|
||||
#include "Layer.h"
|
||||
|
||||
@@ -7,7 +8,10 @@ namespace tkDNN {
|
||||
|
||||
Network::Network() {
|
||||
|
||||
std::cout<<"New NETWORK with CUDNN v"<<float(cudnnGetVersion())/1000<<"\n";
|
||||
float tk_ver = float(tkDNN::getVersion())/1000;
|
||||
float cu_ver = float(cudnnGetVersion())/1000;
|
||||
|
||||
std::cout<<"New NETWORK (tkDNN v"<<tk_ver<<", CUDNN v"<<cu_ver<<")\n";
|
||||
dataType = CUDNN_DATA_FLOAT;
|
||||
tensorFormat = CUDNN_TENSOR_NCHW;
|
||||
|
||||
|
||||
+111
@@ -0,0 +1,111 @@
|
||||
#include <iostream>
|
||||
|
||||
#include "Layer.h"
|
||||
#include "kernels.h"
|
||||
|
||||
namespace tkDNN {
|
||||
|
||||
Pooling::Pooling( Network *net, dataDim_t input_dim,
|
||||
int winH, int winW, int strideH, int strideW, tkdnnPoolingMode_t pool_mode) :
|
||||
Layer(net, input_dim) {
|
||||
|
||||
|
||||
if(winH != strideH || winW != strideW)
|
||||
FatalError("stride pooling not yet implemented");
|
||||
|
||||
this->winH = winH;
|
||||
this->winW = winW;
|
||||
this->strideH = strideH;
|
||||
this->strideW = strideW;
|
||||
this->pool_mode = pool_mode;
|
||||
|
||||
checkCUDNN( cudnnCreatePoolingDescriptor(&poolingDesc) );
|
||||
|
||||
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;
|
||||
|
||||
poolOn3d = false;
|
||||
|
||||
if(l > 1) {
|
||||
poolOn3d = true;
|
||||
|
||||
if(n != 1)
|
||||
FatalError("N value on 3d pool must be 1");
|
||||
|
||||
//use batch as l
|
||||
n = l;
|
||||
}
|
||||
|
||||
checkCUDNN( cudnnSetPooling2dDescriptor(poolingDesc, cudnnPoolingMode_t(pool_mode),
|
||||
winH, winW, 0, 0, strideH, strideW) );
|
||||
|
||||
checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc,
|
||||
net->tensorFormat, net->dataType, n, c, h, w) );
|
||||
|
||||
//get out dim
|
||||
h = h / winH; w = w / winW;
|
||||
|
||||
checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc,
|
||||
net->tensorFormat, net->dataType, n, c, h, w) );
|
||||
|
||||
|
||||
output_dim.n = n;
|
||||
output_dim.c = c;
|
||||
output_dim.h = h;
|
||||
output_dim.w = w;
|
||||
output_dim.l = l;
|
||||
|
||||
checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) );
|
||||
|
||||
//pool on 3d data need transposition at the enter and on the exit
|
||||
//allocate for initial and final transposition
|
||||
if(poolOn3d) {
|
||||
output_dim.n = 1;
|
||||
|
||||
checkCuda( cudaMalloc(&tmpInputData, input_dim.tot()*sizeof(value_type)) );
|
||||
checkCuda( cudaMalloc(&tmpOutputData, output_dim.tot()*sizeof(value_type)) );
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
Pooling::~Pooling() {
|
||||
|
||||
if(poolOn3d) {
|
||||
checkCuda( cudaFree(tmpInputData) );
|
||||
checkCuda( cudaFree(tmpOutputData) );
|
||||
}
|
||||
|
||||
checkCUDNN( cudnnDestroyPoolingDescriptor(poolingDesc) );
|
||||
checkCuda( cudaFree(dstData) );
|
||||
}
|
||||
|
||||
value_type* Pooling::infer(dataDim_t &dim, value_type* srcData) {
|
||||
|
||||
value_type *poolSrc = srcData;
|
||||
value_type *poolDst = dstData;
|
||||
|
||||
if(poolOn3d) {
|
||||
matrixTranspose(net->cublasHandle, srcData, tmpInputData, dim.h*dim.w*dim.c, dim.l);
|
||||
poolSrc = tmpInputData;
|
||||
poolDst = tmpOutputData;
|
||||
}
|
||||
|
||||
value_type alpha = value_type(1);
|
||||
value_type beta = value_type(0);
|
||||
checkCUDNN( cudnnPoolingForward(net->cudnnHandle, poolingDesc,
|
||||
&alpha, srcTensorDesc, poolSrc,
|
||||
&beta, dstTensorDesc, poolDst) );
|
||||
|
||||
//update dim
|
||||
dim = output_dim;
|
||||
|
||||
if(poolOn3d)
|
||||
matrixTranspose(net->cublasHandle, tmpOutputData, dstData, dim.l, dim.h*dim.w*dim.c);
|
||||
|
||||
return dstData;
|
||||
}
|
||||
|
||||
}
|
||||
@@ -9,7 +9,7 @@
|
||||
__global__
|
||||
void activation_elu(value_type *input, value_type *output, int size) {
|
||||
|
||||
int i = threadIdx.x*(blockIdx.x +1);
|
||||
int i = blockDim.x*blockIdx.x + threadIdx.x;
|
||||
|
||||
if(i<size) {
|
||||
value_type k0, k1;
|
||||
@@ -30,6 +30,9 @@ void activation_elu(value_type *input, value_type *output, int size) {
|
||||
*/
|
||||
void activationELUForward(value_type* srcData, value_type* dstData, int size)
|
||||
{
|
||||
activation_elu<<<(size+255)/256, 256>>>(srcData, dstData, size);
|
||||
int blocks = (size+255)/256;
|
||||
int threads = 256;
|
||||
|
||||
activation_elu<<<blocks, threads>>>(srcData, dstData, size);
|
||||
checkCuda( cudaDeviceSynchronize() );
|
||||
}
|
||||
@@ -42,4 +42,25 @@ void resize(int size, value_type **data)
|
||||
if (*data != NULL)
|
||||
checkCuda( cudaFree(*data) );
|
||||
checkCuda( cudaMalloc(data, size*sizeof(value_type)) );
|
||||
}
|
||||
|
||||
void matrixTranspose(cublasHandle_t handle, value_type* srcData, value_type* dstData, int rows, int cols) {
|
||||
|
||||
value_type *A = srcData, *clone = dstData;
|
||||
int m = rows, n= cols;
|
||||
checkCuda( cudaMemcpy(clone, A, m*n*sizeof(value_type), cudaMemcpyDeviceToDevice));
|
||||
|
||||
float const alpha(1.0);
|
||||
float const beta(0.0);
|
||||
checkERROR( cublasSgeam( handle, CUBLAS_OP_T, CUBLAS_OP_N, m, n, &alpha, A, n, &beta, A, m, clone, m ));
|
||||
}
|
||||
|
||||
void matrixMulAdd( cublasHandle_t handle, value_type* srcData, value_type* dstData,
|
||||
value_type* add_vector, int dim, value_type mul) {
|
||||
|
||||
checkCuda( cudaMemcpy(dstData, add_vector, dim*sizeof(value_type), cudaMemcpyDeviceToDevice));
|
||||
|
||||
value_type alpha = mul;
|
||||
checkERROR( cublasSaxpy(handle, dim, &alpha, srcData, 1, dstData, 1));
|
||||
|
||||
}
|
||||
Reference in New Issue
Block a user