   if(Session::GetSession()->gpu) {
     dim3 blocksInGrid(batch_size/BLOCKDIM+1, in_channels/BLOCKDIM+1);
     dim3 threadsPerBlock(BLOCKDIM, BLOCKDIM);
-    FCKernels::ComputeBottomDiffs<Dtype><<<blocksInGrid, threadsPerBlock>>>(top_diff, bottom_diff, W_);
+    FCGPUKernels::ComputeBottomDiffs<Dtype><<<blocksInGrid, threadsPerBlock>>>(top_diff, bottom_diff, W_);
     blocksInGrid = dim3(in_channels/BLOCKDIM+1, out_channels/BLOCKDIM+1);
-    FCKernels::ComputeWDiffs<Dtype><<<blocksInGrid, threadsPerBlock>>>(bottom, top_diff, W_diff_);
-    FCKernels::ComputeBDiffs<Dtype><<<out_channels/BLOCKDIM/BLOCKDIM, BLOCKDIM*BLOCKDIM>>>(top_diff, b_diff_);
+    FCGPUKernels::ComputeWDiffs<Dtype><<<blocksInGrid, threadsPerBlock>>>(bottom, top_diff, W_diff_);
+    FCGPUKernels::ComputeBDiffs<Dtype><<<out_channels/BLOCKDIM/BLOCKDIM, BLOCKDIM*BLOCKDIM>>>(top_diff, b_diff_);
   } else { 
     // compute bottom diffs
     for(int i = 0; i < in_channels; i++) {

#include "initializers/const_initializer.cu"
#include "basics/layer.hpp"
#include "utils/utils.cu"
// #include <stdio.h>

#define BLOCKDIM 32

namespace FCGPUKernels {
  template <class Dtype> 

    sum += b_->at(0,0,0,o);
    top->at(bi,0,0,o) = sum;
  }


  template<class Dtype>
  __global__ void ComputeBottomDiffs(Tensor<Dtype>* top_diff, Tensor<Dtype> * bottom_diff, Tensor<Dtype> * W_) {
    int b = (blockDim.x * blockIdx.x) + threadIdx.x;
    int i = (blockDim.y * blockIdx.y) + threadIdx.y;
    
    size_t batch_size = bottom_diff->GetDims()[0];
    size_t in_channels = bottom_diff->GetDims()[3];
    size_t out_channels = top_diff->GetDims()[3];

    if(b < 0 || b >= batch_size || i < 0 || i >= in_channels) return;
    
    Dtype sum = 0;
    for(int o = 0; o < out_channels; o++) {
      sum += W_->at(0,0,o,i) * top_diff->at(b,0,0,o);
    }
    bottom_diff->at(b,0,0,i) = sum;
  }

  template<class Dtype>
  __global__ void ComputeWDiffs(Tensor<Dtype>* bottom, Tensor<Dtype> * top_diff, Tensor<Dtype> * W_diff_) {
    int i = (blockDim.x * blockIdx.x) + threadIdx.x;
    int o = (blockDim.y * blockIdx.y) + threadIdx.y;
    
    size_t batch_size = bottom->GetDims()[0];
    size_t in_channels = bottom->GetDims()[3];
    size_t out_channels = top_diff->GetDims()[3];

    if(i < 0 || i >= in_channels || o < 0 || o >= out_channels) return;

    Dtype sum_w = 0;
    for(int b = 0; b < batch_size; b++) {
      sum_w += bottom->at(b, 0, 0, i)*top_diff->at(b, 0, 0, o);
    }
    W_diff_->at(0,0,o,i) = sum_w;
  }

  template<class Dtype>
  __global__ void ComputeBDiffs(Tensor<Dtype> * top_diff, Tensor<Dtype>* b_diff_) {
    int o = (blockDim.x * blockIdx.x) + threadIdx.x;
    
    size_t batch_size = top_diff->GetDims()[0];
    size_t out_ch = top_diff->GetDims()[3];

    if(o < 0 || o >= out_ch) return;

    Dtype sum_b = 0;
    for(int b = 0; b < batch_size; b++) {
      sum_b += top_diff->at(b, 0, 0, o);
    }
    b_diff_->at(0,0,0,o) = sum_b;
  }
}



template<class Dtype>
void FC<Dtype>::Forward(const std::vector<Tensor<Dtype>*> &bottoms, const std::vector<Tensor<Dtype>*> &tops) {
  assert(bottoms.size() == 1);
  assert(tops.size() == 1);
  Tensor<Dtype> * bottom = bottoms[0];
  Tensor<Dtype> * top = tops[0];

  if (Session::GetSession()->gpu) {
    size_t batch_size = Session::GetSession()->batch_size;
    dim3 blocksInGrid(batch_size / BLOCKDIM + 1, out_channels / BLOCKDIM + 1);
    dim3 threadsPerBlock(BLOCKDIM, BLOCKDIM);

    if (in_channels*2 < 384) {
      FCGPUKernels::ForwardGPUShared<<<blocksInGrid,threadsPerBlock, 2*in_channels*BLOCKDIM*sizeof(Dtype)>>>(bottom, top, W_, b_);
    } else {
      FCGPUKernels::ForwardGPU<<<blocksInGrid,threadsPerBlock>>>(bottom, top, W_, b_);
    }
  } else {
    for(int b = 0; b < bottom->GetDims()[0]; b++) {
      for(int o = 0; o < out_channels; o++) {
        Dtype sum = 0;
        for(int i = 0; i < in_channels; i++) {
          sum += bottom->at(b, 0, 0, i) * W_->at(0,0,o,i);
        }
        sum += b_->at(0,0,0,o);
        top->at(b,0,0,o) = sum;
      }
    }
  }
}


template<class Dtype>
void FC<Dtype>::GetTopsDims(const std::vector<size_t*> &bottoms_dims, 
                   const std::vector<size_t*> &tops_dims) {
  assert(bottoms_dims.size());
  assert(tops_dims.size());
  size_t * b_dims = bottoms_dims[0];
  assert(b_dims[1] == 1);
  assert(b_dims[2] == 1);
  size_t * t_dims = tops_dims[0];
  t_dims[0] = b_dims[0];
  t_dims[1] = 1;
  t_dims[2] = 1;
  t_dims[3] = out_channels;
}


template <class Dtype>
void FC<Dtype>::InitParams() {
  if (initializer_!=NULL) {
    initializer_->Initialize(W_, b_, Session::GetSession()->gpu);
  } else {
    ConstInitializer<Dtype>(1.0, 1.0).Initialize(W_, b_, Session::GetSession()->gpu);
  }
}

template <class Dtype>
FC<Dtype>::FC(size_t in_channels, size_t out_channels, Initializer<Dtype>* initializer = NULL):
    in_channels(in_channels), out_channels(out_channels), initializer_(initializer), W_diff_(NULL), b_diff_(NULL) {
  size_t w_dims[4] = {1, 1, out_channels, in_channels};
  size_t b_dims[4] = {1, 1, 1, out_channels};
  if (Session::GetSession()->gpu) {
    W_ = Tensor<Dtype>::CreateTensorGPU(w_dims);
    b_ = Tensor<Dtype>::CreateTensorGPU(b_dims);
  } else {
    W_ = Tensor<Dtype>::CreateTensorCPU(w_dims);
    b_ = Tensor<Dtype>::CreateTensorCPU(b_dims);
  }
  InitParams();
}


FC<Dtype>::~FC() {
  if (Session::GetSession()->gpu) {
    if (W_!= NULL) {
      cudaFree(W_);
      W_ = NULL;
    }
    if (b_ != NULL) {
      cudaFree(b_);
      b_ = NULL;
    }
    if (W_diff_ != NULL) {
      cudaFree(W_diff_);
      W_diff_ = NULL;
    }
    if (b_diff_ != NULL) {
      cudaFree(b_diff_);
      b_diff_ = NULL;
    }
  } else {
    if(W_ != NULL) {
      delete W_;
      W_ = NULL;
    }
    if(b_ != NULL) {
      delete b_;
      b_ = NULL;
    }
    if(W_diff_ != NULL) {
      delete W_diff_;
      W_diff_ = NULL;
    }
    if(b_diff_ != NULL) {
      delete b_diff_; 
      b_diff_ = NULL;
    }
  }
}


template<class Dtype>
