You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@singa.apache.org by ka...@apache.org on 2016/06/27 14:11:48 UTC

[1/6] incubator-singa git commit: SINGA-204 Support the training of feed-forward neural nets

Repository: incubator-singa
Updated Branches:
  refs/heads/dev dd08f4130 -> 4db968c2e


SINGA-204 Support the training of feed-forward neural nets

Draft FeedForwardNet including functions for
1. training, training for one mini-batch
2. evaluation, evaluation for one mini-batch

Draf the Alexnet model for Cifar10 dataset,
1. Cifar10 class for reading the dataset
2. alexnet.cc for creating the FeedForwardNet and conducting the training.


Project: http://git-wip-us.apache.org/repos/asf/incubator-singa/repo
Commit: http://git-wip-us.apache.org/repos/asf/incubator-singa/commit/d826b2e8
Tree: http://git-wip-us.apache.org/repos/asf/incubator-singa/tree/d826b2e8
Diff: http://git-wip-us.apache.org/repos/asf/incubator-singa/diff/d826b2e8

Branch: refs/heads/dev
Commit: d826b2e856b778b02b74be1659e990820538cffb
Parents: dd08f41
Author: Wei Wang <wa...@comp.nus.edu.sg>
Authored: Wed Jun 22 18:05:46 2016 +0800
Committer: Wei Wang <wa...@comp.nus.edu.sg>
Committed: Sat Jun 25 16:15:59 2016 +0800

----------------------------------------------------------------------
 examples/cifar10/alexnet.cc            | 144 +++++++++++++++
 examples/cifar10/cifar10.cc            |  98 +++++++++++
 examples/cifar10/download_data.py      |  30 ++++
 include/singa/core/tensor.h            |   9 +-
 include/singa/model/feed_forward_net.h | 133 ++++++++++++++
 include/singa/model/layer.h            |   3 +
 include/singa/utils/singleton.h        |  22 ++-
 include/singa/utils/string.h           |   9 +
 src/core/tensor/tensor.cc              |  18 +-
 src/model/feed_forward_net.cc          | 263 ++++++++++++++++++++++++++++
 10 files changed, 708 insertions(+), 21 deletions(-)
----------------------------------------------------------------------


http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/d826b2e8/examples/cifar10/alexnet.cc
----------------------------------------------------------------------
diff --git a/examples/cifar10/alexnet.cc b/examples/cifar10/alexnet.cc
new file mode 100644
index 0000000..2917dd2
--- /dev/null
+++ b/examples/cifar10/alexnet.cc
@@ -0,0 +1,144 @@
+/************************************************************
+*
+* Licensed to the Apache Software Foundation (ASF) under one
+* or more contributor license agreements.  See the NOTICE file
+* distributed with this work for additional information
+* regarding copyright ownership.  The ASF licenses this file
+* to you under the Apache License, Version 2.0 (the
+* "License"); you may not use this file except in compliance
+* with the License.  You may obtain a copy of the License at
+*
+*   http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing,
+* software distributed under the License is distributed on an
+* "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+* KIND, either express or implied.  See the License for the
+* specific language governing permissions and limitations
+* under the License.
+*
+*************************************************************/
+#include "./cifar10.h"
+#include "singa/model/feed_forward_net.h"
+#include "singa/model/optimizer.h"
+#include "singa/model/initializer.h"
+
+namespace singa {
+
+LayerConf GenConvConf(string name, int nb_filter, int kernel, int stride,
+                      int pad) {
+  LayerConf conf;
+  conf.set_name(name);
+  conf.set_type("CudnnConvolution");
+  ConvolutionConf *conv = conf.mutable_convolution_conf();
+  conv->set_num_output(nb_filter);
+  conv->set_kernel_size(kernel);
+  conv->set_stride(stride);
+  conv->set_pad(pad);
+
+  FillerConf *weight = conv->mutable_weight_filler();
+  weight->set_type("Xavier");
+  return conf;
+}
+
+LayerConf GenPoolingConf(string name, bool max_pool, int kernel, int stride, int pad) {
+  LayerConf conf;
+  conf.set_name(name);
+  conf.set_type("CudnnPooling");
+  PoolingConf *pool = conf.mutable_pooling_conf();
+  pool->set_kernel_size(kernel);
+  pool->set_stride(stride);
+  pool->set_pad(pad);
+  if (!max_pool)
+    pool->set_pool(PoolingConf_AVE);
+  return conf;
+}
+
+LayerConf GenReLUConf(string name) {
+  LayerConf conf;
+  conf.set_name(name);
+  conf.set_type("RELU");
+  return conf;
+}
+
+LayerConf GenDenseConf(string name, int num_output) {
+  LayerConf conf;
+  conf.set_name(name);
+  conf.set_type("Dense");
+  DenseConf *dense = conf->mutable_dense_conf();
+  dense->set_num_output(num_output);
+  FillerConf *weight = conv->mutable_weight_filler();
+  weight->set_type("Xavier");
+  return conf;
+}
+
+LayerConf GenSoftmaxConf(string name) {
+  LayerConf conf;
+  conf.set_name(name);
+  conf.set_type("CudnnSoftmax");
+  return conf;
+}
+
+
+FeedForwordNet CreateNet(Optimizer* opt, Loss* loss, Metric* metric) {
+  FeedForwordNet net;
+  Shape s{3, 32, 32};
+  net.AddLayer(GenConvConf("conv1", 32, 5, 1, 2), &s);
+  net.AddLayer(GenReLUConf("relu1"));
+  net.AddLayer(GenConvConf("pool1", 3, 2, 0));
+  net.AddLayer(GenConvConf("conv2", 32, 5, 1, 2));
+  net.AddLayer(GenReLUConf("relu2"));
+  net.AddLayer(GenConvConf("pool2", 3, 2, 0));
+  net.AddLayer(GenConvConf("conv3", 64, 5, 1, 2));
+  net.AddLayer(GenReLUConf("relu3"));
+  net.AddLayer(GenConvConf("pool3", 3, 2, 0));
+  net.AddLayer(GenDenseConf("ip1", 10));
+  net.AddLayer(GenSoftmaxConf("softmax"));
+
+  OptimizerConf opt_conf;
+  opt_conf.set_momentum(0.9);
+  opt->Setup(opt_conf);
+  net.Compile(true, opt, loss, metric);
+  return net;
+}
+
+void Train(float lr, int num_epoch, string data_dir) {
+  SoftmaxCrossEntropy loss;
+  Accuracy acc;
+  SGD sgd;
+  sgd.SetLearningRate([lr](int step) {return lr;});
+  auto net = CreateNet(&opt, &loss, &metric);
+  Cifar10 data(data_dir);
+  Tensor train_x, tain_y, test_x, test_y;
+  {
+    auto train = data.ReadTrainData();
+    const auto mean = Average(train.first, 0);
+    train_x = SubRow(train.first, mean);
+    auto test = data.ReadTestData();
+    test_x = SubRow(test.first, mean);
+    train_y = train.second;
+    test_y = test.second;
+  }
+  net.Train(100, num_epoch, train_x, train_y, test_x, test_y);
+}
+
+int main(int argc, char** argv) {
+  InitChannel();
+  int pos = ArgPos(argc, argv, "-epoch");
+  int nEpoch = 5;
+  if (pos != -1)
+    nEpoch = atoi(argv[pos + 1]);
+  pos = ArgPos(argc, argv, "-lr");
+  float lr = 0.01;
+  if (pos != -1)
+    lr = atof(argv[pos + 1]);
+  pos = ArgPos(argc, argv, "-data");
+  string data = "cifar-10-batch-bin";
+  if (pos != -1)
+    data = argv[pos + 1];
+
+  LOG(INFO) << "Start training";
+  Train(lr, nEpoch, data);
+  LOG(INFO) << "End training";
+}
+}

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/d826b2e8/examples/cifar10/cifar10.cc
----------------------------------------------------------------------
diff --git a/examples/cifar10/cifar10.cc b/examples/cifar10/cifar10.cc
new file mode 100644
index 0000000..7efc18f
--- /dev/null
+++ b/examples/cifar10/cifar10.cc
@@ -0,0 +1,98 @@
+/************************************************************
+*
+* Licensed to the Apache Software Foundation (ASF) under one
+* or more contributor license agreements.  See the NOTICE file
+* distributed with this work for additional information
+* regarding copyright ownership.  The ASF licenses this file
+* to you under the Apache License, Version 2.0 (the
+* "License"); you may not use this file except in compliance
+* with the License.  You may obtain a copy of the License at
+*
+*   http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing,
+* software distributed under the License is distributed on an
+* "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+* KIND, either express or implied.  See the License for the
+* specific language governing permissions and limitations
+* under the License.
+*
+*************************************************************/
+#include <fstream>
+#include <string>
+#include <cstdint>
+#include <iostream>
+
+using std::string;
+namespace singa {
+/// For reading cifar10 binary data as tensors.
+class Cifar10 {
+ public:
+  /// 'dir_path': path to the folder including the *.bin files
+  Cifar10(string dir_path, bool normalize = true)
+      : dir_path_(dir_path), normalize_(normalize) {}
+
+  /// read all training data into an image Tensor and a label Tensor
+  const std::pair<Tensor, Tensor> ReadTrainData(bool shuffle = false);
+  /// read all test data into an image Tensor and a label Tensor
+  const std::pair<Tensor, Tensor> ReadTestData();
+  /// read data from one file into an image Tensor and a label Tensor
+  const std::pair<Tensor, Tensor> ReadFile(string file, bool shuffle = false);
+
+ private:
+  const int kImageSize = 32;
+  const int kImageVol = 3072;
+  const int kBatchSize = 10000;
+  const int kTrainFiles = 5;
+
+  string dir_path_;
+  bool normalize_;
+};
+
+void read_image(std::ifstream* file, int* label, char* buffer) {
+  char label_char;
+  file->read(&label_char, 1);
+  *label = label_char;
+  file->read(buffer, kImageVol);
+  return;
+}
+const std::pair<Tensor, Tensor> Cifar10::ReadFile(string file,
+                                                  bool shuffle = false) {
+  Tensor images(Shape{kTrainFiles, 3, kImageSize, kImageSize});
+  Tensor labels(Shape{kTrainFiles}, kInt);
+  if (dir_path_.back() != '/') dir_path_.push_back('/');
+  LOG(INFO) << "Reading file " << dir_path_ + file;
+  std::ifstream data_file((dir_path_ + file).c_str(),
+                          std::ios::in | std::ios::binary);
+  CHECK(data_file.is_open()) << "Unable to open file " << file;
+  int label;
+  char image[kImageVol];
+  float float_image[kImageVol];
+  int tmplabels[kBatchSize];
+  for (int itemid = 0; itemid < kBatchSize; ++itemid) {
+    read_image(&data_file, &label, image);
+    for (int i = 0; i < kImageVol; i++)
+      float_image[i] = static_cast<float>(static_cast<int>(image[i]));
+    images.CopyDataFromHostPtr(float_image, kImageVol, itemid * kImageVol);
+    tmplabels[itemid] = label;
+  }
+  labels.CopyDataFromHostPtr(tmplabels, kBatchSize);
+  return std::make_pair(images, labels);
+}
+
+const std::pair<Tensor, Tensor> Cifar10::ReadTrainData(bool shuffle = false) {
+  Tensor images(Shape{kBatchSize * kTrainFiles, 3, kImageSize, kImageSize});
+  Tensor labels(Shape{kBatchSize * kTrainFiles, 3, kImageSize, kImageSize});
+  for (int fileid = 0; fileid < kTrainFiles; ++fileid) {
+    string file = "data_batch_" + std::to_string(fileid + 1) + ".bin";
+    const auto ret = ReadFile(file);
+    CopyDataToFrom(&images, ret.first, ret.first.Size(),
+                   fileid * ret.first.Size());
+    CopyDataToFrom(&labels, ret.second, kBatchSize, fileid * kBatchSize);
+  }
+  return std::make_pair(images, labels);
+}
+const std::pair<Tensor, Tensor> Cifar10::ReadTrainData() {
+  return ReadFile("test_batch.bin");
+}
+}  // namespace singa

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/d826b2e8/examples/cifar10/download_data.py
----------------------------------------------------------------------
diff --git a/examples/cifar10/download_data.py b/examples/cifar10/download_data.py
new file mode 100644
index 0000000..ffb9724
--- /dev/null
+++ b/examples/cifar10/download_data.py
@@ -0,0 +1,30 @@
+#!/usr/bin/env python
+import urllib
+import tarfile
+import os
+import sys
+
+
+def extract_tarfile(filepath):
+    with tarfile.open(filepath, 'r') as f:
+        f.extractall('.')
+
+
+dirpath = 'cifar-10-batches-bin'
+gzfile = 'cifar-10-binary' + '.tar.gz'
+if os.path.exists(dirpath):
+    print 'Directory %s does exist. To redownload the files, '\
+        'remove the existing directory and %s.tar.gz' % (dirpath, dirpath)
+    sys.exit(0)
+
+if os.path.exists(gzfile):
+    print 'The tar file does exist. Extracting it now..'
+    extract_tarfile(gzfile)
+    print 'Finished!'
+    sys.exit(0)
+
+url = 'http://www.cs.toronto.edu/~kriz/cifar-10-binary.tar.gz'
+print 'Downloading CIFAR10 from %s' % (url)
+urllib.urlretrieve(url, gzfile)
+extract_tarfile(gzfile)
+print 'Finished!'

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/d826b2e8/include/singa/core/tensor.h
----------------------------------------------------------------------
diff --git a/include/singa/core/tensor.h b/include/singa/core/tensor.h
index a4f42db..6de5c0c 100644
--- a/include/singa/core/tensor.h
+++ b/include/singa/core/tensor.h
@@ -21,6 +21,7 @@
 
 #include <vector>
 #include <tuple>
+#include <memory>
 
 #include "singa/core/common.h"
 #include "singa/core/device.h"
@@ -127,14 +128,16 @@ class Tensor {
 
   /// For init the tensor values, copy 'num' elements.
   template <typename SType>
-  void CopyDataFromHostPtr(const SType *src, const size_t num);
+  void CopyDataFromHostPtr(const SType *src, const size_t num,
+                           const size_t offset = 0);
 
   /// Copy data from another Tensor which may be on a diff device.
   /// Meta data would not be copied!
   void CopyData(const Tensor &other);
 
-  /// return an exactly the same Tensor with data been deep copied.
-  Tensor Clone() const;
+  /// return an exactly the same Tensor with data been deep copied to the given
+  /// device. If 'device' is nullptr, then clone it one the current device.
+  Tensor Clone(std::shared_ptr<Device> device = nullptr) const;
 
   // Tensor operations
 

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/d826b2e8/include/singa/model/feed_forward_net.h
----------------------------------------------------------------------
diff --git a/include/singa/model/feed_forward_net.h b/include/singa/model/feed_forward_net.h
new file mode 100644
index 0000000..173600b
--- /dev/null
+++ b/include/singa/model/feed_forward_net.h
@@ -0,0 +1,133 @@
+/**
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#ifndef SINGA_MODEL_FEED_FORWARD_NET_H_
+#define SINGA_MODEL_FEED_FORWARD_NET_H_
+#include "singa/model/layer.h"
+
+namespace singa {
+
+/// The feed-forward neural net.
+/// It provides functions for constructing the layers, access layer parameters,
+/// and conducting training, evaluation and prediction.
+class FeedForwardNet {
+ public:
+  FeedForwardNet() = explicit;
+  ~FeedForwardNet();
+
+  /// Add a layer with the assumption that
+  /// 1. this function is called in correct order, i.e., the layers are added
+  ///    following the topological order.
+  /// 2. this layer has already been setup (Setup function is called outside).
+  void Add(Layer *layer);
+
+  // TODO(wangwei) add ConcatenateLayer and SliceLayer
+  // AddConcatenateLayer(vector<Layer*> src, Layer *dst);
+  // AddSliceLayer(Layer* layer, vector<Layer*> dst);
+
+  /// Add a layer by providing its configuration, and setup it.
+  /// Assume the layer is added in corret order.
+  /// For the first layer, 'sample_shape' (the input sample shape) is necessary
+  /// for calling Setup().
+  void Add(const LayerConf &conf, const Shape *sample_shape = nullptr);
+
+  /// Set some fields used for training and evaluating the neural net.
+  /// If the neural net is constructed for evaluation only, then 'opt' is not
+  /// necessary; But for training, both 'opt' and 'loss' are necessary.
+  /// 'shuffle' indicates shuffling training samples within one epoch it is
+  /// valid using Train();
+  void Compile(bool shuffle, Optimizer *opt, Loss *loss, Metric *metric);
+
+  /// Conduct the training giving the training data 'x' and label 'y'.
+  /// Due to memory limit, 'x' and 'y' could not be very large. Hence, it is
+  /// typically used for small training datasets, e.g., cifar10 and MNIST which
+  /// can be stored in main memory.
+  void Train(int batchsize, int nb_epoch, Tensor x, Tensor y);
+  /// Conduct the training giving the training data 'x' and label 'y'.
+  /// 'val_split' is a ratio for splitting (1-'val_split') of training data for
+  /// validation. Validation is performance before every epoch.
+  /// Due to memory limit, 'x' and 'y' could not be very large. Hence, it is
+  /// typically used for small training datasets, e.g., cifar10 and MNIST which
+  /// can be stored in main memory.
+  void Train(int batchsize, int nb_epoch, float val_split, Tensor x, Tensor y);
+  /// Conduct the training given the training and validation data.
+  /// Validation is performance before every epoch.
+  /// Due to memory limit, 'x' and 'y' could not be very large. Hence, it is
+  /// typically used for small training datasets, e.g., cifar10 and MNIST which
+  /// can be stored in main memory.
+  void Train(int batchsize, int nb_epoch, Tensor x, Tensor y, Tensor val_x,
+             Tensor val_y);
+  /// Train the neural net over one batch of training data.
+  Tensor TrainOnBatch(Tensor x, Tensor y);
+
+  /// Evaluate the neural net with given data.
+  /// Returns one tensor for loss values and one tensor for metric values;
+  /// Each sample would have a loss value and a metric value (if 'metic' is set
+  /// in Compile()).'batchsize' is used for controlling the memory footprint.
+  /// It should be smaller than the total number of samples.
+  /// Due to memory limit, 'x' and 'y' could not be very large. Hence, it is
+  /// typically used for small training datasets, e.g., cifar10 and MNIST which
+  /// can be stored in main memory.
+  std::pair<Tensor, Tensor> Evaluate(Tensor x, Tensor y, int batchsize = 128);
+  /// Evaluate the neural net for one batch of data
+  std::pair<Tensor, Tensor> EvaluateOnBatch(Tensor x, Tensor y);
+
+  /// Predict the probability distributation over candicate classes for each
+  /// data sample. 'batchsize' is used for controlling the memory footprint.
+  /// It should be smaller than the total number of samples.
+  /// Due to memory limit, 'x' and 'y' could not be very large. Hence, it is
+  /// typically used for small training datasets, e.g., cifar10 and MNIST which
+  /// can be stored in main memory.
+  Tensor Predict(const Tensor &x, int batchsize = 128);
+  /// Predict for one batch data.
+  Tensor PredictOnBatch(const Tensor &x);
+
+  /// Forward layers one by one using the data batch 'x'.
+  /// Returns the prediction results (from the last layer).
+  Tensor Forward(const Tensor& x);
+  /// Backward layers one by one using the gradient batch 'grad'.
+  /// Returns the parameter gradients.
+  const vector<Tensor> Backward(const Tensor& grad);
+
+  /// Clone the neuaral net by cloning every layer to the given device.
+  /// If 'device' is nullptr, then clone it one the current device.
+  FeedForwardNet Clone(std::shared_ptr<Device> device = nullptr);
+  /// Move the layer data to the given device.
+  void ToDevice(Device *device);
+  /// Set the data type of each layer.
+  void AsType(DataType dtype);
+
+  const vector<Layer *> layers() const { return layers_; }
+  const vector<string> GetParamNames() const;
+  const vector<Tensor *> GetParamValues() const;
+  const vector<Tensor *> GetParamGrads() const;
+
+ protected:
+  vector<Layer *> layers_;
+  Optimizer *opt_;
+  Loss *loss_;
+  Metric *metric_;
+
+  bool shuffle_ = true;
+  Device* device_ = nullptr;
+  DataType dtype_ = kFloat32;
+};
+
+}  /* singa */
+
+
+#endif  // SINGA_MODEL_FEED_FORWARD_NET_H_

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/d826b2e8/include/singa/model/layer.h
----------------------------------------------------------------------
diff --git a/include/singa/model/layer.h b/include/singa/model/layer.h
index 0570b04..79eb069 100644
--- a/include/singa/model/layer.h
+++ b/include/singa/model/layer.h
@@ -149,6 +149,9 @@ class Layer {
     return std::make_pair(input_grad, param_grad);
   }
 
+  /// Clone the layer to the given device. Layer data (e.g., parameters) are
+  /// deep copied. If 'device' is nullptr, then clone it one the current device.
+  virtual Layer* Clone(std::shared_ptr<Device> device);
   /// Move the layer (including its parameters and other internal Tensor) onto
   /// the given device
   virtual void ToDevice(std::shared_ptr<Device> device) {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/d826b2e8/include/singa/utils/singleton.h
----------------------------------------------------------------------
diff --git a/include/singa/utils/singleton.h b/include/singa/utils/singleton.h
index 4cf487e..de831c4 100644
--- a/include/singa/utils/singleton.h
+++ b/include/singa/utils/singleton.h
@@ -7,9 +7,9 @@
 * to you under the Apache License, Version 2.0 (the
 * "License"); you may not use this file except in compliance
 * with the License.  You may obtain a copy of the License at
-* 
+*
 *   http://www.apache.org/licenses/LICENSE-2.0
-* 
+*
 * Unless required by applicable law or agreed to in writing,
 * software distributed under the License is distributed on an
 * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
@@ -22,10 +22,8 @@
 #ifndef SINGA_UTILS_SINGLETON_H_
 #define SINGA_UTILS_SINGLETON_H_
 
-/**
-  * Thread-safe implementation for C++11 according to
-  * http://stackoverflow.com/questions/2576022/efficient-thread-safe-singleton-in-c
-  */
+/// Thread-safe implementation for C++11 according to
+//  http://stackoverflow.com/questions/2576022/efficient-thread-safe-singleton-in-c
 template<typename T>
 class Singleton {
  public:
@@ -35,18 +33,18 @@ class Singleton {
   }
 };
 
-/**
- * Thread Specific Singleton
- *
- * Each thread will have its own data_ storage.
- */
+/// Thread Specific Singleton
+/// Each thread will have its own data_ storage.
+/*
 template<typename T>
 class TSingleton {
  public:
   static T* Instance() {
-    static thread_local T data_;
+    static thread_local T data_;  // thread_local is not available in some
+                                  // compilers
     return &data_;
   }
 };
+*/
 
 #endif  // SINGA_UTILS_SINGLETON_H_

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/d826b2e8/include/singa/utils/string.h
----------------------------------------------------------------------
diff --git a/include/singa/utils/string.h b/include/singa/utils/string.h
index b739afc..cbfb28b 100644
--- a/include/singa/utils/string.h
+++ b/include/singa/utils/string.h
@@ -42,6 +42,15 @@ inline string ToLowerCase(const string& input) {
   return out;
 }
 
+inline int ArgPos(int argc, char** arglist, const char* arg) {
+  for (int i = 0; i < argc; i++) {
+    if (strcmp(arglist[i], arg) == 0) {
+      return i;
+    }
+  }
+  return -1;
+}
+
 /**
  * Tokenize a string.
  *

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/d826b2e8/src/core/tensor/tensor.cc
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor.cc b/src/core/tensor/tensor.cc
index b852a54..ec59aaa 100644
--- a/src/core/tensor/tensor.cc
+++ b/src/core/tensor/tensor.cc
@@ -130,20 +130,24 @@ void Tensor::ToDevice(std::shared_ptr<Device> dst) {
 void Tensor::ToHost() { ToDevice(device_->host()); }
 
 template <typename DType>
-void Tensor::CopyDataFromHostPtr(const DType *src, const size_t num) {
+void Tensor::CopyDataFromHostPtr(const DType *src, const size_t num,
+    const size_t offset) {
   CHECK_EQ(sizeof(DType), SizeOf(data_type_))
       << "data_type is " << DataType_Name(data_type_)
       << " user given type is of size " << sizeof(DType);
   if (src != nullptr) {
-    device_->CopyDataFromHostPtr(block(), src, sizeof(DType) * num, 0);
+    device_->CopyDataFromHostPtr(block(), src, sizeof(DType) * num, offset);
   } else {
     LOG(WARNING) << "Copy data from null host ptr";
   }
 }
 template void Tensor::CopyDataFromHostPtr(const unsigned char *src,
-                                          const size_t num);
-template void Tensor::CopyDataFromHostPtr(const float *src, const size_t num);
-template void Tensor::CopyDataFromHostPtr(const int *src, const size_t num);
+                                          const size_t num,
+                                          const size_t offset);
+template void Tensor::CopyDataFromHostPtr(const float *src, const size_t num,
+                                          const size_t offset);
+template void Tensor::CopyDataFromHostPtr(const int *src, const size_t num,
+                                          const size_t offset);
 
 void Tensor::CopyData(const Tensor &src) {
   CHECK_EQ(Size(), src.Size());
@@ -154,7 +158,9 @@ void Tensor::CopyData(const Tensor &src) {
   }
 }
 
-Tensor Tensor::Clone() const {
+Tensor Tensor::Clone(std::shared_ptr<Device> device) const {
+  if (device == nullptr)
+    device = device_;
   Tensor t(shape_, device_, data_type_);
   t.transpose_ = transpose_;
   t.CopyData(*this);

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/d826b2e8/src/model/feed_forward_net.cc
----------------------------------------------------------------------
diff --git a/src/model/feed_forward_net.cc b/src/model/feed_forward_net.cc
new file mode 100644
index 0000000..f9e6480
--- /dev/null
+++ b/src/model/feed_forward_net.cc
@@ -0,0 +1,263 @@
+/**
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "singa/model/feed_forward_net.h"
+#include "singa/utils/logging.h"
+namespace singa {
+
+~FeedForwardNet::FeedForwardNet() {
+  for (auto layer : layers_)
+    delete layer;
+}
+Layer* FeedForwardNet::Add(Layer* layer) {
+  layers_.push_back(layer);
+  return layer;
+}
+
+Layer* FeedForwardNet::Add(const LayerConf& conf, const Shape* sample_shape) {
+  CHECK(sample_shape != nullptr || layers_.size())
+      << "Must provide the input sample shape for the first layer";
+  Layer* layer = CreateLayer(conf.type());
+  if (sample_shape == nullptr)
+    layer->Setup(layers_.back()->GetOutputSampleShape(), conf);
+  else
+    layer->Setup(*sample_shape, conf);
+  Add(layer);
+  return layer;
+}
+
+const vector<string> FeedForwardNet::GetParamNames() const {
+  vector<string> names;
+  for (auto layer : layers_)
+    for (const auto name : layer->param_names())
+      names.push_back(name);
+  return names;
+}
+const vector<Tensor *> FeedForwardNet::GetParamValues() const {
+  vector<Tensor *> values;
+  for (auto layer : layers_)
+    for (const auto value : layer->param_values())
+      values.push_back(value);
+  return values;
+}
+
+const vector<Tensor *> FeedForwardNet::GetParamSpecs() const {
+  vector<ParamSpec *> specs;
+  for (auto layer : layers_)
+    for (const auto spec : layer->param_specs())
+      specs.push_back(spec);
+  return specs;
+}
+
+void FeedForwardNet::Compile(bool shuffle, Optimizer* opt, Loss* loss,
+                             Metric* metric) {
+  shuffle_ = shuffle;
+  bool train = (opt != nullptr) && (loss != nullptr);
+  bool test = metric != nullptr;
+  CHECK(train || test) << "Must set opt and loss, or set metric";
+  opt_ = opt;
+  loss_ = loss;
+  metric_ = metric;
+}
+
+void FeedForwardNet::ToDevice(std::shared_ptr<Device> device) {
+  for (auto layer: layers_)
+    layer->ToDevice(device);
+  opt_->ToDevice(device);
+  loss_->ToDevice(device);
+  metric_->ToDevice(device);
+}
+
+FeedForwardNet FeedForwardNet::Clone(std::shared_ptr<Device> device) {
+  FeedForwardNet net;
+  /*
+  for (auto layer: layers_)
+    net.layers_.push_back(layer->CloneTo(device));
+  if (opt_ != nullptr)
+    net.opt_ = opt_->CloneTo(device);
+  if (loss_ != nullptr)
+    net.loss_ = loss_.CloneTo(device);
+  if (metric_ != nullptr)
+    net.metric_ = metric_->CloneTo(device);
+  net.shuffle_ = shuffle_;
+  net.device_ = device;
+  net.dtype_ = dtype;
+  */
+}
+
+void FeedForwardNet::AsType(DataType dtype) {
+  LOG(FATAL) << "FeedForwardNet::AsType not implemented";
+}
+
+void FeedForwardNet::Train(int batchsize, int nb_epoch, Tensor x, Tensor y) {
+  CHECK_EQ(x.shape(0), y.shape(0)) << "Diff num of sampels in x and y";
+  int num_extra_samples = x.shape(0) % batchsize;
+  if (num_extra_samples != 0)
+    LOG(WARNING) << "The last " << num_extra_samples << " would not be used";
+  Channel *ch = GetChannel("perf");
+  for (int epoch = 0; epoch < nb_epoch; epoch++) {
+    float loss = 0.0f, metric = 0.0f;
+    int batch = 0;
+    for (; batch < x.shape(0) / batchsize; batch++) {
+      Tesnor bx = x.Slice(batch * batchsize, batch * batchsize + batchsize);
+      Tesnor by = y.Slice(batch * batchsize, batch * batchsize + batchsize);
+      const auto ret = TrainOnBatch(bx, by);
+      loss += ret.first;
+      metric += ret.second;
+    }
+    loss /= batch;
+    metric /= batch;
+    ch->Send("Epoch " + std::to_string(epoch) + ", training loss = " +
+             std::to_string(loss) + ", accuracy = " + std::to_string(metric));
+  }
+}
+
+void FeedForwardNet::Train(int batchsize, int nb_epoch, Tensor x, Tensor y,
+   float val_split) {
+  CHECK_EQ(x.shape(0), y.shape(0)) << "Diff num of sampels in x and y";
+  size_t num_train = x.shape(0) * val_split;
+  const Tensor train_x = CopyRows(x, 0, num_train);
+  const Tensor train_y = CopyRows(y, 0, num_train);
+  const Tensor val_x = CopyRows(x, num_train, x.shape(0));
+  const Tensor val_y = CopyRows(y, num_train, x.shape(0));
+  Train(batchsize, nb_epoch, train_x, train_y, val_x, val_y);
+}
+
+
+void FeedForwardNet::Train(int batchsize, int nb_epoch, Tensor x, Tensor y,
+    const Tensor & val_x, const Tensor &val_y) {
+  CHECK_EQ(x.shape(0), y.shape(0)) << "Diff num of sampels in x and y";
+  int num_extra_samples = x.shape(0) % batchsize;
+  if (num_extra_samples != 0)
+    LOG(WARNING) << "The last " << num_extra_samples << " would not be used";
+  Channel *train_ch = GetChannel("train_perf");
+  Channel *test_ch = GetChannel("test_perf");
+  for (int epoch = 0; epoch < nb_epoch; epoch++) {
+    float loss = 0.0f, metric = 0.0f;
+    int b = 0;
+    for (;b < x.shape(0) / batchsize; b++) {
+      Tesnor bx = CopyRows(x, b * batchsize, b * batchsize + batchsize);
+      Tesnor by = CopyRows(y, b * batchsize, b * batchsize + batchsize);
+      const auto ret = TrainOnBatch(bx, by);
+      loss += ret.first;
+      metric += ret.second;
+    }
+    loss /= batch;
+    metric /= batch;
+    train_ch->Send("Epoch " + std::to_string(epoch) + ", training loss = " +
+             std::to_string(loss) + ", accuracy = " + std::to_string(metric));
+    const auto val_perf = Evaluate(val_x, val_y, batchsize);
+    test_ch->Send("Epoch " + std::to_string(epoch)
+        + ", test loss = " + std::to_string(Average(val_perf.first))
+        + ", metric = " + std::to_string(Average(val_perf.second)));
+  }
+}
+
+const std::pair<float, float> FeedForwardNet::TrainOnBatch(const Tensor x,
+                                                           const Tensor y) {
+  const Tensor fea = Forward(kTrain, bx);
+  float loss = loss->Evaluate(fea, fy);
+  float metric = metric->Evaluate(fea, by);
+  const Tensor grad = loss->Backward();
+  Backward(kTrain, grad);
+  return std::make_pair(loss, metric);
+}
+
+const Tensor FeedForwardNet::Forward(int flag, const Tensor data) {
+  Tensor tmp = data;
+  for (auto layer : layers_) {
+    tmp = layer->Forward(flag, tmp);
+  }
+  return tmp;
+}
+
+cons vector<Tensor> FeedForwardNet::Backward(int flag, const Tensor grad) {
+  vector<Tensor> param_grads;
+  Tensor tmp = grad;
+  for (size_t i = layers_.size() - 1; i >= 0; i--) {
+    auto ret = layers_.at(i)->Backward(flag, tmp);
+    tmp =ret.first;
+    if (ret.second.size())
+      for (const auto x: ret.second)
+        param_grads.push_back(x);
+  }
+  return param_grads;
+}
+
+std::pair<Tensor, Tensor> Evaluate(Tensor x, Tensor y, int batchsize) {
+  CHECK_EQ(x.shape(0), y.shape(0)) << "Diff num of sampels in x and y";
+  CHECK_GE(x.shape(0), batchsize);
+  int num_extra_samples = x.shape(0) % batchsize;
+  int b = 0;
+  Tensor loss(Shape{x.shape(0)}), metric(Shape{x.shape(0)});
+  for (; b < x.shape(0) / batchsize; b++) {
+    int start = b * batchsize, end = start + batchsize;
+    const Tensor bx = CopyRows(x, start, end);
+    const Tensor by = CopyRows(y, start, end);
+    const Tensor fea = Forward(kEval, bx);
+    const auto ret = EvaluateOnBatch(bx, by);
+    CopyDataToFrom(&loss, ret.first, batchsize, start, 0);
+    CopyDataToFrom(&metric, ret.second, batchsize, start, 0);
+  }
+  {
+    int start = x.shape(0) - batchsize, end = x.shape(0);
+    const Tensor bx = CopyRows(x, start, end);
+    const Tensor by = CopyRows(y, start, end);
+    const auto ret = EvaluateOnBatch(bx, by);
+    int dst_offset = x.shape(0) - num_extra_samples;
+    int src_offset = batchsize - num_extra_samples;
+    CopyDataToFrom(&loss, ret.first, num_extra_samples, dst_offset, src_offset);
+    CopyDataToFrom(&metric, ret.second, num_extra_samples, dst_offset,
+                   src_offset);
+  }
+  return std::make_pair(loss, metric);
+}
+
+std::pair<Tensor, Tensor> FeedForwardNet::EvaluateOnBatch(const Tensor& x,
+                                                          const Tensor& y) {
+  const Tensor fea = Forward(kEval, bx);
+  const Tensor m = metric_->Forward(fea, by);
+  const Tensor l = loss_->Forward(fea, by);
+  return std::make_pair(m, l);
+}
+
+const Tensor FeedForwardNet::Predict(const Tensor& x, int batchsize) {
+  CHECK_GE(x.shape(0), batchsize);
+  int num_extra_samples = x.shape(0) % batchsize;
+  const auto outshape = layers_.back().GetOutputSampleShape();
+  Tensor y(Shape{x.shape(0), Product(outshape)}, x.device());
+  for (int b = 0; b < x.shape(0) / batchsize; b++) {
+    int start = b * batchsize, end = start + batchsize;
+    const Tensor bx = CopyRows(x, start, end);
+    CopyDataToFrom(&y, PredictOnBatch(bx), batchsize * y.shape(1),
+                   start * y.shape(1), 0);
+  }
+  if (num_extra_samples > 0) {
+    int start = x.shape(0) - batchsize, end = x.shape(0);
+    const Tensor bx = CopyRows(x, start, end);
+    CopyDataToFrom(&y, PredictOnBatch(bx), num_extra_samples * y.shape(1),
+                   (x.shape(0) - num_extra_samples) * y.shape(1),
+                   (batchsize - num_extra_samples) * y.shape(1));
+  }
+  return y;
+}
+
+const Tensor FeedForwardNet::PredictOnBatch(const Tensor& x) {
+  return Foward(kEval, x);
+}
+}  // namespace singa


[2/6] incubator-singa git commit: SINGA-204 Support the training of feed-forward neural nets

Posted by ka...@apache.org.
SINGA-204 Support the training of feed-forward neural nets

Fix the bug from pre/post increament in Block Inc/Dec reference which
resulted in out-of-memory.  Cifar10 data loading works.


Project: http://git-wip-us.apache.org/repos/asf/incubator-singa/repo
Commit: http://git-wip-us.apache.org/repos/asf/incubator-singa/commit/cf1d8418
Tree: http://git-wip-us.apache.org/repos/asf/incubator-singa/tree/cf1d8418
Diff: http://git-wip-us.apache.org/repos/asf/incubator-singa/diff/cf1d8418

Branch: refs/heads/dev
Commit: cf1d841890842c6cf1573491f4fc9d7e1eca30f4
Parents: d826b2e
Author: Wei Wang <wa...@comp.nus.edu.sg>
Authored: Sat Jun 25 19:34:13 2016 +0800
Committer: Wei Wang <wa...@comp.nus.edu.sg>
Committed: Mon Jun 27 15:27:19 2016 +0800

----------------------------------------------------------------------
 examples/cifar10/alexnet.cc            | 105 +++++++++++------
 examples/cifar10/cifar10.cc            |  98 ----------------
 examples/cifar10/cifar10.h             |  99 ++++++++++++++++
 examples/cifar10/make.sh               |   1 +
 include/singa/core/common.h            |  16 +--
 include/singa/core/tensor.h            |   4 +-
 include/singa/model/feed_forward_net.h |  63 +++++-----
 include/singa/model/layer.h            |   2 +-
 include/singa/model/loss.h             |   4 +-
 include/singa/model/metric.h           |  23 ++++
 include/singa/model/optimizer.h        |   2 +-
 src/CMakeLists.txt                     |   3 +-
 src/core/device/cuda_gpu.cc            |   4 +-
 src/core/tensor/tensor.cc              |  14 ++-
 src/model/feed_forward_net.cc          | 176 ++++++++++++++--------------
 src/model/layer/convolution.cc         |   6 +-
 src/model/layer/convolution.h          |   1 +
 src/model/layer/cudnn_activation.cc    |   1 +
 src/model/layer/cudnn_convolution.cc   |   1 +
 src/model/layer/cudnn_dropout.cc       |   4 +
 src/model/layer/cudnn_dropout.h        |   1 +
 src/model/layer/cudnn_pooling.cc       |   1 +
 src/model/layer/cudnn_softmax.cc       |   1 +
 src/model/layer/dense.cc               |   3 +
 src/model/metric/accuracy.cc           |  62 ++++++++++
 src/model/metric/accuracy.h            |  84 -------------
 test/singa/test_accuracy.cc            |   2 +-
 27 files changed, 419 insertions(+), 362 deletions(-)
----------------------------------------------------------------------


http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/examples/cifar10/alexnet.cc
----------------------------------------------------------------------
diff --git a/examples/cifar10/alexnet.cc b/examples/cifar10/alexnet.cc
index 2917dd2..45d8571 100644
--- a/examples/cifar10/alexnet.cc
+++ b/examples/cifar10/alexnet.cc
@@ -22,7 +22,14 @@
 #include "singa/model/feed_forward_net.h"
 #include "singa/model/optimizer.h"
 #include "singa/model/initializer.h"
-
+#include "singa/model/metric.h"
+#include "singa/utils/channel.h"
+#include "singa/utils/string.h"
+#include "../../src/model/layer/cudnn_convolution.h"
+#include "../../src/model/layer/cudnn_activation.h"
+#include "../../src/model/layer/cudnn_pooling.h"
+#include "../../src/model/layer/dense.h"
+#include "../../src/model/layer/flatten.h"
 namespace singa {
 
 LayerConf GenConvConf(string name, int nb_filter, int kernel, int stride,
@@ -32,9 +39,9 @@ LayerConf GenConvConf(string name, int nb_filter, int kernel, int stride,
   conf.set_type("CudnnConvolution");
   ConvolutionConf *conv = conf.mutable_convolution_conf();
   conv->set_num_output(nb_filter);
-  conv->set_kernel_size(kernel);
-  conv->set_stride(stride);
-  conv->set_pad(pad);
+  conv->add_kernel_size(kernel);
+  conv->add_stride(stride);
+  conv->add_pad(pad);
 
   FillerConf *weight = conv->mutable_weight_filler();
   weight->set_type("Xavier");
@@ -50,7 +57,7 @@ LayerConf GenPoolingConf(string name, bool max_pool, int kernel, int stride, int
   pool->set_stride(stride);
   pool->set_pad(pad);
   if (!max_pool)
-    pool->set_pool(PoolingConf_AVE);
+    pool->set_pool(PoolingConf_PoolMethod_AVE);
   return conf;
 }
 
@@ -65,9 +72,9 @@ LayerConf GenDenseConf(string name, int num_output) {
   LayerConf conf;
   conf.set_name(name);
   conf.set_type("Dense");
-  DenseConf *dense = conf->mutable_dense_conf();
+  DenseConf *dense = conf.mutable_dense_conf();
   dense->set_num_output(num_output);
-  FillerConf *weight = conv->mutable_weight_filler();
+  FillerConf *weight = dense->mutable_weight_filler();
   weight->set_type("Xavier");
   return conf;
 }
@@ -79,22 +86,27 @@ LayerConf GenSoftmaxConf(string name) {
   return conf;
 }
 
-
-FeedForwordNet CreateNet(Optimizer* opt, Loss* loss, Metric* metric) {
-  FeedForwordNet net;
+LayerConf GenFlattenConf(string name) {
+  LayerConf conf;
+  conf.set_name(name);
+  conf.set_type("Flatten");
+  return conf;
+}
+FeedForwardNet CreateNet(Optimizer* opt, Loss<Tensor>* loss, Metric<Tensor>* metric) {
+  FeedForwardNet net;
   Shape s{3, 32, 32};
-  net.AddLayer(GenConvConf("conv1", 32, 5, 1, 2), &s);
-  net.AddLayer(GenReLUConf("relu1"));
-  net.AddLayer(GenConvConf("pool1", 3, 2, 0));
-  net.AddLayer(GenConvConf("conv2", 32, 5, 1, 2));
-  net.AddLayer(GenReLUConf("relu2"));
-  net.AddLayer(GenConvConf("pool2", 3, 2, 0));
-  net.AddLayer(GenConvConf("conv3", 64, 5, 1, 2));
-  net.AddLayer(GenReLUConf("relu3"));
-  net.AddLayer(GenConvConf("pool3", 3, 2, 0));
-  net.AddLayer(GenDenseConf("ip1", 10));
-  net.AddLayer(GenSoftmaxConf("softmax"));
 
+  net.Add(new CudnnConvolution(), GenConvConf("conv1", 32, 5, 1, 2), &s);
+  net.Add(new CudnnActivation(), GenReLUConf("relu1"));
+  net.Add(new CudnnPooling, GenPoolingConf("pool1", true, 3, 2, 0));
+  net.Add(new CudnnConvolution(), GenConvConf("conv2", 32, 5, 1, 2));
+  net.Add(new CudnnActivation(), GenReLUConf("relu2"));
+  net.Add(new CudnnPooling(), GenPoolingConf("pool2", true, 3, 2, 0));
+  net.Add(new CudnnConvolution, GenConvConf("conv3", 64, 5, 1, 2));
+  net.Add(new CudnnActivation(), GenReLUConf("relu3"));
+  net.Add(new CudnnConvolution(), GenConvConf("pool3", true, 3, 2, 0));
+  net.Add(new Flatten(), GenFlattenConf("flat"));
+  net.Add(new Dense(), GenDenseConf("ip1", 10));
   OptimizerConf opt_conf;
   opt_conf.set_momentum(0.9);
   opt->Setup(opt_conf);
@@ -103,42 +115,57 @@ FeedForwordNet CreateNet(Optimizer* opt, Loss* loss, Metric* metric) {
 }
 
 void Train(float lr, int num_epoch, string data_dir) {
-  SoftmaxCrossEntropy loss;
-  Accuracy acc;
-  SGD sgd;
-  sgd.SetLearningRate([lr](int step) {return lr;});
-  auto net = CreateNet(&opt, &loss, &metric);
   Cifar10 data(data_dir);
-  Tensor train_x, tain_y, test_x, test_y;
+  Tensor train_x, train_y, test_x, test_y;
   {
     auto train = data.ReadTrainData();
-    const auto mean = Average(train.first, 0);
-    train_x = SubRow(train.first, mean);
-    auto test = data.ReadTestData();
-    test_x = SubRow(test.first, mean);
+    size_t nsamples = train.first.shape(0);
+    auto matx = Reshape(train.first, Shape{nsamples, train.first.Size() / nsamples});
+    const auto mean = Average(matx, 0);
+    SubRow(mean, &matx);
+    train_x = Reshape(matx, train.first.shape());
     train_y = train.second;
+    auto test = data.ReadTestData();
+    nsamples = test.first.shape(0);
+    auto maty = Reshape(test.first, Shape{nsamples, test.first.Size() / nsamples});
+    SubRow(mean, &maty);
+    test_x = Reshape(maty, test.first.shape());
     test_y = test.second;
   }
-  net.Train(100, num_epoch, train_x, train_y, test_x, test_y);
+  LOG(ERROR) << "creating net";
+  SoftmaxCrossEntropy loss;
+  Accuracy acc;
+  SGD sgd;
+  sgd.SetLearningRateGenerator([lr](int step) {return lr;});
+  auto net = CreateNet(&sgd, &loss, &acc);
+
+  auto cuda = std::make_shared<CudaGPU>();
+  net.ToDevice(cuda);
+
+  train_x.ToDevice(cuda);
+  train_y.ToDevice(cuda);
+  net.Train(50, num_epoch, train_x, train_y); // test_x, test_y);
+}
+
+
 }
 
 int main(int argc, char** argv) {
-  InitChannel();
-  int pos = ArgPos(argc, argv, "-epoch");
+  singa::InitChannel(nullptr);
+  int pos = singa::ArgPos(argc, argv, "-epoch");
   int nEpoch = 5;
   if (pos != -1)
     nEpoch = atoi(argv[pos + 1]);
-  pos = ArgPos(argc, argv, "-lr");
+  pos = singa::ArgPos(argc, argv, "-lr");
   float lr = 0.01;
   if (pos != -1)
     lr = atof(argv[pos + 1]);
-  pos = ArgPos(argc, argv, "-data");
-  string data = "cifar-10-batch-bin";
+  pos = singa::ArgPos(argc, argv, "-data");
+  string data = "cifar-10-batches-bin";
   if (pos != -1)
     data = argv[pos + 1];
 
   LOG(INFO) << "Start training";
-  Train(lr, nEpoch, data);
+  singa::Train(lr, nEpoch, data);
   LOG(INFO) << "End training";
 }
-}

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/examples/cifar10/cifar10.cc
----------------------------------------------------------------------
diff --git a/examples/cifar10/cifar10.cc b/examples/cifar10/cifar10.cc
deleted file mode 100644
index 7efc18f..0000000
--- a/examples/cifar10/cifar10.cc
+++ /dev/null
@@ -1,98 +0,0 @@
-/************************************************************
-*
-* Licensed to the Apache Software Foundation (ASF) under one
-* or more contributor license agreements.  See the NOTICE file
-* distributed with this work for additional information
-* regarding copyright ownership.  The ASF licenses this file
-* to you under the Apache License, Version 2.0 (the
-* "License"); you may not use this file except in compliance
-* with the License.  You may obtain a copy of the License at
-*
-*   http://www.apache.org/licenses/LICENSE-2.0
-*
-* Unless required by applicable law or agreed to in writing,
-* software distributed under the License is distributed on an
-* "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-* KIND, either express or implied.  See the License for the
-* specific language governing permissions and limitations
-* under the License.
-*
-*************************************************************/
-#include <fstream>
-#include <string>
-#include <cstdint>
-#include <iostream>
-
-using std::string;
-namespace singa {
-/// For reading cifar10 binary data as tensors.
-class Cifar10 {
- public:
-  /// 'dir_path': path to the folder including the *.bin files
-  Cifar10(string dir_path, bool normalize = true)
-      : dir_path_(dir_path), normalize_(normalize) {}
-
-  /// read all training data into an image Tensor and a label Tensor
-  const std::pair<Tensor, Tensor> ReadTrainData(bool shuffle = false);
-  /// read all test data into an image Tensor and a label Tensor
-  const std::pair<Tensor, Tensor> ReadTestData();
-  /// read data from one file into an image Tensor and a label Tensor
-  const std::pair<Tensor, Tensor> ReadFile(string file, bool shuffle = false);
-
- private:
-  const int kImageSize = 32;
-  const int kImageVol = 3072;
-  const int kBatchSize = 10000;
-  const int kTrainFiles = 5;
-
-  string dir_path_;
-  bool normalize_;
-};
-
-void read_image(std::ifstream* file, int* label, char* buffer) {
-  char label_char;
-  file->read(&label_char, 1);
-  *label = label_char;
-  file->read(buffer, kImageVol);
-  return;
-}
-const std::pair<Tensor, Tensor> Cifar10::ReadFile(string file,
-                                                  bool shuffle = false) {
-  Tensor images(Shape{kTrainFiles, 3, kImageSize, kImageSize});
-  Tensor labels(Shape{kTrainFiles}, kInt);
-  if (dir_path_.back() != '/') dir_path_.push_back('/');
-  LOG(INFO) << "Reading file " << dir_path_ + file;
-  std::ifstream data_file((dir_path_ + file).c_str(),
-                          std::ios::in | std::ios::binary);
-  CHECK(data_file.is_open()) << "Unable to open file " << file;
-  int label;
-  char image[kImageVol];
-  float float_image[kImageVol];
-  int tmplabels[kBatchSize];
-  for (int itemid = 0; itemid < kBatchSize; ++itemid) {
-    read_image(&data_file, &label, image);
-    for (int i = 0; i < kImageVol; i++)
-      float_image[i] = static_cast<float>(static_cast<int>(image[i]));
-    images.CopyDataFromHostPtr(float_image, kImageVol, itemid * kImageVol);
-    tmplabels[itemid] = label;
-  }
-  labels.CopyDataFromHostPtr(tmplabels, kBatchSize);
-  return std::make_pair(images, labels);
-}
-
-const std::pair<Tensor, Tensor> Cifar10::ReadTrainData(bool shuffle = false) {
-  Tensor images(Shape{kBatchSize * kTrainFiles, 3, kImageSize, kImageSize});
-  Tensor labels(Shape{kBatchSize * kTrainFiles, 3, kImageSize, kImageSize});
-  for (int fileid = 0; fileid < kTrainFiles; ++fileid) {
-    string file = "data_batch_" + std::to_string(fileid + 1) + ".bin";
-    const auto ret = ReadFile(file);
-    CopyDataToFrom(&images, ret.first, ret.first.Size(),
-                   fileid * ret.first.Size());
-    CopyDataToFrom(&labels, ret.second, kBatchSize, fileid * kBatchSize);
-  }
-  return std::make_pair(images, labels);
-}
-const std::pair<Tensor, Tensor> Cifar10::ReadTrainData() {
-  return ReadFile("test_batch.bin");
-}
-}  // namespace singa

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/examples/cifar10/cifar10.h
----------------------------------------------------------------------
diff --git a/examples/cifar10/cifar10.h b/examples/cifar10/cifar10.h
new file mode 100644
index 0000000..261c048
--- /dev/null
+++ b/examples/cifar10/cifar10.h
@@ -0,0 +1,99 @@
+/************************************************************
+*
+* Licensed to the Apache Software Foundation (ASF) under one
+* or more contributor license agreements.  See the NOTICE file
+* distributed with this work for additional information
+* regarding copyright ownership.  The ASF licenses this file
+* to you under the Apache License, Version 2.0 (the
+* "License"); you may not use this file except in compliance
+* with the License.  You may obtain a copy of the License at
+*
+*   http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing,
+* software distributed under the License is distributed on an
+* "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+* KIND, either express or implied.  See the License for the
+* specific language governing permissions and limitations
+* under the License.
+*
+*************************************************************/
+#include <fstream>
+#include <string>
+#include <cstdint>
+#include <iostream>
+#include "singa/core/tensor.h"
+using std::string;
+namespace singa {
+/// For reading cifar10 binary data as tensors.
+class Cifar10 {
+ public:
+  /// 'dir_path': path to the folder including the *.bin files
+  Cifar10(string dir_path, bool normalize = true)
+      : dir_path_(dir_path), normalize_(normalize) {}
+
+  /// read all training data into an image Tensor and a label Tensor
+  const std::pair<Tensor, Tensor> ReadTrainData(bool shuffle = false);
+  /// read all test data into an image Tensor and a label Tensor
+  const std::pair<Tensor, Tensor> ReadTestData();
+  /// read data from one file into an image Tensor and a label Tensor
+  const std::pair<Tensor, Tensor> ReadFile(string file, bool shuffle = false);
+
+  void ReadImage(std::ifstream* file, int* label, char* buffer);
+ private:
+  const size_t kImageSize = 32;
+  const size_t kImageVol = 3072;
+  const size_t kBatchSize = 10000;
+  const size_t kTrainFiles = 1;
+
+  string dir_path_;
+  bool normalize_;
+};
+
+void Cifar10::ReadImage(std::ifstream* file, int* label, char* buffer) {
+  char label_char;
+  file->read(&label_char, 1);
+  *label = static_cast<int>(label_char);
+  file->read(buffer, kImageVol);
+  return;
+}
+const std::pair<Tensor, Tensor> Cifar10::ReadFile(string file, bool shuffle) {
+  Tensor images(Shape{kBatchSize, 3, kImageSize, kImageSize});
+  Tensor labels(Shape{kBatchSize}, kInt);
+  if (dir_path_.back() != '/') dir_path_.push_back('/');
+  LOG(INFO) << "Reading file " << dir_path_ + file;
+  std::ifstream data_file((dir_path_ + file).c_str(),
+                          std::ios::in | std::ios::binary);
+  CHECK(data_file.is_open()) << "Unable to open file " << dir_path_ + file;
+  int label;
+  char image[kImageVol];
+  float float_image[kImageVol];
+  int tmplabels[kBatchSize];
+  for (int itemid = 0; itemid < kBatchSize; ++itemid) {
+    // LOG(INFO) << "reading " << itemid << "-th image";
+    ReadImage(&data_file, &label, image);
+    for (int i = 0; i < kImageVol; i++)
+      float_image[i] = static_cast<float>(static_cast<int>(image[i]));
+    images.CopyDataFromHostPtr(float_image, kImageVol, itemid * kImageVol);
+    tmplabels[itemid] = label;
+  }
+  labels.CopyDataFromHostPtr(tmplabels, kBatchSize);
+  return std::make_pair(images, labels);
+}
+
+const std::pair<Tensor, Tensor> Cifar10::ReadTrainData(bool shuffle) {
+  Tensor images(Shape{kBatchSize * kTrainFiles, 3, kImageSize, kImageSize});
+  Tensor labels(Shape{kBatchSize * kTrainFiles}, kInt);
+  for (int fileid = 0; fileid < kTrainFiles; ++fileid) {
+    string file = "data_batch_" + std::to_string(fileid + 1) + ".bin";
+    const auto ret = ReadFile(file);
+    CopyDataToFrom(&images, ret.first, ret.first.Size(),
+                   fileid * ret.first.Size());
+    CopyDataToFrom(&labels, ret.second, kBatchSize, fileid * kBatchSize);
+  }
+  return std::make_pair(images, labels);
+}
+const std::pair<Tensor, Tensor> Cifar10::ReadTestData() {
+  return ReadFile("test_batch.bin");
+}
+}  // namespace singa

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/examples/cifar10/make.sh
----------------------------------------------------------------------
diff --git a/examples/cifar10/make.sh b/examples/cifar10/make.sh
new file mode 100755
index 0000000..17e4b39
--- /dev/null
+++ b/examples/cifar10/make.sh
@@ -0,0 +1 @@
+g++ -g --std=c++11 alexnet.cc -o alexnet -I../../include -I../../build/include -I/home/wangwei/local/cudnn4/include -I/home/wangwei/local/include -I/usr/local/cuda/include/ -I../../lib/cnmem/include -L../../build/lib/ -lsinga_core -lsinga_model  -lsinga_utils -lcudart -lcublas -lcurand -lcudnn -L/usr/local/cuda/lib64 -L/home/wangwei/local/cudnn4/lib64 ../../build/lib/libproto.a -lprotobuf

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/include/singa/core/common.h
----------------------------------------------------------------------
diff --git a/include/singa/core/common.h b/include/singa/core/common.h
index cb1bdca..691d7d4 100644
--- a/include/singa/core/common.h
+++ b/include/singa/core/common.h
@@ -49,27 +49,29 @@ class Block {
  public:
   Block(void* ptr, size_t size, size_t offset = 0)
       : data_(ptr), size_(size), offset_(offset) {
-    ref_count_ = std::make_shared<std::atomic<int>>(1);
+    ref_count_ = 1; //std::make_shared<std::atomic<int>>(1);
   }
-  Block(void* ptr, size_t size, size_t offset, std::shared_ptr<atomic<int>> ref)
-      : data_(ptr), size_(size), offset_(offset), ref_count_(ref) {}
+//  Block(void* ptr, size_t size, size_t offset, std::shared_ptr<atomic<int>> ref)
+//      : data_(ptr), size_(size), offset_(offset), ref_count_(ref) {}
   void* mutable_data() const { return static_cast<char*>(data_) + offset_; }
   const void* data() const { return static_cast<char*>(data_) + offset_; }
   size_t size() const { return size_; }
   size_t offset() const { return offset_; }
   int IncRefCount() {
-    return (*ref_count_)++;
+    return ++ref_count_;  //(*ref_count_)++;
   }
   int DecRefCount() {
-    return  (*ref_count_)--;
+    return --ref_count_; // (*ref_count_)--;
   }
-  int ref_count() const { return ref_count_->load(); }
+  int ref_count() const { return ref_count_.load(); }
 
  private:
+  Block() {}
   void* data_ = nullptr;
   size_t size_ = 0;
   size_t offset_ = 0;
-  std::shared_ptr<std::atomic<int>> ref_count_ = nullptr;
+  // std::shared_ptr<std::atomic<int>> ref_count_ = nullptr;
+  std::atomic<int> ref_count_;
 };
 
 typedef struct _Context {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/include/singa/core/tensor.h
----------------------------------------------------------------------
diff --git a/include/singa/core/tensor.h b/include/singa/core/tensor.h
index 6de5c0c..3b496d9 100644
--- a/include/singa/core/tensor.h
+++ b/include/singa/core/tensor.h
@@ -96,6 +96,8 @@ class Tensor {
 
   /// return number of total elements
   size_t Size() const {
+    if (block_ == nullptr)
+      return 0u;
     CHECK_EQ(block_->size() % SizeOf(data_type_), 0u);
     return block_->size() / SizeOf(data_type_);
   }
@@ -315,7 +317,7 @@ Tensor Div(const SType x, const Tensor &in);
 template <typename SType>
 void Div(const SType x, const Tensor &in, Tensor *out);
 
-template <typename SType>
+template <typename SType = float>
 SType Sum(const Tensor &in);
 // ============Matrix (row/column) operations==================================
 /// Average elements in the Tensor, currently only support vector and matrix.

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/include/singa/model/feed_forward_net.h
----------------------------------------------------------------------
diff --git a/include/singa/model/feed_forward_net.h b/include/singa/model/feed_forward_net.h
index 173600b..9beeb7a 100644
--- a/include/singa/model/feed_forward_net.h
+++ b/include/singa/model/feed_forward_net.h
@@ -18,7 +18,9 @@
 #ifndef SINGA_MODEL_FEED_FORWARD_NET_H_
 #define SINGA_MODEL_FEED_FORWARD_NET_H_
 #include "singa/model/layer.h"
-
+#include "singa/model/loss.h"
+#include "singa/model/metric.h"
+#include "singa/model/optimizer.h"
 namespace singa {
 
 /// The feed-forward neural net.
@@ -26,14 +28,14 @@ namespace singa {
 /// and conducting training, evaluation and prediction.
 class FeedForwardNet {
  public:
-  FeedForwardNet() = explicit;
+  FeedForwardNet() = default;
   ~FeedForwardNet();
 
   /// Add a layer with the assumption that
   /// 1. this function is called in correct order, i.e., the layers are added
   ///    following the topological order.
   /// 2. this layer has already been setup (Setup function is called outside).
-  void Add(Layer *layer);
+  Layer* Add(Layer* layer);
 
   // TODO(wangwei) add ConcatenateLayer and SliceLayer
   // AddConcatenateLayer(vector<Layer*> src, Layer *dst);
@@ -43,36 +45,34 @@ class FeedForwardNet {
   /// Assume the layer is added in corret order.
   /// For the first layer, 'sample_shape' (the input sample shape) is necessary
   /// for calling Setup().
-  void Add(const LayerConf &conf, const Shape *sample_shape = nullptr);
+  Layer* Add(const LayerConf& conf, const Shape* sample_shape = nullptr);
 
+  Layer* Add(Layer* layer, const LayerConf& conf, const Shape* sample_shape = nullptr);
   /// Set some fields used for training and evaluating the neural net.
   /// If the neural net is constructed for evaluation only, then 'opt' is not
   /// necessary; But for training, both 'opt' and 'loss' are necessary.
   /// 'shuffle' indicates shuffling training samples within one epoch it is
   /// valid using Train();
-  void Compile(bool shuffle, Optimizer *opt, Loss *loss, Metric *metric);
+  void Compile(bool shuffle, Optimizer* opt, Loss<Tensor>* loss,
+               Metric<Tensor>* metric);
 
   /// Conduct the training giving the training data 'x' and label 'y'.
-  /// Due to memory limit, 'x' and 'y' could not be very large. Hence, it is
-  /// typically used for small training datasets, e.g., cifar10 and MNIST which
-  /// can be stored in main memory.
-  void Train(int batchsize, int nb_epoch, Tensor x, Tensor y);
-  /// Conduct the training giving the training data 'x' and label 'y'.
-  /// 'val_split' is a ratio for splitting (1-'val_split') of training data for
+  /// 'val_split' of training data is used for
   /// validation. Validation is performance before every epoch.
   /// Due to memory limit, 'x' and 'y' could not be very large. Hence, it is
   /// typically used for small training datasets, e.g., cifar10 and MNIST which
   /// can be stored in main memory.
-  void Train(int batchsize, int nb_epoch, float val_split, Tensor x, Tensor y);
+  void Train(size_t batchsize, int nb_epoch, const Tensor& x, const Tensor& y,
+             float val_split = 0.0f);
   /// Conduct the training given the training and validation data.
   /// Validation is performance before every epoch.
   /// Due to memory limit, 'x' and 'y' could not be very large. Hence, it is
   /// typically used for small training datasets, e.g., cifar10 and MNIST which
   /// can be stored in main memory.
-  void Train(int batchsize, int nb_epoch, Tensor x, Tensor y, Tensor val_x,
-             Tensor val_y);
+  void Train(size_t batchsize, int nb_epoch, const Tensor& x, const Tensor& y,
+             const Tensor& val_x, const Tensor& val_y);
   /// Train the neural net over one batch of training data.
-  Tensor TrainOnBatch(Tensor x, Tensor y);
+  const std::pair<float, float> TrainOnBatch(const Tensor& x, const Tensor& y);
 
   /// Evaluate the neural net with given data.
   /// Returns one tensor for loss values and one tensor for metric values;
@@ -82,9 +82,10 @@ class FeedForwardNet {
   /// Due to memory limit, 'x' and 'y' could not be very large. Hence, it is
   /// typically used for small training datasets, e.g., cifar10 and MNIST which
   /// can be stored in main memory.
-  std::pair<Tensor, Tensor> Evaluate(Tensor x, Tensor y, int batchsize = 128);
+  std::pair<Tensor, Tensor> Evaluate(const Tensor& x, const Tensor& y,
+                                     size_t batchsize = 128);
   /// Evaluate the neural net for one batch of data
-  std::pair<Tensor, Tensor> EvaluateOnBatch(Tensor x, Tensor y);
+  std::pair<Tensor, Tensor> EvaluateOnBatch(const Tensor& x, const Tensor& y);
 
   /// Predict the probability distributation over candicate classes for each
   /// data sample. 'batchsize' is used for controlling the memory footprint.
@@ -92,35 +93,37 @@ class FeedForwardNet {
   /// Due to memory limit, 'x' and 'y' could not be very large. Hence, it is
   /// typically used for small training datasets, e.g., cifar10 and MNIST which
   /// can be stored in main memory.
-  Tensor Predict(const Tensor &x, int batchsize = 128);
+  const Tensor Predict(const Tensor& x, size_t batchsize = 128);
   /// Predict for one batch data.
-  Tensor PredictOnBatch(const Tensor &x);
+  const Tensor PredictOnBatch(const Tensor& x);
 
   /// Forward layers one by one using the data batch 'x'.
   /// Returns the prediction results (from the last layer).
-  Tensor Forward(const Tensor& x);
+  const Tensor Forward(int flag, const Tensor& x);
   /// Backward layers one by one using the gradient batch 'grad'.
   /// Returns the parameter gradients.
-  const vector<Tensor> Backward(const Tensor& grad);
+  const vector<Tensor> Backward(int flag, const Tensor& grad);
 
   /// Clone the neuaral net by cloning every layer to the given device.
   /// If 'device' is nullptr, then clone it one the current device.
-  FeedForwardNet Clone(std::shared_ptr<Device> device = nullptr);
+  FeedForwardNet Clone(std::shared_ptr<Device> device);
   /// Move the layer data to the given device.
-  void ToDevice(Device *device);
+  void ToDevice(std::shared_ptr<Device> device);
+  void ToHost() { ToDevice(defaultDevice); }
   /// Set the data type of each layer.
   void AsType(DataType dtype);
 
-  const vector<Layer *> layers() const { return layers_; }
+  const vector<Layer*> layers() const { return layers_; }
   const vector<string> GetParamNames() const;
-  const vector<Tensor *> GetParamValues() const;
-  const vector<Tensor *> GetParamGrads() const;
+  const vector<ParamSpec> GetParamSpecs() const;
+  const vector<Tensor*> GetParamValues() const;
+  const vector<Tensor*> GetParamGrads() const;
 
  protected:
-  vector<Layer *> layers_;
-  Optimizer *opt_;
-  Loss *loss_;
-  Metric *metric_;
+  vector<Layer*> layers_;
+  Optimizer* opt_;
+  Loss<Tensor>* loss_;
+  Metric<Tensor>* metric_;
 
   bool shuffle_ = true;
   Device* device_ = nullptr;

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/include/singa/model/layer.h
----------------------------------------------------------------------
diff --git a/include/singa/model/layer.h b/include/singa/model/layer.h
index 79eb069..ce8007c 100644
--- a/include/singa/model/layer.h
+++ b/include/singa/model/layer.h
@@ -151,7 +151,7 @@ class Layer {
 
   /// Clone the layer to the given device. Layer data (e.g., parameters) are
   /// deep copied. If 'device' is nullptr, then clone it one the current device.
-  virtual Layer* Clone(std::shared_ptr<Device> device);
+  // virtual Layer* Clone(std::shared_ptr<Device> device);
   /// Move the layer (including its parameters and other internal Tensor) onto
   /// the given device
   virtual void ToDevice(std::shared_ptr<Device> device) {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/include/singa/model/loss.h
----------------------------------------------------------------------
diff --git a/include/singa/model/loss.h b/include/singa/model/loss.h
index 79abace..41ec701 100644
--- a/include/singa/model/loss.h
+++ b/include/singa/model/loss.h
@@ -37,6 +37,7 @@ class Loss {
     Setup(loss);
   }
 	virtual ~Loss(){};
+  virtual void ToDevice(std::shared_ptr<Device> device) {}
   /// Set meta fields from user configurations.
   virtual void Setup(const LossConf& conf) {}
 
@@ -48,7 +49,8 @@ class Loss {
   /// It calls Forward() internally. The calling pattern should be
   /// [Evaluate|Forward] Backward.
   float Evaluate(const Tensor& prediction, const T& target) {
-    const Tensor& loss = Forward(prediction, target);
+    Tensor loss = Forward(prediction, target);
+    loss.ToHost();
     return Sum<float>(loss) / (1.0f * loss.Size());
   }
 

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/include/singa/model/metric.h
----------------------------------------------------------------------
diff --git a/include/singa/model/metric.h b/include/singa/model/metric.h
index b99ff0d..d013fa4 100644
--- a/include/singa/model/metric.h
+++ b/include/singa/model/metric.h
@@ -33,6 +33,7 @@ class Metric {
  public:
   // TODO(wangwei) call Setup using a default MetricConf.
   Metric() = default;
+  virtual void ToDevice(std::shared_ptr<Device> device) {}
   void Setup(const string& conf) {
     MetricConf metric;
     metric.ParseFromString(conf);
@@ -51,6 +52,28 @@ class Metric {
     return Sum<float>(metric) / (1.0f * metric.Size());
   }
 };
+/// Compute the accuray of the prediction, which is matched against the
+/// ground truth labels.
+/// TODO(wangwei) consider multi-label cases.
+class Accuracy : public Metric<Tensor> {
+ public:
+  /// Set meta fields from user configurations.
+  void Setup(const MetricConf& conf) override { top_k_ = conf.top_k(); }
+
+  /// Check the prediction against the target (ground truth) for each data
+  /// sample. The returned Tensor has a float value for each sample, 0 for wrong
+  /// and 1 for correct. Users can call Sum(const Tensor&) / Tensor::Size() to
+  /// get the accuracy.
+  Tensor Forward(const Tensor& prediction, const Tensor& target);
+
+ private:
+  /// \copydoc Match(const Tensor&, const Tensor&);
+  Tensor Match(const Tensor& prediction, const vector<int>& target);
+  /// If the ground truth label is in the top k predicted labels, then the
+  /// prediction is correct.
+  size_t top_k_ = 1;
+};
+
 
 }  // namespace singa
 

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/include/singa/model/optimizer.h
----------------------------------------------------------------------
diff --git a/include/singa/model/optimizer.h b/include/singa/model/optimizer.h
index f912668..a268126 100644
--- a/include/singa/model/optimizer.h
+++ b/include/singa/model/optimizer.h
@@ -155,7 +155,7 @@ class Regularizer {
 };
 
 // =============Vallina SGD with Momentum=====================================
-class SGD : Optimizer {
+class SGD : public Optimizer {
  public:
   void Setup(const OptimizerConf& conf);
   /// Apply the updating algorithm.

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/CMakeLists.txt
----------------------------------------------------------------------
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index c51d454..af09799 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -36,7 +36,7 @@ IF (USE_CUDA)
     SET(FLAGS_BACKUP ${CMAKE_CXX_FLAGS})
     SET(CMAKE_CXX_FLAGS "")
     IF (CMAKE_BUILD_TYPE MATCHES DEBUG)
-        CUDA_COMPILE(cuda_objs SHARED ${cuda_source} 
+        CUDA_COMPILE(cuda_objs SHARED ${cuda_source}
             OPTIONS "-Xcompiler -fPIC -G -g")
     ELSE (CMAKE_BUILD_TYPE MATCHES  DEBUG)
         CUDA_COMPILE(cuda_objs SHARED ${cuda_source} OPTIONS "-Xcompiler -fPIC")
@@ -57,6 +57,7 @@ AUX_SOURCE_DIRECTORY(model model_source)
 AUX_SOURCE_DIRECTORY(model/layer model_source)
 AUX_SOURCE_DIRECTORY(model/optimizer model_source)
 AUX_SOURCE_DIRECTORY(model/loss model_source)
+AUX_SOURCE_DIRECTORY(model/metric model_source)
 #MESSAGE(STATUS "MODEL ${model_source}")
 ADD_LIBRARY(singa_model SHARED ${model_source})
 TARGET_LINK_LIBRARIES(singa_model ${SINGA_LINKER_LIBS})

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/core/device/cuda_gpu.cc
----------------------------------------------------------------------
diff --git a/src/core/device/cuda_gpu.cc b/src/core/device/cuda_gpu.cc
index 5879c58..5f6ac17 100644
--- a/src/core/device/cuda_gpu.cc
+++ b/src/core/device/cuda_gpu.cc
@@ -121,7 +121,7 @@ void CudaGPU::CopyToFrom(void* dst, const void* src, size_t nBytes,
   // cudaMemcpyAsync(dst, src, nBytes,cudaMemcpyDefault, ctx_.stream);
 }
 
-/// Allocate cpu memory.
+/// Allocate gpu memory.
 void* CudaGPU::Malloc(int size) {
   void* ptr = nullptr;
   if (size > 0) {
@@ -132,7 +132,7 @@ void* CudaGPU::Malloc(int size) {
   return ptr;
 }
 
-/// Free cpu memory.
+/// Free gpu memory.
 void CudaGPU::Free(void* ptr) {
   if (ptr != nullptr) {
     // CUDA_CHECK(cudaFree(ptr));

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/core/tensor/tensor.cc
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor.cc b/src/core/tensor/tensor.cc
index ec59aaa..898cdc6 100644
--- a/src/core/tensor/tensor.cc
+++ b/src/core/tensor/tensor.cc
@@ -118,7 +118,8 @@ void Tensor::ToDevice(std::shared_ptr<Device> dst) {
   // TODO(wangwei) the comparison is very strict. May compare against device ID?
   if (device_ != dst) {
     Tensor tmp(shape_, dst, data_type_);
-    tmp.CopyData(*this);
+    if (block_ != nullptr && Size())
+      tmp.CopyData(*this);
     if (block_ != nullptr && block_->DecRefCount() == 0)
       device_->FreeBlock(block_);
     block_ = tmp.block_;
@@ -136,7 +137,8 @@ void Tensor::CopyDataFromHostPtr(const DType *src, const size_t num,
       << "data_type is " << DataType_Name(data_type_)
       << " user given type is of size " << sizeof(DType);
   if (src != nullptr) {
-    device_->CopyDataFromHostPtr(block(), src, sizeof(DType) * num, offset);
+    device_->CopyDataFromHostPtr(block(), src, sizeof(DType) * num,
+        sizeof(DType) * offset);
   } else {
     LOG(WARNING) << "Copy data from null host ptr";
   }
@@ -637,13 +639,13 @@ Tensor ConcatenateColumns(const vector<Tensor> &in) {
   return out;
 }
 Tensor CopyRows(const Tensor &in, const size_t start, const size_t end) {
-  CHECK_EQ(in.nDim(), 2u);
   CHECK_LT(start, end);
   CHECK_GE(in.shape(0), end);
-  Shape s;
-  s = Shape{end - start, in.shape(1)};
+  Shape s = in.shape();
+  s[0] = end - start;
+  size_t sample_size = in.Size() / in.shape(0);
   Tensor out(s, in.device(), in.data_type());
-  CopyDataToFrom(&out, in, out.Size(), 0, start * out.shape(1));
+  CopyDataToFrom(&out, in, out.Size(), 0, start * sample_size);
   return out;
 }
 Tensor CopyColumns(const Tensor &in, const size_t start, const size_t end) {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/feed_forward_net.cc
----------------------------------------------------------------------
diff --git a/src/model/feed_forward_net.cc b/src/model/feed_forward_net.cc
index f9e6480..a24d36a 100644
--- a/src/model/feed_forward_net.cc
+++ b/src/model/feed_forward_net.cc
@@ -18,11 +18,11 @@
 
 #include "singa/model/feed_forward_net.h"
 #include "singa/utils/logging.h"
+#include "singa/utils/channel.h"
 namespace singa {
 
-~FeedForwardNet::FeedForwardNet() {
-  for (auto layer : layers_)
-    delete layer;
+FeedForwardNet::~FeedForwardNet() {
+  for (auto layer : layers_) delete layer;
 }
 Layer* FeedForwardNet::Add(Layer* layer) {
   layers_.push_back(layer);
@@ -32,7 +32,12 @@ Layer* FeedForwardNet::Add(Layer* layer) {
 Layer* FeedForwardNet::Add(const LayerConf& conf, const Shape* sample_shape) {
   CHECK(sample_shape != nullptr || layers_.size())
       << "Must provide the input sample shape for the first layer";
-  Layer* layer = CreateLayer(conf.type());
+  Layer* layer = nullptr;  // TODO(wangwei) use CreateLayer(conf.type());
+  Add(layer, conf, sample_shape);
+  return layer;
+}
+
+Layer* FeedForwardNet::Add(Layer* layer, const LayerConf& conf, const Shape* sample_shape) {
   if (sample_shape == nullptr)
     layer->Setup(layers_.back()->GetOutputSampleShape(), conf);
   else
@@ -44,28 +49,25 @@ Layer* FeedForwardNet::Add(const LayerConf& conf, const Shape* sample_shape) {
 const vector<string> FeedForwardNet::GetParamNames() const {
   vector<string> names;
   for (auto layer : layers_)
-    for (const auto name : layer->param_names())
-      names.push_back(name);
+    for (const auto name : layer->param_names()) names.push_back(name);
   return names;
 }
-const vector<Tensor *> FeedForwardNet::GetParamValues() const {
-  vector<Tensor *> values;
+const vector<Tensor*> FeedForwardNet::GetParamValues() const {
+  vector<Tensor*> values;
   for (auto layer : layers_)
-    for (const auto value : layer->param_values())
-      values.push_back(value);
+    for (const auto value : layer->param_values()) values.push_back(value);
   return values;
 }
 
-const vector<Tensor *> FeedForwardNet::GetParamSpecs() const {
-  vector<ParamSpec *> specs;
+const vector<ParamSpec> FeedForwardNet::GetParamSpecs() const {
+  vector<ParamSpec> specs;
   for (auto layer : layers_)
-    for (const auto spec : layer->param_specs())
-      specs.push_back(spec);
+    for (const auto spec : layer->param_specs()) specs.push_back(spec);
   return specs;
 }
 
-void FeedForwardNet::Compile(bool shuffle, Optimizer* opt, Loss* loss,
-                             Metric* metric) {
+void FeedForwardNet::Compile(bool shuffle, Optimizer* opt, Loss<Tensor>* loss,
+                             Metric<Tensor>* metric) {
   shuffle_ = shuffle;
   bool train = (opt != nullptr) && (loss != nullptr);
   bool test = metric != nullptr;
@@ -73,14 +75,17 @@ void FeedForwardNet::Compile(bool shuffle, Optimizer* opt, Loss* loss,
   opt_ = opt;
   loss_ = loss;
   metric_ = metric;
+  // init params and register them to sgd
 }
 
 void FeedForwardNet::ToDevice(std::shared_ptr<Device> device) {
   for (auto layer: layers_)
     layer->ToDevice(device);
+  /*
   opt_->ToDevice(device);
   loss_->ToDevice(device);
   metric_->ToDevice(device);
+  */
 }
 
 FeedForwardNet FeedForwardNet::Clone(std::shared_ptr<Device> device) {
@@ -98,118 +103,110 @@ FeedForwardNet FeedForwardNet::Clone(std::shared_ptr<Device> device) {
   net.device_ = device;
   net.dtype_ = dtype;
   */
+  return net;
 }
 
 void FeedForwardNet::AsType(DataType dtype) {
   LOG(FATAL) << "FeedForwardNet::AsType not implemented";
 }
 
-void FeedForwardNet::Train(int batchsize, int nb_epoch, Tensor x, Tensor y) {
-  CHECK_EQ(x.shape(0), y.shape(0)) << "Diff num of sampels in x and y";
-  int num_extra_samples = x.shape(0) % batchsize;
-  if (num_extra_samples != 0)
-    LOG(WARNING) << "The last " << num_extra_samples << " would not be used";
-  Channel *ch = GetChannel("perf");
-  for (int epoch = 0; epoch < nb_epoch; epoch++) {
-    float loss = 0.0f, metric = 0.0f;
-    int batch = 0;
-    for (; batch < x.shape(0) / batchsize; batch++) {
-      Tesnor bx = x.Slice(batch * batchsize, batch * batchsize + batchsize);
-      Tesnor by = y.Slice(batch * batchsize, batch * batchsize + batchsize);
-      const auto ret = TrainOnBatch(bx, by);
-      loss += ret.first;
-      metric += ret.second;
-    }
-    loss /= batch;
-    metric /= batch;
-    ch->Send("Epoch " + std::to_string(epoch) + ", training loss = " +
-             std::to_string(loss) + ", accuracy = " + std::to_string(metric));
-  }
-}
-
-void FeedForwardNet::Train(int batchsize, int nb_epoch, Tensor x, Tensor y,
-   float val_split) {
+void FeedForwardNet::Train(size_t batchsize, int nb_epoch, const Tensor& x,
+                           const Tensor& y, float val_split) {
   CHECK_EQ(x.shape(0), y.shape(0)) << "Diff num of sampels in x and y";
   size_t num_train = x.shape(0) * val_split;
-  const Tensor train_x = CopyRows(x, 0, num_train);
-  const Tensor train_y = CopyRows(y, 0, num_train);
-  const Tensor val_x = CopyRows(x, num_train, x.shape(0));
-  const Tensor val_y = CopyRows(y, num_train, x.shape(0));
-  Train(batchsize, nb_epoch, train_x, train_y, val_x, val_y);
+  if (val_split == 0.0f) {
+    Tensor dummy;
+    Train(batchsize, nb_epoch, x, y, dummy, dummy);
+  } else {
+    const Tensor train_x = CopyRows(x, 0, num_train);
+    const Tensor train_y = CopyRows(y, 0, num_train);
+    const Tensor test_x = CopyRows(x, num_train, x.shape(0));
+    const Tensor test_y = CopyRows(y, num_train, y.shape(0));
+    Train(batchsize, nb_epoch, train_x, train_y, test_x, test_y);
+  }
 }
 
-
-void FeedForwardNet::Train(int batchsize, int nb_epoch, Tensor x, Tensor y,
-    const Tensor & val_x, const Tensor &val_y) {
+void FeedForwardNet::Train(size_t batchsize, int nb_epoch, const Tensor& x,
+                           const Tensor& y, const Tensor& val_x,
+                           const Tensor& val_y) {
+  InitNetParams();
   CHECK_EQ(x.shape(0), y.shape(0)) << "Diff num of sampels in x and y";
   int num_extra_samples = x.shape(0) % batchsize;
   if (num_extra_samples != 0)
     LOG(WARNING) << "The last " << num_extra_samples << " would not be used";
-  Channel *train_ch = GetChannel("train_perf");
-  Channel *test_ch = GetChannel("test_perf");
+  Channel* train_ch = GetChannel("train_perf");
+  train_ch->EnableDestStderr(true);
+  Channel* val_ch = GetChannel("val_perf");
   for (int epoch = 0; epoch < nb_epoch; epoch++) {
     float loss = 0.0f, metric = 0.0f;
-    int b = 0;
-    for (;b < x.shape(0) / batchsize; b++) {
-      Tesnor bx = CopyRows(x, b * batchsize, b * batchsize + batchsize);
-      Tesnor by = CopyRows(y, b * batchsize, b * batchsize + batchsize);
+    size_t b = 0;
+    for (; b < x.shape(0) / batchsize; b++) {
+      const Tensor bx = CopyRows(x, b * batchsize, b * batchsize + batchsize);
+      const Tensor by = CopyRows(y, b * batchsize, b * batchsize + batchsize);
       const auto ret = TrainOnBatch(bx, by);
       loss += ret.first;
       metric += ret.second;
     }
-    loss /= batch;
-    metric /= batch;
+    loss /= b;
+    metric /= b;
     train_ch->Send("Epoch " + std::to_string(epoch) + ", training loss = " +
-             std::to_string(loss) + ", accuracy = " + std::to_string(metric));
-    const auto val_perf = Evaluate(val_x, val_y, batchsize);
-    test_ch->Send("Epoch " + std::to_string(epoch)
-        + ", test loss = " + std::to_string(Average(val_perf.first))
-        + ", metric = " + std::to_string(Average(val_perf.second)));
+                   std::to_string(loss) + ", accuracy = " +
+                   std::to_string(metric));
+    if (val_x.Size() && val_y.Size()) {
+      const auto val_perf = Evaluate(val_x, val_y, batchsize);
+      val_ch->Send("Epoch " + std::to_string(epoch) + ", val loss = " +
+                   std::to_string(Sum(val_perf.first) / val_y.Size()) +
+                   ", metric = " +
+                   std::to_string(Sum(val_perf.second) / val_y.Size()));
+    }
   }
 }
 
-const std::pair<float, float> FeedForwardNet::TrainOnBatch(const Tensor x,
-                                                           const Tensor y) {
-  const Tensor fea = Forward(kTrain, bx);
-  float loss = loss->Evaluate(fea, fy);
-  float metric = metric->Evaluate(fea, by);
-  const Tensor grad = loss->Backward();
-  Backward(kTrain, grad);
+const std::pair<float, float> FeedForwardNet::TrainOnBatch(const Tensor& x,
+                                                           const Tensor& y) {
+  int flag = kTrain;
+  const Tensor fea = Forward(flag, x);
+  float loss = loss_->Evaluate(fea, y);
+  float metric = metric_->Evaluate(fea, y);
+  const Tensor grad = loss_->Backward();
+  const auto grads = Backward(kTrain, grad);
   return std::make_pair(loss, metric);
 }
 
-const Tensor FeedForwardNet::Forward(int flag, const Tensor data) {
-  Tensor tmp = data;
+const Tensor FeedForwardNet::Forward(int flag, const Tensor& data) {
+  Tensor input = data, output;
   for (auto layer : layers_) {
-    tmp = layer->Forward(flag, tmp);
+//    LOG(INFO) << layer->name();
+    output = layer->Forward(flag, input);
+    input = output;
   }
-  return tmp;
+  return output;
 }
 
-cons vector<Tensor> FeedForwardNet::Backward(int flag, const Tensor grad) {
+const vector<Tensor> FeedForwardNet::Backward(int flag, const Tensor& grad) {
   vector<Tensor> param_grads;
   Tensor tmp = grad;
-  for (size_t i = layers_.size() - 1; i >= 0; i--) {
+  for (int i = layers_.size() - 1; i >= 0; i--) {
+ //   LOG(INFO) << layers_.at(i)->name();
     auto ret = layers_.at(i)->Backward(flag, tmp);
-    tmp =ret.first;
+    tmp = ret.first;
     if (ret.second.size())
-      for (const auto x: ret.second)
-        param_grads.push_back(x);
+      for (const auto x : ret.second) param_grads.push_back(x);
   }
   return param_grads;
 }
 
-std::pair<Tensor, Tensor> Evaluate(Tensor x, Tensor y, int batchsize) {
+std::pair<Tensor, Tensor> FeedForwardNet::Evaluate(const Tensor& x,
+                                                   const Tensor& y,
+                                                   size_t batchsize) {
   CHECK_EQ(x.shape(0), y.shape(0)) << "Diff num of sampels in x and y";
   CHECK_GE(x.shape(0), batchsize);
   int num_extra_samples = x.shape(0) % batchsize;
-  int b = 0;
   Tensor loss(Shape{x.shape(0)}), metric(Shape{x.shape(0)});
-  for (; b < x.shape(0) / batchsize; b++) {
+  for (size_t b = 0; b < x.shape(0) / batchsize; b++) {
     int start = b * batchsize, end = start + batchsize;
     const Tensor bx = CopyRows(x, start, end);
     const Tensor by = CopyRows(y, start, end);
-    const Tensor fea = Forward(kEval, bx);
     const auto ret = EvaluateOnBatch(bx, by);
     CopyDataToFrom(&loss, ret.first, batchsize, start, 0);
     CopyDataToFrom(&metric, ret.second, batchsize, start, 0);
@@ -230,18 +227,19 @@ std::pair<Tensor, Tensor> Evaluate(Tensor x, Tensor y, int batchsize) {
 
 std::pair<Tensor, Tensor> FeedForwardNet::EvaluateOnBatch(const Tensor& x,
                                                           const Tensor& y) {
-  const Tensor fea = Forward(kEval, bx);
-  const Tensor m = metric_->Forward(fea, by);
-  const Tensor l = loss_->Forward(fea, by);
+  int flag = kEval;
+  const Tensor fea = Forward(flag, x);
+  const Tensor m = metric_->Forward(fea, y);
+  const Tensor l = loss_->Forward(fea, y);
   return std::make_pair(m, l);
 }
 
-const Tensor FeedForwardNet::Predict(const Tensor& x, int batchsize) {
+const Tensor FeedForwardNet::Predict(const Tensor& x, size_t batchsize) {
   CHECK_GE(x.shape(0), batchsize);
   int num_extra_samples = x.shape(0) % batchsize;
-  const auto outshape = layers_.back().GetOutputSampleShape();
+  const auto outshape = layers_.back()->GetOutputSampleShape();
   Tensor y(Shape{x.shape(0), Product(outshape)}, x.device());
-  for (int b = 0; b < x.shape(0) / batchsize; b++) {
+  for (size_t b = 0; b < x.shape(0) / batchsize; b++) {
     int start = b * batchsize, end = start + batchsize;
     const Tensor bx = CopyRows(x, start, end);
     CopyDataToFrom(&y, PredictOnBatch(bx), batchsize * y.shape(1),
@@ -258,6 +256,6 @@ const Tensor FeedForwardNet::Predict(const Tensor& x, int batchsize) {
 }
 
 const Tensor FeedForwardNet::PredictOnBatch(const Tensor& x) {
-  return Foward(kEval, x);
+  return Forward(kEval, x);
 }
 }  // namespace singa

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/layer/convolution.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/convolution.cc b/src/model/layer/convolution.cc
index c27960d..e4991a4 100644
--- a/src/model/layer/convolution.cc
+++ b/src/model/layer/convolution.cc
@@ -112,5 +112,9 @@ const std::pair<Tensor, vector<Tensor>> Convolution::Backward(
 
   return std::make_pair(input_grad, param_grad);
 }
-
+void Convolution::ToDevice(std::shared_ptr<Device> device) {
+  Layer::ToDevice(device);
+  weight_.ToDevice(device);
+  bias_.ToDevice(device);
+}
 }  // namespace singa

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/layer/convolution.h
----------------------------------------------------------------------
diff --git a/src/model/layer/convolution.h b/src/model/layer/convolution.h
index 7ea5712..0e0b160 100644
--- a/src/model/layer/convolution.h
+++ b/src/model/layer/convolution.h
@@ -44,6 +44,7 @@ class Convolution : public Layer {
   const std::pair<Tensor, vector<Tensor>> Backward(int flag,
                                                    const Tensor &grad) override;
 
+  void ToDevice(std::shared_ptr<Device> device) override;
   size_t kernel_w() const { return kernel_w_; }
   size_t kernel_h() const { return kernel_h_; }
   size_t pad_w() const { return pad_w_; }

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/layer/cudnn_activation.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/cudnn_activation.cc b/src/model/layer/cudnn_activation.cc
index 4216fea..df728ce 100644
--- a/src/model/layer/cudnn_activation.cc
+++ b/src/model/layer/cudnn_activation.cc
@@ -54,6 +54,7 @@ void CudnnActivation::InitCudnn(size_t size, DataType dtype) {
 }
 
 const Tensor CudnnActivation::Forward(int flag, const Tensor& input) {
+  CHECK(buf_.empty());
   auto size = input.Size();
   DataType dtype = input.data_type();
   if (!has_init_cudnn_) {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/layer/cudnn_convolution.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/cudnn_convolution.cc b/src/model/layer/cudnn_convolution.cc
index d5ac2a3..eb507b2 100644
--- a/src/model/layer/cudnn_convolution.cc
+++ b/src/model/layer/cudnn_convolution.cc
@@ -156,6 +156,7 @@ void CudnnConvolution::InitCudnn(const Tensor &input) {
 }
 
 const Tensor CudnnConvolution::Forward(int flag, const Tensor &input) {
+  CHECK(buf_.empty());
   CHECK_EQ(input.device()->lang(), kCuda);
   CHECK_EQ(input.nDim(), 4u);
   if (flag & kTrain) buf_.push(input);  // buffer the input for backward

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/layer/cudnn_dropout.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/cudnn_dropout.cc b/src/model/layer/cudnn_dropout.cc
index 2e2e12b..f9b9dbf 100644
--- a/src/model/layer/cudnn_dropout.cc
+++ b/src/model/layer/cudnn_dropout.cc
@@ -106,6 +106,10 @@ const std::pair<Tensor, vector<Tensor>> CudnnDropout::Backward(
   }
   return std::make_pair(dx, param_grad);
 }
+void CudnnDropout::ToDevice(std::shared_ptr<Device> device) {
+  Dropout::ToDevice(device);
+  state.ToDevice(device);
+}
 }  // namespace singa
 #endif  // CUDNN_VERSION_MAJOR>=5
 #endif  // USE_CUDNN

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/layer/cudnn_dropout.h
----------------------------------------------------------------------
diff --git a/src/model/layer/cudnn_dropout.h b/src/model/layer/cudnn_dropout.h
index 6809653..9e0cb9e 100644
--- a/src/model/layer/cudnn_dropout.h
+++ b/src/model/layer/cudnn_dropout.h
@@ -42,6 +42,7 @@ class CudnnDropout : public Dropout {
   const std::pair<Tensor, vector<Tensor>> Backward(int flag,
                                                    const Tensor& grad) override;
 
+  void ToDevice(std::shared_ptr<Device> device) override;
  private:
   /// Init cudnn related data structures.
   void InitCudnn(int size, DataType dtype, std::shared_ptr<Device> dev,

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/layer/cudnn_pooling.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/cudnn_pooling.cc b/src/model/layer/cudnn_pooling.cc
index 6d7a5b1..e49a1ec 100644
--- a/src/model/layer/cudnn_pooling.cc
+++ b/src/model/layer/cudnn_pooling.cc
@@ -78,6 +78,7 @@ void CudnnPooling::InitCudnn(const Tensor &input) {
 }
 
 const Tensor CudnnPooling::Forward(int flag, const Tensor &input) {
+  CHECK(buf_.empty());
   CHECK_EQ(input.device()->lang(), kCuda);
   CHECK_EQ(input.nDim(), 4u);
   size_t batchsize = input.shape(0);

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/layer/cudnn_softmax.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/cudnn_softmax.cc b/src/model/layer/cudnn_softmax.cc
index 77eab98..1d9e0b8 100644
--- a/src/model/layer/cudnn_softmax.cc
+++ b/src/model/layer/cudnn_softmax.cc
@@ -57,6 +57,7 @@ void CudnnSoftmax::InitCudnn(Shape shape, DataType dtype) {
 }
 
 const Tensor CudnnSoftmax::Forward(int flag, const Tensor& input) {
+  CHECK(buf_.empty());
   auto shape = input.shape();
   DataType dtype = input.data_type();
   if (!has_init_cudnn_) {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/layer/dense.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/dense.cc b/src/model/layer/dense.cc
index bad26a8..c6a9f8a 100644
--- a/src/model/layer/dense.cc
+++ b/src/model/layer/dense.cc
@@ -45,7 +45,9 @@ void Dense::Setup(const Shape& in_sample, const LayerConf &conf) {
 
 /// \copydoc Layer::Forward(int flag, const Tensor&)
 const Tensor Dense::Forward(int flag, const Tensor &input) {
+  CHECK(buf_.empty());
   Tensor output;
+  CHECK_EQ(input.nDim(), 2);
   if (transpose_)  // use the transposed version of weight_ for computing
     output = Mult(input, weight_);
   else
@@ -81,6 +83,7 @@ const std::pair<Tensor, vector<Tensor>> Dense::Backward(int flag,
 }
 
 void Dense::ToDevice(std::shared_ptr<Device> device) {
+  Layer::ToDevice(device);
   weight_.ToDevice(device);
   bias_.ToDevice(device);
 }

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/metric/accuracy.cc
----------------------------------------------------------------------
diff --git a/src/model/metric/accuracy.cc b/src/model/metric/accuracy.cc
new file mode 100644
index 0000000..1b667b1
--- /dev/null
+++ b/src/model/metric/accuracy.cc
@@ -0,0 +1,62 @@
+/**
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "singa/model/metric.h"
+#include <algorithm>
+namespace singa {
+
+Tensor Accuracy::Match(const Tensor& predict, const vector<int>& target) {
+  Tensor prediction(predict.shape());
+  prediction.CopyData(predict);
+  size_t batchsize = target.size();
+  size_t nb_classes = prediction.Size() / batchsize;
+  // each row of prediction is the prob distribution for one sample
+  CHECK_EQ(prediction.shape().at(0), batchsize);
+  // TODO(wangwei) CloneToDevice(host);
+  const float* prob = prediction.data<float>();
+  float* score = new float[batchsize];
+  for (size_t b = 0; b < batchsize; b++) {
+    vector<std::pair<float, int>> prob_class;
+    for (size_t c = 0; c < nb_classes; c++) {
+      prob_class.push_back(std::make_pair(prob[b * nb_classes + c], c));
+    }
+    std::partial_sort(prob_class.begin(), prob_class.begin() + top_k_,
+                      prob_class.end(), std::greater<std::pair<float, int>>());
+
+    for (size_t k = 0; k < top_k_; k++)
+      if (prob_class.at(k).second == target.at(b)) score[b] = 1;
+  }
+  Tensor ret(Shape{batchsize});
+  ret.CopyDataFromHostPtr(score, batchsize);
+  return ret;
+}
+
+// TODO(wangwei) consider multi-label cases, where target is of shape
+// nb_samples * nb_classes
+Tensor Accuracy::Forward(const Tensor& prediction, const Tensor& t) {
+  Tensor target(t.shape(), t.data_type());
+  target.CopyData(t);
+  vector<int> target_vec;
+  // TODO(wangwei) copy target to host.
+  const int* target_value = target.data<int>();
+  for (size_t i = 0; i < target.Size(); i++)
+    target_vec.push_back(target_value[i]);
+  return Match(prediction, target_vec);
+}
+
+}  // namespace singa

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/src/model/metric/accuracy.h
----------------------------------------------------------------------
diff --git a/src/model/metric/accuracy.h b/src/model/metric/accuracy.h
deleted file mode 100644
index 69bd96b..0000000
--- a/src/model/metric/accuracy.h
+++ /dev/null
@@ -1,84 +0,0 @@
-/**
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements.  See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership.  The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License.  You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#ifndef SINGA_MODEL_METRIC_ACCURACY_H_
-#define SINGA_MODEL_METRIC_ACCURACY_H_
-#include "singa/model/metric.h"
-#include <algorithm>
-namespace singa {
-
-/// Compute the accuray of the prediction, which is matched against the
-/// ground truth labels.
-/// TODO(wangwei) consider multi-label cases.
-class Accuracy : public Metric<Tensor> {
- public:
-  /// Set meta fields from user configurations.
-  void Setup(const MetricConf& conf) override { top_k_ = conf.top_k(); }
-
-  /// Check the prediction against the target (ground truth) for each data
-  /// sample. The returned Tensor has a float value for each sample, 0 for wrong
-  /// and 1 for correct. Users can call Sum(const Tensor&) / Tensor::Size() to
-  /// get the accuracy.
-  Tensor Forward(const Tensor& prediction, const Tensor& target);
-
- private:
-  /// \copydoc Match(const Tensor&, const Tensor&);
-  Tensor Match(const Tensor& prediction, const vector<int>& target);
-  /// If the ground truth label is in the top k predicted labels, then the
-  /// prediction is correct.
-  size_t top_k_ = 1;
-};
-
-Tensor Accuracy::Match(const Tensor& prediction, const vector<int>& target) {
-  size_t batchsize = target.size();
-  size_t nb_classes = prediction.Size() / batchsize;
-  // each row of prediction is the prob distribution for one sample
-  CHECK_EQ(prediction.shape().at(0), batchsize);
-  // TODO(wangwei) CloneToDevice(host);
-  const float* prob = prediction.data<float>();
-  float* score = new float[batchsize];
-  for (size_t b = 0; b < batchsize; b++) {
-    vector<std::pair<float, int>> prob_class;
-    for (size_t c = 0; c < nb_classes; c++) {
-      prob_class.push_back(std::make_pair(prob[b * nb_classes + c], c));
-    }
-    std::partial_sort(prob_class.begin(), prob_class.begin() + top_k_,
-                      prob_class.end(), std::greater<std::pair<float, int>>());
-
-    for (size_t k = 0; k < top_k_; k++)
-      if (prob_class.at(k).second == target.at(b)) score[b] = 1;
-  }
-  Tensor ret(Shape{batchsize});
-  ret.CopyDataFromHostPtr(score, batchsize);
-  return ret;
-}
-
-// TODO(wangwei) consider multi-label cases, where target is of shape
-// nb_samples * nb_classes
-Tensor Accuracy::Forward(const Tensor& prediction, const Tensor& target) {
-  vector<int> target_vec;
-  // TODO(wangwei) copy target to host.
-  const int* target_value = target.data<int>();
-  for (size_t i = 0; i < target.Size(); i++)
-    target_vec.push_back(target_value[i]);
-  return Match(prediction, target_vec);
-}
-
-}  // namespace singa
-
-#endif  // SINGA_MODEL_METRIC_ACCURACY_H_

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/cf1d8418/test/singa/test_accuracy.cc
----------------------------------------------------------------------
diff --git a/test/singa/test_accuracy.cc b/test/singa/test_accuracy.cc
index dc7719b..4ff14c0 100644
--- a/test/singa/test_accuracy.cc
+++ b/test/singa/test_accuracy.cc
@@ -20,7 +20,7 @@
 *************************************************************/
 
 #include "gtest/gtest.h"
-#include "../src/model/metric/accuracy.h"
+#include "singa/model/metric.h"
 
 TEST(Accuracy, Compute) {
   singa::Accuracy acc;


[5/6] incubator-singa git commit: SINGA-204 Support the training of feed-forward neural nets

Posted by ka...@apache.org.
SINGA-204 Support the training of feed-forward neural nets

Implement Alexnet model for Cifar10 https://code.google.com/p/cuda-convnet/
The accuracy of the test data is 0.82 (the same as reported by in the above link).
NOTE:
1. do not convert from uint_8 data to float directly. Instead, cast
the data to uint8_t and then to float.
2. it is necessary to substract the max value for softmax. Numeric error
(nan) is likely to happen otherwise.


Project: http://git-wip-us.apache.org/repos/asf/incubator-singa/repo
Commit: http://git-wip-us.apache.org/repos/asf/incubator-singa/commit/97648a60
Tree: http://git-wip-us.apache.org/repos/asf/incubator-singa/tree/97648a60
Diff: http://git-wip-us.apache.org/repos/asf/incubator-singa/diff/97648a60

Branch: refs/heads/dev
Commit: 97648a60a9ff9feb55ef1a8b86e5372837e0b4f8
Parents: cf1d841
Author: Wei Wang <wa...@comp.nus.edu.sg>
Authored: Mon Jun 27 01:21:59 2016 +0800
Committer: Wei Wang <wa...@comp.nus.edu.sg>
Committed: Mon Jun 27 19:28:40 2016 +0800

----------------------------------------------------------------------
 examples/cifar10/alexnet.cc             | 136 ++++++++++++++++++---------
 examples/cifar10/cifar10.h              |  17 ++--
 examples/cifar10/make.sh                |   2 +-
 include/singa/core/common.h             |  12 ++-
 include/singa/core/tensor.h             |   6 +-
 include/singa/model/feed_forward_net.h  |   9 +-
 include/singa/model/initializer.h       |  27 +++++-
 include/singa/model/loss.h              |  12 +--
 include/singa/model/metric.h            |   2 +-
 include/singa/model/optimizer.h         |  18 ++--
 include/singa/utils/string.h            |  11 +++
 src/core/tensor/math_kernel.cu          |  22 +++++
 src/core/tensor/math_kernel.h           |   2 +
 src/core/tensor/tensor.cc               |  69 +++++++++-----
 src/core/tensor/tensor_math.h           |   5 +
 src/core/tensor/tensor_math_cpp.h       |  14 +++
 src/core/tensor/tensor_math_cuda.h      |   9 ++
 src/model/feed_forward_net.cc           | 105 ++++++++++++++-------
 src/model/layer/cudnn_convolution.cc    |   4 +-
 src/model/layer/cudnn_dropout.cc        |   2 +-
 src/model/layer/dense.cc                |   4 +-
 src/model/loss/mse.cc                   |   5 +-
 src/model/loss/softmax_cross_entropy.cc |  11 ++-
 src/model/metric/accuracy.cc            |   2 +
 src/model/optimizer/optimizer.cc        |  30 +++++-
 src/model/optimizer/sgd.cc              |   1 +
 src/proto/model.proto                   |   5 +
 test/singa/test_cross_entropy.cc        |   8 +-
 test/singa/test_dense.cc                |   2 +-
 test/singa/test_mse.cc                  |   8 +-
 test/singa/test_tensor_math.cc          |  11 ++-
 31 files changed, 408 insertions(+), 163 deletions(-)
----------------------------------------------------------------------


http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/examples/cifar10/alexnet.cc
----------------------------------------------------------------------
diff --git a/examples/cifar10/alexnet.cc b/examples/cifar10/alexnet.cc
index 45d8571..81713f1 100644
--- a/examples/cifar10/alexnet.cc
+++ b/examples/cifar10/alexnet.cc
@@ -28,12 +28,13 @@
 #include "../../src/model/layer/cudnn_convolution.h"
 #include "../../src/model/layer/cudnn_activation.h"
 #include "../../src/model/layer/cudnn_pooling.h"
+#include "../../src/model/layer/cudnn_lrn.h"
 #include "../../src/model/layer/dense.h"
 #include "../../src/model/layer/flatten.h"
 namespace singa {
 
 LayerConf GenConvConf(string name, int nb_filter, int kernel, int stride,
-                      int pad) {
+                      int pad, float std) {
   LayerConf conf;
   conf.set_name(name);
   conf.set_type("CudnnConvolution");
@@ -42,13 +43,23 @@ LayerConf GenConvConf(string name, int nb_filter, int kernel, int stride,
   conv->add_kernel_size(kernel);
   conv->add_stride(stride);
   conv->add_pad(pad);
+  conv->set_bias_term(true);
 
-  FillerConf *weight = conv->mutable_weight_filler();
-  weight->set_type("Xavier");
+  ParamSpec *wspec = conf.add_param();
+  wspec->set_name(name + "_weight");
+  auto wfill = wspec->mutable_filler();
+  wfill->set_type("Gaussian");
+  wfill->set_std(std);
+
+  ParamSpec *bspec = conf.add_param();
+  bspec->set_name(name + "_bias");
+  bspec->set_lr_mult(2);
+//  bspec->set_decay_mult(0);
   return conf;
 }
 
-LayerConf GenPoolingConf(string name, bool max_pool, int kernel, int stride, int pad) {
+LayerConf GenPoolingConf(string name, bool max_pool, int kernel, int stride,
+                         int pad) {
   LayerConf conf;
   conf.set_name(name);
   conf.set_type("CudnnPooling");
@@ -56,8 +67,7 @@ LayerConf GenPoolingConf(string name, bool max_pool, int kernel, int stride, int
   pool->set_kernel_size(kernel);
   pool->set_stride(stride);
   pool->set_pad(pad);
-  if (!max_pool)
-    pool->set_pool(PoolingConf_PoolMethod_AVE);
+  if (!max_pool) pool->set_pool(PoolingConf_PoolMethod_AVE);
   return conf;
 }
 
@@ -68,21 +78,38 @@ LayerConf GenReLUConf(string name) {
   return conf;
 }
 
-LayerConf GenDenseConf(string name, int num_output) {
+LayerConf GenDenseConf(string name, int num_output, float std, float wd) {
   LayerConf conf;
   conf.set_name(name);
   conf.set_type("Dense");
   DenseConf *dense = conf.mutable_dense_conf();
   dense->set_num_output(num_output);
-  FillerConf *weight = dense->mutable_weight_filler();
-  weight->set_type("Xavier");
+  FillerConf *bias = dense->mutable_bias_filler();
+
+  ParamSpec *wspec = conf.add_param();
+  wspec->set_name(name + "_weight");
+  wspec->set_decay_mult(wd);
+
+  auto wfill = wspec->mutable_filler();
+  wfill->set_type("Gaussian");
+  wfill->set_std(std);
+
+  ParamSpec *bspec = conf.add_param();
+  bspec->set_name(name + "_bias");
+  bspec->set_lr_mult(2);
+  bspec->set_decay_mult(0);
+
   return conf;
 }
 
-LayerConf GenSoftmaxConf(string name) {
+LayerConf GenLRNConf(string name) {
   LayerConf conf;
   conf.set_name(name);
-  conf.set_type("CudnnSoftmax");
+  conf.set_type("CudnnLRN");
+  LRNConf *lrn = conf.mutable_lrn_conf();
+  lrn->set_local_size(3);
+  lrn->set_alpha(5e-05);
+  lrn->set_beta(0.75);
   return conf;
 }
 
@@ -92,25 +119,25 @@ LayerConf GenFlattenConf(string name) {
   conf.set_type("Flatten");
   return conf;
 }
-FeedForwardNet CreateNet(Optimizer* opt, Loss<Tensor>* loss, Metric<Tensor>* metric) {
+
+FeedForwardNet CreateNet() {
   FeedForwardNet net;
   Shape s{3, 32, 32};
 
-  net.Add(new CudnnConvolution(), GenConvConf("conv1", 32, 5, 1, 2), &s);
+  net.Add(new CudnnConvolution(), GenConvConf("conv1", 32, 5, 1, 2, 0.0001),
+          &s);
+  net.Add(new CudnnPooling(), GenPoolingConf("pool1", true, 3, 2, 1));
   net.Add(new CudnnActivation(), GenReLUConf("relu1"));
-  net.Add(new CudnnPooling, GenPoolingConf("pool1", true, 3, 2, 0));
-  net.Add(new CudnnConvolution(), GenConvConf("conv2", 32, 5, 1, 2));
+  net.Add(new CudnnLRN(), GenLRNConf("lrn1"));
+  net.Add(new CudnnConvolution(), GenConvConf("conv2", 32, 5, 1, 2, 0.01));
   net.Add(new CudnnActivation(), GenReLUConf("relu2"));
-  net.Add(new CudnnPooling(), GenPoolingConf("pool2", true, 3, 2, 0));
-  net.Add(new CudnnConvolution, GenConvConf("conv3", 64, 5, 1, 2));
+  net.Add(new CudnnPooling(), GenPoolingConf("pool2", false, 3, 2, 1));
+  net.Add(new CudnnLRN(), GenLRNConf("lrn2"));
+  net.Add(new CudnnConvolution, GenConvConf("conv3", 64, 5, 1, 2, 0.01));
   net.Add(new CudnnActivation(), GenReLUConf("relu3"));
-  net.Add(new CudnnConvolution(), GenConvConf("pool3", true, 3, 2, 0));
+  net.Add(new CudnnPooling(), GenPoolingConf("pool3", false, 3, 2, 1));
   net.Add(new Flatten(), GenFlattenConf("flat"));
-  net.Add(new Dense(), GenDenseConf("ip1", 10));
-  OptimizerConf opt_conf;
-  opt_conf.set_momentum(0.9);
-  opt->Setup(opt_conf);
-  net.Compile(true, opt, loss, metric);
+  net.Add(new Dense(), GenDenseConf("ip", 10, 0.01, 250));
   return net;
 }
 
@@ -120,50 +147,65 @@ void Train(float lr, int num_epoch, string data_dir) {
   {
     auto train = data.ReadTrainData();
     size_t nsamples = train.first.shape(0);
-    auto matx = Reshape(train.first, Shape{nsamples, train.first.Size() / nsamples});
-    const auto mean = Average(matx, 0);
-    SubRow(mean, &matx);
-    train_x = Reshape(matx, train.first.shape());
+    auto mtrain =
+        Reshape(train.first, Shape{nsamples, train.first.Size() / nsamples});
+    const Tensor& mean = Average(mtrain, 0);
+    SubRow(mean, &mtrain);
+    train_x = Reshape(mtrain, train.first.shape());
     train_y = train.second;
     auto test = data.ReadTestData();
     nsamples = test.first.shape(0);
-    auto maty = Reshape(test.first, Shape{nsamples, test.first.Size() / nsamples});
-    SubRow(mean, &maty);
-    test_x = Reshape(maty, test.first.shape());
+    auto mtest =
+        Reshape(test.first, Shape{nsamples, test.first.Size() / nsamples});
+    SubRow(mean, &mtest);
+    test_x = Reshape(mtest, test.first.shape());
     test_y = test.second;
   }
-  LOG(ERROR) << "creating net";
+  CHECK_EQ(train_x.shape(0), train_y.shape(0));
+  CHECK_EQ(test_x.shape(0), test_y.shape(0));
+  LOG(INFO) << "Training samples = " << train_y.shape(0)
+            << " Test samples = " << test_y.shape(0);
+  auto net = CreateNet();
+  SGD sgd;
+  OptimizerConf opt_conf;
+  opt_conf.set_momentum(0.9);
+  auto reg = opt_conf.mutable_regularizer();
+  reg->set_coefficient(0.004);
+  sgd.Setup(opt_conf);
+  sgd.SetLearningRateGenerator([lr](int step) {
+    if (step <= 120)
+      return 0.001;
+    else if (step <= 130)
+      return 0.0001;
+    else if (step <= 140)
+      return 0.00001;
+  });
+
   SoftmaxCrossEntropy loss;
   Accuracy acc;
-  SGD sgd;
-  sgd.SetLearningRateGenerator([lr](int step) {return lr;});
-  auto net = CreateNet(&sgd, &loss, &acc);
+  net.Compile(true, &sgd, &loss, &acc);
 
   auto cuda = std::make_shared<CudaGPU>();
   net.ToDevice(cuda);
-
   train_x.ToDevice(cuda);
   train_y.ToDevice(cuda);
-  net.Train(50, num_epoch, train_x, train_y); // test_x, test_y);
+  test_x.ToDevice(cuda);
+  test_y.ToDevice(cuda);
+  net.Train(100, num_epoch, train_x, train_y, test_x, test_y);
 }
-
-
 }
 
-int main(int argc, char** argv) {
+int main(int argc, char **argv) {
   singa::InitChannel(nullptr);
   int pos = singa::ArgPos(argc, argv, "-epoch");
-  int nEpoch = 5;
-  if (pos != -1)
-    nEpoch = atoi(argv[pos + 1]);
+  int nEpoch = 1;
+  if (pos != -1) nEpoch = atoi(argv[pos + 1]);
   pos = singa::ArgPos(argc, argv, "-lr");
-  float lr = 0.01;
-  if (pos != -1)
-    lr = atof(argv[pos + 1]);
+  float lr = 0.001;
+  if (pos != -1) lr = atof(argv[pos + 1]);
   pos = singa::ArgPos(argc, argv, "-data");
   string data = "cifar-10-batches-bin";
-  if (pos != -1)
-    data = argv[pos + 1];
+  if (pos != -1) data = argv[pos + 1];
 
   LOG(INFO) << "Start training";
   singa::Train(lr, nEpoch, data);

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/examples/cifar10/cifar10.h
----------------------------------------------------------------------
diff --git a/examples/cifar10/cifar10.h b/examples/cifar10/cifar10.h
index 261c048..5af54e2 100644
--- a/examples/cifar10/cifar10.h
+++ b/examples/cifar10/cifar10.h
@@ -29,25 +29,24 @@ namespace singa {
 class Cifar10 {
  public:
   /// 'dir_path': path to the folder including the *.bin files
-  Cifar10(string dir_path, bool normalize = true)
-      : dir_path_(dir_path), normalize_(normalize) {}
+  Cifar10(string dir_path, bool normalize = true) : dir_path_(dir_path) {}
 
   /// read all training data into an image Tensor and a label Tensor
-  const std::pair<Tensor, Tensor> ReadTrainData(bool shuffle = false);
+  const std::pair<Tensor, Tensor> ReadTrainData();
   /// read all test data into an image Tensor and a label Tensor
   const std::pair<Tensor, Tensor> ReadTestData();
   /// read data from one file into an image Tensor and a label Tensor
-  const std::pair<Tensor, Tensor> ReadFile(string file, bool shuffle = false);
+  const std::pair<Tensor, Tensor> ReadFile(string file);
 
   void ReadImage(std::ifstream* file, int* label, char* buffer);
+
  private:
   const size_t kImageSize = 32;
   const size_t kImageVol = 3072;
   const size_t kBatchSize = 10000;
-  const size_t kTrainFiles = 1;
+  const size_t kTrainFiles = 5;
 
   string dir_path_;
-  bool normalize_;
 };
 
 void Cifar10::ReadImage(std::ifstream* file, int* label, char* buffer) {
@@ -57,7 +56,7 @@ void Cifar10::ReadImage(std::ifstream* file, int* label, char* buffer) {
   file->read(buffer, kImageVol);
   return;
 }
-const std::pair<Tensor, Tensor> Cifar10::ReadFile(string file, bool shuffle) {
+const std::pair<Tensor, Tensor> Cifar10::ReadFile(string file) {
   Tensor images(Shape{kBatchSize, 3, kImageSize, kImageSize});
   Tensor labels(Shape{kBatchSize}, kInt);
   if (dir_path_.back() != '/') dir_path_.push_back('/');
@@ -73,7 +72,7 @@ const std::pair<Tensor, Tensor> Cifar10::ReadFile(string file, bool shuffle) {
     // LOG(INFO) << "reading " << itemid << "-th image";
     ReadImage(&data_file, &label, image);
     for (int i = 0; i < kImageVol; i++)
-      float_image[i] = static_cast<float>(static_cast<int>(image[i]));
+      float_image[i] = static_cast<float>(static_cast<uint8_t>(image[i]));
     images.CopyDataFromHostPtr(float_image, kImageVol, itemid * kImageVol);
     tmplabels[itemid] = label;
   }
@@ -81,7 +80,7 @@ const std::pair<Tensor, Tensor> Cifar10::ReadFile(string file, bool shuffle) {
   return std::make_pair(images, labels);
 }
 
-const std::pair<Tensor, Tensor> Cifar10::ReadTrainData(bool shuffle) {
+const std::pair<Tensor, Tensor> Cifar10::ReadTrainData() {
   Tensor images(Shape{kBatchSize * kTrainFiles, 3, kImageSize, kImageSize});
   Tensor labels(Shape{kBatchSize * kTrainFiles}, kInt);
   for (int fileid = 0; fileid < kTrainFiles; ++fileid) {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/examples/cifar10/make.sh
----------------------------------------------------------------------
diff --git a/examples/cifar10/make.sh b/examples/cifar10/make.sh
index 17e4b39..5a41612 100755
--- a/examples/cifar10/make.sh
+++ b/examples/cifar10/make.sh
@@ -1 +1 @@
-g++ -g --std=c++11 alexnet.cc -o alexnet -I../../include -I../../build/include -I/home/wangwei/local/cudnn4/include -I/home/wangwei/local/include -I/usr/local/cuda/include/ -I../../lib/cnmem/include -L../../build/lib/ -lsinga_core -lsinga_model  -lsinga_utils -lcudart -lcublas -lcurand -lcudnn -L/usr/local/cuda/lib64 -L/home/wangwei/local/cudnn4/lib64 ../../build/lib/libproto.a -lprotobuf
+g++ -g --std=c++11 alexnet.cc -o alexnet -I../../include -I../../build/include -I/home/wangwei/local/cudnn5/include -I/home/wangwei/local/include -I/usr/local/cuda/include/ -I../../lib/cnmem/include -L../../build/lib/ -lsinga_core -lsinga_model  -lsinga_utils -lcudart -lcublas -lcurand -lcudnn -L/home/wangwei/local/cudnn5/lib64 -L/usr/local/cuda/lib64 ../../build/lib/libproto.a -lprotobuf

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/include/singa/core/common.h
----------------------------------------------------------------------
diff --git a/include/singa/core/common.h b/include/singa/core/common.h
index 691d7d4..0a52a05 100644
--- a/include/singa/core/common.h
+++ b/include/singa/core/common.h
@@ -49,19 +49,20 @@ class Block {
  public:
   Block(void* ptr, size_t size, size_t offset = 0)
       : data_(ptr), size_(size), offset_(offset) {
-    ref_count_ = 1; //std::make_shared<std::atomic<int>>(1);
+    ref_count_ = 1;  // std::make_shared<std::atomic<int>>(1);
   }
-//  Block(void* ptr, size_t size, size_t offset, std::shared_ptr<atomic<int>> ref)
-//      : data_(ptr), size_(size), offset_(offset), ref_count_(ref) {}
+  // Disabled as it is not used currently.
+  // Block(void* ptr, size_t size, size_t offset, std::shared_ptr<atomic<int>>
+  //  ref) : data_(ptr), size_(size), offset_(offset), ref_count_(ref) {}
   void* mutable_data() const { return static_cast<char*>(data_) + offset_; }
   const void* data() const { return static_cast<char*>(data_) + offset_; }
   size_t size() const { return size_; }
   size_t offset() const { return offset_; }
   int IncRefCount() {
-    return ++ref_count_;  //(*ref_count_)++;
+    return ++ref_count_;  // Note do not use ref_count_++;
   }
   int DecRefCount() {
-    return --ref_count_; // (*ref_count_)--;
+    return --ref_count_;
   }
   int ref_count() const { return ref_count_.load(); }
 
@@ -70,6 +71,7 @@ class Block {
   void* data_ = nullptr;
   size_t size_ = 0;
   size_t offset_ = 0;
+  // Disabled as it is not used currently.
   // std::shared_ptr<std::atomic<int>> ref_count_ = nullptr;
   std::atomic<int> ref_count_;
 };

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/include/singa/core/tensor.h
----------------------------------------------------------------------
diff --git a/include/singa/core/tensor.h b/include/singa/core/tensor.h
index 3b496d9..15e7b7f 100644
--- a/include/singa/core/tensor.h
+++ b/include/singa/core/tensor.h
@@ -128,7 +128,8 @@ class Tensor {
   template <typename SType>
   void SetValue(const SType x);
 
-  /// For init the tensor values, copy 'num' elements.
+  /// For init the tensor values, copy 'num' elements from 'src' to the internal
+  /// memory with 'offset' (elements).
   template <typename SType>
   void CopyDataFromHostPtr(const SType *src, const size_t num,
                            const size_t offset = 0);
@@ -180,6 +181,9 @@ class Tensor {
   template <typename SType>
   Tensor &operator/=(const SType x);
 
+  /// Return average L1 norm
+  float L1() const;
+  /// Return average L2 norm
   float L2() const;
 
  protected:

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/include/singa/model/feed_forward_net.h
----------------------------------------------------------------------
diff --git a/include/singa/model/feed_forward_net.h b/include/singa/model/feed_forward_net.h
index 9beeb7a..87d2179 100644
--- a/include/singa/model/feed_forward_net.h
+++ b/include/singa/model/feed_forward_net.h
@@ -29,12 +29,14 @@ namespace singa {
 class FeedForwardNet {
  public:
   FeedForwardNet() = default;
+  /// Delete all layers.
   ~FeedForwardNet();
 
   /// Add a layer with the assumption that
   /// 1. this function is called in correct order, i.e., the layers are added
   ///    following the topological order.
   /// 2. this layer has already been setup (Setup function is called outside).
+  /// The layer will be freed in the destructor of FeedForwardNet.
   Layer* Add(Layer* layer);
 
   // TODO(wangwei) add ConcatenateLayer and SliceLayer
@@ -47,7 +49,9 @@ class FeedForwardNet {
   /// for calling Setup().
   Layer* Add(const LayerConf& conf, const Shape* sample_shape = nullptr);
 
-  Layer* Add(Layer* layer, const LayerConf& conf, const Shape* sample_shape = nullptr);
+  /// Add a layer, and call its Setup function.
+  Layer* Add(Layer* layer, const LayerConf& conf,
+             const Shape* sample_shape = nullptr);
   /// Set some fields used for training and evaluating the neural net.
   /// If the neural net is constructed for evaluation only, then 'opt' is not
   /// necessary; But for training, both 'opt' and 'loss' are necessary.
@@ -72,7 +76,8 @@ class FeedForwardNet {
   void Train(size_t batchsize, int nb_epoch, const Tensor& x, const Tensor& y,
              const Tensor& val_x, const Tensor& val_y);
   /// Train the neural net over one batch of training data.
-  const std::pair<float, float> TrainOnBatch(const Tensor& x, const Tensor& y);
+  const std::pair<float, float> TrainOnBatch(int epoch, const Tensor& x,
+                                             const Tensor& y);
 
   /// Evaluate the neural net with given data.
   /// Returns one tensor for loss values and one tensor for metric values;

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/include/singa/model/initializer.h
----------------------------------------------------------------------
diff --git a/include/singa/model/initializer.h b/include/singa/model/initializer.h
index 302fc97..5692185 100644
--- a/include/singa/model/initializer.h
+++ b/include/singa/model/initializer.h
@@ -21,8 +21,8 @@
 #include <string>
 #include "singa/core/tensor.h"
 #include "singa/proto/model.pb.h"
+#include "singa/utils/string.h"
 namespace singa {
-namespace init {
 /// Base class for initializing parameter values.
 using InitializerConf = FillerConf;
 class Initializer {
@@ -40,6 +40,7 @@ class Initializer {
   virtual void Fill(Tensor* t) = 0;
 };
 
+namespace init {
 class Constant : public Initializer {
 public:
   Constant() = default;
@@ -76,7 +77,7 @@ public:
   void Fill(Tensor* t) override { singa::Gaussian(mean_, std_, t); }
 
  private:
-  float mean_ = 0, std_ = 0.01;
+  float mean_ = 0, std_ = 1;
 };
 
 /// Ref: [Bengio and Glorot 2010] Understanding the difficulty of training deep
@@ -86,6 +87,7 @@ public:
   void Fill(Tensor* t) override {
     CHECK_EQ(t->nDim(), 2u);
     float scale = sqrt(6.0f / (t->shape(0) + t->shape(1)));
+    LOG(INFO) << "xavier scale " << scale;
     singa::Uniform(-scale, scale, t);
   }
 };
@@ -100,6 +102,27 @@ class MSRA : public Initializer {
     singa::Gaussian(0.0f, std, t);
   }
 };
+
 }  // namespace init
+
+/// TODO(wangwei) create the initializers from factory like that for Layer.
+std::shared_ptr<Initializer> CreateInitializer(const InitializerConf& conf) {
+  std::shared_ptr<Initializer> init;
+  if (ToLowerCase(conf.type()) == "constant") {
+    init = std::make_shared<init::Constant>();
+  } else if (ToLowerCase(conf.type()) == "uniform") {
+    init = std::make_shared<init::Uniform>();
+  } else if (ToLowerCase(conf.type()) == "gaussian") {
+    init = std::make_shared<init::Gaussian>();
+  } else if (ToLowerCase(conf.type()) == "xavier") {
+    init = std::make_shared<init::Xavier>();
+  } else if (ToLowerCase(conf.type()) == "msra") {
+    init = std::make_shared<init::MSRA>();
+  } else {
+    LOG(FATAL) << "Unknown initialization type: " << conf.type();
+  }
+  init->Setup(conf);
+  return init;
+}
 }  // namespace singa
 #endif  // SINGA_MODEL_INITIALIZER_H_

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/include/singa/model/loss.h
----------------------------------------------------------------------
diff --git a/include/singa/model/loss.h b/include/singa/model/loss.h
index 41ec701..f400768 100644
--- a/include/singa/model/loss.h
+++ b/include/singa/model/loss.h
@@ -43,13 +43,13 @@ class Loss {
 
   /// Compute the loss values for each sample/instance given the prediction
   /// and the target.
-  virtual Tensor Forward(const Tensor& prediction, const T& target) = 0;
+  virtual Tensor Forward(int flag, const Tensor& prediction, const T& target) = 0;
 
   /// Average loss values for all samples in the mini-batch
   /// It calls Forward() internally. The calling pattern should be
   /// [Evaluate|Forward] Backward.
-  float Evaluate(const Tensor& prediction, const T& target) {
-    Tensor loss = Forward(prediction, target);
+  float Evaluate(int flag, const Tensor& prediction, const T& target) {
+    Tensor loss = Forward(flag, prediction, target);
     loss.ToHost();
     return Sum<float>(loss) / (1.0f * loss.Size());
   }
@@ -68,7 +68,7 @@ class MSE : public Loss<Tensor> {
   /// and the target, which is 0.5/||prediction-target||^2
   /// Users can call Average(const Tensor&) to get the average
   /// loss value over all samples in the batch.
-  Tensor Forward(const Tensor& prediction, const Tensor& target) override;
+  Tensor Forward(int flag, const Tensor& prediction, const Tensor& target) override;
 
   /// Compute the gradients of the loss values w.r.t. the prediction,
   /// which is (prediction-target)/batchsize
@@ -90,7 +90,7 @@ class SoftmaxCrossEntropy : public Loss<Tensor> {
   /// from Softmax(prediction).
   /// Users can call Average(const Tensor&) to get the average
   /// loss value over all samples in the batch.
-  Tensor Forward(const Tensor& prediction, const Tensor& target) override;
+  Tensor Forward(int flag, const Tensor& prediction, const Tensor& target) override;
 
   /// Compute the gradients of the loss values w.r.t. the prediction,
   /// which is: p[idx] - 1 if idx is the truth category's index; else,
@@ -106,5 +106,3 @@ class SoftmaxCrossEntropy : public Loss<Tensor> {
 }  // namespace singa
 
 #endif  // SINGA_MODEL_LOSS_H_
-
-

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/include/singa/model/metric.h
----------------------------------------------------------------------
diff --git a/include/singa/model/metric.h b/include/singa/model/metric.h
index d013fa4..b100435 100644
--- a/include/singa/model/metric.h
+++ b/include/singa/model/metric.h
@@ -48,7 +48,7 @@ class Metric {
 
   /// Comptue the metric value averaged over all samples (in a batch)
   float Evaluate(const Tensor& prediction, const T& target) {
-    const Tensor& metric = Forward(prediction, target);
+    const Tensor metric = Forward(prediction, target);
     return Sum<float>(metric) / (1.0f * metric.Size());
   }
 };

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/include/singa/model/optimizer.h
----------------------------------------------------------------------
diff --git a/include/singa/model/optimizer.h b/include/singa/model/optimizer.h
index a268126..2ec68fe 100644
--- a/include/singa/model/optimizer.h
+++ b/include/singa/model/optimizer.h
@@ -41,7 +41,7 @@ class Regularizer;
 class Optimizer {
  public:
   Optimizer() = default;
-
+  virtual ~Optimizer();
   /// Setup the optimzier using configurations from serialized string (for
   /// binding languages).
   void Setup(const string& str) {
@@ -51,7 +51,7 @@ class Optimizer {
   }
 
   /// Setup the meta fields of the optimizer
-  virtual void Setup(const OptimizerConf& conf) {}
+  virtual void Setup(const OptimizerConf& conf);
   /// Register the parameter, e.g., create Constraint and Regularizers.
   /// If there is no constraint or regularizer, then no need to register the
   /// parameter.
@@ -76,15 +76,21 @@ class Optimizer {
   void SetLearningRateGenerator(function<float(int)> func) {
     learning_rate_generator_ = func;
   }
-  /// Since Optimizer base layer has pure virtual function, a virtual
-  /// deconstructor is needed.
-  virtual ~Optimizer() = default;
+  float GetLearningRate(int step) {
+    if (learning_rate_generator_)
+      return learning_rate_generator_(step);
+    else
+      return 0;
+  }
 
  protected:
   function<float(int)> learning_rate_generator_;
   std::unordered_map<std::string, float> learning_rate_multplier_;
+  std::unordered_map<std::string, float> weight_decay_multplier_;
   std::unordered_map<std::string, Constraint*> constraints_;
   std::unordered_map<std::string, Regularizer*> regularizers_;
+  Constraint* constraint_ = nullptr;
+  Regularizer* regularizer_ = nullptr;
 };
 
 /// Apply constraints for parameters (gradient).
@@ -141,7 +147,7 @@ class Regularizer {
   /// e.g., clip each gradient if it is too large w.r.t the threshold,
   /// \ref
   /// https://www.reddit.com/r/MachineLearning/comments/31b6x8/gradient_clipping_rnns/
-  void Apply(int step, Tensor* grad, Tensor* value);
+  void Apply(int step, Tensor* grad, Tensor* value, float scale = 1.0f);
   /// Apply the regularizer for multiple parameter objects together.
   /// \ref https://github.com/Lasagne/Lasagne/blob/master/lasagne/updates.py
   void Apply(int step, const vector<Tensor*>& grads,

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/include/singa/utils/string.h
----------------------------------------------------------------------
diff --git a/include/singa/utils/string.h b/include/singa/utils/string.h
index cbfb28b..b4c7c24 100644
--- a/include/singa/utils/string.h
+++ b/include/singa/utils/string.h
@@ -51,6 +51,17 @@ inline int ArgPos(int argc, char** arglist, const char* arg) {
   return -1;
 }
 
+template<typename T>
+inline std::string VecToStr(const std::vector<T> & in) {
+  std::string out = "(";
+
+  for (auto x : in) {
+    out += std::to_string(x) + ", ";
+  }
+  out += ")";
+  return out;
+}
+
 /**
  * Tokenize a string.
  *

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/core/tensor/math_kernel.cu
----------------------------------------------------------------------
diff --git a/src/core/tensor/math_kernel.cu b/src/core/tensor/math_kernel.cu
index 4135ab8..13005af 100644
--- a/src/core/tensor/math_kernel.cu
+++ b/src/core/tensor/math_kernel.cu
@@ -265,6 +265,19 @@ __global__ void KernelLT(const size_t num, const float *in, const float x,
     out[idx] = in[idx] < x ? 1.0f : 0.0f;
   }
 }
+
+__global__ void KernelRowMax(const size_t nrow, const size_t ncol, const float *inPtr,
+    float *outPtr) {
+  for (size_t idx = blockIdx.x * blockDim.x + threadIdx.x; idx < nrow;
+       idx += blockDim.x * gridDim.x) {
+    int offset = idx * ncol;
+    float maxval = inPtr[offset];
+    for (size_t k = 1; k < ncol; k++) {
+      maxval = max(maxval, inPtr[offset + k]);
+    }
+    outPtr[idx] = maxval;
+  }
+}
 __global__ void KernelComputeCrossEntropy(const size_t batchsize,
                                           const size_t dim, const float *p,
                                           const int *t, float *loss) {
@@ -286,6 +299,9 @@ __global__ void KernelSoftmaxCrossEntropyBwd(const size_t batchsize,
     grad[pos] = p[pos] - 1.0f;  // TODO(wangwei) Consider p and grad are diff
   }
 }
+
+
+
 // ********************************
 // Functions call kernels
 // ********************************
@@ -421,6 +437,12 @@ void SoftmaxCrossEntropyBwd(size_t batchsize, const size_t dim, const float *p,
   KernelSoftmaxCrossEntropyBwd <<<ceil(batchsize / CU1DBLOCKF), CU1DBLOCKF>>>
       (batchsize, dim, p, t, grad);
 }
+
+void RowMax(const size_t nrow, const size_t ncol, const float *inPtr,
+    float *outPtr, cudaStream_t stream) {
+  KernelRowMax <<<ceil(nrow / CU1DBLOCKF), CU1DBLOCKF>>>(nrow, ncol, inPtr, outPtr);
+}
+
 /*
 void square_grad(int n, const float *in, float *out, cudaStream_t s) {
   kernel_square_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (in, out, n);

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/core/tensor/math_kernel.h
----------------------------------------------------------------------
diff --git a/src/core/tensor/math_kernel.h b/src/core/tensor/math_kernel.h
index d4087e5..63b0d82 100644
--- a/src/core/tensor/math_kernel.h
+++ b/src/core/tensor/math_kernel.h
@@ -98,6 +98,8 @@ void SoftmaxCrossEntropyBwd(const size_t batchsize, const size_t dim,
                             const float *p, const int *t, float *grad,
                             cudaStream_t stream);
 
+void RowMax(const size_t nrow, const size_t ncol, const float *inPtr,
+    float *outPtr, cudaStream_t stream);
 }  // cuda
 
 }  // namespace singa

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/core/tensor/tensor.cc
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor.cc b/src/core/tensor/tensor.cc
index 898cdc6..b07a23c 100644
--- a/src/core/tensor/tensor.cc
+++ b/src/core/tensor/tensor.cc
@@ -42,7 +42,8 @@ Tensor::Tensor(Shape &&shape, DataType dtype)
   device_ = defaultDevice;
   block_ = device_->NewBlock(Product(shape_) * SizeOf(data_type_));
 }
-Tensor::Tensor(const Shape &shape, std::shared_ptr<Device> device, DataType dtype)
+Tensor::Tensor(const Shape &shape, std::shared_ptr<Device> device,
+               DataType dtype)
     : data_type_(dtype), device_(device), shape_(shape) {
   block_ = device_->NewBlock(Product(shape_) * SizeOf(data_type_));
 }
@@ -68,11 +69,10 @@ Tensor::Tensor(Tensor &&in)
   in.block_ = nullptr;
 }
 
-void Tensor::SetBlock(Block* block) {
+void Tensor::SetBlock(Block *block) {
   LOG(WARNING) << "Pls avoid using this function, which may have side-effect.";
   if (block_ != nullptr)
-    if (block_->DecRefCount())
-      device_->FreeBlock(block_);
+    if (block_->DecRefCount()) device_->FreeBlock(block_);
   block_ = block;
 }
 
@@ -118,8 +118,7 @@ void Tensor::ToDevice(std::shared_ptr<Device> dst) {
   // TODO(wangwei) the comparison is very strict. May compare against device ID?
   if (device_ != dst) {
     Tensor tmp(shape_, dst, data_type_);
-    if (block_ != nullptr && Size())
-      tmp.CopyData(*this);
+    if (block_ != nullptr && Size()) tmp.CopyData(*this);
     if (block_ != nullptr && block_->DecRefCount() == 0)
       device_->FreeBlock(block_);
     block_ = tmp.block_;
@@ -132,13 +131,13 @@ void Tensor::ToHost() { ToDevice(device_->host()); }
 
 template <typename DType>
 void Tensor::CopyDataFromHostPtr(const DType *src, const size_t num,
-    const size_t offset) {
+                                 const size_t offset) {
   CHECK_EQ(sizeof(DType), SizeOf(data_type_))
       << "data_type is " << DataType_Name(data_type_)
       << " user given type is of size " << sizeof(DType);
   if (src != nullptr) {
     device_->CopyDataFromHostPtr(block(), src, sizeof(DType) * num,
-        sizeof(DType) * offset);
+                                 sizeof(DType) * offset);
   } else {
     LOG(WARNING) << "Copy data from null host ptr";
   }
@@ -161,8 +160,7 @@ void Tensor::CopyData(const Tensor &src) {
 }
 
 Tensor Tensor::Clone(std::shared_ptr<Device> device) const {
-  if (device == nullptr)
-    device = device_;
+  if (device == nullptr) device = device_;
   Tensor t(shape_, device_, data_type_);
   t.transpose_ = transpose_;
   t.CopyData(*this);
@@ -244,8 +242,6 @@ GenUnaryScalarArgMemberFn(operator+=, Add);
 GenUnaryScalarArgMemberFn(operator*=, EltwiseMult);
 GenUnaryScalarArgMemberFn(operator/=, Div);
 
-
-
 // ====================Tensor Operations=======================================
 void CopyDataToFrom(Tensor *dst, const Tensor &src, const size_t num,
                     const size_t dst_offset, const size_t src_offset) {
@@ -336,6 +332,18 @@ void CopyDataToFrom(Tensor *dst, const Tensor &src, const size_t num,
   } while (0)
 
 // =============Element-wise operations====================================
+float Tensor::L1() const {
+  float nrm = 0.0f;
+  TYPE_LANG_SWITCH(data_type_, DType, device_->lang(), Lang, {
+    device_->Exec([&nrm, this](Context *ctx) {
+      DType ret;
+      Asum<DType, Lang>(this->Size(), this->block(), &ret, ctx);
+      nrm = TypeCast<DType, float>(ret);
+    }, {this->block()}, {});
+  });
+  return nrm / Size();
+}
+
 /// L2 norm, Do not use Nrm2 (name conflict).
 float Tensor::L2() const {
   float nrm = 0.0f;
@@ -346,8 +354,10 @@ float Tensor::L2() const {
       nrm = TypeCast<DType, float>(ret);
     }, {this->block()}, {});
   });
-  return nrm;
+  return nrm / Size();
 }
+
+
 template <typename SType>
 void Tensor::SetValue(const SType x) {
   CHECK_EQ(sizeof(SType), SizeOf(data_type_));
@@ -525,18 +535,35 @@ Tensor SoftMax(const Tensor &in) {
   return out;
 }
 
+Tensor RowMax(const Tensor &in) {
+  Tensor ret({in.shape(0)}, in.device(), in.data_type());
+  TYPE_LANG_SWITCH(in.data_type(), DType, in.device()->lang(), Lang, {
+    in.device()->Exec([in, ret](Context *ctx) {
+      size_t nrow = 1;
+      if (in.nDim() > 1) nrow = in.shape(0);
+      size_t ncol = in.Size() / nrow;
+      RowMax<DType, Lang>(nrow, ncol, in.block(), ret.block(), ctx);
+    }, {in.block()}, {ret.block()});
+  });
+  return ret;
+}
+
 void SoftMax(const Tensor &in, Tensor *out) {
   CHECK_LE(in.nDim(), 2u);
-  Exp(in, out);
+  out->CopyData(in);
   size_t nrow = 1, ncol = in.Size(), size = ncol;
   if (in.nDim() == 2u) {
     nrow = in.shape(0);
     ncol = size / nrow;
     out->Reshape(Shape{nrow, ncol});
   }
-  Tensor sum(Shape{nrow}, in.device(), in.data_type());
-  SumColumns(*out, &sum);
-  DivColumn(sum, out);
+  Tensor tmp = RowMax(*out);
+  SubColumn(tmp, out);
+  Exp(*out, out);
+
+  SumColumns(*out, &tmp);
+  DivColumn(tmp, out);
+  out->Reshape(in.shape());
 }
 
 void AddColumn(const Tensor &v, Tensor *M) { AddColumn(1, 1, v, M); }
@@ -582,8 +609,8 @@ void AddRow(const SType alpha, const SType beta, const Tensor &v, Tensor *M) {
     Mult(alpha, one, vmat, beta, M);
   }
 }
-template
-void AddRow(const float alpha, const float beta, const Tensor &v, Tensor *M);
+template void AddRow(const float alpha, const float beta, const Tensor &v,
+                     Tensor *M);
 
 /// Divide column 'v' by each column of matrix M; write results into 'out'
 void DivColumn(const Tensor &v, Tensor *M) {
@@ -699,7 +726,7 @@ void MultRow(const Tensor &v, Tensor *M) {
   });
 }
 
-Tensor SliceRows(const Tensor& in, const size_t start, const size_t end) {
+Tensor SliceRows(const Tensor &in, const size_t start, const size_t end) {
   LOG(FATAL) << "Tensor::SliceRows is not implemented";
   Tensor ret;
   /*
@@ -788,6 +815,7 @@ void Gaussian(const SType mean, const SType std, Tensor *out) {
 template void Gaussian<float>(const float mean, const float std, Tensor *out);
 
 // ================Blas operations============================================
+
 template <typename SType>
 void Axpy(const SType alpha, const Tensor &in, Tensor *out) {
   TYPE_LANG_SWITCH(in.data_type(), DType, in.device()->lang(), Lang, {
@@ -869,5 +897,4 @@ void SoftmaxCrossEntropyBwd(const Tensor &t, Tensor *p) {
   });
 }
 
-
 }  // namespace singa

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/core/tensor/tensor_math.h
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor_math.h b/src/core/tensor/tensor_math.h
index 57ccb88..7732dd2 100644
--- a/src/core/tensor/tensor_math.h
+++ b/src/core/tensor/tensor_math.h
@@ -339,6 +339,11 @@ void SoftmaxCrossEntropyBwd(const size_t batchsize, const size_t dim,
   LOG(FATAL) << "Not Implemented";
 }
 
+template <typename DType, typename Lang>
+void RowMax(const size_t nrow, const size_t ncol, const Block *in,
+    const Block *ret, Context* ctx) {
+  LOG(FATAL) << "Not Implemented";
+}
 // **************************************
 // Matrix functions
 // **************************************

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/core/tensor/tensor_math_cpp.h
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor_math_cpp.h b/src/core/tensor/tensor_math_cpp.h
index 4717b5f..3e0c8ad 100644
--- a/src/core/tensor/tensor_math_cpp.h
+++ b/src/core/tensor/tensor_math_cpp.h
@@ -549,6 +549,20 @@ void SoftmaxCrossEntropyBwd<float, lang::Cpp>(const size_t batchsize,
   }
 }
 
+template <>
+void RowMax<float, lang::Cpp>(const size_t nrow, const size_t ncol,
+                              const Block *in, const Block *out, Context *ctx) {
+  const float *inPtr = static_cast<const float *>(in->data());
+  float *outPtr = static_cast<float *>(out->mutable_data());
+  for (size_t r = 0; r < nrow; r++) {
+    int offset = r * ncol;
+    float maxval = inPtr[offset];
+    for (size_t c = 1; c < ncol; c++)
+      maxval = std::max(maxval, inPtr[offset + c]);
+    outPtr[r] = maxval;
+  }
+}
+
 // =========Matrix operations ================================================
 /*
 template <>

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/core/tensor/tensor_math_cuda.h
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor_math_cuda.h b/src/core/tensor/tensor_math_cuda.h
index 67ee861..43bfa1b 100644
--- a/src/core/tensor/tensor_math_cuda.h
+++ b/src/core/tensor/tensor_math_cuda.h
@@ -421,6 +421,15 @@ void SoftmaxCrossEntropyBwd<float, lang::Cuda>(const size_t batchsize,
   cuda::SoftmaxCrossEntropyBwd(batchsize, dim, pPtr, tPtr, gradPtr,
                                ctx->stream);
 }
+
+template <>
+void RowMax<float, lang::Cuda>(const size_t nrow, const size_t ncol,
+                               const Block* in, const Block* out,
+                               Context* ctx) {
+  const float* inPtr = static_cast<const float*>(in->data());
+  float* outPtr = static_cast<float*>(out->mutable_data());
+  cuda::RowMax(nrow, ncol, inPtr, outPtr, ctx->stream);
+}
 }  // namespace singa
 
 #endif  // USE_CUDA

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/model/feed_forward_net.cc
----------------------------------------------------------------------
diff --git a/src/model/feed_forward_net.cc b/src/model/feed_forward_net.cc
index a24d36a..30d030a 100644
--- a/src/model/feed_forward_net.cc
+++ b/src/model/feed_forward_net.cc
@@ -1,22 +1,26 @@
-/**
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements.  See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership.  The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License.  You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
+/************************************************************
+*
+* Licensed to the Apache Software Foundation (ASF) under one
+* or more contributor license agreements.  See the NOTICE file
+* distributed with this work for additional information
+* regarding copyright ownership.  The ASF licenses this file
+* to you under the Apache License, Version 2.0 (the
+* "License"); you may not use this file except in compliance
+* with the License.  You may obtain a copy of the License at
+*
+*   http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing,
+* software distributed under the License is distributed on an
+* "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+* KIND, either express or implied.  See the License for the
+* specific language governing permissions and limitations
+* under the License.
+*
+*************************************************************/
 
 #include "singa/model/feed_forward_net.h"
+#include "singa/model/initializer.h"
 #include "singa/utils/logging.h"
 #include "singa/utils/channel.h"
 namespace singa {
@@ -37,12 +41,15 @@ Layer* FeedForwardNet::Add(const LayerConf& conf, const Shape* sample_shape) {
   return layer;
 }
 
-Layer* FeedForwardNet::Add(Layer* layer, const LayerConf& conf, const Shape* sample_shape) {
+Layer* FeedForwardNet::Add(Layer* layer, const LayerConf& conf,
+                           const Shape* sample_shape) {
+  CHECK(conf.has_name()) << "Must set layer name";
   if (sample_shape == nullptr)
     layer->Setup(layers_.back()->GetOutputSampleShape(), conf);
   else
     layer->Setup(*sample_shape, conf);
   Add(layer);
+  LOG(INFO) << layer->name() << VecToStr(layer->GetOutputSampleShape());
   return layer;
 }
 
@@ -75,12 +82,19 @@ void FeedForwardNet::Compile(bool shuffle, Optimizer* opt, Loss<Tensor>* loss,
   opt_ = opt;
   loss_ = loss;
   metric_ = metric;
-  // init params and register them to sgd
+  const auto specs = GetParamSpecs();
+  const auto params = GetParamValues();
+  CHECK_EQ(specs.size(), params.size());
+  for (size_t k = 0; k < specs.size(); k++) {
+    opt_->Register(specs[k].name(), specs[k]);
+    auto init = CreateInitializer(specs[k].filler());
+    init->Fill(params[k]);
+    LOG(INFO) << specs[k].name() << " : " << params[k]->L1();
+  }
 }
 
 void FeedForwardNet::ToDevice(std::shared_ptr<Device> device) {
-  for (auto layer: layers_)
-    layer->ToDevice(device);
+  for (auto layer : layers_) layer->ToDevice(device);
   /*
   opt_->ToDevice(device);
   loss_->ToDevice(device);
@@ -129,21 +143,27 @@ void FeedForwardNet::Train(size_t batchsize, int nb_epoch, const Tensor& x,
 void FeedForwardNet::Train(size_t batchsize, int nb_epoch, const Tensor& x,
                            const Tensor& y, const Tensor& val_x,
                            const Tensor& val_y) {
-  InitNetParams();
   CHECK_EQ(x.shape(0), y.shape(0)) << "Diff num of sampels in x and y";
   int num_extra_samples = x.shape(0) % batchsize;
   if (num_extra_samples != 0)
-    LOG(WARNING) << "The last " << num_extra_samples << " would not be used";
+    LOG(WARNING) << "Pls set batchsize to make num_total_samples "
+      << "% batchsize == 0. Otherwise, the last " << num_extra_samples
+      << " samples would not be used";
   Channel* train_ch = GetChannel("train_perf");
   train_ch->EnableDestStderr(true);
   Channel* val_ch = GetChannel("val_perf");
+  val_ch->EnableDestStderr(true);
+  std::vector<size_t> index;
+  for (size_t i = 0; i < x.shape(0) / batchsize; i++) index.push_back(i);
   for (int epoch = 0; epoch < nb_epoch; epoch++) {
+    if (shuffle_) std::random_shuffle(index.begin(), index.end());
     float loss = 0.0f, metric = 0.0f;
     size_t b = 0;
     for (; b < x.shape(0) / batchsize; b++) {
-      const Tensor bx = CopyRows(x, b * batchsize, b * batchsize + batchsize);
-      const Tensor by = CopyRows(y, b * batchsize, b * batchsize + batchsize);
-      const auto ret = TrainOnBatch(bx, by);
+      size_t idx = index[b];
+      const Tensor bx = CopyRows(x, idx * batchsize, (idx + 1) * batchsize);
+      const Tensor by = CopyRows(y, idx * batchsize, (idx + 1) * batchsize);
+      const auto ret = TrainOnBatch(epoch, bx, by);
       loss += ret.first;
       metric += ret.second;
     }
@@ -162,21 +182,27 @@ void FeedForwardNet::Train(size_t batchsize, int nb_epoch, const Tensor& x,
   }
 }
 
-const std::pair<float, float> FeedForwardNet::TrainOnBatch(const Tensor& x,
+const std::pair<float, float> FeedForwardNet::TrainOnBatch(int epoch,
+                                                           const Tensor& x,
                                                            const Tensor& y) {
   int flag = kTrain;
   const Tensor fea = Forward(flag, x);
-  float loss = loss_->Evaluate(fea, y);
+  float loss = loss_->Evaluate(flag, fea, y);
   float metric = metric_->Evaluate(fea, y);
   const Tensor grad = loss_->Backward();
-  const auto grads = Backward(kTrain, grad);
+  auto grads = Backward(kTrain, grad / static_cast<float>(x.shape(0)));
+  auto names = GetParamNames();
+  auto values = GetParamValues();
+  for (size_t k = 0; k < grads.size(); k++) {
+    opt_->Apply(epoch, names[k], &grads[k], values.at(k));
+  }
   return std::make_pair(loss, metric);
 }
 
 const Tensor FeedForwardNet::Forward(int flag, const Tensor& data) {
   Tensor input = data, output;
   for (auto layer : layers_) {
-//    LOG(INFO) << layer->name();
+//    LOG(INFO) << layer->name() << ": " << input.L1();
     output = layer->Forward(flag, input);
     input = output;
   }
@@ -185,13 +211,22 @@ const Tensor FeedForwardNet::Forward(int flag, const Tensor& data) {
 
 const vector<Tensor> FeedForwardNet::Backward(int flag, const Tensor& grad) {
   vector<Tensor> param_grads;
+  std::stack<Tensor> buf;
   Tensor tmp = grad;
   for (int i = layers_.size() - 1; i >= 0; i--) {
- //   LOG(INFO) << layers_.at(i)->name();
+ //   LOG(INFO) << layers_.at(i)->name() << " : " << tmp.L1();
     auto ret = layers_.at(i)->Backward(flag, tmp);
     tmp = ret.first;
-    if (ret.second.size())
-      for (const auto x : ret.second) param_grads.push_back(x);
+    if (ret.second.size()) {
+      for (int k = ret.second.size() - 1; k >= 0; k--) {
+        buf.push(ret.second[k]);
+ //       LOG(INFO) <<  "      " << buf.top().L1();
+      }
+    }
+  }
+  while (!buf.empty()) {
+    param_grads.push_back(buf.top());
+    buf.pop();
   }
   return param_grads;
 }
@@ -229,9 +264,9 @@ std::pair<Tensor, Tensor> FeedForwardNet::EvaluateOnBatch(const Tensor& x,
                                                           const Tensor& y) {
   int flag = kEval;
   const Tensor fea = Forward(flag, x);
+  const Tensor l = loss_->Forward(flag, fea, y);
   const Tensor m = metric_->Forward(fea, y);
-  const Tensor l = loss_->Forward(fea, y);
-  return std::make_pair(m, l);
+  return std::make_pair(l, m);
 }
 
 const Tensor FeedForwardNet::Predict(const Tensor& x, size_t batchsize) {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/model/layer/cudnn_convolution.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/cudnn_convolution.cc b/src/model/layer/cudnn_convolution.cc
index eb507b2..33f2cf8 100644
--- a/src/model/layer/cudnn_convolution.cc
+++ b/src/model/layer/cudnn_convolution.cc
@@ -72,8 +72,8 @@ void CudnnConvolution::InitCudnn(const Tensor &input) {
       num_filters_, conv_height_, conv_width_));
   if (bias_term_)
     CUDNN_CHECK(cudnnSetTensor4dDescriptor(bias_desc_, CUDNN_TENSOR_NCHW,
-                                           GetCudnnDataType(dtype), 1, 1,
-                                           num_filters_, 1));
+                                           GetCudnnDataType(dtype), 1, num_filters_, 1,
+                                           1));
   CUDNN_CHECK(cudnnSetConvolution2dDescriptor(conv_desc_, pad_h_, pad_w_,
                                               stride_h_, stride_w_, 1, 1,
                                               CUDNN_CROSS_CORRELATION));

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/model/layer/cudnn_dropout.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/cudnn_dropout.cc b/src/model/layer/cudnn_dropout.cc
index f9b9dbf..ab83226 100644
--- a/src/model/layer/cudnn_dropout.cc
+++ b/src/model/layer/cudnn_dropout.cc
@@ -108,7 +108,7 @@ const std::pair<Tensor, vector<Tensor>> CudnnDropout::Backward(
 }
 void CudnnDropout::ToDevice(std::shared_ptr<Device> device) {
   Dropout::ToDevice(device);
-  state.ToDevice(device);
+  state_.ToDevice(device);
 }
 }  // namespace singa
 #endif  // CUDNN_VERSION_MAJOR>=5

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/model/layer/dense.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/dense.cc b/src/model/layer/dense.cc
index c6a9f8a..338409c 100644
--- a/src/model/layer/dense.cc
+++ b/src/model/layer/dense.cc
@@ -41,13 +41,15 @@ void Dense::Setup(const Shape& in_sample, const LayerConf &conf) {
   bias_.Reshape(Shape{hdim_});
   param_values_.push_back(&weight_);
   param_values_.push_back(&bias_);
+  for (auto specs: conf.param())
+    param_specs_.push_back(specs);
 }
 
 /// \copydoc Layer::Forward(int flag, const Tensor&)
 const Tensor Dense::Forward(int flag, const Tensor &input) {
   CHECK(buf_.empty());
   Tensor output;
-  CHECK_EQ(input.nDim(), 2);
+  CHECK_EQ(input.nDim(), 2u);
   if (transpose_)  // use the transposed version of weight_ for computing
     output = Mult(input, weight_);
   else

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/model/loss/mse.cc
----------------------------------------------------------------------
diff --git a/src/model/loss/mse.cc b/src/model/loss/mse.cc
index a4bbb72..6e19059 100644
--- a/src/model/loss/mse.cc
+++ b/src/model/loss/mse.cc
@@ -20,7 +20,7 @@
 
 namespace singa {
 
-Tensor MSE::Forward(const Tensor& prediction, const Tensor& target) {
+Tensor MSE::Forward(int flag, const Tensor& prediction, const Tensor& target) {
   CHECK(buf_.empty()) << "Do not call Forward successively for more than twice."
                       << " The calling pattern is [Forward|Evaluate] Backward";
   Tensor t = prediction - target;
@@ -28,7 +28,8 @@ Tensor MSE::Forward(const Tensor& prediction, const Tensor& target) {
   if (t.nDim() > 1) batchsize = t.shape().at(0);
   size_t dim = t.Size() / batchsize;
   t.Reshape(Shape{batchsize, dim});
-  buf_.push(t);
+  if (kTrain & flag)
+    buf_.push(t);
   // TODO(wangwei) use CastType for operator/
   return Sum(Square(t), 1) * 0.5f;
 }

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/model/loss/softmax_cross_entropy.cc
----------------------------------------------------------------------
diff --git a/src/model/loss/softmax_cross_entropy.cc b/src/model/loss/softmax_cross_entropy.cc
index bed3348..3411fbe 100644
--- a/src/model/loss/softmax_cross_entropy.cc
+++ b/src/model/loss/softmax_cross_entropy.cc
@@ -21,7 +21,7 @@
 
 namespace singa {
 
-Tensor SoftmaxCrossEntropy::Forward(const Tensor& prediction,
+Tensor SoftmaxCrossEntropy::Forward(int flag, const Tensor& prediction,
                                     const Tensor& target) {
   CHECK(buf_.empty()) << "Do not call Forward successively for more than twice."
                       << " The calling pattern is [Forward|Evaluate] Backward";
@@ -30,13 +30,17 @@ Tensor SoftmaxCrossEntropy::Forward(const Tensor& prediction,
   size_t dim = prediction.Size() / batchsize;
   const Tensor& input = Reshape(prediction, Shape{batchsize, dim});
   Tensor prob = SoftMax(input);
+  // LOG(INFO) << "prob: " << prob.L2();
 
   // buffer intermediate data
-  buf_.push(prob);
-  buf_.push(target);
+  if (flag & kTrain) {
+    buf_.push(prob);
+    buf_.push(target);
+  }
   Tensor loss(Shape{batchsize}, prob.device(), prob.data_type());
 
   ComputeCrossEntropy(prob, target, &loss);
+
   return loss;
 }
 
@@ -50,4 +54,3 @@ Tensor SoftmaxCrossEntropy::Backward() {
 }
 }  // namespace singa
 
-

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/model/metric/accuracy.cc
----------------------------------------------------------------------
diff --git a/src/model/metric/accuracy.cc b/src/model/metric/accuracy.cc
index 1b667b1..a95486b 100644
--- a/src/model/metric/accuracy.cc
+++ b/src/model/metric/accuracy.cc
@@ -30,6 +30,7 @@ Tensor Accuracy::Match(const Tensor& predict, const vector<int>& target) {
   // TODO(wangwei) CloneToDevice(host);
   const float* prob = prediction.data<float>();
   float* score = new float[batchsize];
+  memset(score, 0, batchsize * sizeof(float));
   for (size_t b = 0; b < batchsize; b++) {
     vector<std::pair<float, int>> prob_class;
     for (size_t c = 0; c < nb_classes; c++) {
@@ -43,6 +44,7 @@ Tensor Accuracy::Match(const Tensor& predict, const vector<int>& target) {
   }
   Tensor ret(Shape{batchsize});
   ret.CopyDataFromHostPtr(score, batchsize);
+  delete score;
   return ret;
 }
 

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/model/optimizer/optimizer.cc
----------------------------------------------------------------------
diff --git a/src/model/optimizer/optimizer.cc b/src/model/optimizer/optimizer.cc
index c9e7a72..9be47c8 100644
--- a/src/model/optimizer/optimizer.cc
+++ b/src/model/optimizer/optimizer.cc
@@ -21,6 +21,17 @@
 
 namespace singa {
 
+Optimizer::~Optimizer() {
+  for (auto entry : regularizers_) delete entry.second;
+  for (auto entry : constraints_) delete entry.second;
+  if (constraint_ != nullptr) delete constraint_;
+  if (regularizer_ != nullptr) delete regularizer_;
+}
+void Optimizer::Setup(const OptimizerConf& conf) {
+  if (conf.has_regularizer())
+    regularizer_ = new Regularizer(conf.regularizer());
+  if (conf.has_constraint()) constraint_ = new Constraint(conf.constraint());
+}
 void Optimizer::Register(const string& name, const ParamSpec& specs) {
   if (specs.has_constraint()) {
     CHECK(constraints_.find(name) == constraints_.end())
@@ -32,6 +43,11 @@ void Optimizer::Register(const string& name, const ParamSpec& specs) {
         << "Parameter with name = " << name << " has already registered";
     regularizers_[name] = new Regularizer(specs.regularizer());
   }
+  if (specs.has_decay_mult()) {
+    CHECK(weight_decay_multplier_.find(name) == weight_decay_multplier_.end())
+        << "Parameter with name = " << name << " has already registered";
+    weight_decay_multplier_[name] = specs.decay_mult();
+  }
   if (specs.has_lr_mult()) {
     CHECK(learning_rate_multplier_.find(name) == learning_rate_multplier_.end())
         << "Parameter with name = " << name << " has already registered";
@@ -47,10 +63,18 @@ void Optimizer::Register(const string& name, const ParamSpec& specs) {
 void Optimizer::Apply(int step, const string& name, Tensor* grad,
                       Tensor* param) {
   // TODO(wangwei) need to consider the order of constraint and regularizer
-  if (regularizers_.find(name) != regularizers_.end())
+  if (regularizers_.find(name) != regularizers_.end()) {
     regularizers_.at(name)->Apply(step, param, grad);
+  } else if (regularizer_ != nullptr) {
+    float scale = 1.0f;
+    if (weight_decay_multplier_.find(name) != weight_decay_multplier_.end())
+      scale = weight_decay_multplier_.at(name);
+    regularizer_->Apply(step, param, grad, scale);
+  }
   if (constraints_.find(name) != constraints_.end())
     constraints_.at(name)->Apply(step, param, grad);
+  else if (constraint_ != nullptr)
+    constraint_->Apply(step, param, grad);
   float lr = learning_rate_generator_(step);
   if (learning_rate_multplier_.find(name) != learning_rate_multplier_.end())
     lr *= learning_rate_multplier_.at(name);
@@ -62,9 +86,9 @@ void Regularizer::Setup(const RegularizerConf& conf) {
   coefficient_ = conf.coefficient();
 }
 
-void Regularizer::Apply(int step, Tensor* value, Tensor* grad) {
+void Regularizer::Apply(int step, Tensor* value, Tensor* grad, float scale) {
   if (type_ == "L2" || type_ == "l2") {
-    (*grad) -= (*value) * coefficient_;
+    Axpy(coefficient_ * scale, *value, grad);
   } else {
     CHECK(type_ == "NotSet") << "Unknown regularizer type = " << type_;
   }

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/model/optimizer/sgd.cc
----------------------------------------------------------------------
diff --git a/src/model/optimizer/sgd.cc b/src/model/optimizer/sgd.cc
index a5c66a1..71071ff 100644
--- a/src/model/optimizer/sgd.cc
+++ b/src/model/optimizer/sgd.cc
@@ -22,6 +22,7 @@
 namespace singa {
 
 void SGD::Setup(const OptimizerConf& conf) {
+  Optimizer::Setup(conf);
   if (conf.has_momentum()) {
     float m = conf.momentum();
     SetMomentumGenerator([m](int step) { return m; });

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/src/proto/model.proto
----------------------------------------------------------------------
diff --git a/src/proto/model.proto b/src/proto/model.proto
index c06deec..b1318d9 100644
--- a/src/proto/model.proto
+++ b/src/proto/model.proto
@@ -89,6 +89,11 @@ message OptimizerConf {
 
   // delta is used to avoid dividing zero
   optional float delta = 6 [default = 1e-8];
+
+  // global regularizer lower priority than ParamSpec regularizer
+  optional RegularizerConf regularizer = 10;
+  // global constraint lower priority than ParamSpec constraint
+  optional ConstraintConf constraint = 11;
 }
 
 message ConstraintConf {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/test/singa/test_cross_entropy.cc
----------------------------------------------------------------------
diff --git a/test/singa/test_cross_entropy.cc b/test/singa/test_cross_entropy.cc
index d73591f..c7fa2fb 100644
--- a/test/singa/test_cross_entropy.cc
+++ b/test/singa/test_cross_entropy.cc
@@ -44,7 +44,7 @@ TEST_F(TestSoftmaxCrossEntropy, CppForward) {
   t.CopyDataFromHostPtr(tdat, 2);
 
   singa::SoftmaxCrossEntropy cross_entropy;
-  const Tensor& loss = cross_entropy.Forward(p, t);
+  const Tensor& loss = cross_entropy.Forward(singa::kEval, p, t);
   auto ldat = loss.data<float>();
 
   const float result_test = -log(0.25);
@@ -58,7 +58,7 @@ TEST_F(TestSoftmaxCrossEntropy, CppBackward) {
   t.CopyDataFromHostPtr(tdat, 2);
 
   singa::SoftmaxCrossEntropy cross_entropy;
-  cross_entropy.Forward(p, t);
+  cross_entropy.Forward(singa::kTrain, p, t);
   const Tensor& grad = cross_entropy.Backward();
 
   auto gdat = grad.data<float>();
@@ -82,7 +82,7 @@ TEST_F(TestSoftmaxCrossEntropy, CudaForward) {
   p.CopyDataFromHostPtr(pdat, 8);
   t.CopyDataFromHostPtr(tdat, 2);
 
-  Tensor loss = cross_entropy.Forward(p, t);
+  Tensor loss = cross_entropy.Forward(singa::kEval, p, t);
   loss.ToHost();
   auto ldat = loss.data<float>();
 
@@ -99,7 +99,7 @@ TEST_F(TestSoftmaxCrossEntropy, CudaBackward) {
   p.CopyDataFromHostPtr(pdat, 8);
   t.CopyDataFromHostPtr(tdat, 2);
 
-  cross_entropy.Forward(p, t);
+  cross_entropy.Forward(singa::kTrain, p, t);
   Tensor grad = cross_entropy.Backward();
 
   grad.ToHost();

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/test/singa/test_dense.cc
----------------------------------------------------------------------
diff --git a/test/singa/test_dense.cc b/test/singa/test_dense.cc
index 363fb6e..e80384f 100644
--- a/test/singa/test_dense.cc
+++ b/test/singa/test_dense.cc
@@ -207,7 +207,7 @@ TEST(Dense, BackwardCuda) {
   singa::Tensor grad(singa::Shape{batchsize, hdim}, cuda);
   grad.CopyDataFromHostPtr(dy, batchsize * hdim);
 
-  const auto ret = dense.Backward(singa::kTrain, grad);
+  auto ret = dense.Backward(singa::kTrain, grad);
   singa::Tensor in_grad = ret.first;
   singa::Tensor dweight = ret.second.at(0);
   singa::Tensor dbias = ret.second.at(1);

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/test/singa/test_mse.cc
----------------------------------------------------------------------
diff --git a/test/singa/test_mse.cc b/test/singa/test_mse.cc
index 788652f..640caf4 100644
--- a/test/singa/test_mse.cc
+++ b/test/singa/test_mse.cc
@@ -42,7 +42,7 @@ class TestMSE : public ::testing::Test {
 #ifdef USE_CBLAS
 TEST_F(TestMSE, CppForward) {
   singa::MSE mse;
-  const Tensor& loss = mse.Forward(p, t);
+  const Tensor& loss = mse.Forward(singa::kEval, p, t);
   auto ldat = loss.data<float>();
 
   for (size_t i = 0, k = 0; i < loss.Size(); i++) {
@@ -57,7 +57,7 @@ TEST_F(TestMSE, CppForward) {
 
 TEST_F(TestMSE, CppBackward) {
   singa::MSE mse;
-  mse.Forward(p, t);
+  mse.Forward(singa::kTrain, p, t);
   const Tensor& grad = mse.Backward();
 
   auto gdat = grad.data<float>();
@@ -72,7 +72,7 @@ TEST_F(TestMSE, CudaForward) {
   auto dev = std::make_shared<singa::CudaGPU>();
   p.ToDevice(dev);
   t.ToDevice(dev);
-  Tensor loss = mse->Forward(p, t);
+  Tensor loss = mse->Forward(singa::kEval, p, t);
 
   loss.ToHost();
   auto ldat = loss.data<float>();
@@ -94,7 +94,7 @@ TEST_F(TestMSE, CudaBackward) {
   auto dev = std::make_shared<singa::CudaGPU>();
   p.ToDevice(dev);
   t.ToDevice(dev);
-  mse.Forward(p, t);
+  mse.Forward(singa::kTrain, p, t);
   Tensor grad = mse.Backward();
   grad.ToHost();
   auto gdat = grad.data<float>();

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/97648a60/test/singa/test_tensor_math.cc
----------------------------------------------------------------------
diff --git a/test/singa/test_tensor_math.cc b/test/singa/test_tensor_math.cc
index f8d0351..2803a23 100644
--- a/test/singa/test_tensor_math.cc
+++ b/test/singa/test_tensor_math.cc
@@ -346,7 +346,7 @@ TEST_F(TestTensorMath, L2Cpp) {
   float l2 = a.L2();
   float target = 0.0f;
   for (size_t i = 0; i < a.Size(); i++) target += dat1[i] * dat1[i];
-  EXPECT_FLOAT_EQ(l2, sqrt(target));
+  EXPECT_FLOAT_EQ(l2, sqrt(target) / a.Size());
 }
 TEST_F(TestTensorMath, MultCpp) {
   const float x[4] = {1.0f, 2.0f, 3.0f, 4.0f};
@@ -480,13 +480,16 @@ TEST_F(TestTensorMath, MultRowCpp) {
 
 TEST_F(TestTensorMath, SumRowsCpp) {
   Tensor t(Shape{2});
-  d.CopyDataFromHostPtr(dat1, 6);
+  float dat[6];
+  for (int i = 0; i < 6; i ++)
+    dat[i] = (float)rand()/(float)(RAND_MAX/ 10);
+  d.CopyDataFromHostPtr(dat, 6);
   SumRows(d, &t);
   const float *tptr = t.data<float>();
   for (int i = 0; i < 2; i++) {
     float tmp = 0;
     for (int j = 0; j < 3; j++) {
-      tmp += dat1[j * 2 + i];
+      tmp += dat[j * 2 + i];
     }
     EXPECT_FLOAT_EQ(tptr[i], tmp);
   }
@@ -514,7 +517,7 @@ TEST_F(TestTensorMath, L2Cuda) {
   float l2 = t.L2();
   float target = 0.0f;
   for (size_t i = 0; i < t.Size(); i++) target += dat1[i] * dat1[i];
-  EXPECT_FLOAT_EQ(l2, sqrt(target));
+  EXPECT_FLOAT_EQ(l2, sqrt(target) / t.Size());
 }
 TEST_F(TestTensorMath, MultCuda) {
   const float x[4] = {1.0f, 2.0f, 3.0f, 4.0f};


[4/6] incubator-singa git commit: SINGA-204 Support the training of feed-forward neural nets

Posted by ka...@apache.org.
SINGA-204 Support the training of feed-forward neural nets

Add cmake supporting for examples(cifar/alexnet).


Project: http://git-wip-us.apache.org/repos/asf/incubator-singa/repo
Commit: http://git-wip-us.apache.org/repos/asf/incubator-singa/commit/fce5d813
Tree: http://git-wip-us.apache.org/repos/asf/incubator-singa/tree/fce5d813
Diff: http://git-wip-us.apache.org/repos/asf/incubator-singa/diff/fce5d813

Branch: refs/heads/dev
Commit: fce5d8137d50f6ad6579a53137cf832bf2c9509c
Parents: 71eb059
Author: xiezl <xi...@comp.nus.edu.sg>
Authored: Mon Jun 27 18:49:10 2016 +0800
Committer: xiezl <xi...@comp.nus.edu.sg>
Committed: Mon Jun 27 18:49:10 2016 +0800

----------------------------------------------------------------------
 CMakeLists.txt          | 1 +
 examples/CMakeLists.txt | 8 ++++++++
 2 files changed, 9 insertions(+)
----------------------------------------------------------------------


http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/fce5d813/CMakeLists.txt
----------------------------------------------------------------------
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 7a5caf3..21d25cf 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -43,3 +43,4 @@ ADD_SUBDIRECTORY(lib/cnmem)
 LIST(APPEND SINGA_LINKER_LIBS cnmem)
 ADD_SUBDIRECTORY(src)
 ADD_SUBDIRECTORY(test)
+ADD_SUBDIRECTORY(examples)

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/fce5d813/examples/CMakeLists.txt
----------------------------------------------------------------------
diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
new file mode 100644
index 0000000..b06f434
--- /dev/null
+++ b/examples/CMakeLists.txt
@@ -0,0 +1,8 @@
+INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR})
+INCLUDE_DIRECTORIES(${CMAKE_BINARY_DIR}/include)
+
+AUX_SOURCE_DIRECTORY(cifar10 cifar_source)
+
+ADD_EXECUTABLE(alexnet ${cifar_source})
+ADD_DEPENDENCIES(alexnet singa_core singa_model singa_utils)
+TARGET_LINK_LIBRARIES(alexnet singa_core singa_utils singa_model protobuf ${SINGA_LIBKER_LIBS})


[6/6] incubator-singa git commit: SINGA-204 Support the training of feed-forward neural nets

Posted by ka...@apache.org.
SINGA-204 Support the training of feed-forward neural nets

Merge commits from #xiezl (zhongle) for adding CMake for examples.
To test the cifar10 exmaple,
1. mkdir build && cd build
2. cmake .. && make
3. cd ../examples/cifar10
4. export LD_LIBRARY_PATH=../../build/lib:$LD_LIBRARY_PATH
5. ./run.sh


Project: http://git-wip-us.apache.org/repos/asf/incubator-singa/repo
Commit: http://git-wip-us.apache.org/repos/asf/incubator-singa/commit/4db968c2
Tree: http://git-wip-us.apache.org/repos/asf/incubator-singa/tree/4db968c2
Diff: http://git-wip-us.apache.org/repos/asf/incubator-singa/diff/4db968c2

Branch: refs/heads/dev
Commit: 4db968c2ec003a051118cf4b49f6ae10cfbc054f
Parents: 97648a6 fce5d81
Author: Wei Wang <wa...@comp.nus.edu.sg>
Authored: Mon Jun 27 19:55:46 2016 +0800
Committer: Wei Wang <wa...@comp.nus.edu.sg>
Committed: Mon Jun 27 19:55:46 2016 +0800

----------------------------------------------------------------------
 CMakeLists.txt                       | 3 ++-
 examples/CMakeLists.txt              | 8 ++++++++
 examples/cifar10/alexnet.cc          | 6 ++----
 examples/cifar10/cifar10.h           | 6 +++---
 examples/cifar10/make.sh             | 1 -
 examples/cifar10/run.sh              | 2 ++
 src/model/feed_forward_net.cc        | 4 +++-
 src/model/layer/cudnn_convolution.cc | 4 ++--
 8 files changed, 22 insertions(+), 12 deletions(-)
----------------------------------------------------------------------


http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4db968c2/CMakeLists.txt
----------------------------------------------------------------------
diff --cc CMakeLists.txt
index 7a5caf3,21d25cf..ff94d4e
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@@ -20,7 -20,7 +20,7 @@@ INCLUDE_DIRECTORIES(${SINGA_INCLUDE_DIR
  
  OPTION(USE_CBLAS "Use CBlas libs" ON)
  OPTION(USE_CUDA "Use Cuda libs" ON)
--OPTION(USE_CUDNN "Use Cudnn libs" OFF)
++OPTION(USE_CUDNN "Use Cudnn libs" ON)
  OPTION(USE_OPENCV "Use opencv" OFF)
  OPTION(USE_LMDB "Use LMDB libs" OFF)
  OPTION(USE_PYTHON "Generate py wrappers" OFF)

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4db968c2/examples/cifar10/alexnet.cc
----------------------------------------------------------------------
diff --cc examples/cifar10/alexnet.cc
index 81713f1,d6541a3..5175f49
--- a/examples/cifar10/alexnet.cc
+++ b/examples/cifar10/alexnet.cc
@@@ -84,12 -84,12 +84,10 @@@ LayerConf GenDenseConf(string name, in
    conf.set_type("Dense");
    DenseConf *dense = conf.mutable_dense_conf();
    dense->set_num_output(num_output);
--  FillerConf *bias = dense->mutable_bias_filler();
  
    ParamSpec *wspec = conf.add_param();
    wspec->set_name(name + "_weight");
    wspec->set_decay_mult(wd);
--
    auto wfill = wspec->mutable_filler();
    wfill->set_type("Gaussian");
    wfill->set_std(std);
@@@ -155,16 -155,14 +153,16 @@@ void Train(float lr, int num_epoch, str
      train_y = train.second;
      auto test = data.ReadTestData();
      nsamples = test.first.shape(0);
 -    auto maty =
 +    auto mtest =
          Reshape(test.first, Shape{nsamples, test.first.Size() / nsamples});
 -    SubRow(mean, &maty);
 -    test_x = Reshape(maty, test.first.shape());
 +    SubRow(mean, &mtest);
 +    test_x = Reshape(mtest, test.first.shape());
      test_y = test.second;
    }
 +  CHECK_EQ(train_x.shape(0), train_y.shape(0));
 +  CHECK_EQ(test_x.shape(0), test_y.shape(0));
    LOG(INFO) << "Training samples = " << train_y.shape(0)
-             << " Test samples = " << test_y.shape(0);
 -            << " Test samples =" << test_y.shape(0);
++            << ", Test samples = " << test_y.shape(0);
    auto net = CreateNet();
    SGD sgd;
    OptimizerConf opt_conf;
@@@ -177,10 -175,9 +175,10 @@@
        return 0.001;
      else if (step <= 130)
        return 0.0001;
--    else if (step <= 140)
++    else
        return 0.00001;
    });
 +
    SoftmaxCrossEntropy loss;
    Accuracy acc;
    net.Compile(true, &sgd, &loss, &acc);

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4db968c2/examples/cifar10/cifar10.h
----------------------------------------------------------------------
diff --cc examples/cifar10/cifar10.h
index 5af54e2,7f10153..d2b9225
--- a/examples/cifar10/cifar10.h
+++ b/examples/cifar10/cifar10.h
@@@ -68,11 -70,11 +68,11 @@@ const std::pair<Tensor, Tensor> Cifar10
    char image[kImageVol];
    float float_image[kImageVol];
    int tmplabels[kBatchSize];
--  for (int itemid = 0; itemid < kBatchSize; ++itemid) {
++  for (size_t itemid = 0; itemid < kBatchSize; ++itemid) {
      // LOG(INFO) << "reading " << itemid << "-th image";
      ReadImage(&data_file, &label, image);
--    for (int i = 0; i < kImageVol; i++)
 -      float_image[i] = static_cast<float>(static_cast<int>(image[i]));
++    for (size_t i = 0; i < kImageVol; i++)
 +      float_image[i] = static_cast<float>(static_cast<uint8_t>(image[i]));
      images.CopyDataFromHostPtr(float_image, kImageVol, itemid * kImageVol);
      tmplabels[itemid] = label;
    }
@@@ -80,10 -82,10 +80,10 @@@
    return std::make_pair(images, labels);
  }
  
 -const std::pair<Tensor, Tensor> Cifar10::ReadTrainData(bool shuffle) {
 +const std::pair<Tensor, Tensor> Cifar10::ReadTrainData() {
    Tensor images(Shape{kBatchSize * kTrainFiles, 3, kImageSize, kImageSize});
    Tensor labels(Shape{kBatchSize * kTrainFiles}, kInt);
--  for (int fileid = 0; fileid < kTrainFiles; ++fileid) {
++  for (size_t fileid = 0; fileid < kTrainFiles; ++fileid) {
      string file = "data_batch_" + std::to_string(fileid + 1) + ".bin";
      const auto ret = ReadFile(file);
      CopyDataToFrom(&images, ret.first, ret.first.Size(),

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4db968c2/examples/cifar10/make.sh
----------------------------------------------------------------------
diff --cc examples/cifar10/make.sh
index 5a41612,5a41612..0000000
deleted file mode 100755,100755
--- a/examples/cifar10/make.sh
+++ /dev/null
@@@ -1,1 -1,1 +1,0 @@@
--g++ -g --std=c++11 alexnet.cc -o alexnet -I../../include -I../../build/include -I/home/wangwei/local/cudnn5/include -I/home/wangwei/local/include -I/usr/local/cuda/include/ -I../../lib/cnmem/include -L../../build/lib/ -lsinga_core -lsinga_model  -lsinga_utils -lcudart -lcublas -lcurand -lcudnn -L/home/wangwei/local/cudnn5/lib64 -L/usr/local/cuda/lib64 ../../build/lib/libproto.a -lprotobuf

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4db968c2/examples/cifar10/run.sh
----------------------------------------------------------------------
diff --cc examples/cifar10/run.sh
index 0000000,0000000..c01ec18
new file mode 100755
--- /dev/null
+++ b/examples/cifar10/run.sh
@@@ -1,0 -1,0 +1,2 @@@
++#!/usr/bin/env sh
++../../build/bin/alexnet -epoch 140

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4db968c2/src/model/feed_forward_net.cc
----------------------------------------------------------------------
diff --cc src/model/feed_forward_net.cc
index 30d030a,e682918..ebbe00c
--- a/src/model/feed_forward_net.cc
+++ b/src/model/feed_forward_net.cc
@@@ -202,8 -201,8 +203,9 @@@ const std::pair<float, float> FeedForwa
  const Tensor FeedForwardNet::Forward(int flag, const Tensor& data) {
    Tensor input = data, output;
    for (auto layer : layers_) {
 +//    LOG(INFO) << layer->name() << ": " << input.L1();
      output = layer->Forward(flag, input);
+     // LOG(INFO) << layer->name() << ": " << output.L2();
      input = output;
    }
    return output;

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4db968c2/src/model/layer/cudnn_convolution.cc
----------------------------------------------------------------------
diff --cc src/model/layer/cudnn_convolution.cc
index 33f2cf8,3dca28a..82cf4e5
--- a/src/model/layer/cudnn_convolution.cc
+++ b/src/model/layer/cudnn_convolution.cc
@@@ -72,8 -72,8 +72,8 @@@ void CudnnConvolution::InitCudnn(const 
        num_filters_, conv_height_, conv_width_));
    if (bias_term_)
      CUDNN_CHECK(cudnnSetTensor4dDescriptor(bias_desc_, CUDNN_TENSOR_NCHW,
-                                            GetCudnnDataType(dtype), 1, num_filters_, 1,
-                                            1));
 -                                           GetCudnnDataType(dtype), 1, 1, 1,
 -                                           num_filters_));
++                                           GetCudnnDataType(dtype), 1,
++                                           num_filters_, 1, 1));
    CUDNN_CHECK(cudnnSetConvolution2dDescriptor(conv_desc_, pad_h_, pad_w_,
                                                stride_h_, stride_w_, 1, 1,
                                                CUDNN_CROSS_CORRELATION));


[3/6] incubator-singa git commit: SINGA-204 Support the training of feed-forward neural nets

Posted by ka...@apache.org.
SINGA-204 Support the training of feed-forward neural nets

Implement Alexnet model for Cifar10 https://code.google.com/p/cuda-convnet/
But the test accuracy is low 0.72 (which should be 0.82).


Project: http://git-wip-us.apache.org/repos/asf/incubator-singa/repo
Commit: http://git-wip-us.apache.org/repos/asf/incubator-singa/commit/71eb059c
Tree: http://git-wip-us.apache.org/repos/asf/incubator-singa/tree/71eb059c
Diff: http://git-wip-us.apache.org/repos/asf/incubator-singa/diff/71eb059c

Branch: refs/heads/dev
Commit: 71eb059cd13ea41e74195c7c115f927aaf143490
Parents: cf1d841
Author: Wei Wang <wa...@comp.nus.edu.sg>
Authored: Mon Jun 27 01:21:59 2016 +0800
Committer: Wei Wang <wa...@comp.nus.edu.sg>
Committed: Mon Jun 27 15:29:05 2016 +0800

----------------------------------------------------------------------
 examples/cifar10/alexnet.cc             | 123 ++++++++++++++++++---------
 examples/cifar10/cifar10.h              |   3 +-
 examples/cifar10/make.sh                |   2 +-
 include/singa/core/tensor.h             |   1 +
 include/singa/model/feed_forward_net.h  |   2 +-
 include/singa/model/initializer.h       |  26 +++++-
 include/singa/model/loss.h              |  12 ++-
 include/singa/model/metric.h            |   2 +-
 include/singa/model/optimizer.h         |  18 ++--
 include/singa/utils/string.h            |  11 +++
 src/core/tensor/math_kernel.cu          |  22 +++++
 src/core/tensor/math_kernel.h           |   2 +
 src/core/tensor/tensor.cc               |  69 ++++++++++-----
 src/core/tensor/tensor_math.h           |   5 ++
 src/core/tensor/tensor_math_cpp.h       |  14 +++
 src/core/tensor/tensor_math_cuda.h      |   9 ++
 src/model/feed_forward_net.cc           | 104 ++++++++++++++--------
 src/model/layer/cudnn_convolution.cc    |   5 +-
 src/model/layer/cudnn_dropout.cc        |   2 +-
 src/model/layer/dense.cc                |   4 +-
 src/model/loss/mse.cc                   |   5 +-
 src/model/loss/softmax_cross_entropy.cc |  11 ++-
 src/model/metric/accuracy.cc            |   1 +
 src/model/optimizer/optimizer.cc        |  30 ++++++-
 src/model/optimizer/sgd.cc              |   1 +
 src/proto/model.proto                   |   5 ++
 test/singa/test_cross_entropy.cc        |   8 +-
 test/singa/test_dense.cc                |   2 +-
 test/singa/test_mse.cc                  |   8 +-
 test/singa/test_tensor_math.cc          |   4 +-
 30 files changed, 370 insertions(+), 141 deletions(-)
----------------------------------------------------------------------


http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/examples/cifar10/alexnet.cc
----------------------------------------------------------------------
diff --git a/examples/cifar10/alexnet.cc b/examples/cifar10/alexnet.cc
index 45d8571..d6541a3 100644
--- a/examples/cifar10/alexnet.cc
+++ b/examples/cifar10/alexnet.cc
@@ -28,12 +28,13 @@
 #include "../../src/model/layer/cudnn_convolution.h"
 #include "../../src/model/layer/cudnn_activation.h"
 #include "../../src/model/layer/cudnn_pooling.h"
+#include "../../src/model/layer/cudnn_lrn.h"
 #include "../../src/model/layer/dense.h"
 #include "../../src/model/layer/flatten.h"
 namespace singa {
 
 LayerConf GenConvConf(string name, int nb_filter, int kernel, int stride,
-                      int pad) {
+                      int pad, float std) {
   LayerConf conf;
   conf.set_name(name);
   conf.set_type("CudnnConvolution");
@@ -42,13 +43,23 @@ LayerConf GenConvConf(string name, int nb_filter, int kernel, int stride,
   conv->add_kernel_size(kernel);
   conv->add_stride(stride);
   conv->add_pad(pad);
+  conv->set_bias_term(true);
 
-  FillerConf *weight = conv->mutable_weight_filler();
-  weight->set_type("Xavier");
+  ParamSpec *wspec = conf.add_param();
+  wspec->set_name(name + "_weight");
+  auto wfill = wspec->mutable_filler();
+  wfill->set_type("Gaussian");
+  wfill->set_std(std);
+
+  ParamSpec *bspec = conf.add_param();
+  bspec->set_name(name + "_bias");
+  bspec->set_lr_mult(2);
+//  bspec->set_decay_mult(0);
   return conf;
 }
 
-LayerConf GenPoolingConf(string name, bool max_pool, int kernel, int stride, int pad) {
+LayerConf GenPoolingConf(string name, bool max_pool, int kernel, int stride,
+                         int pad) {
   LayerConf conf;
   conf.set_name(name);
   conf.set_type("CudnnPooling");
@@ -56,8 +67,7 @@ LayerConf GenPoolingConf(string name, bool max_pool, int kernel, int stride, int
   pool->set_kernel_size(kernel);
   pool->set_stride(stride);
   pool->set_pad(pad);
-  if (!max_pool)
-    pool->set_pool(PoolingConf_PoolMethod_AVE);
+  if (!max_pool) pool->set_pool(PoolingConf_PoolMethod_AVE);
   return conf;
 }
 
@@ -68,21 +78,38 @@ LayerConf GenReLUConf(string name) {
   return conf;
 }
 
-LayerConf GenDenseConf(string name, int num_output) {
+LayerConf GenDenseConf(string name, int num_output, float std, float wd) {
   LayerConf conf;
   conf.set_name(name);
   conf.set_type("Dense");
   DenseConf *dense = conf.mutable_dense_conf();
   dense->set_num_output(num_output);
-  FillerConf *weight = dense->mutable_weight_filler();
-  weight->set_type("Xavier");
+  FillerConf *bias = dense->mutable_bias_filler();
+
+  ParamSpec *wspec = conf.add_param();
+  wspec->set_name(name + "_weight");
+  wspec->set_decay_mult(wd);
+
+  auto wfill = wspec->mutable_filler();
+  wfill->set_type("Gaussian");
+  wfill->set_std(std);
+
+  ParamSpec *bspec = conf.add_param();
+  bspec->set_name(name + "_bias");
+  bspec->set_lr_mult(2);
+  bspec->set_decay_mult(0);
+
   return conf;
 }
 
-LayerConf GenSoftmaxConf(string name) {
+LayerConf GenLRNConf(string name) {
   LayerConf conf;
   conf.set_name(name);
-  conf.set_type("CudnnSoftmax");
+  conf.set_type("CudnnLRN");
+  LRNConf *lrn = conf.mutable_lrn_conf();
+  lrn->set_local_size(3);
+  lrn->set_alpha(5e-05);
+  lrn->set_beta(0.75);
   return conf;
 }
 
@@ -92,25 +119,25 @@ LayerConf GenFlattenConf(string name) {
   conf.set_type("Flatten");
   return conf;
 }
-FeedForwardNet CreateNet(Optimizer* opt, Loss<Tensor>* loss, Metric<Tensor>* metric) {
+
+FeedForwardNet CreateNet() {
   FeedForwardNet net;
   Shape s{3, 32, 32};
 
-  net.Add(new CudnnConvolution(), GenConvConf("conv1", 32, 5, 1, 2), &s);
+  net.Add(new CudnnConvolution(), GenConvConf("conv1", 32, 5, 1, 2, 0.0001),
+          &s);
   net.Add(new CudnnActivation(), GenReLUConf("relu1"));
-  net.Add(new CudnnPooling, GenPoolingConf("pool1", true, 3, 2, 0));
-  net.Add(new CudnnConvolution(), GenConvConf("conv2", 32, 5, 1, 2));
+  net.Add(new CudnnPooling(), GenPoolingConf("pool1", true, 3, 2, 1));
+  net.Add(new CudnnLRN(), GenLRNConf("lrn1"));
+  net.Add(new CudnnConvolution(), GenConvConf("conv2", 32, 5, 1, 2, 0.01));
   net.Add(new CudnnActivation(), GenReLUConf("relu2"));
-  net.Add(new CudnnPooling(), GenPoolingConf("pool2", true, 3, 2, 0));
-  net.Add(new CudnnConvolution, GenConvConf("conv3", 64, 5, 1, 2));
+  net.Add(new CudnnPooling(), GenPoolingConf("pool2", false, 3, 2, 1));
+  net.Add(new CudnnLRN(), GenLRNConf("lrn2"));
+  net.Add(new CudnnConvolution, GenConvConf("conv3", 64, 5, 1, 2, 0.01));
   net.Add(new CudnnActivation(), GenReLUConf("relu3"));
-  net.Add(new CudnnConvolution(), GenConvConf("pool3", true, 3, 2, 0));
+  net.Add(new CudnnPooling(), GenPoolingConf("pool3", false, 3, 2, 1));
   net.Add(new Flatten(), GenFlattenConf("flat"));
-  net.Add(new Dense(), GenDenseConf("ip1", 10));
-  OptimizerConf opt_conf;
-  opt_conf.set_momentum(0.9);
-  opt->Setup(opt_conf);
-  net.Compile(true, opt, loss, metric);
+  net.Add(new Dense(), GenDenseConf("ip", 10, 0.01, 250));
   return net;
 }
 
@@ -120,50 +147,62 @@ void Train(float lr, int num_epoch, string data_dir) {
   {
     auto train = data.ReadTrainData();
     size_t nsamples = train.first.shape(0);
-    auto matx = Reshape(train.first, Shape{nsamples, train.first.Size() / nsamples});
+    auto matx =
+        Reshape(train.first, Shape{nsamples, train.first.Size() / nsamples});
     const auto mean = Average(matx, 0);
     SubRow(mean, &matx);
     train_x = Reshape(matx, train.first.shape());
     train_y = train.second;
     auto test = data.ReadTestData();
     nsamples = test.first.shape(0);
-    auto maty = Reshape(test.first, Shape{nsamples, test.first.Size() / nsamples});
+    auto maty =
+        Reshape(test.first, Shape{nsamples, test.first.Size() / nsamples});
     SubRow(mean, &maty);
     test_x = Reshape(maty, test.first.shape());
     test_y = test.second;
   }
-  LOG(ERROR) << "creating net";
+  LOG(INFO) << "Training samples = " << train_y.shape(0)
+            << " Test samples =" << test_y.shape(0);
+  auto net = CreateNet();
+  SGD sgd;
+  OptimizerConf opt_conf;
+  opt_conf.set_momentum(0.9);
+  auto reg = opt_conf.mutable_regularizer();
+  reg->set_coefficient(0.004);
+  sgd.Setup(opt_conf);
+  sgd.SetLearningRateGenerator([lr](int step) {
+    if (step <= 120)
+      return 0.001;
+    else if (step <= 130)
+      return 0.0001;
+    else if (step <= 140)
+      return 0.00001;
+  });
   SoftmaxCrossEntropy loss;
   Accuracy acc;
-  SGD sgd;
-  sgd.SetLearningRateGenerator([lr](int step) {return lr;});
-  auto net = CreateNet(&sgd, &loss, &acc);
-
+  net.Compile(true, &sgd, &loss, &acc);
   auto cuda = std::make_shared<CudaGPU>();
   net.ToDevice(cuda);
 
   train_x.ToDevice(cuda);
   train_y.ToDevice(cuda);
-  net.Train(50, num_epoch, train_x, train_y); // test_x, test_y);
+  test_x.ToDevice(cuda);
+  test_y.ToDevice(cuda);
+  net.Train(100, num_epoch, train_x, train_y, test_x, test_y);
 }
-
-
 }
 
-int main(int argc, char** argv) {
+int main(int argc, char **argv) {
   singa::InitChannel(nullptr);
   int pos = singa::ArgPos(argc, argv, "-epoch");
-  int nEpoch = 5;
-  if (pos != -1)
-    nEpoch = atoi(argv[pos + 1]);
+  int nEpoch = 140;
+  if (pos != -1) nEpoch = atoi(argv[pos + 1]);
   pos = singa::ArgPos(argc, argv, "-lr");
-  float lr = 0.01;
-  if (pos != -1)
-    lr = atof(argv[pos + 1]);
+  float lr = 0.001;
+  if (pos != -1) lr = atof(argv[pos + 1]);
   pos = singa::ArgPos(argc, argv, "-data");
   string data = "cifar-10-batches-bin";
-  if (pos != -1)
-    data = argv[pos + 1];
+  if (pos != -1) data = argv[pos + 1];
 
   LOG(INFO) << "Start training";
   singa::Train(lr, nEpoch, data);

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/examples/cifar10/cifar10.h
----------------------------------------------------------------------
diff --git a/examples/cifar10/cifar10.h b/examples/cifar10/cifar10.h
index 261c048..7f10153 100644
--- a/examples/cifar10/cifar10.h
+++ b/examples/cifar10/cifar10.h
@@ -40,11 +40,12 @@ class Cifar10 {
   const std::pair<Tensor, Tensor> ReadFile(string file, bool shuffle = false);
 
   void ReadImage(std::ifstream* file, int* label, char* buffer);
+
  private:
   const size_t kImageSize = 32;
   const size_t kImageVol = 3072;
   const size_t kBatchSize = 10000;
-  const size_t kTrainFiles = 1;
+  const size_t kTrainFiles = 5;
 
   string dir_path_;
   bool normalize_;

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/examples/cifar10/make.sh
----------------------------------------------------------------------
diff --git a/examples/cifar10/make.sh b/examples/cifar10/make.sh
index 17e4b39..5a41612 100755
--- a/examples/cifar10/make.sh
+++ b/examples/cifar10/make.sh
@@ -1 +1 @@
-g++ -g --std=c++11 alexnet.cc -o alexnet -I../../include -I../../build/include -I/home/wangwei/local/cudnn4/include -I/home/wangwei/local/include -I/usr/local/cuda/include/ -I../../lib/cnmem/include -L../../build/lib/ -lsinga_core -lsinga_model  -lsinga_utils -lcudart -lcublas -lcurand -lcudnn -L/usr/local/cuda/lib64 -L/home/wangwei/local/cudnn4/lib64 ../../build/lib/libproto.a -lprotobuf
+g++ -g --std=c++11 alexnet.cc -o alexnet -I../../include -I../../build/include -I/home/wangwei/local/cudnn5/include -I/home/wangwei/local/include -I/usr/local/cuda/include/ -I../../lib/cnmem/include -L../../build/lib/ -lsinga_core -lsinga_model  -lsinga_utils -lcudart -lcublas -lcurand -lcudnn -L/home/wangwei/local/cudnn5/lib64 -L/usr/local/cuda/lib64 ../../build/lib/libproto.a -lprotobuf

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/include/singa/core/tensor.h
----------------------------------------------------------------------
diff --git a/include/singa/core/tensor.h b/include/singa/core/tensor.h
index 3b496d9..18aa7ef 100644
--- a/include/singa/core/tensor.h
+++ b/include/singa/core/tensor.h
@@ -180,6 +180,7 @@ class Tensor {
   template <typename SType>
   Tensor &operator/=(const SType x);
 
+  float L1() const;
   float L2() const;
 
  protected:

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/include/singa/model/feed_forward_net.h
----------------------------------------------------------------------
diff --git a/include/singa/model/feed_forward_net.h b/include/singa/model/feed_forward_net.h
index 9beeb7a..1ca417c 100644
--- a/include/singa/model/feed_forward_net.h
+++ b/include/singa/model/feed_forward_net.h
@@ -72,7 +72,7 @@ class FeedForwardNet {
   void Train(size_t batchsize, int nb_epoch, const Tensor& x, const Tensor& y,
              const Tensor& val_x, const Tensor& val_y);
   /// Train the neural net over one batch of training data.
-  const std::pair<float, float> TrainOnBatch(const Tensor& x, const Tensor& y);
+  const std::pair<float, float> TrainOnBatch(int epoch, const Tensor& x, const Tensor& y);
 
   /// Evaluate the neural net with given data.
   /// Returns one tensor for loss values and one tensor for metric values;

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/include/singa/model/initializer.h
----------------------------------------------------------------------
diff --git a/include/singa/model/initializer.h b/include/singa/model/initializer.h
index 302fc97..7024f70 100644
--- a/include/singa/model/initializer.h
+++ b/include/singa/model/initializer.h
@@ -21,8 +21,8 @@
 #include <string>
 #include "singa/core/tensor.h"
 #include "singa/proto/model.pb.h"
+#include "singa/utils/string.h"
 namespace singa {
-namespace init {
 /// Base class for initializing parameter values.
 using InitializerConf = FillerConf;
 class Initializer {
@@ -40,6 +40,7 @@ class Initializer {
   virtual void Fill(Tensor* t) = 0;
 };
 
+namespace init {
 class Constant : public Initializer {
 public:
   Constant() = default;
@@ -76,7 +77,7 @@ public:
   void Fill(Tensor* t) override { singa::Gaussian(mean_, std_, t); }
 
  private:
-  float mean_ = 0, std_ = 0.01;
+  float mean_ = 0, std_ = 1;
 };
 
 /// Ref: [Bengio and Glorot 2010] Understanding the difficulty of training deep
@@ -86,6 +87,7 @@ public:
   void Fill(Tensor* t) override {
     CHECK_EQ(t->nDim(), 2u);
     float scale = sqrt(6.0f / (t->shape(0) + t->shape(1)));
+    LOG(INFO) << "xavier scale " << scale;
     singa::Uniform(-scale, scale, t);
   }
 };
@@ -100,6 +102,26 @@ class MSRA : public Initializer {
     singa::Gaussian(0.0f, std, t);
   }
 };
+
 }  // namespace init
+
+std::shared_ptr<Initializer> CreateInitializer(const InitializerConf& conf) {
+  std::shared_ptr<Initializer> init;
+  if (ToLowerCase(conf.type()) == "constant") {
+    init = std::make_shared<init::Constant>();
+  } else if (ToLowerCase(conf.type()) == "uniform") {
+    init = std::make_shared<init::Uniform>();
+  } else if (ToLowerCase(conf.type()) == "gaussian") {
+    init = std::make_shared<init::Gaussian>();
+  } else if (ToLowerCase(conf.type()) == "xavier") {
+    init = std::make_shared<init::Xavier>();
+  } else if (ToLowerCase(conf.type()) == "msra") {
+    init = std::make_shared<init::MSRA>();
+  } else {
+    LOG(FATAL) << "Unknown initialization type: " << conf.type();
+  }
+  init->Setup(conf);
+  return init;
+}
 }  // namespace singa
 #endif  // SINGA_MODEL_INITIALIZER_H_

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/include/singa/model/loss.h
----------------------------------------------------------------------
diff --git a/include/singa/model/loss.h b/include/singa/model/loss.h
index 41ec701..f400768 100644
--- a/include/singa/model/loss.h
+++ b/include/singa/model/loss.h
@@ -43,13 +43,13 @@ class Loss {
 
   /// Compute the loss values for each sample/instance given the prediction
   /// and the target.
-  virtual Tensor Forward(const Tensor& prediction, const T& target) = 0;
+  virtual Tensor Forward(int flag, const Tensor& prediction, const T& target) = 0;
 
   /// Average loss values for all samples in the mini-batch
   /// It calls Forward() internally. The calling pattern should be
   /// [Evaluate|Forward] Backward.
-  float Evaluate(const Tensor& prediction, const T& target) {
-    Tensor loss = Forward(prediction, target);
+  float Evaluate(int flag, const Tensor& prediction, const T& target) {
+    Tensor loss = Forward(flag, prediction, target);
     loss.ToHost();
     return Sum<float>(loss) / (1.0f * loss.Size());
   }
@@ -68,7 +68,7 @@ class MSE : public Loss<Tensor> {
   /// and the target, which is 0.5/||prediction-target||^2
   /// Users can call Average(const Tensor&) to get the average
   /// loss value over all samples in the batch.
-  Tensor Forward(const Tensor& prediction, const Tensor& target) override;
+  Tensor Forward(int flag, const Tensor& prediction, const Tensor& target) override;
 
   /// Compute the gradients of the loss values w.r.t. the prediction,
   /// which is (prediction-target)/batchsize
@@ -90,7 +90,7 @@ class SoftmaxCrossEntropy : public Loss<Tensor> {
   /// from Softmax(prediction).
   /// Users can call Average(const Tensor&) to get the average
   /// loss value over all samples in the batch.
-  Tensor Forward(const Tensor& prediction, const Tensor& target) override;
+  Tensor Forward(int flag, const Tensor& prediction, const Tensor& target) override;
 
   /// Compute the gradients of the loss values w.r.t. the prediction,
   /// which is: p[idx] - 1 if idx is the truth category's index; else,
@@ -106,5 +106,3 @@ class SoftmaxCrossEntropy : public Loss<Tensor> {
 }  // namespace singa
 
 #endif  // SINGA_MODEL_LOSS_H_
-
-

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/include/singa/model/metric.h
----------------------------------------------------------------------
diff --git a/include/singa/model/metric.h b/include/singa/model/metric.h
index d013fa4..b100435 100644
--- a/include/singa/model/metric.h
+++ b/include/singa/model/metric.h
@@ -48,7 +48,7 @@ class Metric {
 
   /// Comptue the metric value averaged over all samples (in a batch)
   float Evaluate(const Tensor& prediction, const T& target) {
-    const Tensor& metric = Forward(prediction, target);
+    const Tensor metric = Forward(prediction, target);
     return Sum<float>(metric) / (1.0f * metric.Size());
   }
 };

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/include/singa/model/optimizer.h
----------------------------------------------------------------------
diff --git a/include/singa/model/optimizer.h b/include/singa/model/optimizer.h
index a268126..2ec68fe 100644
--- a/include/singa/model/optimizer.h
+++ b/include/singa/model/optimizer.h
@@ -41,7 +41,7 @@ class Regularizer;
 class Optimizer {
  public:
   Optimizer() = default;
-
+  virtual ~Optimizer();
   /// Setup the optimzier using configurations from serialized string (for
   /// binding languages).
   void Setup(const string& str) {
@@ -51,7 +51,7 @@ class Optimizer {
   }
 
   /// Setup the meta fields of the optimizer
-  virtual void Setup(const OptimizerConf& conf) {}
+  virtual void Setup(const OptimizerConf& conf);
   /// Register the parameter, e.g., create Constraint and Regularizers.
   /// If there is no constraint or regularizer, then no need to register the
   /// parameter.
@@ -76,15 +76,21 @@ class Optimizer {
   void SetLearningRateGenerator(function<float(int)> func) {
     learning_rate_generator_ = func;
   }
-  /// Since Optimizer base layer has pure virtual function, a virtual
-  /// deconstructor is needed.
-  virtual ~Optimizer() = default;
+  float GetLearningRate(int step) {
+    if (learning_rate_generator_)
+      return learning_rate_generator_(step);
+    else
+      return 0;
+  }
 
  protected:
   function<float(int)> learning_rate_generator_;
   std::unordered_map<std::string, float> learning_rate_multplier_;
+  std::unordered_map<std::string, float> weight_decay_multplier_;
   std::unordered_map<std::string, Constraint*> constraints_;
   std::unordered_map<std::string, Regularizer*> regularizers_;
+  Constraint* constraint_ = nullptr;
+  Regularizer* regularizer_ = nullptr;
 };
 
 /// Apply constraints for parameters (gradient).
@@ -141,7 +147,7 @@ class Regularizer {
   /// e.g., clip each gradient if it is too large w.r.t the threshold,
   /// \ref
   /// https://www.reddit.com/r/MachineLearning/comments/31b6x8/gradient_clipping_rnns/
-  void Apply(int step, Tensor* grad, Tensor* value);
+  void Apply(int step, Tensor* grad, Tensor* value, float scale = 1.0f);
   /// Apply the regularizer for multiple parameter objects together.
   /// \ref https://github.com/Lasagne/Lasagne/blob/master/lasagne/updates.py
   void Apply(int step, const vector<Tensor*>& grads,

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/include/singa/utils/string.h
----------------------------------------------------------------------
diff --git a/include/singa/utils/string.h b/include/singa/utils/string.h
index cbfb28b..b4c7c24 100644
--- a/include/singa/utils/string.h
+++ b/include/singa/utils/string.h
@@ -51,6 +51,17 @@ inline int ArgPos(int argc, char** arglist, const char* arg) {
   return -1;
 }
 
+template<typename T>
+inline std::string VecToStr(const std::vector<T> & in) {
+  std::string out = "(";
+
+  for (auto x : in) {
+    out += std::to_string(x) + ", ";
+  }
+  out += ")";
+  return out;
+}
+
 /**
  * Tokenize a string.
  *

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/core/tensor/math_kernel.cu
----------------------------------------------------------------------
diff --git a/src/core/tensor/math_kernel.cu b/src/core/tensor/math_kernel.cu
index 4135ab8..13005af 100644
--- a/src/core/tensor/math_kernel.cu
+++ b/src/core/tensor/math_kernel.cu
@@ -265,6 +265,19 @@ __global__ void KernelLT(const size_t num, const float *in, const float x,
     out[idx] = in[idx] < x ? 1.0f : 0.0f;
   }
 }
+
+__global__ void KernelRowMax(const size_t nrow, const size_t ncol, const float *inPtr,
+    float *outPtr) {
+  for (size_t idx = blockIdx.x * blockDim.x + threadIdx.x; idx < nrow;
+       idx += blockDim.x * gridDim.x) {
+    int offset = idx * ncol;
+    float maxval = inPtr[offset];
+    for (size_t k = 1; k < ncol; k++) {
+      maxval = max(maxval, inPtr[offset + k]);
+    }
+    outPtr[idx] = maxval;
+  }
+}
 __global__ void KernelComputeCrossEntropy(const size_t batchsize,
                                           const size_t dim, const float *p,
                                           const int *t, float *loss) {
@@ -286,6 +299,9 @@ __global__ void KernelSoftmaxCrossEntropyBwd(const size_t batchsize,
     grad[pos] = p[pos] - 1.0f;  // TODO(wangwei) Consider p and grad are diff
   }
 }
+
+
+
 // ********************************
 // Functions call kernels
 // ********************************
@@ -421,6 +437,12 @@ void SoftmaxCrossEntropyBwd(size_t batchsize, const size_t dim, const float *p,
   KernelSoftmaxCrossEntropyBwd <<<ceil(batchsize / CU1DBLOCKF), CU1DBLOCKF>>>
       (batchsize, dim, p, t, grad);
 }
+
+void RowMax(const size_t nrow, const size_t ncol, const float *inPtr,
+    float *outPtr, cudaStream_t stream) {
+  KernelRowMax <<<ceil(nrow / CU1DBLOCKF), CU1DBLOCKF>>>(nrow, ncol, inPtr, outPtr);
+}
+
 /*
 void square_grad(int n, const float *in, float *out, cudaStream_t s) {
   kernel_square_grad <<<ceil(n / CU1DBLOCKF), CU1DBLOCKF>>> (in, out, n);

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/core/tensor/math_kernel.h
----------------------------------------------------------------------
diff --git a/src/core/tensor/math_kernel.h b/src/core/tensor/math_kernel.h
index d4087e5..63b0d82 100644
--- a/src/core/tensor/math_kernel.h
+++ b/src/core/tensor/math_kernel.h
@@ -98,6 +98,8 @@ void SoftmaxCrossEntropyBwd(const size_t batchsize, const size_t dim,
                             const float *p, const int *t, float *grad,
                             cudaStream_t stream);
 
+void RowMax(const size_t nrow, const size_t ncol, const float *inPtr,
+    float *outPtr, cudaStream_t stream);
 }  // cuda
 
 }  // namespace singa

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/core/tensor/tensor.cc
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor.cc b/src/core/tensor/tensor.cc
index 898cdc6..b07a23c 100644
--- a/src/core/tensor/tensor.cc
+++ b/src/core/tensor/tensor.cc
@@ -42,7 +42,8 @@ Tensor::Tensor(Shape &&shape, DataType dtype)
   device_ = defaultDevice;
   block_ = device_->NewBlock(Product(shape_) * SizeOf(data_type_));
 }
-Tensor::Tensor(const Shape &shape, std::shared_ptr<Device> device, DataType dtype)
+Tensor::Tensor(const Shape &shape, std::shared_ptr<Device> device,
+               DataType dtype)
     : data_type_(dtype), device_(device), shape_(shape) {
   block_ = device_->NewBlock(Product(shape_) * SizeOf(data_type_));
 }
@@ -68,11 +69,10 @@ Tensor::Tensor(Tensor &&in)
   in.block_ = nullptr;
 }
 
-void Tensor::SetBlock(Block* block) {
+void Tensor::SetBlock(Block *block) {
   LOG(WARNING) << "Pls avoid using this function, which may have side-effect.";
   if (block_ != nullptr)
-    if (block_->DecRefCount())
-      device_->FreeBlock(block_);
+    if (block_->DecRefCount()) device_->FreeBlock(block_);
   block_ = block;
 }
 
@@ -118,8 +118,7 @@ void Tensor::ToDevice(std::shared_ptr<Device> dst) {
   // TODO(wangwei) the comparison is very strict. May compare against device ID?
   if (device_ != dst) {
     Tensor tmp(shape_, dst, data_type_);
-    if (block_ != nullptr && Size())
-      tmp.CopyData(*this);
+    if (block_ != nullptr && Size()) tmp.CopyData(*this);
     if (block_ != nullptr && block_->DecRefCount() == 0)
       device_->FreeBlock(block_);
     block_ = tmp.block_;
@@ -132,13 +131,13 @@ void Tensor::ToHost() { ToDevice(device_->host()); }
 
 template <typename DType>
 void Tensor::CopyDataFromHostPtr(const DType *src, const size_t num,
-    const size_t offset) {
+                                 const size_t offset) {
   CHECK_EQ(sizeof(DType), SizeOf(data_type_))
       << "data_type is " << DataType_Name(data_type_)
       << " user given type is of size " << sizeof(DType);
   if (src != nullptr) {
     device_->CopyDataFromHostPtr(block(), src, sizeof(DType) * num,
-        sizeof(DType) * offset);
+                                 sizeof(DType) * offset);
   } else {
     LOG(WARNING) << "Copy data from null host ptr";
   }
@@ -161,8 +160,7 @@ void Tensor::CopyData(const Tensor &src) {
 }
 
 Tensor Tensor::Clone(std::shared_ptr<Device> device) const {
-  if (device == nullptr)
-    device = device_;
+  if (device == nullptr) device = device_;
   Tensor t(shape_, device_, data_type_);
   t.transpose_ = transpose_;
   t.CopyData(*this);
@@ -244,8 +242,6 @@ GenUnaryScalarArgMemberFn(operator+=, Add);
 GenUnaryScalarArgMemberFn(operator*=, EltwiseMult);
 GenUnaryScalarArgMemberFn(operator/=, Div);
 
-
-
 // ====================Tensor Operations=======================================
 void CopyDataToFrom(Tensor *dst, const Tensor &src, const size_t num,
                     const size_t dst_offset, const size_t src_offset) {
@@ -336,6 +332,18 @@ void CopyDataToFrom(Tensor *dst, const Tensor &src, const size_t num,
   } while (0)
 
 // =============Element-wise operations====================================
+float Tensor::L1() const {
+  float nrm = 0.0f;
+  TYPE_LANG_SWITCH(data_type_, DType, device_->lang(), Lang, {
+    device_->Exec([&nrm, this](Context *ctx) {
+      DType ret;
+      Asum<DType, Lang>(this->Size(), this->block(), &ret, ctx);
+      nrm = TypeCast<DType, float>(ret);
+    }, {this->block()}, {});
+  });
+  return nrm / Size();
+}
+
 /// L2 norm, Do not use Nrm2 (name conflict).
 float Tensor::L2() const {
   float nrm = 0.0f;
@@ -346,8 +354,10 @@ float Tensor::L2() const {
       nrm = TypeCast<DType, float>(ret);
     }, {this->block()}, {});
   });
-  return nrm;
+  return nrm / Size();
 }
+
+
 template <typename SType>
 void Tensor::SetValue(const SType x) {
   CHECK_EQ(sizeof(SType), SizeOf(data_type_));
@@ -525,18 +535,35 @@ Tensor SoftMax(const Tensor &in) {
   return out;
 }
 
+Tensor RowMax(const Tensor &in) {
+  Tensor ret({in.shape(0)}, in.device(), in.data_type());
+  TYPE_LANG_SWITCH(in.data_type(), DType, in.device()->lang(), Lang, {
+    in.device()->Exec([in, ret](Context *ctx) {
+      size_t nrow = 1;
+      if (in.nDim() > 1) nrow = in.shape(0);
+      size_t ncol = in.Size() / nrow;
+      RowMax<DType, Lang>(nrow, ncol, in.block(), ret.block(), ctx);
+    }, {in.block()}, {ret.block()});
+  });
+  return ret;
+}
+
 void SoftMax(const Tensor &in, Tensor *out) {
   CHECK_LE(in.nDim(), 2u);
-  Exp(in, out);
+  out->CopyData(in);
   size_t nrow = 1, ncol = in.Size(), size = ncol;
   if (in.nDim() == 2u) {
     nrow = in.shape(0);
     ncol = size / nrow;
     out->Reshape(Shape{nrow, ncol});
   }
-  Tensor sum(Shape{nrow}, in.device(), in.data_type());
-  SumColumns(*out, &sum);
-  DivColumn(sum, out);
+  Tensor tmp = RowMax(*out);
+  SubColumn(tmp, out);
+  Exp(*out, out);
+
+  SumColumns(*out, &tmp);
+  DivColumn(tmp, out);
+  out->Reshape(in.shape());
 }
 
 void AddColumn(const Tensor &v, Tensor *M) { AddColumn(1, 1, v, M); }
@@ -582,8 +609,8 @@ void AddRow(const SType alpha, const SType beta, const Tensor &v, Tensor *M) {
     Mult(alpha, one, vmat, beta, M);
   }
 }
-template
-void AddRow(const float alpha, const float beta, const Tensor &v, Tensor *M);
+template void AddRow(const float alpha, const float beta, const Tensor &v,
+                     Tensor *M);
 
 /// Divide column 'v' by each column of matrix M; write results into 'out'
 void DivColumn(const Tensor &v, Tensor *M) {
@@ -699,7 +726,7 @@ void MultRow(const Tensor &v, Tensor *M) {
   });
 }
 
-Tensor SliceRows(const Tensor& in, const size_t start, const size_t end) {
+Tensor SliceRows(const Tensor &in, const size_t start, const size_t end) {
   LOG(FATAL) << "Tensor::SliceRows is not implemented";
   Tensor ret;
   /*
@@ -788,6 +815,7 @@ void Gaussian(const SType mean, const SType std, Tensor *out) {
 template void Gaussian<float>(const float mean, const float std, Tensor *out);
 
 // ================Blas operations============================================
+
 template <typename SType>
 void Axpy(const SType alpha, const Tensor &in, Tensor *out) {
   TYPE_LANG_SWITCH(in.data_type(), DType, in.device()->lang(), Lang, {
@@ -869,5 +897,4 @@ void SoftmaxCrossEntropyBwd(const Tensor &t, Tensor *p) {
   });
 }
 
-
 }  // namespace singa

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/core/tensor/tensor_math.h
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor_math.h b/src/core/tensor/tensor_math.h
index 57ccb88..7732dd2 100644
--- a/src/core/tensor/tensor_math.h
+++ b/src/core/tensor/tensor_math.h
@@ -339,6 +339,11 @@ void SoftmaxCrossEntropyBwd(const size_t batchsize, const size_t dim,
   LOG(FATAL) << "Not Implemented";
 }
 
+template <typename DType, typename Lang>
+void RowMax(const size_t nrow, const size_t ncol, const Block *in,
+    const Block *ret, Context* ctx) {
+  LOG(FATAL) << "Not Implemented";
+}
 // **************************************
 // Matrix functions
 // **************************************

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/core/tensor/tensor_math_cpp.h
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor_math_cpp.h b/src/core/tensor/tensor_math_cpp.h
index 4717b5f..3e0c8ad 100644
--- a/src/core/tensor/tensor_math_cpp.h
+++ b/src/core/tensor/tensor_math_cpp.h
@@ -549,6 +549,20 @@ void SoftmaxCrossEntropyBwd<float, lang::Cpp>(const size_t batchsize,
   }
 }
 
+template <>
+void RowMax<float, lang::Cpp>(const size_t nrow, const size_t ncol,
+                              const Block *in, const Block *out, Context *ctx) {
+  const float *inPtr = static_cast<const float *>(in->data());
+  float *outPtr = static_cast<float *>(out->mutable_data());
+  for (size_t r = 0; r < nrow; r++) {
+    int offset = r * ncol;
+    float maxval = inPtr[offset];
+    for (size_t c = 1; c < ncol; c++)
+      maxval = std::max(maxval, inPtr[offset + c]);
+    outPtr[r] = maxval;
+  }
+}
+
 // =========Matrix operations ================================================
 /*
 template <>

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/core/tensor/tensor_math_cuda.h
----------------------------------------------------------------------
diff --git a/src/core/tensor/tensor_math_cuda.h b/src/core/tensor/tensor_math_cuda.h
index 67ee861..43bfa1b 100644
--- a/src/core/tensor/tensor_math_cuda.h
+++ b/src/core/tensor/tensor_math_cuda.h
@@ -421,6 +421,15 @@ void SoftmaxCrossEntropyBwd<float, lang::Cuda>(const size_t batchsize,
   cuda::SoftmaxCrossEntropyBwd(batchsize, dim, pPtr, tPtr, gradPtr,
                                ctx->stream);
 }
+
+template <>
+void RowMax<float, lang::Cuda>(const size_t nrow, const size_t ncol,
+                               const Block* in, const Block* out,
+                               Context* ctx) {
+  const float* inPtr = static_cast<const float*>(in->data());
+  float* outPtr = static_cast<float*>(out->mutable_data());
+  cuda::RowMax(nrow, ncol, inPtr, outPtr, ctx->stream);
+}
 }  // namespace singa
 
 #endif  // USE_CUDA

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/model/feed_forward_net.cc
----------------------------------------------------------------------
diff --git a/src/model/feed_forward_net.cc b/src/model/feed_forward_net.cc
index a24d36a..e682918 100644
--- a/src/model/feed_forward_net.cc
+++ b/src/model/feed_forward_net.cc
@@ -1,22 +1,26 @@
-/**
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements.  See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership.  The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License.  You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
+/************************************************************
+*
+* Licensed to the Apache Software Foundation (ASF) under one
+* or more contributor license agreements.  See the NOTICE file
+* distributed with this work for additional information
+* regarding copyright ownership.  The ASF licenses this file
+* to you under the Apache License, Version 2.0 (the
+* "License"); you may not use this file except in compliance
+* with the License.  You may obtain a copy of the License at
+*
+*   http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing,
+* software distributed under the License is distributed on an
+* "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+* KIND, either express or implied.  See the License for the
+* specific language governing permissions and limitations
+* under the License.
+*
+*************************************************************/
 
 #include "singa/model/feed_forward_net.h"
+#include "singa/model/initializer.h"
 #include "singa/utils/logging.h"
 #include "singa/utils/channel.h"
 namespace singa {
@@ -37,12 +41,15 @@ Layer* FeedForwardNet::Add(const LayerConf& conf, const Shape* sample_shape) {
   return layer;
 }
 
-Layer* FeedForwardNet::Add(Layer* layer, const LayerConf& conf, const Shape* sample_shape) {
+Layer* FeedForwardNet::Add(Layer* layer, const LayerConf& conf,
+                           const Shape* sample_shape) {
+  CHECK(conf.has_name()) << "Must set layer name";
   if (sample_shape == nullptr)
     layer->Setup(layers_.back()->GetOutputSampleShape(), conf);
   else
     layer->Setup(*sample_shape, conf);
   Add(layer);
+  LOG(INFO) << layer->name() << VecToStr(layer->GetOutputSampleShape());
   return layer;
 }
 
@@ -75,12 +82,19 @@ void FeedForwardNet::Compile(bool shuffle, Optimizer* opt, Loss<Tensor>* loss,
   opt_ = opt;
   loss_ = loss;
   metric_ = metric;
-  // init params and register them to sgd
+  const auto specs = GetParamSpecs();
+  const auto params = GetParamValues();
+  CHECK_EQ(specs.size(), params.size());
+  for (size_t k = 0; k < specs.size(); k++) {
+    opt_->Register(specs[k].name(), specs[k]);
+    auto init = CreateInitializer(specs[k].filler());
+    init->Fill(params[k]);
+    LOG(INFO) << specs[k].name() << " : " << params[k]->L1();
+  }
 }
 
 void FeedForwardNet::ToDevice(std::shared_ptr<Device> device) {
-  for (auto layer: layers_)
-    layer->ToDevice(device);
+  for (auto layer : layers_) layer->ToDevice(device);
   /*
   opt_->ToDevice(device);
   loss_->ToDevice(device);
@@ -129,7 +143,6 @@ void FeedForwardNet::Train(size_t batchsize, int nb_epoch, const Tensor& x,
 void FeedForwardNet::Train(size_t batchsize, int nb_epoch, const Tensor& x,
                            const Tensor& y, const Tensor& val_x,
                            const Tensor& val_y) {
-  InitNetParams();
   CHECK_EQ(x.shape(0), y.shape(0)) << "Diff num of sampels in x and y";
   int num_extra_samples = x.shape(0) % batchsize;
   if (num_extra_samples != 0)
@@ -137,13 +150,18 @@ void FeedForwardNet::Train(size_t batchsize, int nb_epoch, const Tensor& x,
   Channel* train_ch = GetChannel("train_perf");
   train_ch->EnableDestStderr(true);
   Channel* val_ch = GetChannel("val_perf");
+  val_ch->EnableDestStderr(true);
+  std::vector<size_t> index;
+  for (size_t i = 0; i < x.shape(0) / batchsize; i++) index.push_back(i);
   for (int epoch = 0; epoch < nb_epoch; epoch++) {
+    if (shuffle_) std::random_shuffle(index.begin(), index.end());
     float loss = 0.0f, metric = 0.0f;
     size_t b = 0;
     for (; b < x.shape(0) / batchsize; b++) {
-      const Tensor bx = CopyRows(x, b * batchsize, b * batchsize + batchsize);
-      const Tensor by = CopyRows(y, b * batchsize, b * batchsize + batchsize);
-      const auto ret = TrainOnBatch(bx, by);
+      size_t idx = index[b];
+      const Tensor bx = CopyRows(x, idx * batchsize, (idx + 1) * batchsize);
+      const Tensor by = CopyRows(y, idx * batchsize, (idx + 1) * batchsize);
+      const auto ret = TrainOnBatch(epoch, bx, by);
       loss += ret.first;
       metric += ret.second;
     }
@@ -151,7 +169,8 @@ void FeedForwardNet::Train(size_t batchsize, int nb_epoch, const Tensor& x,
     metric /= b;
     train_ch->Send("Epoch " + std::to_string(epoch) + ", training loss = " +
                    std::to_string(loss) + ", accuracy = " +
-                   std::to_string(metric));
+                   std::to_string(metric) + ", lr = " +
+                   std::to_string(opt_->GetLearningRate(epoch)));
     if (val_x.Size() && val_y.Size()) {
       const auto val_perf = Evaluate(val_x, val_y, batchsize);
       val_ch->Send("Epoch " + std::to_string(epoch) + ", val loss = " +
@@ -162,22 +181,28 @@ void FeedForwardNet::Train(size_t batchsize, int nb_epoch, const Tensor& x,
   }
 }
 
-const std::pair<float, float> FeedForwardNet::TrainOnBatch(const Tensor& x,
+const std::pair<float, float> FeedForwardNet::TrainOnBatch(int epoch,
+                                                           const Tensor& x,
                                                            const Tensor& y) {
   int flag = kTrain;
   const Tensor fea = Forward(flag, x);
-  float loss = loss_->Evaluate(fea, y);
+  float loss = loss_->Evaluate(flag, fea, y);
   float metric = metric_->Evaluate(fea, y);
   const Tensor grad = loss_->Backward();
-  const auto grads = Backward(kTrain, grad);
+  auto grads = Backward(kTrain, grad / static_cast<float>(x.shape(0)));
+  auto names = GetParamNames();
+  auto values = GetParamValues();
+  for (size_t k = 0; k < grads.size(); k++) {
+    opt_->Apply(epoch, names[k], &grads[k], values.at(k));
+  }
   return std::make_pair(loss, metric);
 }
 
 const Tensor FeedForwardNet::Forward(int flag, const Tensor& data) {
   Tensor input = data, output;
   for (auto layer : layers_) {
-//    LOG(INFO) << layer->name();
     output = layer->Forward(flag, input);
+    // LOG(INFO) << layer->name() << ": " << output.L2();
     input = output;
   }
   return output;
@@ -185,13 +210,22 @@ const Tensor FeedForwardNet::Forward(int flag, const Tensor& data) {
 
 const vector<Tensor> FeedForwardNet::Backward(int flag, const Tensor& grad) {
   vector<Tensor> param_grads;
+  std::stack<Tensor> buf;
   Tensor tmp = grad;
   for (int i = layers_.size() - 1; i >= 0; i--) {
- //   LOG(INFO) << layers_.at(i)->name();
+    // LOG(INFO) << layers_.at(i)->name() << " : " << tmp.L2();
     auto ret = layers_.at(i)->Backward(flag, tmp);
     tmp = ret.first;
-    if (ret.second.size())
-      for (const auto x : ret.second) param_grads.push_back(x);
+    if (ret.second.size()) {
+      for (int k = ret.second.size() - 1; k >= 0; k--) {
+        buf.push(ret.second[k]);
+        // LOG(INFO) <<  "      " << buf.top().L2();
+      }
+    }
+  }
+  while (!buf.empty()) {
+    param_grads.push_back(buf.top());
+    buf.pop();
   }
   return param_grads;
 }
@@ -230,8 +264,8 @@ std::pair<Tensor, Tensor> FeedForwardNet::EvaluateOnBatch(const Tensor& x,
   int flag = kEval;
   const Tensor fea = Forward(flag, x);
   const Tensor m = metric_->Forward(fea, y);
-  const Tensor l = loss_->Forward(fea, y);
-  return std::make_pair(m, l);
+  const Tensor l = loss_->Forward(flag, fea, y);
+  return std::make_pair(l, m);
 }
 
 const Tensor FeedForwardNet::Predict(const Tensor& x, size_t batchsize) {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/model/layer/cudnn_convolution.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/cudnn_convolution.cc b/src/model/layer/cudnn_convolution.cc
index eb507b2..3dca28a 100644
--- a/src/model/layer/cudnn_convolution.cc
+++ b/src/model/layer/cudnn_convolution.cc
@@ -72,8 +72,8 @@ void CudnnConvolution::InitCudnn(const Tensor &input) {
       num_filters_, conv_height_, conv_width_));
   if (bias_term_)
     CUDNN_CHECK(cudnnSetTensor4dDescriptor(bias_desc_, CUDNN_TENSOR_NCHW,
-                                           GetCudnnDataType(dtype), 1, 1,
-                                           num_filters_, 1));
+                                           GetCudnnDataType(dtype), 1, 1, 1,
+                                           num_filters_));
   CUDNN_CHECK(cudnnSetConvolution2dDescriptor(conv_desc_, pad_h_, pad_w_,
                                               stride_h_, stride_w_, 1, 1,
                                               CUDNN_CROSS_CORRELATION));
@@ -244,6 +244,7 @@ const std::pair<Tensor, vector<Tensor>> CudnnConvolution::Backward(
   }, {grad.block(), weight_.block()}, {dx.block(), workspace_.block()});
   param_grad.push_back(dw);
   param_grad.push_back(db);
+  LOG(INFO) << "bias nrm " << db.L1();
   return std::make_pair(dx, param_grad);
 }
 

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/model/layer/cudnn_dropout.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/cudnn_dropout.cc b/src/model/layer/cudnn_dropout.cc
index f9b9dbf..ab83226 100644
--- a/src/model/layer/cudnn_dropout.cc
+++ b/src/model/layer/cudnn_dropout.cc
@@ -108,7 +108,7 @@ const std::pair<Tensor, vector<Tensor>> CudnnDropout::Backward(
 }
 void CudnnDropout::ToDevice(std::shared_ptr<Device> device) {
   Dropout::ToDevice(device);
-  state.ToDevice(device);
+  state_.ToDevice(device);
 }
 }  // namespace singa
 #endif  // CUDNN_VERSION_MAJOR>=5

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/model/layer/dense.cc
----------------------------------------------------------------------
diff --git a/src/model/layer/dense.cc b/src/model/layer/dense.cc
index c6a9f8a..338409c 100644
--- a/src/model/layer/dense.cc
+++ b/src/model/layer/dense.cc
@@ -41,13 +41,15 @@ void Dense::Setup(const Shape& in_sample, const LayerConf &conf) {
   bias_.Reshape(Shape{hdim_});
   param_values_.push_back(&weight_);
   param_values_.push_back(&bias_);
+  for (auto specs: conf.param())
+    param_specs_.push_back(specs);
 }
 
 /// \copydoc Layer::Forward(int flag, const Tensor&)
 const Tensor Dense::Forward(int flag, const Tensor &input) {
   CHECK(buf_.empty());
   Tensor output;
-  CHECK_EQ(input.nDim(), 2);
+  CHECK_EQ(input.nDim(), 2u);
   if (transpose_)  // use the transposed version of weight_ for computing
     output = Mult(input, weight_);
   else

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/model/loss/mse.cc
----------------------------------------------------------------------
diff --git a/src/model/loss/mse.cc b/src/model/loss/mse.cc
index a4bbb72..6e19059 100644
--- a/src/model/loss/mse.cc
+++ b/src/model/loss/mse.cc
@@ -20,7 +20,7 @@
 
 namespace singa {
 
-Tensor MSE::Forward(const Tensor& prediction, const Tensor& target) {
+Tensor MSE::Forward(int flag, const Tensor& prediction, const Tensor& target) {
   CHECK(buf_.empty()) << "Do not call Forward successively for more than twice."
                       << " The calling pattern is [Forward|Evaluate] Backward";
   Tensor t = prediction - target;
@@ -28,7 +28,8 @@ Tensor MSE::Forward(const Tensor& prediction, const Tensor& target) {
   if (t.nDim() > 1) batchsize = t.shape().at(0);
   size_t dim = t.Size() / batchsize;
   t.Reshape(Shape{batchsize, dim});
-  buf_.push(t);
+  if (kTrain & flag)
+    buf_.push(t);
   // TODO(wangwei) use CastType for operator/
   return Sum(Square(t), 1) * 0.5f;
 }

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/model/loss/softmax_cross_entropy.cc
----------------------------------------------------------------------
diff --git a/src/model/loss/softmax_cross_entropy.cc b/src/model/loss/softmax_cross_entropy.cc
index bed3348..3411fbe 100644
--- a/src/model/loss/softmax_cross_entropy.cc
+++ b/src/model/loss/softmax_cross_entropy.cc
@@ -21,7 +21,7 @@
 
 namespace singa {
 
-Tensor SoftmaxCrossEntropy::Forward(const Tensor& prediction,
+Tensor SoftmaxCrossEntropy::Forward(int flag, const Tensor& prediction,
                                     const Tensor& target) {
   CHECK(buf_.empty()) << "Do not call Forward successively for more than twice."
                       << " The calling pattern is [Forward|Evaluate] Backward";
@@ -30,13 +30,17 @@ Tensor SoftmaxCrossEntropy::Forward(const Tensor& prediction,
   size_t dim = prediction.Size() / batchsize;
   const Tensor& input = Reshape(prediction, Shape{batchsize, dim});
   Tensor prob = SoftMax(input);
+  // LOG(INFO) << "prob: " << prob.L2();
 
   // buffer intermediate data
-  buf_.push(prob);
-  buf_.push(target);
+  if (flag & kTrain) {
+    buf_.push(prob);
+    buf_.push(target);
+  }
   Tensor loss(Shape{batchsize}, prob.device(), prob.data_type());
 
   ComputeCrossEntropy(prob, target, &loss);
+
   return loss;
 }
 
@@ -50,4 +54,3 @@ Tensor SoftmaxCrossEntropy::Backward() {
 }
 }  // namespace singa
 
-

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/model/metric/accuracy.cc
----------------------------------------------------------------------
diff --git a/src/model/metric/accuracy.cc b/src/model/metric/accuracy.cc
index 1b667b1..ffda938 100644
--- a/src/model/metric/accuracy.cc
+++ b/src/model/metric/accuracy.cc
@@ -30,6 +30,7 @@ Tensor Accuracy::Match(const Tensor& predict, const vector<int>& target) {
   // TODO(wangwei) CloneToDevice(host);
   const float* prob = prediction.data<float>();
   float* score = new float[batchsize];
+  memset(score, 0, batchsize * sizeof(float));
   for (size_t b = 0; b < batchsize; b++) {
     vector<std::pair<float, int>> prob_class;
     for (size_t c = 0; c < nb_classes; c++) {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/model/optimizer/optimizer.cc
----------------------------------------------------------------------
diff --git a/src/model/optimizer/optimizer.cc b/src/model/optimizer/optimizer.cc
index c9e7a72..9be47c8 100644
--- a/src/model/optimizer/optimizer.cc
+++ b/src/model/optimizer/optimizer.cc
@@ -21,6 +21,17 @@
 
 namespace singa {
 
+Optimizer::~Optimizer() {
+  for (auto entry : regularizers_) delete entry.second;
+  for (auto entry : constraints_) delete entry.second;
+  if (constraint_ != nullptr) delete constraint_;
+  if (regularizer_ != nullptr) delete regularizer_;
+}
+void Optimizer::Setup(const OptimizerConf& conf) {
+  if (conf.has_regularizer())
+    regularizer_ = new Regularizer(conf.regularizer());
+  if (conf.has_constraint()) constraint_ = new Constraint(conf.constraint());
+}
 void Optimizer::Register(const string& name, const ParamSpec& specs) {
   if (specs.has_constraint()) {
     CHECK(constraints_.find(name) == constraints_.end())
@@ -32,6 +43,11 @@ void Optimizer::Register(const string& name, const ParamSpec& specs) {
         << "Parameter with name = " << name << " has already registered";
     regularizers_[name] = new Regularizer(specs.regularizer());
   }
+  if (specs.has_decay_mult()) {
+    CHECK(weight_decay_multplier_.find(name) == weight_decay_multplier_.end())
+        << "Parameter with name = " << name << " has already registered";
+    weight_decay_multplier_[name] = specs.decay_mult();
+  }
   if (specs.has_lr_mult()) {
     CHECK(learning_rate_multplier_.find(name) == learning_rate_multplier_.end())
         << "Parameter with name = " << name << " has already registered";
@@ -47,10 +63,18 @@ void Optimizer::Register(const string& name, const ParamSpec& specs) {
 void Optimizer::Apply(int step, const string& name, Tensor* grad,
                       Tensor* param) {
   // TODO(wangwei) need to consider the order of constraint and regularizer
-  if (regularizers_.find(name) != regularizers_.end())
+  if (regularizers_.find(name) != regularizers_.end()) {
     regularizers_.at(name)->Apply(step, param, grad);
+  } else if (regularizer_ != nullptr) {
+    float scale = 1.0f;
+    if (weight_decay_multplier_.find(name) != weight_decay_multplier_.end())
+      scale = weight_decay_multplier_.at(name);
+    regularizer_->Apply(step, param, grad, scale);
+  }
   if (constraints_.find(name) != constraints_.end())
     constraints_.at(name)->Apply(step, param, grad);
+  else if (constraint_ != nullptr)
+    constraint_->Apply(step, param, grad);
   float lr = learning_rate_generator_(step);
   if (learning_rate_multplier_.find(name) != learning_rate_multplier_.end())
     lr *= learning_rate_multplier_.at(name);
@@ -62,9 +86,9 @@ void Regularizer::Setup(const RegularizerConf& conf) {
   coefficient_ = conf.coefficient();
 }
 
-void Regularizer::Apply(int step, Tensor* value, Tensor* grad) {
+void Regularizer::Apply(int step, Tensor* value, Tensor* grad, float scale) {
   if (type_ == "L2" || type_ == "l2") {
-    (*grad) -= (*value) * coefficient_;
+    Axpy(coefficient_ * scale, *value, grad);
   } else {
     CHECK(type_ == "NotSet") << "Unknown regularizer type = " << type_;
   }

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/model/optimizer/sgd.cc
----------------------------------------------------------------------
diff --git a/src/model/optimizer/sgd.cc b/src/model/optimizer/sgd.cc
index a5c66a1..71071ff 100644
--- a/src/model/optimizer/sgd.cc
+++ b/src/model/optimizer/sgd.cc
@@ -22,6 +22,7 @@
 namespace singa {
 
 void SGD::Setup(const OptimizerConf& conf) {
+  Optimizer::Setup(conf);
   if (conf.has_momentum()) {
     float m = conf.momentum();
     SetMomentumGenerator([m](int step) { return m; });

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/src/proto/model.proto
----------------------------------------------------------------------
diff --git a/src/proto/model.proto b/src/proto/model.proto
index c06deec..b1318d9 100644
--- a/src/proto/model.proto
+++ b/src/proto/model.proto
@@ -89,6 +89,11 @@ message OptimizerConf {
 
   // delta is used to avoid dividing zero
   optional float delta = 6 [default = 1e-8];
+
+  // global regularizer lower priority than ParamSpec regularizer
+  optional RegularizerConf regularizer = 10;
+  // global constraint lower priority than ParamSpec constraint
+  optional ConstraintConf constraint = 11;
 }
 
 message ConstraintConf {

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/test/singa/test_cross_entropy.cc
----------------------------------------------------------------------
diff --git a/test/singa/test_cross_entropy.cc b/test/singa/test_cross_entropy.cc
index d73591f..c7fa2fb 100644
--- a/test/singa/test_cross_entropy.cc
+++ b/test/singa/test_cross_entropy.cc
@@ -44,7 +44,7 @@ TEST_F(TestSoftmaxCrossEntropy, CppForward) {
   t.CopyDataFromHostPtr(tdat, 2);
 
   singa::SoftmaxCrossEntropy cross_entropy;
-  const Tensor& loss = cross_entropy.Forward(p, t);
+  const Tensor& loss = cross_entropy.Forward(singa::kEval, p, t);
   auto ldat = loss.data<float>();
 
   const float result_test = -log(0.25);
@@ -58,7 +58,7 @@ TEST_F(TestSoftmaxCrossEntropy, CppBackward) {
   t.CopyDataFromHostPtr(tdat, 2);
 
   singa::SoftmaxCrossEntropy cross_entropy;
-  cross_entropy.Forward(p, t);
+  cross_entropy.Forward(singa::kTrain, p, t);
   const Tensor& grad = cross_entropy.Backward();
 
   auto gdat = grad.data<float>();
@@ -82,7 +82,7 @@ TEST_F(TestSoftmaxCrossEntropy, CudaForward) {
   p.CopyDataFromHostPtr(pdat, 8);
   t.CopyDataFromHostPtr(tdat, 2);
 
-  Tensor loss = cross_entropy.Forward(p, t);
+  Tensor loss = cross_entropy.Forward(singa::kEval, p, t);
   loss.ToHost();
   auto ldat = loss.data<float>();
 
@@ -99,7 +99,7 @@ TEST_F(TestSoftmaxCrossEntropy, CudaBackward) {
   p.CopyDataFromHostPtr(pdat, 8);
   t.CopyDataFromHostPtr(tdat, 2);
 
-  cross_entropy.Forward(p, t);
+  cross_entropy.Forward(singa::kTrain, p, t);
   Tensor grad = cross_entropy.Backward();
 
   grad.ToHost();

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/test/singa/test_dense.cc
----------------------------------------------------------------------
diff --git a/test/singa/test_dense.cc b/test/singa/test_dense.cc
index 363fb6e..e80384f 100644
--- a/test/singa/test_dense.cc
+++ b/test/singa/test_dense.cc
@@ -207,7 +207,7 @@ TEST(Dense, BackwardCuda) {
   singa::Tensor grad(singa::Shape{batchsize, hdim}, cuda);
   grad.CopyDataFromHostPtr(dy, batchsize * hdim);
 
-  const auto ret = dense.Backward(singa::kTrain, grad);
+  auto ret = dense.Backward(singa::kTrain, grad);
   singa::Tensor in_grad = ret.first;
   singa::Tensor dweight = ret.second.at(0);
   singa::Tensor dbias = ret.second.at(1);

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/test/singa/test_mse.cc
----------------------------------------------------------------------
diff --git a/test/singa/test_mse.cc b/test/singa/test_mse.cc
index 788652f..640caf4 100644
--- a/test/singa/test_mse.cc
+++ b/test/singa/test_mse.cc
@@ -42,7 +42,7 @@ class TestMSE : public ::testing::Test {
 #ifdef USE_CBLAS
 TEST_F(TestMSE, CppForward) {
   singa::MSE mse;
-  const Tensor& loss = mse.Forward(p, t);
+  const Tensor& loss = mse.Forward(singa::kEval, p, t);
   auto ldat = loss.data<float>();
 
   for (size_t i = 0, k = 0; i < loss.Size(); i++) {
@@ -57,7 +57,7 @@ TEST_F(TestMSE, CppForward) {
 
 TEST_F(TestMSE, CppBackward) {
   singa::MSE mse;
-  mse.Forward(p, t);
+  mse.Forward(singa::kTrain, p, t);
   const Tensor& grad = mse.Backward();
 
   auto gdat = grad.data<float>();
@@ -72,7 +72,7 @@ TEST_F(TestMSE, CudaForward) {
   auto dev = std::make_shared<singa::CudaGPU>();
   p.ToDevice(dev);
   t.ToDevice(dev);
-  Tensor loss = mse->Forward(p, t);
+  Tensor loss = mse->Forward(singa::kEval, p, t);
 
   loss.ToHost();
   auto ldat = loss.data<float>();
@@ -94,7 +94,7 @@ TEST_F(TestMSE, CudaBackward) {
   auto dev = std::make_shared<singa::CudaGPU>();
   p.ToDevice(dev);
   t.ToDevice(dev);
-  mse.Forward(p, t);
+  mse.Forward(singa::kTrain, p, t);
   Tensor grad = mse.Backward();
   grad.ToHost();
   auto gdat = grad.data<float>();

http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/71eb059c/test/singa/test_tensor_math.cc
----------------------------------------------------------------------
diff --git a/test/singa/test_tensor_math.cc b/test/singa/test_tensor_math.cc
index f8d0351..2a0df0d 100644
--- a/test/singa/test_tensor_math.cc
+++ b/test/singa/test_tensor_math.cc
@@ -346,7 +346,7 @@ TEST_F(TestTensorMath, L2Cpp) {
   float l2 = a.L2();
   float target = 0.0f;
   for (size_t i = 0; i < a.Size(); i++) target += dat1[i] * dat1[i];
-  EXPECT_FLOAT_EQ(l2, sqrt(target));
+  EXPECT_FLOAT_EQ(l2, sqrt(target) / a.Size());
 }
 TEST_F(TestTensorMath, MultCpp) {
   const float x[4] = {1.0f, 2.0f, 3.0f, 4.0f};
@@ -514,7 +514,7 @@ TEST_F(TestTensorMath, L2Cuda) {
   float l2 = t.L2();
   float target = 0.0f;
   for (size_t i = 0; i < t.Size(); i++) target += dat1[i] * dat1[i];
-  EXPECT_FLOAT_EQ(l2, sqrt(target));
+  EXPECT_FLOAT_EQ(l2, sqrt(target) / t.Size());
 }
 TEST_F(TestTensorMath, MultCuda) {
   const float x[4] = {1.0f, 2.0f, 3.0f, 4.0f};