From b94931f9f79b8fbef19dec11489ad367277a63a0 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 1 Aug 2017 16:08:56 +0200 Subject: [PATCH] yolo layers --- CMakeLists.txt | 8 ++- include/Layer.h | 100 +++++++++++++++++++++++------ include/Network.h | 1 - include/kernels.h | 9 ++- include/utils.h | 3 +- src/Activation.cpp | 57 +++++++++------- src/Conv2d.cpp | 31 ++++++--- src/Dense.cpp | 4 +- src/LayerWgs.cpp | 32 +++++++-- src/Network.cpp | 2 +- src/Region.cpp | 60 +++++++++++++++++ src/Reorg.cpp | 35 ++++++++++ src/Route.cpp | 62 ++++++++++++++++++ src/kernels/activation_elu.cu | 2 +- src/kernels/activation_leaky.cu | 29 +++++++++ src/kernels/activation_logistic.cu | 26 ++++++++ src/kernels/reorg.cu | 49 ++++++++++++++ src/kernels/softmax.cu | 43 +++++++++++++ src/utils.cpp | 30 ++++++++- tests/mnist/test.cpp | 12 ++-- tests/test/test.cpp | 11 ++-- 21 files changed, 522 insertions(+), 84 deletions(-) create mode 100644 src/Region.cpp create mode 100644 src/Reorg.cpp create mode 100644 src/Route.cpp create mode 100644 src/kernels/activation_leaky.cu create mode 100644 src/kernels/activation_logistic.cu create mode 100644 src/kernels/reorg.cu create mode 100644 src/kernels/softmax.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index fefc792..71f8908 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -5,12 +5,16 @@ project (tkDNN) find_package(CUDA QUIET REQUIRED) cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS}) -cuda_add_library(kernels SHARED src/kernels/activation_elu.cu) +cuda_add_library(kernels SHARED src/kernels/activation_elu.cu + src/kernels/activation_leaky.cu + src/kernels/activation_logistic.cu + src/kernels/reorg.cu + src/kernels/softmax.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/Flatten.cpp src/MulAdd.cpp src/Pooling.cpp src/Softmax.cpp - src/Network.cpp src/utils.cpp) + src/Route.cpp src/Reorg.cpp src/Region.cpp src/Network.cpp src/utils.cpp) target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn) add_executable(test_simple tests/test/test.cpp) diff --git a/include/Layer.h b/include/Layer.h index 98aead8..5e96ad5 100644 --- a/include/Layer.h +++ b/include/Layer.h @@ -49,10 +49,12 @@ public: } dataDim_t input_dim, output_dim; + value_type *dstData; //where results will be putted protected: Network *net; cudnnTensorDescriptor_t srcTensorDesc, dstTensorDesc; + }; @@ -64,15 +66,21 @@ class LayerWgs : public Layer { public: LayerWgs(Network *net, dataDim_t input_dim, int inputs, int outputs, int kh, int kw, int kt, - const char* fname_weights, const char* fname_bias); + const char* fname_weights, bool batchnorm = false); virtual ~LayerWgs(); protected: int inputs, outputs; - std::string weights_path, bias_path; + std::string weights_path; value_type *data_h, *data_d; value_type *bias_h, *bias_d; + + //batchnorm + bool batchnorm; + value_type *scales_h, *scales_d; + value_type *mean_h, *mean_d; + value_type *variance_h, *variance_d; }; @@ -83,31 +91,35 @@ class Dense : public LayerWgs { public: Dense(Network *net, dataDim_t in_dim, int out_ch, - const char* fname_weights, const char* fname_bias); + const char* fname_weights); virtual ~Dense(); virtual value_type* infer(dataDim_t &dim, value_type* srcData); - -protected: - value_type *dstData; //where results will be putted }; +/** + Avaible activation functions +*/ +typedef enum { + ACTIVATION_ELU = 100, + ACTIVATION_LEAKY = 101 +} tkdnnActivationMode_t; + /** Activation layer (it doesnt need weigths) */ class Activation : public Layer { public: - Activation(Network *net, dataDim_t input_dim, cudnnActivationMode_t act_mode); + Activation(Network *net, dataDim_t input_dim, int act_mode); virtual ~Activation(); virtual value_type* infer(dataDim_t &dim, value_type* srcData); protected: - cudnnActivationMode_t act_mode; + int act_mode; cudnnActivationDescriptor_t activDesc; - value_type *dstData; //where results will be putted }; @@ -117,15 +129,15 @@ protected: class Conv2d : public LayerWgs { public: - Conv2d(Network *net, dataDim_t in_dim, int out_ch, - int kernelH, int kernelW, int strideH, int strideW, - const char* fname_weights, const char* fname_bias); + Conv2d( Network *net, dataDim_t in_dim, int out_ch, + int kernelH, int kernelW, int strideH, int strideW, + int paddingH, int paddingW, + const char* fname_weights, bool batchnorm = false); virtual ~Conv2d(); virtual value_type* infer(dataDim_t &dim, value_type* srcData); protected: - value_type *dstData; //where results will be putted int kernelH, kernelW, strideH, strideW; cudnnFilterDescriptor_t filterDesc; @@ -149,9 +161,6 @@ public: virtual ~Flatten(); virtual value_type* infer(dataDim_t &dim, value_type* srcData); - -protected: - value_type *dstData; //where results will be putted }; @@ -169,7 +178,7 @@ public: protected: value_type mul, add; - value_type *dstData, *add_vector; //where results will be putted + value_type *add_vector; }; @@ -203,7 +212,7 @@ protected: int winH, winW; int strideH, strideW; tkdnnPoolingMode_t pool_mode; - value_type *dstData, *tmpInputData, *tmpOutputData; //where results will be putted + value_type *tmpInputData, *tmpOutputData; bool poolOn3d; }; @@ -216,11 +225,64 @@ public: Softmax(Network *net, dataDim_t input_dim); virtual ~Softmax(); + virtual value_type* infer(dataDim_t &dim, value_type* srcData); +}; + +/** + Route layer + Merge a list of layers +*/ +class Route : public Layer { + +public: + Route(Network *net, int *layers_id, int layers_n); + virtual ~Route(); + + virtual value_type* infer(dataDim_t &dim, value_type* srcData); + +public: + Layer **layers; //ids of layers to be merged + int layers_n; //number of layers +}; + + +/** + Reorg layer + Mantain same dimension but change C*H*W distribution +*/ +class Reorg : public Layer { + +public: + Reorg(Network *net, dataDim_t input_dim, int stride); + virtual ~Reorg(); + virtual value_type* infer(dataDim_t &dim, value_type* srcData); protected: - value_type *dstData; //where results will be putted + int stride; }; +/** + Region layer + Mantain same dimension but change C*H*W distribution +*/ +class Region : public Layer { + +public: + Region(Network *net, dataDim_t input_dim, + int classes, int coords, int num, float thresh); + virtual ~Region(); + + virtual value_type* infer(dataDim_t &dim, value_type* srcData); + +protected: + int classes, coords, num; + float thresh; + + int entry_index(int batch, int location, int entry); +}; + + + } #endif //LAYER_H diff --git a/include/Network.h b/include/Network.h index fb898a6..76cb6e1 100644 --- a/include/Network.h +++ b/include/Network.h @@ -27,7 +27,6 @@ public: cudnnHandle_t cudnnHandle; cublasHandle_t cublasHandle; -private: Layer* layers[MAX_LAYERS]; //contains layers of the net int num_layers; //current number of layers }; diff --git a/include/kernels.h b/include/kernels.h index 4085ef7..3901099 100644 --- a/include/kernels.h +++ b/include/kernels.h @@ -1,3 +1,10 @@ #include "utils.h" +#include "Layer.h" -void activationELUForward(value_type* srcData, value_type* dstData, int size); \ No newline at end of file +void activationELUForward(value_type* srcData, value_type* dstData, int size); +void activationLEAKYForward(value_type* srcData, value_type* dstData, int size); +void activationLOGISTICForward(value_type* srcData, value_type* dstData, int size); + +void reorgForward(value_type* srcData, value_type* dstData, tkDNN::dataDim_t dim, int stride); +void softmaxForward(float *input, int n, int batch, int batch_offset, + int groups, int group_offset, int stride, float temp, float *output); diff --git a/include/utils.h b/include/utils.h index ed50073..c68d9e5 100644 --- a/include/utils.h +++ b/include/utils.h @@ -62,7 +62,8 @@ } \ } -void readBinaryFile(const char* fname, int size, value_type** data_h, value_type** data_d); +void readBinaryFile(const char* fname, int size, value_type** data_h, value_type** data_d, int seek = 0); +int checkResult(int size, value_type *data_d, value_type *correct_d); void printDeviceVector(int size, value_type* vec_d); void resize(int size, value_type **data); diff --git a/src/Activation.cpp b/src/Activation.cpp index 23007ae..f5580f3 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -5,52 +5,61 @@ namespace tkDNN { -Activation::Activation(Network *net, dataDim_t input_dim, cudnnActivationMode_t act_mode) : +Activation::Activation(Network *net, dataDim_t input_dim, int act_mode) : Layer(net, input_dim) { this->act_mode = act_mode; checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) ); - checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc, - net->tensorFormat, - net->dataType, - input_dim.n*input_dim.l, - input_dim.c, - input_dim.h, input_dim.w) ); - checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc, + if(int(act_mode) < 100) { + + checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc, net->tensorFormat, net->dataType, 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.l, + input_dim.c, + input_dim.h, input_dim.w) ); - checkCUDNN( cudnnCreateActivationDescriptor(&activDesc) ); - checkCUDNN( cudnnSetActivationDescriptor(activDesc, - act_mode, - CUDNN_PROPAGATE_NAN, - 0.0) ); + checkCUDNN( cudnnCreateActivationDescriptor(&activDesc) ); + checkCUDNN( cudnnSetActivationDescriptor(activDesc, + (cudnnActivationMode_t) act_mode, + CUDNN_PROPAGATE_NAN, + 0.0) ); + } } Activation::~Activation() { checkCuda( cudaFree(dstData) ); - checkCUDNN( cudnnDestroyActivationDescriptor(activDesc) ); + if(int(act_mode) < 100) + checkCUDNN( cudnnDestroyActivationDescriptor(activDesc) ); } value_type* Activation::infer(dataDim_t &dim, value_type* srcData) { - value_type alpha = value_type(1); - value_type beta = value_type(0); - checkCUDNN( cudnnActivationForward(net->cudnnHandle, - activDesc, - &alpha, - srcTensorDesc, - srcData, - &beta, - dstTensorDesc, - dstData) ); + if(act_mode == ACTIVATION_LEAKY) { + activationLEAKYForward(srcData, dstData, dim.tot()); + + } else { + 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; } diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index 3b0e816..0ce544a 100644 --- a/src/Conv2d.cpp +++ b/src/Conv2d.cpp @@ -6,10 +6,11 @@ namespace tkDNN { Conv2d::Conv2d( Network *net, dataDim_t in_dim, int out_ch, int kernelH, int kernelW, int strideH, int strideW, - const char* fname_weights, const char* fname_bias) : + int paddingH, int paddingW, + const char* fname_weights, bool batchnorm) : LayerWgs(net, in_dim, in_dim.c, out_ch, kernelH, kernelW, 1, - fname_weights, fname_bias) { + fname_weights, batchnorm) { this->kernelH = kernelH; this->kernelW = kernelW; @@ -33,7 +34,7 @@ Conv2d::Conv2d( Network *net, dataDim_t in_dim, int out_ch, kernelH, kernelW) ); checkCUDNN( cudnnSetConvolution2dDescriptor(convDesc, - 0,0, // padding + paddingH, paddingW, // padding strideH, strideW, // stride 1,1, // upscale CUDNN_CROSS_CORRELATION) ); @@ -100,13 +101,23 @@ value_type* Conv2d::infer(dataDim_t &dim, value_type* srcData) { data_d, convDesc, algo, workSpace, ws_sizeInBytes, &beta, dstTensorDesc, dstData) ); - // bias - alpha = value_type(1); - beta = value_type(1); - checkCUDNN( cudnnAddTensor(net->cudnnHandle, - &alpha, biasTensorDesc, bias_d, - &beta, dstTensorDesc, dstData) ); - + if(!batchnorm) { + // bias + alpha = value_type(1); + beta = value_type(1); + checkCUDNN( cudnnAddTensor(net->cudnnHandle, + &alpha, biasTensorDesc, bias_d, + &beta, dstTensorDesc, dstData) ); + } else { + float one = 1; + float zero = 0; + cudnnBatchNormalizationForwardInference(net->cudnnHandle, + CUDNN_BATCHNORM_SPATIAL, &one, &zero, + dstTensorDesc, dstData, dstTensorDesc, + dstData, biasTensorDesc, //same tensor descriptor as bias + scales_d, bias_d, mean_d, variance_d, + CUDNN_BN_MIN_EPSILON); + } //update data dimensions dim = output_dim; diff --git a/src/Dense.cpp b/src/Dense.cpp index 1ad8965..4f8cc05 100644 --- a/src/Dense.cpp +++ b/src/Dense.cpp @@ -5,8 +5,8 @@ namespace tkDNN { Dense::Dense(Network *net, dataDim_t in_dim, - int out_ch, const char* fname_weights, const char* fname_bias) : - LayerWgs(net, in_dim, in_dim.tot(), out_ch, 1, 1, 1, fname_weights, fname_bias) { + int out_ch, const char* fname_weights) : + LayerWgs(net, in_dim, in_dim.tot(), out_ch, 1, 1, 1, fname_weights) { output_dim.n = 1; output_dim.c = out_ch; diff --git a/src/LayerWgs.cpp b/src/LayerWgs.cpp index b4c6c8b..84bb39b 100644 --- a/src/LayerWgs.cpp +++ b/src/LayerWgs.cpp @@ -6,16 +6,27 @@ namespace tkDNN { LayerWgs::LayerWgs(Network *net, dataDim_t in_dim, int inputs, int outputs, int kh, int kw, int kl, - const char* fname_weights, const char* fname_bias) : Layer(net, in_dim) { + const char* fname_weights, bool batchnorm) : Layer(net, in_dim) { this->inputs = inputs; this->outputs = outputs; this->weights_path = std::string(fname_weights); - this->bias_path = std::string(fname_bias); - + std::cout<<"Reading weights: I="<batchnorm = batchnorm; + if(batchnorm) { + seek += outputs*4; + readBinaryFile(weights_path.c_str(), outputs, &scales_h, &scales_d, seek); + seek += outputs*4; + readBinaryFile(weights_path.c_str(), outputs, &mean_h, &mean_d, seek); + seek += outputs*4; + readBinaryFile(weights_path.c_str(), outputs, &variance_h, &variance_d, seek); + } } LayerWgs::~LayerWgs() { @@ -24,6 +35,15 @@ LayerWgs::~LayerWgs() { delete [] bias_h; checkCuda( cudaFree(data_d) ); checkCuda( cudaFree(bias_d) ); + + if(batchnorm) { + delete [] scales_h; + delete [] mean_h; + delete [] variance_h; + checkCuda( cudaFree(scales_d) ); + checkCuda( cudaFree(mean_d) ); + checkCuda( cudaFree(variance_d) ); + } } -} \ No newline at end of file +} diff --git a/src/Network.cpp b/src/Network.cpp index 24aa84b..6e705e6 100644 --- a/src/Network.cpp +++ b/src/Network.cpp @@ -32,7 +32,7 @@ value_type* Network::infer(dataDim_t &dim, value_type* data) { //do infer for every layer for(int i=0; iinfer(dim, data); - + return data; } diff --git a/src/Region.cpp b/src/Region.cpp new file mode 100644 index 0000000..0a032cf --- /dev/null +++ b/src/Region.cpp @@ -0,0 +1,60 @@ +#include + +#include "Layer.h" +#include "kernels.h" + +namespace tkDNN { + +Region::Region(Network *net, dataDim_t input_dim, + int classes, int coords, int num, float thresh) : + Layer(net, input_dim) { + + this->classes = classes; + this->coords = coords; + this->num = num; + this->thresh = thresh; + + // same + output_dim.n = input_dim.n; + output_dim.c = input_dim.c; + output_dim.h = input_dim.h; + output_dim.w = input_dim.w; + output_dim.l = input_dim.l; + + checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) ); +} + +Region::~Region() { + checkCuda( cudaFree(dstData) ); +} + +int Region::entry_index(int batch, int location, int entry) { + int n = location / (input_dim.w*input_dim.h); + int loc = location % (input_dim.w*input_dim.h); + return batch*output_dim.tot() + n*input_dim.w*input_dim.h*(coords+classes+1) + entry*input_dim.w*input_dim.h + loc; +} + +value_type* Region::infer(dataDim_t &dim, value_type* srcData) { + + checkCuda( cudaMemcpy(dstData, srcData, dim.tot()*sizeof(value_type), cudaMemcpyDeviceToDevice)); + + for (int b = 0; b < dim.n; ++b){ + for(int n = 0; n < num; ++n){ + int index = entry_index(b, n*dim.w*dim.h, 0); + activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h); + + index = entry_index(b, n*dim.w*dim.h, coords); + activationLOGISTICForward(srcData + index, dstData + index, dim.w*dim.h); + } + } + + //softmax start + int index = entry_index(0, 0, coords + 1); + softmaxForward(srcData + index, classes, output_dim.n*num, output_dim.tot()/num, + output_dim.w*output_dim.h, 1, output_dim.w*output_dim.h, 1, dstData + index); + + dim = output_dim; + return dstData; +} + +} diff --git a/src/Reorg.cpp b/src/Reorg.cpp new file mode 100644 index 0000000..543457e --- /dev/null +++ b/src/Reorg.cpp @@ -0,0 +1,35 @@ +#include + +#include "Layer.h" +#include "kernels.h" + +namespace tkDNN { + +Reorg::Reorg(Network *net, dataDim_t input_dim, int stride) : + Layer(net, input_dim) { + + this->stride = stride; + + output_dim.n = input_dim.n; + output_dim.c = input_dim.c*stride*stride; + output_dim.h = input_dim.h/stride; + output_dim.w = input_dim.w/stride; + output_dim.l = input_dim.l; + + checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) ); +} + +Reorg::~Reorg() { + + checkCuda( cudaFree(dstData) ); +} + +value_type* Reorg::infer(dataDim_t &dim, value_type* srcData) { + + reorgForward(srcData, dstData, dim, stride); + + dim = output_dim; + return dstData; +} + +} diff --git a/src/Route.cpp b/src/Route.cpp new file mode 100644 index 0000000..643cdea --- /dev/null +++ b/src/Route.cpp @@ -0,0 +1,62 @@ +#include + +#include "Layer.h" +#include "kernels.h" + +namespace tkDNN { + +Route::Route(Network *net, int *layers_id, int layers_n) : + Layer(net, dataDim_t()) { + + this->layers_n = layers_n; + + //get layers + layers = new Layer*[layers_n]; + for(int i=0; ilayers[layers_id[i]]; + + + //get dims + output_dim.l = 1; + output_dim.c = 0; + for(int i=0; ioutput_dim.w; + output_dim.h = layers[i]->output_dim.h; + } else { + if( layers[i]->output_dim.w != output_dim.w || + layers[i]->output_dim.h != output_dim.h ) + FatalError("Route Output dim missmatch"); + } + output_dim.c += layers[i]->output_dim.c; + } + + input_dim = output_dim; + + checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) ); +} + +Route::~Route() { + + checkCuda( cudaFree(dstData) ); +} + +value_type* Route::infer(dataDim_t &dim, value_type* srcData) { + + + int offset = 0; + for(int i=0; idstData; + int in_dim = layers[i]->input_dim.tot(); + checkCuda( cudaMemcpy(dstData + offset, input, in_dim*sizeof(value_type), cudaMemcpyDeviceToDevice)); + offset += in_dim; + } + + //update data dimensions + dim = output_dim; + + return dstData; +} + +} \ No newline at end of file diff --git a/src/kernels/activation_elu.cu b/src/kernels/activation_elu.cu index a80e3ae..073d5cf 100644 --- a/src/kernels/activation_elu.cu +++ b/src/kernels/activation_elu.cu @@ -35,4 +35,4 @@ void activationELUForward(value_type* srcData, value_type* dstData, int size) activation_elu<<>>(srcData, dstData, size); checkCuda( cudaDeviceSynchronize() ); -} \ No newline at end of file +} diff --git a/src/kernels/activation_leaky.cu b/src/kernels/activation_leaky.cu new file mode 100644 index 0000000..d22ca0c --- /dev/null +++ b/src/kernels/activation_leaky.cu @@ -0,0 +1,29 @@ +#include "kernels.h" + +__global__ +void activation_leaky(value_type *input, value_type *output, int size) { + + int i = blockDim.x*blockIdx.x + threadIdx.x; + + if(i0) + output[i] = input[i]; + else + output[i] = 0.1f*input[i]; + } + } + + +/** + ELU activation function +*/ +void activationLEAKYForward(value_type* srcData, value_type* dstData, int size) +{ + int blocks = (size+255)/256; + int threads = 256; + + activation_leaky<<>>(srcData, dstData, size); + checkCuda( cudaDeviceSynchronize() ); +} + + diff --git a/src/kernels/activation_logistic.cu b/src/kernels/activation_logistic.cu new file mode 100644 index 0000000..23ace0c --- /dev/null +++ b/src/kernels/activation_logistic.cu @@ -0,0 +1,26 @@ +#include "kernels.h" + +__global__ +void activation_logistic(value_type *input, value_type *output, int size) { + + int i = blockDim.x*blockIdx.x + threadIdx.x; + + if(i>>(srcData, dstData, size); + checkCuda( cudaDeviceSynchronize() ); +} + + diff --git a/src/kernels/reorg.cu b/src/kernels/reorg.cu new file mode 100644 index 0000000..37abc24 --- /dev/null +++ b/src/kernels/reorg.cu @@ -0,0 +1,49 @@ +#include "kernels.h" + +__global__ void reorg_kernel(int N, float *x, int w, int h, int c, int batch, int stride, int forward, float *out) +{ + int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; + if(i >= N) return; + int in_index = i; + int in_w = i%w; + i = i/w; + int in_h = i%h; + i = i/h; + int in_c = i%c; + i = i/c; + int b = i%batch; + + int out_c = c/(stride*stride); + + int c2 = in_c % out_c; + int offset = in_c / out_c; + int w2 = in_w*stride + offset % stride; + int h2 = in_h*stride + offset / stride; + //printf("%d\n", offset); + int out_index = w2 + w*stride*(h2 + h*stride*(c2 + out_c*b)); + + // printf("%d %d %d\n", w2, h2, c2); + //printf("%d %d\n", in_index, out_index); + //if(out_index >= N || out_index < 0) printf("bad bad bad \n"); + + if(forward) out[out_index] = x[in_index]; + else out[in_index] = x[out_index]; + //if(forward) out[1] = x[1]; + //else out[0] = x[0]; +} + +/** + reorg function function +*/ +void reorgForward(value_type* srcData, value_type* dstData, tkDNN::dataDim_t dim, int stride) +{ + int size = dim.tot(); + + int blocks = (size+255)/256; + int threads = 256; + + reorg_kernel<<>>(size, srcData, dim.w, dim.h, dim.c, dim.n, stride, false, dstData); + checkCuda( cudaDeviceSynchronize() ); +} + + diff --git a/src/kernels/softmax.cu b/src/kernels/softmax.cu new file mode 100644 index 0000000..5ba112b --- /dev/null +++ b/src/kernels/softmax.cu @@ -0,0 +1,43 @@ +#include "kernels.h" + +__device__ void softmax_device(float *input, int n, float temp, int stride, float *output) +{ + int i; + float sum = 0; + float largest = -INFINITY; + for(i = 0; i < n; ++i){ + int val = input[i*stride]; + largest = (val>largest) ? val : largest; + } + for(i = 0; i < n; ++i){ + float e = exp(input[i*stride]/temp - largest/temp); + sum += e; + output[i*stride] = e; + } + for(i = 0; i < n; ++i){ + output[i*stride] /= sum; + } +} + +__global__ void softmax_kernel(float *input, int n, int batch, int batch_offset, int groups, int group_offset, int stride, float temp, float *output) +{ + int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; + if (id >= batch*groups) return; + int b = id / groups; + int g = id % groups; + softmax_device(input + b*batch_offset + g*group_offset, n, temp, stride, output + b*batch_offset + g*group_offset); +} + +/** + softmax function +*/ +void softmaxForward(float *input, int n, int batch, int batch_offset, + int groups, int group_offset, int stride, float temp, float *output) +{ + int size = groups*batch; + int blocks = (size+255)/256; + int threads = 256; + + softmax_kernel<<>>(input, n, batch, batch_offset, groups, group_offset, stride, temp, output); + checkCuda( cudaDeviceSynchronize() ); +} diff --git a/src/utils.cpp b/src/utils.cpp index 9414395..3806b5c 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -1,6 +1,6 @@ #include "utils.h" -void readBinaryFile(const char* fname, int size, value_type** data_h, value_type** data_d) +void readBinaryFile(const char* fname, int size, value_type** data_h, value_type** data_d, int seek) { std::ifstream dataFile (fname, std::ios::in | std::ios::binary); std::stringstream error_s; @@ -9,6 +9,11 @@ void readBinaryFile(const char* fname, int size, value_type** data_h, value_type error_s << "Error opening file " << fname; FatalError(error_s.str()); } + + if(seek != 0) { + dataFile.seekg(seek, dataFile.cur); + } + int size_b = size*sizeof(value_type); *data_h = new value_type[size]; if (!dataFile.read ((char*) *data_h, size_b)) @@ -37,6 +42,29 @@ void printDeviceVector(int size, value_type* vec_d) delete [] vec; } +int checkResult(int size, value_type *data_d, value_type *correct_d) { + + value_type *data_h, *correct_h; + data_h = new value_type[size]; + correct_h = new value_type[size]; + cudaDeviceSynchronize(); + cudaMemcpy(data_h, data_d, size*sizeof(value_type), cudaMemcpyDeviceToHost); + cudaMemcpy(correct_h, correct_d, size*sizeof(value_type), cudaMemcpyDeviceToHost); + + int diffs = 0; + for(int i=0; i 0.0001) { + diffs += 1; + printf("%d\n", i); + } + } + + delete [] data_h; + delete [] correct_h; + + return diffs; +} + void resize(int size, value_type **data) { if (*data != NULL) diff --git a/tests/mnist/test.cpp b/tests/mnist/test.cpp index 7462423..699b440 100644 --- a/tests/mnist/test.cpp +++ b/tests/mnist/test.cpp @@ -3,13 +3,9 @@ const char *input_bin = "../tests/mnist/input.bin"; const char *c0_bin = "../tests/mnist/layers/Convolution0.bin"; -const char *c0_bias_bin = "../tests/mnist/layers/Convolution0.bias.bin"; const char *c1_bin = "../tests/mnist/layers/Convolution1.bin"; -const char *c1_bias_bin = "../tests/mnist/layers/Convolution1.bias.bin"; const char *d2_bin = "../tests/mnist/layers/InnerProduct2.bin"; -const char *d2_bias_bin = "../tests/mnist/layers/InnerProduct2.bias.bin"; const char *d3_bin = "../tests/mnist/layers/InnerProduct3.bin"; -const char *d3_bias_bin = "../tests/mnist/layers/InnerProduct3.bias.bin"; const char *output_bin = "../tests/mnist/output.bin"; int main() { @@ -18,13 +14,13 @@ int main() { tkDNN::Network net; tkDNN::dataDim_t dim(1, 1, 28, 28, 1); tkDNN::Layer *l; - l = new tkDNN::Conv2d (&net, dim, 20, 5, 5, 1, 1, c0_bin, c0_bias_bin); + l = new tkDNN::Conv2d (&net, dim, 20, 5, 5, 1, 1, 1, 1, c0_bin); l = new tkDNN::Pooling (&net, l->output_dim, 2, 2, 2, 2, tkDNN::POOLING_MAX); - l = new tkDNN::Conv2d (&net, l->output_dim, 50, 5, 5, 1, 1, c1_bin, c1_bias_bin); + l = new tkDNN::Conv2d (&net, l->output_dim, 50, 5, 5, 1, 1, 1, 1, c1_bin); l = new tkDNN::Pooling (&net, l->output_dim, 2, 2, 2, 2, tkDNN::POOLING_MAX); - l = new tkDNN::Dense (&net, l->output_dim, 500, d2_bin, d2_bias_bin); + l = new tkDNN::Dense (&net, l->output_dim, 500, d2_bin); l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU); - l = new tkDNN::Dense (&net, l->output_dim, 10, d3_bin, d3_bias_bin); + l = new tkDNN::Dense (&net, l->output_dim, 10, d3_bin); l = new tkDNN::Softmax (&net, l->output_dim); // Load input diff --git a/tests/test/test.cpp b/tests/test/test.cpp index df4ca47..da0f1b3 100644 --- a/tests/test/test.cpp +++ b/tests/test/test.cpp @@ -3,11 +3,8 @@ const char *input_bin = "../tests/test/input.bin"; const char *c0_bin = "../tests/test/layers/conv0.bin"; -const char *c0_bias_bin = "../tests/test/layers/conv0.bias.bin"; const char *c1_bin = "../tests/test/layers/conv1.bin"; -const char *c1_bias_bin = "../tests/test/layers/conv1.bias.bin"; const char *d2_bin = "../tests/test/layers/dense2.bin"; -const char *d2_bias_bin = "../tests/test/layers/dense2.bias.bin"; const char *output_bin = "../tests/test/output.bin"; int main() { @@ -16,12 +13,12 @@ int main() { tkDNN::Network net; tkDNN::dataDim_t dim(1, 1, 10, 10, 1); tkDNN::Layer *l; - l = new tkDNN::Conv2d (&net, dim, 2, 4, 4, 2, 2, c0_bin, c0_bias_bin); + l = new tkDNN::Conv2d (&net, dim, 2, 4, 4, 2, 2, 1, 1, c0_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::Conv2d (&net, l->output_dim, 4, 2, 2, 1, 1, 1, 1, c1_bin); l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU); - l = new tkDNN::Flatten (&net, l->output_dim); - l = new tkDNN::Dense (&net, l->output_dim, 4, d2_bin, d2_bias_bin); + l = new tkDNN::Flatten (&net, l->output_dim); + l = new tkDNN::Dense (&net, l->output_dim, 4, d2_bin); l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU); // Load input