yolo layers

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)

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

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
 };

include/kernels.h

@@ -1,3 +1,10 @@
 #include "utils.h"
+#include "Layer.h"
 
-void activationELUForward(value_type* srcData, value_type* dstData, int size);
+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);

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);
 

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;
 }
 

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;
 

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;

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="<<inputs<<" O="<<outputs<<" KERNEL="<<kh<<"x"<<kw<<"x"<<kl<<"\n";
-    readBinaryFile(weights_path.c_str(), inputs*outputs*kh*kw*kl, &data_h, &data_d);
-    readBinaryFile(bias_path.c_str(), outputs, &bias_h, &bias_d);
+    int seek = 0;
+    readBinaryFile(weights_path.c_str(), inputs*outputs*kh*kw*kl, &data_h, &data_d, seek);
+    seek += inputs*outputs*kh*kw*kl*4;
+    readBinaryFile(weights_path.c_str(), outputs, &bias_h, &bias_d, seek);
+    
+    this->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) );
+    }
 }
 
-}
+}

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; i<num_layers; i++)
         data = layers[i]->infer(dim, data);
-
+        
     return data;
 }
 

src/Region.cpp

@@ -0,0 +1,60 @@
+#include <iostream>
+
+#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;
+}
+
+}

src/Reorg.cpp

@@ -0,0 +1,35 @@
+#include <iostream>
+
+#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;
+}
+
+}

src/Route.cpp

@@ -0,0 +1,62 @@
+#include <iostream>
+
+#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; i<layers_n; i++) 
+        layers[i] = net->layers[layers_id[i]]; 
+        
+
+    //get dims
+    output_dim.l = 1;
+    output_dim.c = 0;
+    for(int i=0; i<layers_n; i++) {
+
+        if(i==0) {
+            output_dim.w = layers[i]->output_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; i<layers_n; i++) {
+        value_type *input = layers[i]->dstData;
+        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;
+}
+
+}

src/kernels/activation_elu.cu

@@ -35,4 +35,4 @@ void activationELUForward(value_type* srcData, value_type* dstData, int size)
     
     activation_elu<<<blocks, threads>>>(srcData, dstData, size);
     checkCuda( cudaDeviceSynchronize() );
-}
+}

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(i<size) {    
+        if (input[i]>0)
+            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<<<blocks, threads>>>(srcData, dstData, size);
+    checkCuda( cudaDeviceSynchronize() );
+}
+
+

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<size) {    
+        output[i] =  1.0f/(1.0f + exp(-input[i]));;
+    }
+ }
+
+
+/**
+    LOGISTIC activation function
+*/
+void activationLOGISTICForward(value_type* srcData, value_type* dstData, int size)
+{
+    int blocks = (size+255)/256;
+    int threads = 256;
+    
+    activation_logistic<<<blocks, threads>>>(srcData, dstData, size);
+    checkCuda( cudaDeviceSynchronize() );
+}
+
+

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<<<blocks, threads>>>(size, srcData, dim.w, dim.h, dim.c, dim.n, stride, false, dstData);
+    checkCuda( cudaDeviceSynchronize() );
+}
+
+

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<<<blocks, threads>>>(input, n, batch, batch_offset, groups, group_offset, stride, temp, output);
+    checkCuda( cudaDeviceSynchronize() );
+}

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<size; i++) {
+        if(fabs(data_h[i] - correct_h[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)

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

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