Merge branch 'develop' into maxseq

doc/design/model_format.md

@@ -12,24 +12,22 @@ The topology is saved as a plain text in a detailed self-contain protobuf file.
 
 The parameters are saved as a binary file. As we all know, the protobuf message has a limit of [64M size](https://developers.google.com/protocol-buffers/docs/reference/cpp/google.protobuf.io.coded_stream#CodedInputStream.SetTotalBytesLimit.details). We have done a [benchmark experiment](https://github.com/PaddlePaddle/Paddle/pull/4610), which shows that protobuf is not fit for the task.
 
-As a result, we design a particular format for tensor serialization. By default, an arbitrary tensor in Paddle is a [LoDTensor](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/lod_tensor.md), and has a description information proto of [LoDTensorDesc](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/framework.proto#L99). We save the DescProto as the byte string header. It contains all the necessary information, such as the `dims`, the `name` of the tensor, and the `LoD` information in [LoDTensor](https://github.com/PaddlePaddle/Paddle/blob/1c0a4c901c9fc881d120249c703b15d1c50dae7d/paddle/framework/lod_tensor.md). A tensor stores values in a continuous memory buffer. For speed we dump the raw memory to disk and save it as the byte string content. So, the binary format of one tensor is, 
-
-|HeaderLength|ContentLength|**LoDTensorDesc**|**TensorValue**|
+As a result, we design a particular format for tensor serialization. By default, an arbitrary tensor in Paddle is a [LoDTensor](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/lod_tensor.md), and has a description information proto of [LoDTensorDesc](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/framework.proto#L99). We save the DescProto as the byte string header. It contains all the necessary information, such as the `dims`, and the `LoD` information in [LoDTensor](https://github.com/PaddlePaddle/Paddle/blob/1c0a4c901c9fc881d120249c703b15d1c50dae7d/paddle/framework/lod_tensor.md). A tensor stores values in a continuous memory buffer. For speed we dump the raw memory to disk and save it as the byte string content. So, the binary format of one tensor is, 
 
 The table below shows a tensor's byte view in detail. Note that all the signed values are written in the little-endian format.
 
-```text
-[offset] [type]              [description] 
-0004     4 bytes integer      HeaderLength, the length of LoDTensorDesc
-0008     4 bytes integer      ContentLength, the length of LodTensor Buffer
-0009     1 bytes char         TensorDesc
-00010    1 bytes char         TensorDesc
-...
-00100    1 bytes char         TensorValue
-00101    1 bytes char         TensorValue
-00102    1 bytes char         TensorValue              ..
-...
-```
+|field name  | type | description |
+| --- | --- | --- |
+| version | uint32_t | Version of saved file. Always 0 now. |
+| tensor desc length | uint32_t | TensorDesc(Protobuf message) length in bytes. |
+| tensor desc | void* | TensorDesc protobuf binary message |
+| tensor data | void* | Tensor's data in binary format. The length of `tensor_data` is decided by `TensorDesc.dims()` and `TensorDesc.data_type()` |
+| lod_level | uint64_t | Level of LoD |
+| length of lod[0] | uint64_t | [Optional] length of lod[0] in bytes. |
+| data of lod[0] | uint64_t*  | [Optional] lod[0].data() |
+| ... | ... | ... |
+
+
 
 ## Summary
 

go/cmd/pserver/pserver.go

@@ -67,7 +67,7 @@ func main() {
 		cp, err = pserver.LoadCheckpoint(e, idx)
 		if err != nil {
 			if err == pserver.ErrCheckpointNotFound {
-				log.Info("Could not find the pserver checkpoint.")
+				log.Info("load checkpoint error", "error", err)
 			} else {
 				panic(err)
 			}
@@ -99,7 +99,7 @@ func main() {
 	candy.Must(err)
 
 	go func() {
-		log.Info("starting pserver", log.Ctx{"port": *port})
+		log.Info("serving pserver", log.Ctx{"port": *port})
 		err = http.Serve(l, nil)
 		candy.Must(err)
 	}()

go/pserver/optimizer.go

@@ -71,9 +71,15 @@ func newOptimizer(paramWithConfigs ParameterWithConfig, State []byte) *optimizer
 		cstate = unsafe.Pointer(&s[0])
 	}
 
+	var cptr (*C.uchar)
+	if len(c) > 0 {
+		cptr = (*C.uchar)(&c[0])
+	} else {
+		log.Error("empty config", "param name", paramWithConfigs.Param.Name)
+	}
 	o.config = c
 	o.opt = C.paddle_create_optimizer(
-		(*C.uchar)(&c[0]),
+		cptr,
 		C.int(len(c)),
 		C.paddle_element_type(p.ElementType),
 		cbuffer,

go/pserver/service.go

@@ -17,12 +17,11 @@ package pserver
 import (
 	"bufio"
 	"bytes"
-	"crypto/md5"
 	"encoding/gob"
-	"encoding/hex"
 	"encoding/json"
 	"errors"
 	"fmt"
+	"hash/crc32"
 	"io/ioutil"
 	"os"
 	"path"
@@ -40,7 +39,7 @@ type ElementType int
 
 // ErrCheckpointNotFound indicates that the pserver checkpoint could
 // not be found.
-var ErrCheckpointNotFound = errors.New("checkpoint not found")
+var ErrCheckpointNotFound = errors.New("checkpoint not found in etcd")
 
 // RPC error message.
 const (
@@ -76,7 +75,7 @@ type ParameterWithConfig struct {
 type checkpointMeta struct {
 	UUID      string `json:"uuid"`
 	Path      string `json:"path"`
-	MD5       string `json:"md5"`
+	CRC32     uint32 `json:"crc32"`
 	Timestamp int64  `json:"timestamp"`
 }
 
@@ -92,7 +91,7 @@ type Service struct {
 	idx                int
 	checkpointInterval time.Duration
 	checkpointPath     string
-	client             *EtcdClient
+	client             KVStore
 
 	mu     sync.Mutex
 	optMap map[string]*optimizer
@@ -104,7 +103,12 @@ type parameterCheckpoint struct {
 	State []byte
 }
 
-func loadMeta(e *EtcdClient, idx int) (meta checkpointMeta, err error) {
+type KVStore interface {
+	GetKey(key string, timeout time.Duration) ([]byte, error)
+	PutKey(key string, value []byte, timeout time.Duration, withLease bool) error
+}
+
+func loadMeta(e KVStore, idx int) (meta checkpointMeta, err error) {
 	v, err := e.GetKey(PsCheckpoint+strconv.Itoa(idx), 3*time.Second)
 	if err != nil {
 		return
@@ -123,7 +127,7 @@ func loadMeta(e *EtcdClient, idx int) (meta checkpointMeta, err error) {
 }
 
 // LoadCheckpoint loads checkpoint from file.
-func LoadCheckpoint(e *EtcdClient, idx int) (Checkpoint, error) {
+func LoadCheckpoint(e KVStore, idx int) (Checkpoint, error) {
 	log.Info("Loading checkpoint", "pserver index", idx)
 	defer traceTime(time.Now(), "load checkpoint")
 
@@ -137,11 +141,8 @@ func LoadCheckpoint(e *EtcdClient, idx int) (Checkpoint, error) {
 		return nil, err
 	}
 
-	// TODO(helin): change MD5 to CRC since CRC is better for file
-	// checksum in our use case (emphasize speed over security).
-	h := md5.New()
-	md5 := hex.EncodeToString(h.Sum(content))
-	if md5 != cpMeta.MD5 {
+	crc32 := crc32.ChecksumIEEE(content)
+	if crc32 != cpMeta.CRC32 {
 		return nil, errors.New(WrongChecksum)
 	}
 
@@ -150,12 +151,13 @@ func LoadCheckpoint(e *EtcdClient, idx int) (Checkpoint, error) {
 	if err = dec.Decode(&cp); err != nil {
 		return nil, err
 	}
+
 	return cp, nil
 }
 
 // NewService creates a new service, will bypass etcd registration if no
 // endpoints specified. It will recovery from checkpoint file if a exists a specified checkpoint.
-func NewService(idx int, interval time.Duration, path string, client *EtcdClient, cp Checkpoint) (*Service, error) {
+func NewService(idx int, interval time.Duration, path string, client KVStore, cp Checkpoint) (*Service, error) {
 	s := &Service{
 		idx:                idx,
 		checkpointInterval: interval,
@@ -173,6 +175,7 @@ func NewService(idx int, interval time.Duration, path string, client *EtcdClient
 			}
 			s.optMap[p.Param.Name] = newOptimizer(p, item.State)
 		}
+		close(s.initialized)
 	}
 	return s, nil
 }
@@ -221,7 +224,7 @@ func (s *Service) FinishInitParams(_ int, _ *int) error {
 		for range t {
 			err := s.checkpoint()
 			if err != nil {
-				log.Error("finish init params error", log.Ctx{"error": err})
+				log.Error("checkpoint error", log.Ctx{"error": err})
 			}
 		}
 	}()
@@ -274,6 +277,7 @@ func (s *Service) GetParam(name string, parameter *Parameter) error {
 	parameter.Name = name
 	parameter.ElementType = opt.elementType
 	parameter.Content = opt.GetWeights()
+
 	log.Info("sending parameter to the trainer", "name", parameter.Name, "size", len(parameter.Content), "type", parameter.ElementType)
 	return nil
 }
@@ -354,20 +358,29 @@ func (s *Service) checkpoint() (err error) {
 
 	oldMeta, err := loadMeta(s.client, s.idx)
 	if err == ErrCheckpointNotFound {
-		log.Info("Do not have existing checkpoint.")
+		log.Info("old meta not found, skip removing old meta")
 		err = nil
+	} else if err == nil {
+		log.Info("removing old meta")
+		if oldMeta.Path != "" {
+			rmErr := os.Remove(oldMeta.Path)
+			if rmErr != nil {
+				// log error, but still treat checkpoint as
+				// successful.
+				log.Error("remove old meta file error", log.Ctx{"error": rmErr})
+			}
+		}
 	}
 
 	if err != nil {
 		return
 	}
 
-	h := md5.New()
-	md5 := hex.EncodeToString(h.Sum(buf.Bytes()))
+	crc32 := crc32.ChecksumIEEE(buf.Bytes())
 	cpMeta := checkpointMeta{
 		UUID:      id,
 		Timestamp: time.Now().UnixNano(),
-		MD5:       md5,
+		CRC32:     crc32,
 		Path:      p,
 	}
 
@@ -381,14 +394,5 @@ func (s *Service) checkpoint() (err error) {
 		return
 	}
 
-	if oldMeta.Path != "" {
-		rmErr := os.Remove(oldMeta.Path)
-		if rmErr != nil {
-			// log error, but still treat checkpoint as
-			// successful.
-			log.Error("remove old meta file error", log.Ctx{"error": rmErr})
-		}
-	}
-
 	return
 }

go/pserver/service_internal_test.go

+package pserver
+
+import (
+	"bytes"
+	"encoding/binary"
+	"fmt"
+	"testing"
+	"time"
+
+	"github.com/stretchr/testify/assert"
+)
+
+const testDir = "./test_data"
+
+type myKV struct {
+	m map[string][]byte
+}
+
+func (m *myKV) GetKey(key string, timeout time.Duration) ([]byte, error) {
+	if m.m == nil {
+		m.m = make(map[string][]byte)
+	}
+	return m.m[key], nil
+}
+
+func (m *myKV) PutKey(key string, value []byte, timeout time.Duration, withLease bool) error {
+	if m.m == nil {
+		m.m = make(map[string][]byte)
+	}
+	m.m[key] = value
+	return nil
+}
+
+func TestCheckpoint(t *testing.T) {
+	kv := &myKV{}
+	s, err := NewService(0, time.Hour, testDir, kv, nil)
+	assert.Nil(t, err)
+	err = s.checkpoint()
+	assert.Nil(t, err)
+	_, err = LoadCheckpoint(kv, 0)
+	assert.Nil(t, err)
+}
+
+func float32ToByte(f float32) []byte {
+	var buf bytes.Buffer
+	err := binary.Write(&buf, binary.LittleEndian, f)
+	if err != nil {
+		fmt.Println("binary.Write failed:", err)
+	}
+	return buf.Bytes()
+}
+
+func TestCheckpointWithData(t *testing.T) {
+	kv := &myKV{}
+	s, err := NewService(0, time.Hour, testDir, kv, nil)
+	assert.Nil(t, err)
+
+	var content []byte
+	for i := 0; i < 50000; i++ {
+		content = append(content, float32ToByte(float32(i))...)
+	}
+
+	p1 := Parameter{Name: "p1", ElementType: 1, Content: content}
+	err = s.InitParam(ParameterWithConfig{Param: p1}, nil)
+	assert.Nil(t, err)
+
+	err = s.FinishInitParams(0, nil)
+	assert.Nil(t, err)
+
+	var p2 Parameter
+	err = s.GetParam(p1.Name, &p2)
+	assert.Nil(t, err)
+	assert.Equal(t, p1, p2)
+
+	err = s.checkpoint()
+	assert.Nil(t, err)
+	cp, err := LoadCheckpoint(kv, 0)
+	assert.Nil(t, err)
+	s1, err := NewService(0, time.Hour, testDir, kv, cp)
+	assert.Nil(t, err)
+
+	var p3 Parameter
+	err = s1.GetParam(p1.Name, &p3)
+	assert.Nil(t, err)
+	assert.Equal(t, p1, p3)
+}

go/pserver/service_test.go

@@ -178,7 +178,3 @@ func TestBlockUntilInitialized(t *testing.T) {
 
 	wg.Wait()
 }
-
-func TestCheckpointSpeed(t *testing.T) {
-	//TODO(zhihong): test speed
-}

paddle/capi/gradient_machine.cpp

@@ -64,12 +64,18 @@ paddle_error paddle_gradient_machine_create_for_inference_with_parameters(
   modelConfigProtobuf.resize(modelConfigSize);
   is.read(&modelConfigProtobuf[0], modelConfigSize);
   paddle::TrainerConfig config;
+  paddle::ModelConfig modelConfig;
   if (!config.ParseFromString(modelConfigProtobuf) || !config.IsInitialized()) {
-    return kPD_PROTOBUF_ERROR;
+    if (!modelConfig.ParseFromString(modelConfigProtobuf) ||
+        !modelConfig.IsInitialized()) {
+      return kPD_PROTOBUF_ERROR;
+    }
+  } else {
+    modelConfig = config.model_config();
   }
   auto ptr = new paddle::capi::CGradientMachine();
   ptr->machine.reset(paddle::GradientMachine::create(
-      config.model_config(), CREATE_MODE_TESTING, {paddle::PARAMETER_VALUE}));
+      modelConfig, CREATE_MODE_TESTING, {paddle::PARAMETER_VALUE}));
   std::vector<paddle::ParameterPtr>& parameters = ptr->machine->getParameters();
   for (auto& para : parameters) {
     para->load(is);

paddle/framework/CMakeLists.txt

 # ddim lib
 proto_library(framework_proto SRCS framework.proto)
-proto_library(saver_proto SRCS framework.proto saver.proto)
 
 cc_library(ddim SRCS ddim.cc DEPS eigen3)
 cc_test(ddim_test SRCS ddim_test.cc DEPS ddim)
@@ -10,7 +9,7 @@ cc_library(tensor SRCS tensor.cc DEPS ddim place paddle_memory device_context)
 cc_test(tensor_test SRCS tensor_test.cc DEPS tensor)
 cc_test(eigen_test SRCS eigen_test.cc DEPS tensor)
 
-cc_library(lod_tensor SRCS lod_tensor.cc DEPS ddim place tensor saver_proto framework_proto)
+cc_library(lod_tensor SRCS lod_tensor.cc DEPS ddim place tensor framework_proto)
 cc_test(lod_tensor_test SRCS lod_tensor_test.cc DEPS lod_tensor paddle_memory)
 nv_test(lod_tensor_gpu_test SRCS lod_tensor_test.cu DEPS lod_tensor)
 
@@ -27,7 +26,7 @@ cc_test(op_proto_maker_test SRCS op_proto_maker_test.cc DEPS op_proto_maker)
 cc_library(op_info SRCS op_info.cc DEPS attribute framework_proto)
 cc_library(operator SRCS operator.cc DEPS op_info device_context tensor scope glog)
 cc_test(operator_test SRCS operator_test.cc DEPS operator op_registry)
-cc_library(proto_desc SRCS var_desc.cc op_desc.cc block_desc.cc program_desc.cc DEPS attribute ddim op_info operator)
+cc_library(proto_desc SRCS var_desc.cc op_desc.cc block_desc.cc program_desc.cc DEPS attribute ddim op_info operator glog)
 
 cc_library(op_registry SRCS op_registry.cc DEPS op_proto_maker op_info operator glog proto_desc)
 cc_test(op_registry_test SRCS op_registry_test.cc DEPS op_registry)
@@ -43,7 +42,7 @@ add_custom_command(TARGET framework_py_proto POST_BUILD
     WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR})
 
 cc_library(backward SRCS backward.cc DEPS net_op)
-cc_test(backward_test SRCS backward_test.cc DEPS backward recurrent_op device_context)
+cc_test(backward_test SRCS backward_test.cc DEPS backward recurrent_op device_context fill_constant_op)
 
 cc_library(executor SRCS executor.cc DEPS op_registry device_context scope framework_proto backward glog)
 

paddle/framework/backward.cc

@@ -315,6 +315,7 @@ static void CreateGradVarInBlock(
                      return false; /* not break */
                    });
     if (need_infer_shape) {
+      ops[op_index]->InferVarType(block_desc);
       ops[op_index]->InferShape(*block_desc);
     }
   }
@@ -452,11 +453,16 @@ ParamGradInfoMap AppendBackward(
   std::transform(target_shape_desc.begin(), target_shape_desc.end(),
                  std::back_inserter(target_shape),
                  [](int64_t dim) { return static_cast<int>(dim); });
+  VLOG(3) << "backward from loss=" << target.Name()
+          << " data_type=" << target.GetDataType();
   std::unique_ptr<OpDescBind> fill_one_op(
       new OpDescBind("fill_constant", {}, {{"Out", {fill_one_op_out}}},
                      {{"shape", target_shape},
                       {"value", static_cast<float>(1.0)},
-                      {"data_type", framework::DataType::FP32}}));
+                      {"data_type", target.GetDataType()}}));
+  // infer var type of fill_one_op
+  fill_one_op->InferVarType(root_block);
+
   root_block->AppendAllocatedOp(std::move(fill_one_op));
   size_t forward_op_num = root_block->OpSize();
   size_t forward_block_num = program_desc.Size();
@@ -475,8 +481,7 @@ ParamGradInfoMap AppendBackward(
   std::unordered_map<std::string, GradVarInfo> retv;
 
   auto var = root_block->Var(fill_one_op_out);
-  // FIXME(qiao) infer the data type
-  var->SetDataType(framework::DataType::FP32);
+  var->SetDataType(target.GetDataType());
   var->SetShape(target.Shape());
   auto& target_grad = retv[target.Name()];
   target_grad.name_ = fill_one_op_out;

paddle/framework/backward_test.cc

@@ -21,6 +21,8 @@
 #include "paddle/framework/var_desc.h"
 #include "paddle/operators/net_op.h"
 
+USE_OP(fill_constant);
+
 namespace paddle {
 namespace framework {
 

paddle/framework/block_desc.cc

@@ -120,6 +120,17 @@ BlockDesc *BlockDescBind::Proto() {
   Flush();
   return desc_;
 }
+
+BlockDescBind::BlockDescBind(ProgramDescBind *prog, BlockDesc *desc)
+    : prog_(prog), desc_(desc), need_update_(false) {
+  for (const VarDesc &var_desc : desc_->vars()) {
+    vars_[var_desc.name()].reset(new VarDescBind(var_desc));
+  }
+  for (const OpDesc &op_desc : desc_->ops()) {
+    ops_.emplace_back(new OpDescBind(op_desc, prog));
+  }
+}
+
 BlockDescBind::BlockDescBind(const BlockDescBind &other, BlockDesc *desc,
                              ProgramDescBind *prog)
     : prog_(prog), desc_(desc) {

paddle/framework/block_desc.h

@@ -36,8 +36,7 @@ class ProgramDescBind;
 
 class BlockDescBind {
  public:
-  BlockDescBind(ProgramDescBind *prog, BlockDesc *desc)
-      : prog_(prog), desc_(desc), need_update_(false) {}
+  BlockDescBind(ProgramDescBind *prog, BlockDesc *desc);
 
   BlockDescBind(const BlockDescBind &other, BlockDesc *desc,
                 ProgramDescBind *prog);

paddle/framework/data_type.h

@@ -15,6 +15,7 @@
 #pragma once
 #include <typeindex>
 #include "paddle/framework/framework.pb.h"
+#include "paddle/platform/enforce.h"
 
 namespace paddle {
 namespace framework {

paddle/framework/details/op_registry.h

@@ -28,7 +28,8 @@ enum OpInfoFillType {
   kOperator = 0,
   kOpProtoAndCheckerMaker = 1,
   kGradOpDescMaker = 2,
-  kVarTypeInference = 3
+  kVarTypeInference = 3,
+  kShapeInference = 4
 };
 
 template <typename T>
@@ -42,7 +43,10 @@ struct OpInfoFillTypeID {
                              ? kGradOpDescMaker
                              : (std::is_base_of<VarTypeInference, T>::value
                                     ? kVarTypeInference
-                                    : static_cast<OpInfoFillType>(-1))));
+                                    : (std::is_base_of<InferShapeBase, T>::value
+                                           ? kShapeInference
+                                           : static_cast<OpInfoFillType>(
+                                                 -1)))));
   }
 };
 
@@ -121,6 +125,16 @@ struct OpInfoFiller<T, kVarTypeInference> {
   }
 };
 
+template <typename T>
+struct OpInfoFiller<T, kShapeInference> {
+  void operator()(const char* op_type, OpInfo* info) const {
+    info->infer_shape_ = [](InferShapeContext* ctx) {
+      T inference;
+      inference(ctx);
+    };
+  }
+};
+
 }  // namespace details
 
 }  // namespace framework

paddle/framework/executor.cc

@@ -20,6 +20,7 @@ limitations under the License. */
 #include <set>
 #include <vector>
 
+#include "paddle/framework/feed_fetch_type.h"
 #include "paddle/framework/lod_tensor.h"
 #include "paddle/framework/op_registry.h"
 #include "paddle/framework/scope.h"
@@ -56,6 +57,22 @@ Executor::~Executor() {
   }
 }
 
+static void CreateTensor(Variable* var, VarDesc::VarType var_type) {
+  if (var_type == VarDesc::LOD_TENSOR) {
+    var->GetMutable<LoDTensor>();
+  } else if (var_type == VarDesc::SELECTED_ROWS) {
+    var->GetMutable<SelectedRows>();
+  } else if (var_type == VarDesc::FEED_MINIBATCH) {
+    var->GetMutable<FeedFetchList>();
+  } else if (var_type == VarDesc::FETCH_LIST) {
+    var->GetMutable<FeedFetchList>();
+  } else {
+    PADDLE_THROW(
+        "Variable type must be "
+        "LoDTensor/SelectedRows/FEED_MINIBATCH/FETCH_LIST.");
+  }
+}
+
 void Executor::Run(const ProgramDesc& pdesc, Scope* scope, int block_id) {
   // TODO(tonyyang-svail):
   //    - only runs on the first device (i.e. no interdevice communication)
@@ -69,10 +86,12 @@ void Executor::Run(const ProgramDesc& pdesc, Scope* scope, int block_id) {
   for (auto& var : block.vars()) {
     if (var.persistable()) {
       auto* ptr = scope->Var(var.name());
+      CreateTensor(ptr, var.type());
       VLOG(3) << "Create Variable " << var.name()
               << " global, which pointer is " << ptr;
     } else {
       auto* ptr = local_scope.Var(var.name());
+      CreateTensor(ptr, var.type());
       VLOG(3) << "Create Variable " << var.name()
               << " locally, which pointer is " << ptr;
     }

paddle/framework/lod_tensor.cc

@@ -13,7 +13,6 @@
    limitations under the License. */
 
 #include "paddle/framework/lod_tensor.h"
-#include "paddle/framework/saver.pb.h"
 
 #include "paddle/memory/memcpy.h"
 #include "paddle/memory/memory.h"
@@ -136,141 +135,5 @@ void LoDTensor::ShrinkInLevel(size_t level, size_t elem_begin,
   PADDLE_ENFORCE_LT(begin, end, "Cannot shrink, the result tensor is empty.");
   ShareDataWith(Slice(begin, end));
 }
-
-std::string LoDTensor::SerializeToString() const {
-  LoDTensorProto desc;
-
-  // set data_type
-  if (this->type() == typeid(int8_t)) desc.set_data_type(DataType::BOOL);
-  if (this->type() == typeid(int16_t)) desc.set_data_type(DataType::INT16);
-  if (this->type() == typeid(int32_t)) desc.set_data_type(DataType::INT32);
-  if (this->type() == typeid(int64_t)) desc.set_data_type(DataType::INT64);
-  // FIXME(dzh): there is no fp16 in standard c++
-
-  if (this->type() == typeid(float))  // NOLINT
-    desc.set_data_type(DataType::FP32);
-  if (this->type() == typeid(double))  // NOLINT
-    desc.set_data_type(DataType::FP64);
-
-  for (int i = 0; i < dims().size(); ++i) {
-    desc.add_dims(dims()[i]);
-  }
-
-  // set lod information
-  desc.set_lod_level(this->NumLevels());
-  for (size_t i = 0; i < this->NumLevels(); ++i) {
-    LoDInfo* lod = desc.add_levels();
-    for (size_t j = 0; j < lod_[i].size(); ++j) {
-      lod->add_level(lod_[i][j]);
-    }
-  }
-
-  desc.set_version(0);
-
-  std::string desc_bytes = desc.SerializeAsString();
-
-  // FIXME(dzh) : implement fix chunk size buffer.
-  size_t DESC_SIZE = desc_bytes.size();
-  size_t DATA_SIZE = holder_->size() - offset_;
-
-  const size_t BUFFER_SIZE = DESC_SIZE + DATA_SIZE + 2 * sizeof(size_t);
-  char* buffer =
-      static_cast<char*>(memory::Alloc(platform::CPUPlace(), BUFFER_SIZE));
-
-  // format: desc_size data_size, desc_bytes, data_bytes.
-  platform::CPUPlace src_place;
-  platform::CPUPlace dst_place;
-
-  memory::Copy(dst_place, buffer, src_place, &BUFFER_SIZE, sizeof(size_t));
-  memory::Copy(dst_place, buffer + sizeof(size_t), src_place, &DESC_SIZE,
-               sizeof(size_t));
-  memory::Copy(dst_place, buffer + sizeof(size_t) * 2, src_place,
-               desc_bytes.c_str(), desc_bytes.size());
-
-  PADDLE_ENFORCE(this->numel() != 0, "Serialize a empty Tensor!");
-
-  platform::Place place = holder_->place();
-  int element_width = holder_->size() / this->numel();
-
-  if (platform::is_cpu_place(place)) {
-    memory::Copy(dst_place, buffer + sizeof(size_t) * 2 + desc_bytes.size(),
-                 boost::get<platform::CPUPlace>(place),
-                 static_cast<char*>(holder_->ptr()) + offset_ / element_width,
-                 DATA_SIZE);
-  }
-#ifdef PADDLE_WITH_GPU
-  if (platform::is_gpu_place(place)) {
-    memory::Copy(dst_place, buffer + sizeof(size_t) * 2 + desc_bytes.size(),
-                 boost::get<platform::GPUPlace>(place),
-                 static_cast<char*>(holder_->ptr()) + offset_ / element_width,
-                 DATA_SIZE);
-  }
-#endif
-
-  std::string ret(buffer, BUFFER_SIZE);
-  memory::Free(platform::CPUPlace(), buffer);
-  return ret;
-}
-
-void LoDTensor::DeserializeFromString(const std::string& s,
-                                      const platform::Place& dst_place) {
-  size_t DESC_SIZE, BUFFER_SIZE;
-  platform::CPUPlace src_place;
-
-  memory::Copy(src_place, &BUFFER_SIZE, src_place, s.c_str(), sizeof(size_t));
-  memory::Copy(src_place, &DESC_SIZE, src_place, s.c_str() + sizeof(size_t),
-               sizeof(size_t));
-
-  const size_t DATA_SIZE = BUFFER_SIZE - DESC_SIZE - sizeof(size_t) * 2;
-
-  // parse LoDTensorDesc
-  LoDTensorProto desc;
-  desc.ParseFromArray(s.c_str() + sizeof(size_t) * 2, DESC_SIZE);
-
-  std::vector<int64_t> dims;
-  std::copy(desc.dims().begin(), desc.dims().end(), std::back_inserter(dims));
-  this->Resize(make_ddim(dims));
-
-  // parse data type
-  void* ptr = nullptr;
-  if (desc.data_type() == DataType::BOOL)
-    ptr = this->mutable_data<bool>(dst_place);
-  if (desc.data_type() == DataType::INT16)
-    ptr = this->mutable_data<int16_t>(dst_place);
-  if (desc.data_type() == DataType::INT32)
-    ptr = this->mutable_data<int32_t>(dst_place);
-  if (desc.data_type() == DataType::INT64)
-    ptr = this->mutable_data<int64_t>(dst_place);
-  // FIXME(dzh): there is no fp16 in standard c++
-
-  if (desc.data_type() == DataType::FP32)
-    ptr = this->mutable_data<float>(dst_place);
-  if (desc.data_type() == DataType::FP64)
-    ptr = this->mutable_data<double>(dst_place);
-
-  LoD lod;
-  std::vector<size_t> levels;
-  for (int i = 0; i < desc.levels().size(); ++i) {
-    auto current_level = desc.levels()[i].level();
-    std::copy(current_level.begin(), current_level.end(),
-              std::back_inserter(levels));
-    lod.emplace_back(levels);
-    levels.clear();
-  }
-
-  this->set_lod(lod);
-
-  if (platform::is_cpu_place(dst_place)) {
-    memory::Copy(boost::get<platform::CPUPlace>(dst_place), ptr, src_place,
-                 s.c_str() + sizeof(size_t) * 2 + DESC_SIZE, DATA_SIZE);
-  }
-#ifdef PADDLE_WITH_GPU
-  if (platform::is_gpu_place(dst_place)) {
-    memory::Copy(boost::get<platform::GPUPlace>(dst_place), ptr, src_place,
-                 s.c_str() + sizeof(size_t) * 2 + DESC_SIZE, DATA_SIZE);
-  }
-#endif
-}
-
 }  // namespace framework
 }  // namespace paddle

paddle/framework/lod_tensor.h

@@ -85,7 +85,9 @@ class LoDTensor : public Tensor {
 
   void set_lod(const LoD& lod) { lod_ = lod; }
 
-  LoD lod() const { return lod_; }
+  const LoD& lod() const { return lod_; }
+
+  LoD* mutable_lod() { return &lod_; }
 
   /*
    * Get the start offset and end offset of an  element from LoD.
@@ -139,27 +141,6 @@ class LoDTensor : public Tensor {
    */
   void ShrinkInLevel(size_t level, size_t elem_begin, size_t elem_end);
 
-  /**
-   *  @brief Serialize tensor to char bytes.
-   *  Please check model_format.md for the format detail.
-   *  NOTE: GPUTensor will copy data to cpu implicitly.
-   *  @return return string
-   */
-
-  // FIXME(dzh) : Currently, this interface should only be used in
-  // save/restore model and checkpoint. ParameterServer do not use shape
-  // information to do the optimization, as a result, when we serialize
-  // parameter/gradient to string, we should serialize the tensor
-  // to string in the ps trainer instead of LoDTensor.
-  std::string SerializeToString() const;
-
-  /**
-   *  @brief Deserialize char bytes to tensor.
-   *  @return return string
-   */
-  void DeserializeFromString(const std::string& s,
-                             const platform::Place& dst_place);
-
  private:
   LoD lod_;
 };

paddle/framework/lod_tensor_test.cc

@@ -144,21 +144,5 @@ TEST(LodExpand, test) {
   }
 }
 
-TEST_F(LoDTensorTester, SerializeDeserialize) {
-  LoDTensor new_lod_tensor = lod_tensor_;
-  float* src_ptr = lod_tensor_.data<float>();
-  std::string s = lod_tensor_.SerializeToString();
-  LoDTensor dst;
-  dst.DeserializeFromString(s, platform::CPUPlace());
-  float* dst_ptr = dst.data<float>();
-  for (int i = 0; i < kLodTensorSize; ++i) {
-    EXPECT_EQ(dst_ptr[i], src_ptr[i]);
-  }
-
-  ASSERT_EQ(dst.NumElements(0), 2UL);
-  ASSERT_EQ(dst.NumElements(1), 3UL);
-  ASSERT_EQ(dst.NumElements(2), 8UL);
-}
-
 }  // namespace framework
 }  // namespace paddle

paddle/framework/lod_tensor_test.cu

@@ -47,31 +47,4 @@ TEST(LoDTensor, LoDInGPU) {
   for (size_t i = 0; i < src_lod[0].size(); ++i) {
     CHECK_EQ(lod[0].data()[i], src_lod[0].data()[i] * 2);
   }
-}
-
-TEST(LoDTensor, SerializeDeserialize) {
-  paddle::framework::LoDTensor lod_tensor;
-  paddle::platform::GPUPlace place(0);
-
-  paddle::framework::LoD src_lod;
-  src_lod.push_back(std::vector<size_t>{0, 2, 4, 6, 8, 10, 12, 14});
-
-  lod_tensor.Resize({14, 16});
-  lod_tensor.mutable_data<float>(place);
-
-  lod_tensor.set_lod(src_lod);
-  CHECK_EQ(lod_tensor.lod_element(0, 2).first, 4UL);
-  CHECK_EQ(lod_tensor.lod_element(0, 4).first, 8UL);
-
-  test<<<1, 8>>>(src_lod[0].data(), src_lod[0].size());
-  cudaDeviceSynchronize();
-
-  std::string s = lod_tensor.SerializeToString();
-  paddle::framework::LoDTensor dst;
-  dst.DeserializeFromString(s, place);
-  paddle::framework::LoD dst_lod = dst.lod();
-
-  for (size_t i = 0; i < dst_lod[0].size(); ++i) {
-    CHECK_EQ(src_lod[0].data()[i], dst_lod[0].data()[i] * 2);
-  }
-}
+}
\ No newline at end of file

paddle/framework/op_desc.cc

@@ -14,9 +14,13 @@ limitations under the License. */
 
 #include "paddle/framework/op_desc.h"
 #include <functional>
+#include <mutex>
 #include <unordered_map>
 #include "paddle/framework/block_desc.h"
 #include "paddle/framework/operator.h"
+#include "paddle/framework/program_desc.h"
+
+#include "glog/logging.h"
 
 namespace paddle {
 namespace framework {
@@ -24,16 +28,47 @@ namespace framework {
 OpDescBind::OpDescBind(const std::string &type, const VariableNameMap &inputs,
                        const VariableNameMap &outputs,
                        const AttributeMap &attrs) {
-  op_desc_.set_type(type);
+  desc_.set_type(type);
   inputs_ = inputs;
   outputs_ = outputs;
   attrs_ = attrs;
   need_update_ = true;
 }
 
+OpDescBind::OpDescBind(const OpDesc &desc, ProgramDescBind *prog)
+    : desc_(desc), need_update_(false) {
+  // restore inputs_
+  int input_size = desc_.inputs_size();
+  for (int i = 0; i < input_size; ++i) {
+    const OpDesc::Var &var = desc_.inputs(i);
+    std::vector<std::string> &args = inputs_[var.parameter()];
+    int argu_size = var.arguments_size();
+    args.reserve(argu_size);
+    for (int j = 0; j < argu_size; ++j) {
+      args.push_back(var.arguments(j));
+    }
+  }
+  // restore outputs_
+  int output_size = desc_.outputs_size();
+  for (int i = 0; i < output_size; ++i) {
+    const OpDesc::Var &var = desc_.outputs(i);
+    std::vector<std::string> &args = outputs_[var.parameter()];
+    int argu_size = var.arguments_size();
+    args.reserve(argu_size);
+    for (int j = 0; j < argu_size; ++j) {
+      args.push_back(var.arguments(j));
+    }
+  }
+  // restore attrs_
+  for (const OpDesc::Attr &attr : desc_.attrs()) {
+    std::string attr_name = attr.name();
+    attrs_[attr_name] = GetAttrValue(attr, prog->Proto());
+  }
+}
+
 OpDesc *OpDescBind::Proto() {
   Flush();
-  return &op_desc_;
+  return &desc_;
 }
 
 const std::vector<std::string> &OpDescBind::Input(
@@ -167,23 +202,23 @@ struct SetAttrDescVisitor : public boost::static_visitor<void> {
 
 void OpDescBind::Flush() {
   if (need_update_) {
-    this->op_desc_.mutable_inputs()->Clear();
+    this->desc_.mutable_inputs()->Clear();
     for (auto &ipt : inputs_) {
-      auto *input = op_desc_.add_inputs();
+      auto *input = desc_.add_inputs();
       input->set_parameter(ipt.first);
       VectorToRepeated(ipt.second, input->mutable_arguments());
     }
 
-    this->op_desc_.mutable_outputs()->Clear();
+    this->desc_.mutable_outputs()->Clear();
     for (auto &opt : outputs_) {
-      auto *output = op_desc_.add_outputs();
+      auto *output = desc_.add_outputs();
       output->set_parameter(opt.first);
       VectorToRepeated(opt.second, output->mutable_arguments());
     }
 
-    this->op_desc_.mutable_attrs()->Clear();
+    this->desc_.mutable_attrs()->Clear();
     for (auto &attr : attrs_) {
-      auto *attr_desc = op_desc_.add_attrs();
+      auto *attr_desc = desc_.add_attrs();
       attr_desc->set_name(attr.first);
       attr_desc->set_type(
           static_cast<framework::AttrType>(attr.second.which() - 1));
@@ -195,26 +230,26 @@ void OpDescBind::Flush() {
   }
 }
 
-using InferShapeFuncMap =
-    std::unordered_map<std::string /*op_type*/,
-                       std::function<void(InferShapeContext *)>>;
-
-static InferShapeFuncMap &InferShapeFuncs() {
-  static InferShapeFuncMap *g_map = nullptr;
-  if (g_map == nullptr) {
-    g_map = new InferShapeFuncMap();
-    auto &info_map = OpInfoMap::Instance();
-    // all registered kernels
-    for (auto &pair : OperatorWithKernel::AllOpKernels()) {
-      auto &info = info_map.Get(pair.first);
-      // use empty type here to avoid runtime checks.
+static std::once_flag init_infer_shape_funcs;
+
+static void InitInferShapeFuncs() {
+  std::call_once(init_infer_shape_funcs, [] {
+    auto &map = OpInfoMap::Instance();
+    auto &info_map = *map.mutable_map();
+
+    for (auto &kern_pair : OperatorWithKernel::AllOpKernels()) {
+      auto op_type = kern_pair.first;
+      auto &op_info = info_map.at(op_type);
       auto op =
-          static_cast<OperatorWithKernel *>(info.Creator()("", {}, {}, {}));
-      g_map->insert(
-          {pair.first, [op](InferShapeContext *ctx) { op->InferShape(ctx); }});
+          static_cast<OperatorWithKernel *>(op_info.Creator()("", {}, {}, {}));
+      if (op_info.infer_shape_) {  // infer_shape has been registered.
+        continue;
+      }
+      op_info.infer_shape_ = [op](InferShapeContext *ctx) {
+        op->InferShape(ctx);
+      };
     }
-  }
-  return *g_map;
+  });
 }
 
 void OpDescBind::CheckAttrs() {
@@ -230,13 +265,13 @@ void OpDescBind::CheckAttrs() {
 }
 
 void OpDescBind::InferShape(const BlockDescBind &block) const {
-  auto &funcs = InferShapeFuncs();
-  auto it = funcs.find(this->Type());
-  if (it == funcs.end()) {
-    PADDLE_THROW("Operator %s has not been registered", this->Type());
-  }
+  VLOG(3) << "CompileTime infer shape on " << Type();
+  InitInferShapeFuncs();
+  auto &infer_shape = OpInfoMap::Instance().Get(this->Type()).infer_shape_;
+  PADDLE_ENFORCE(static_cast<bool>(infer_shape),
+                 "%s's infer_shape has not been registered", this->Type());
   CompileTimeInferShapeContext ctx(*this, block);
-  it->second(&ctx);
+  infer_shape(&ctx);
 }
 
 void OpDescBind::InferVarType(BlockDescBind *block) const {

paddle/framework/op_desc.h

@@ -24,6 +24,7 @@ namespace paddle {
 namespace framework {
 
 class BlockDescBind;
+class ProgramDescBind;
 
 class OpDescBind {
  public:
@@ -32,11 +33,13 @@ class OpDescBind {
   OpDescBind(const std::string &type, const VariableNameMap &inputs,
              const VariableNameMap &outputs, const AttributeMap &attrs);
 
+  OpDescBind(const OpDesc &desc, ProgramDescBind *prog);
+
   OpDesc *Proto();
 
-  std::string Type() const { return op_desc_.type(); }
+  std::string Type() const { return desc_.type(); }
 
-  void SetType(const std::string &type) { op_desc_.set_type(type); }
+  void SetType(const std::string &type) { desc_.set_type(type); }
 
   const std::vector<std::string> &Input(const std::string &name) const;
 
@@ -117,7 +120,7 @@ class OpDescBind {
     return ret_val;
   }
 
-  OpDesc op_desc_;
+  OpDesc desc_;
   VariableNameMap inputs_;
   VariableNameMap outputs_;
   AttributeMap attrs_;

paddle/framework/op_info.h

@@ -25,12 +25,19 @@
 namespace paddle {
 namespace framework {
 
+class InferShapeBase {
+ public:
+  virtual ~InferShapeBase() = default;
+  virtual void operator()(InferShapeContext*) const = 0;
+};
+
 struct OpInfo {
   OpCreator creator_;
   GradOpMakerFN grad_op_maker_;
   OpProto* proto_{nullptr};
   OpAttrChecker* checker_{nullptr};
   InferVarTypeFN infer_var_type_;
+  InferShapeFN infer_shape_;
 
   bool HasOpProtoAndChecker() const {
     return proto_ != nullptr && checker_ != nullptr;
@@ -87,13 +94,13 @@ class OpInfoMap {
     }
   }
 
-  const std::unordered_map<std::string, const OpInfo>& map() const {
-    return map_;
-  }
+  const std::unordered_map<std::string, OpInfo>& map() const { return map_; }
+
+  std::unordered_map<std::string, OpInfo>* mutable_map() { return &map_; }
 
  private:
   OpInfoMap() = default;
-  std::unordered_map<std::string, const OpInfo> map_;
+  std::unordered_map<std::string, OpInfo> map_;
 
   DISABLE_COPY_AND_ASSIGN(OpInfoMap);
 };

paddle/framework/operator.cc

@@ -33,24 +33,6 @@ ExecutionContext::GetEigenDevice<platform::GPUPlace, Eigen::GpuDevice>() const {
 }
 #endif
 
-const Tensor* GetTensorFromVar(const Variable* var) {
-  if (var->IsType<LoDTensor>()) {
-    return &var->Get<LoDTensor>();
-  }
-  PADDLE_ENFORCE(var->IsType<Tensor>(),
-                 "The Input must be LoDTensor or Tensor.");
-  return &var->Get<Tensor>();
-}
-
-Tensor* GetTensorFromVar(Variable* var) {
-  if (var->IsType<LoDTensor>()) {
-    return var->GetMutable<LoDTensor>();
-  }
-  PADDLE_ENFORCE(var->IsType<Tensor>(),
-                 "The Input must be LoDTensor or Tensor.");
-  return var->GetMutable<Tensor>();
-}
-
 std::string OperatorBase::Input(const std::string& name) const {
   auto& ins = Inputs(name);
   PADDLE_ENFORCE_LE(ins.size(), 1UL,
@@ -204,6 +186,30 @@ void OperatorBase::GenerateTemporaryNames() {
   }
 }
 
+static const Tensor* GetTensorFromVar(const Variable* var) {
+  const Tensor* t = nullptr;
+  if (var->IsType<LoDTensor>()) {
+    t = &(var->Get<LoDTensor>());
+  } else if (var->IsType<SelectedRows>()) {
+    t = &(var->Get<SelectedRows>().value());
+  } else {
+    PADDLE_THROW("Variable type must be LoDTensor/SelectedRows.");
+  }
+  return t;
+}
+
+static Tensor* GetMutableTensorFromVar(Variable* var) {
+  Tensor* t = nullptr;
+  if (var->IsType<LoDTensor>()) {
+    t = var->GetMutable<LoDTensor>();
+  } else if (var->IsType<SelectedRows>()) {
+    t = var->GetMutable<SelectedRows>()->mutable_value();
+  } else {
+    PADDLE_THROW("Variable type must be LoDTensor/SelectedRows.");
+  }
+  return t;
+}
+
 template <>
 const Tensor* ExecutionContext::Input<Tensor>(const std::string& name) const {
   auto* var = InputVar(name);
@@ -227,7 +233,7 @@ const std::vector<const Tensor*> ExecutionContext::MultiInput<Tensor>(
 template <>
 Tensor* ExecutionContext::Output<Tensor>(const std::string& name) const {
   auto var = OutputVar(name);
-  return var == nullptr ? nullptr : var->GetMutable<LoDTensor>();
+  return var == nullptr ? nullptr : GetMutableTensorFromVar(var);
 }
 
 template <>
@@ -240,7 +246,7 @@ std::vector<Tensor*> ExecutionContext::MultiOutput<Tensor>(
                  [&](const std::string& sub_name) {
                    auto var = scope_.FindVar(sub_name);
                    return var == nullptr ? nullptr
-                                         : var->GetMutable<LoDTensor>();
+                                         : GetMutableTensorFromVar(var);
                  });
   return res;
 }

paddle/framework/operator.h

@@ -28,6 +28,7 @@ limitations under the License. */
 #include "paddle/framework/lod_tensor.h"
 #include "paddle/framework/op_info.h"
 #include "paddle/framework/scope.h"
+#include "paddle/framework/selected_rows.h"
 #include "paddle/framework/shape_inference.h"
 #include "paddle/framework/tensor.h"
 #include "paddle/platform/device_context.h"
@@ -60,9 +61,6 @@ inline std::string GradVarName(const std::string& var_name) {
 class OperatorBase;
 class ExecutionContext;
 
-extern const Tensor* GetTensorFromVar(const Variable* var);
-extern Tensor* GetTensorFromVar(Variable* var);
-
 /**
  * OperatorBase has the basic element that Net will call to do computation.
  * Only CreateOperator from OpRegistry will new Operator directly. User
@@ -414,7 +412,9 @@ class CompileTimeInferShapeContext : public InferShapeContext {
 
  private:
   DDim GetDim(const std::string& name) const override {
-    return framework::make_ddim(block_.FindVarRecursive(name)->Shape());
+    auto var = block_.FindVarRecursive(name);
+    PADDLE_ENFORCE(var != nullptr, "Cannot find variable %s", name);
+    return framework::make_ddim(var->Shape());
   }
 
   void SetDim(const std::string& name, const DDim& dim) override {
@@ -511,28 +511,26 @@ class RuntimeInferShapeContext : public InferShapeContext {
   }
 
  private:
-  template <bool Allocate>
-  Tensor* GetTensor(const std::string& name) const {
-    Tensor* t = nullptr;
-    auto* var = scope_.FindVar(name);
-    if (!var->IsType<LoDTensor>() && !var->IsType<Tensor>()) {
-      if (Allocate) {
-        t = var->GetMutable<LoDTensor>();
-      } else {
-        PADDLE_THROW("Variable(%s) should be tensor", name);
-      }
+  DDim GetDim(const std::string& name) const override {
+    Variable* var = scope_.FindVar(name);
+    if (var->IsType<LoDTensor>()) {
+      return var->Get<LoDTensor>().dims();
+    } else if (var->IsType<SelectedRows>()) {
+      return var->Get<SelectedRows>().GetCompleteDims();
     } else {
-      t = GetTensorFromVar(scope_.FindVar(name));
+      PADDLE_THROW("Variable type must be LoDTensor/SelectedRows.");
     }
-    return t;
-  }
-
-  DDim GetDim(const std::string& name) const override {
-    return GetTensor<false>(name)->dims();
   }
 
   void SetDim(const std::string& name, const DDim& dim) override {
-    GetTensor<true>(name)->Resize(dim);
+    Variable* var = scope_.FindVar(name);
+    if (var->IsType<LoDTensor>()) {
+      var->GetMutable<LoDTensor>()->Resize(dim);
+    } else if (var->IsType<SelectedRows>()) {
+      var->GetMutable<SelectedRows>()->set_height(dim[0]);
+    } else {
+      PADDLE_THROW("Variable type must be LoDTensor/SelectedRows.");
+    }
   }
 
   const OperatorBase& op_;
@@ -638,7 +636,9 @@ class OperatorWithKernel : public OperatorBase {
                        });
   }
 
-  virtual void InferShape(InferShapeContext* ctx) const = 0;
+  virtual void InferShape(InferShapeContext* ctx) const {
+    OpInfoMap::Instance().Get(Type()).infer_shape_(ctx);
+  }
 
  protected:
   // indicate kernel DataType by input data. Defaultly all input data must be
@@ -655,11 +655,14 @@ class OperatorWithKernel : public OperatorBase {
             t = &var->Get<Tensor>();
           } else if (var->IsType<LoDTensor>()) {
             t = &var->Get<LoDTensor>();
+          } else if (var->IsType<SelectedRows>()) {
+            t = &(var->Get<SelectedRows>().value());
           }
           if (t != nullptr) {
             int tmp = static_cast<int>(ToDataType(t->type()));
+            VLOG(3) << "Input " << ipt_name << " with data_type " << tmp;
             PADDLE_ENFORCE(tmp == data_type || data_type == -1,
-                           "DataType of Paddle Op must be same.");
+                           "DataType of Paddle Op %s must be same.", Type());
             data_type = tmp;
           }
         }

paddle/framework/operator_test.cc

@@ -237,12 +237,12 @@ TEST(OpKernel, multi_inputs) {
 
   paddle::platform::CPUDeviceContext cpu_device_context;
   paddle::framework::Scope scope;
-  scope.Var("x0")->GetMutable<Tensor>();
-  scope.Var("x1")->GetMutable<Tensor>();
-  scope.Var("x2")->GetMutable<Tensor>();
-  scope.Var("k0")->GetMutable<Tensor>();
-  scope.Var("y0")->GetMutable<Tensor>();
-  scope.Var("y1")->GetMutable<Tensor>();
+  scope.Var("x0")->GetMutable<LoDTensor>();
+  scope.Var("x1")->GetMutable<LoDTensor>();
+  scope.Var("x2")->GetMutable<LoDTensor>();
+  scope.Var("k0")->GetMutable<LoDTensor>();
+  scope.Var("y0")->GetMutable<LoDTensor>();
+  scope.Var("y1")->GetMutable<LoDTensor>();
 
   auto op = paddle::framework::OpRegistry::CreateOp(op_desc, nullptr);
   op->Run(scope, cpu_device_context);

paddle/framework/program_desc.cc

@@ -19,9 +19,9 @@ namespace paddle {
 namespace framework {
 
 BlockDescBind *ProgramDescBind::AppendBlock(const BlockDescBind &parent) {
-  auto *b = prog_.add_blocks();
+  auto *b = desc_.add_blocks();
   b->set_parent_idx(parent.ID());
-  b->set_idx(prog_.blocks_size() - 1);
+  b->set_idx(desc_.blocks_size() - 1);
   blocks_.emplace_back(new BlockDescBind(this, b));
   return blocks_.back().get();
 }
@@ -30,23 +30,32 @@ ProgramDesc *ProgramDescBind::Proto() {
   for (auto &block : blocks_) {
     block->Flush();
   }
-  return &prog_;
+  return &desc_;
 }
 
 ProgramDescBind::ProgramDescBind() {
-  auto *block = prog_.mutable_blocks()->Add();
+  auto *block = desc_.mutable_blocks()->Add();
   block->set_idx(kRootBlockIndex);
   block->set_parent_idx(kNoneBlockIndex);
   blocks_.emplace_back(new BlockDescBind(this, block));
 }
 
 ProgramDescBind::ProgramDescBind(const ProgramDescBind &o) {
-  prog_ = o.prog_;
+  desc_ = o.desc_;
 
-  for (int i = 0; i < prog_.blocks_size(); ++i) {
-    auto *block = prog_.mutable_blocks(i);
+  for (int i = 0; i < desc_.blocks_size(); ++i) {
+    auto *block = desc_.mutable_blocks(i);
     blocks_.emplace_back(new BlockDescBind(*o.blocks_[i], block, this));
   }
 }
+
+ProgramDescBind::ProgramDescBind(const std::string &binary_str) {
+  PADDLE_ENFORCE(desc_.ParseFromString(binary_str),
+                 "Fail to parse program_desc from binary string.");
+  for (auto &block_desc : *desc_.mutable_blocks()) {
+    blocks_.emplace_back(new BlockDescBind(this, &block_desc));
+  }
+}
+
 }  // namespace framework
 }  // namespace paddle

paddle/framework/program_desc.h

@@ -31,6 +31,8 @@ class ProgramDescBind {
 
   ProgramDescBind(const ProgramDescBind &o);
 
+  explicit ProgramDescBind(const std::string &binary_str);
+
   BlockDescBind *AppendBlock(const BlockDescBind &parent);
 
   BlockDescBind *Block(size_t idx) { return blocks_[idx].get(); }
@@ -40,7 +42,7 @@ class ProgramDescBind {
   ProgramDesc *Proto();
 
  private:
-  ProgramDesc prog_;
+  ProgramDesc desc_;
 
   std::vector<std::unique_ptr<BlockDescBind>> blocks_;
 };

paddle/framework/program_desc_test.cc

@@ -59,7 +59,7 @@ TEST(ProgramDesc, copy_ctor) {
   };
 
   ASSERT_EQ(global_block->LocalVarNames(), global_block_copy->LocalVarNames());
-  ASSERT_EQ(3, global_block_copy->LocalVarNames().size());
+  ASSERT_EQ(3UL, global_block_copy->LocalVarNames().size());
   assert_same_var("X", x);
   assert_same_var("Y", y);
   assert_same_var("Out", out);
@@ -79,5 +79,67 @@ TEST(ProgramDesc, copy_ctor) {
   // Not check block's protostr are same it because the order of vars could be
   // different and it is correct.
 }
+
+TEST(ProgramDescBind, serialize_and_deserialize) {
+  ProgramDescBind program_origin;
+  auto* global_block = program_origin.Block(0);
+  auto* x = global_block->Var("X");
+  x->SetType(VarDesc_VarType_LOD_TENSOR);
+  x->SetLoDLevel(0);
+  x->SetDataType(FP32);
+  x->SetShape({1000, 784});
+
+  auto* y = global_block->Var("Y");
+  y->SetType(VarDesc_VarType_LOD_TENSOR);
+  y->SetLoDLevel(0);
+  y->SetDataType(FP32);
+  y->SetShape({784, 100});
+
+  auto* op = global_block->AppendOp();
+  op->SetType("mul");
+  op->SetInput("X", {x->Name()});
+  op->SetInput("Y", {y->Name()});
+
+  auto* out = global_block->Var("Out");
+  out->SetType(VarDesc_VarType_LOD_TENSOR);
+  op->SetOutput("Y", {out->Name()});
+
+  std::string binary_str;
+  program_origin.Proto()->SerializeToString(&binary_str);
+
+  ProgramDescBind program_restored(binary_str);
+  auto* global_block_restored = program_restored.Block(0);
+  ASSERT_NE(global_block, global_block_restored);
+
+  auto assert_same_var = [&](const std::string& name, VarDescBind* var_before) {
+    ASSERT_TRUE(global_block_restored->HasVar(name));
+    auto* restored = global_block_restored->Var(name);
+    ASSERT_NE(restored, var_before);
+    ASSERT_EQ(restored->Name(), var_before->Name());
+    ASSERT_EQ(restored->GetType(), var_before->GetType());
+    ASSERT_EQ(restored->Shape(), var_before->Shape());
+    ASSERT_EQ(restored->Proto()->SerializeAsString(),
+              var_before->Proto()->SerializeAsString());
+  };
+
+  ASSERT_EQ(global_block->LocalVarNames(),
+            global_block_restored->LocalVarNames());
+  ASSERT_EQ(3UL, global_block_restored->LocalVarNames().size());
+  assert_same_var("X", x);
+  assert_same_var("Y", y);
+  assert_same_var("Out", out);
+
+  for (size_t i = 0; i < global_block->OpSize(); ++i) {
+    auto op_origin = global_block->Op(i);
+    auto op_restored = global_block->Op(i);
+
+    ASSERT_EQ(op_origin->Type(), op_restored->Type());
+    ASSERT_EQ(op_origin->Inputs(), op_restored->Inputs());
+    ASSERT_EQ(op_origin->Outputs(), op_restored->Outputs());
+
+    ASSERT_EQ(op_restored->Proto()->SerializeAsString(),
+              op_origin->Proto()->SerializeAsString());
+  }
+}
 }  // namespace framework
 }  // namespace paddle

paddle/framework/selected_rows.h

@@ -23,7 +23,10 @@ class SelectedRows {
     value_.reset(new Tensor());
   }
 
-  SelectedRows() { value_.reset(new Tensor()); }
+  SelectedRows() {
+    height_ = 0;
+    value_.reset(new Tensor());
+  }
 
   platform::Place place() const { return value_->place(); }
 
@@ -37,6 +40,8 @@ class SelectedRows {
 
   const Vector<int64_t>& rows() const { return rows_; }
 
+  Vector<int64_t>* mutable_rows() { return &rows_; }
+
   void set_rows(const Vector<int64_t>& rows) { rows_ = rows; }
 
   DDim GetCompleteDims() const {

paddle/framework/tensor.h

@@ -132,6 +132,8 @@ class Tensor {
 
   std::type_index type() const { return holder_->type(); }
 
+  size_t memory_size() const;
+
  private:
   inline void check_memory_size() const;
 

paddle/framework/tensor_array.cc

@@ -254,13 +254,12 @@ LoDTensor TensorArray::LodPackTwo(const LoDTensor& pre, const LoDTensor& cur,
 void TensorArray::LodUnpack(const LoDTensor& source, size_t level) {
   PADDLE_ENFORCE_EQ(level, source.NumLevels() - 1,
                     "only the lowest LoD level supports unpack.");
-  int non_empty_instances = -1;
+  const size_t non_empty_instances = source.dims()[0];
   size_t index = 0;
   Vector<size_t> lowest_lod_level;
   lowest_lod_level.push_back(index);
 
-  for (size_t step = 0; non_empty_instances > 0 || non_empty_instances == -1;
-       step++) {
+  for (size_t step = 0; step < non_empty_instances; step++) {
     size_t num_instances = 0;
     for (size_t id = 0; id < source.NumElements(level); id++) {
       auto instance = source;

paddle/framework/tensor_impl.h

@@ -62,12 +62,16 @@ inline void Tensor::check_memory_size() const {
   PADDLE_ENFORCE_NOT_NULL(
       holder_, "Tensor holds no memory. Call Tensor::mutable_data first.");
   PADDLE_ENFORCE_GE(
-      holder_->size(), numel() * SizeOfType(type()) + offset_,
+      holder_->size(), memory_size() + offset_,
       "Tensor's dims_ is out of bound. Call Tensor::mutable_data "
       "first to re-allocate memory.\n"
       "or maybe the required data-type mismatches the data already stored.");
 }
 
+inline size_t Tensor::memory_size() const {
+  return holder_ == nullptr ? 0UL : numel() * SizeOfType(type());
+}
+
 template <typename T>
 inline const T* Tensor::data() const {
   check_memory_size();

paddle/framework/type_defs.h

@@ -28,6 +28,8 @@ class OperatorBase;
 class OpDescBind;
 class BlockDescBind;
 class BlockDesc;
+class InferShapeContext;
+
 using VariableNameMap = std::map<std::string, std::vector<std::string>>;
 
 // The order should be as same as framework.proto
@@ -49,5 +51,7 @@ using GradOpMakerFN = std::function<std::vector<std::unique_ptr<OpDescBind>>(
 using InferVarTypeFN = std::function<void(const OpDescBind& /*op_desc*/,
                                           BlockDescBind* /*block*/)>;
 
+using InferShapeFN = std::function<void(InferShapeContext*)>;
+
 }  // namespace framework
 }  // namespace paddle

paddle/framework/var_desc.h

@@ -59,6 +59,8 @@ class VarDescBind {
     desc_.set_type(VarDesc::LOD_TENSOR);
   }
 
+  explicit VarDescBind(const VarDesc &desc) : desc_(desc) {}
+
   VarDesc *Proto() { return &desc_; }
 
   std::string Name() const { return desc_.name(); }

paddle/framework/variable.h

@@ -46,6 +46,8 @@ class Variable {
            std::type_index(typeid(T)) == std::type_index(holder_->Type());
   }
 
+  void Clear() { holder_.reset(); }
+
  private:
   struct Placeholder {
     virtual ~Placeholder() {}

paddle/memory/memcpy.h

@@ -54,6 +54,5 @@ void Copy(DstPlace, void* dst, SrcPlace, const void* src, size_t num,
           cudaStream_t stream);
 
 #endif
-
 }  // namespace memory
 }  // namespace paddle

paddle/operators/CMakeLists.txt

@@ -82,7 +82,7 @@ function(op_library TARGET)
         # It's enough to just adding one operator to pybind
         file(APPEND ${pybind_file} "USE_OP(sigmoid);\n")
     endif()
-    
+
     # reduce_op contains several operators
     if ("${TARGET}" STREQUAL "reduce_op")
         set(pybind_flag 1)
@@ -123,6 +123,7 @@ set(DEPS_OPS
     sum_op
     pool_op
     pool_with_index_op
+    sequence_conv_op
     lstm_op)
 
 
@@ -131,9 +132,10 @@ op_library(recurrent_op SRCS recurrent_op.cc rnn/recurrent_op_utils.cc
 op_library(cond_op SRCS cond_op.cc DEPS framework_proto tensor operator net_op)
 op_library(cross_entropy_op DEPS cross_entropy)
 op_library(softmax_with_cross_entropy_op DEPS cross_entropy softmax)
-op_library(sum_op DEPS net_op)
+op_library(sum_op DEPS net_op selected_rows_functor)
 op_library(pool_op DEPS pooling)
 op_library(pool_with_index_op DEPS pooling)
+op_library(sequence_conv_op DEPS context_project)
 op_library(lstm_op DEPS sequence2batch lstm_compute)
 
 list(REMOVE_ITEM GENERAL_OPS ${DEPS_OPS})
@@ -148,3 +150,4 @@ cc_test(net_op_test SRCS net_op_test.cc DEPS net_op)
 cc_test(scatter_test SRCS scatter_test.cc DEPS tensor)
 cc_test(strided_memcpy_test SRCS strided_memcpy_test.cc DEPS tensor paddle_memory)
 cc_test(dynamic_recurrent_op_test SRCS dynamic_recurrent_op_test.cc DEPS dynamic_recurrent_op recurrent_op tensor_array)
+cc_test(save_load_op_test SRCS save_load_op_test.cc DEPS save_op load_op)

paddle/operators/activation_op.cc

@@ -446,12 +446,16 @@ REGISTER_OP(thresholded_relu, ops::ActivationOp,
 REGISTER_OP(hard_sigmoid, ops::ActivationOp, ops::HardSigmoidOpMaker<float>,
             hard_sigmoid_grad, ops::ActivationOpGrad);
 
-#define REGISTER_ACTIVATION_CPU_KERNEL(act_type, functor, grad_functor)        \
-  REGISTER_OP_CPU_KERNEL(                                                      \
-      act_type,                                                                \
-      ops::ActivationKernel<paddle::platform::CPUPlace, ops::functor<float>>); \
-  REGISTER_OP_CPU_KERNEL(act_type##_grad,                                      \
-                         ops::ActivationGradKernel<paddle::platform::CPUPlace, \
-                                                   ops::grad_functor<float>>);
+#define REGISTER_ACTIVATION_CPU_KERNEL(act_type, functor, grad_functor)       \
+  REGISTER_OP_CPU_KERNEL(                                                     \
+      act_type,                                                               \
+      ops::ActivationKernel<paddle::platform::CPUPlace, ops::functor<float>>, \
+      ops::ActivationKernel<paddle::platform::CPUPlace,                       \
+                            ops::functor<double>>);                           \
+  REGISTER_OP_CPU_KERNEL(                                                     \
+      act_type##_grad, ops::ActivationGradKernel<paddle::platform::CPUPlace,  \
+                                                 ops::grad_functor<float>>,   \
+      ops::ActivationGradKernel<paddle::platform::CPUPlace,                   \
+                                ops::grad_functor<double>>);
 
 FOR_EACH_KERNEL_FUNCTOR(REGISTER_ACTIVATION_CPU_KERNEL);

paddle/operators/activation_op.cu

@@ -17,12 +17,16 @@
 
 namespace ops = paddle::operators;
 
-#define REGISTER_ACTIVATION_GPU_KERNEL(act_type, functor, grad_functor)        \
-  REGISTER_OP_GPU_KERNEL(                                                      \
-      act_type,                                                                \
-      ops::ActivationKernel<paddle::platform::GPUPlace, ops::functor<float>>); \
-  REGISTER_OP_GPU_KERNEL(act_type##_grad,                                      \
-                         ops::ActivationGradKernel<paddle::platform::GPUPlace, \
-                                                   ops::grad_functor<float>>);
+#define REGISTER_ACTIVATION_GPU_KERNEL(act_type, functor, grad_functor)       \
+  REGISTER_OP_GPU_KERNEL(                                                     \
+      act_type,                                                               \
+      ops::ActivationKernel<paddle::platform::GPUPlace, ops::functor<float>>, \
+      ops::ActivationKernel<paddle::platform::GPUPlace,                       \
+                            ops::functor<double>>);                           \
+  REGISTER_OP_GPU_KERNEL(                                                     \
+      act_type##_grad, ops::ActivationGradKernel<paddle::platform::GPUPlace,  \
+                                                 ops::grad_functor<float>>,   \
+      ops::ActivationGradKernel<paddle::platform::GPUPlace,                   \
+                                ops::grad_functor<double>>);
 
 FOR_EACH_KERNEL_FUNCTOR(REGISTER_ACTIVATION_GPU_KERNEL);

paddle/operators/activation_op.h

@@ -210,8 +210,8 @@ struct HardShrinkFunctor : public BaseActivationFunctor<T> {
   }
   template <typename Device, typename X, typename Y>
   void operator()(Device d, X x, Y y) const {
-    auto temp1 = (x < (threshold * -1)).template cast<T>().eval();
-    auto temp2 = (x > threshold).template cast<T>().eval();
+    auto temp1 = (x < static_cast<T>(threshold * -1)).template cast<T>().eval();
+    auto temp2 = (x > static_cast<T>(threshold)).template cast<T>().eval();
     y.device(d) = x * (temp1 + temp2);
   }
 };
@@ -226,8 +226,8 @@ struct HardShrinkGradFunctor : public BaseActivationFunctor<T> {
 
   template <typename Device, typename X, typename Y, typename dY, typename dX>
   void operator()(Device d, X x, Y y, dY dy, dX dx) const {
-    auto temp1 = (x < (threshold * -1)).template cast<T>().eval();
-    auto temp2 = (x > threshold).template cast<T>().eval();
+    auto temp1 = (x < static_cast<T>(threshold * -1)).template cast<T>().eval();
+    auto temp2 = (x > static_cast<T>(threshold)).template cast<T>().eval();
     dx.device(d) = dy * (temp1 + temp2).template cast<T>();
   }
 };
@@ -243,9 +243,10 @@ struct SoftShrinkFunctor : public BaseActivationFunctor<T> {
 
   template <typename Device, typename X, typename Y>
   void operator()(Device d, X x, Y y) const {
-    auto temp1 = (x > lambda).template cast<T>().eval();
-    auto temp2 = (x < -lambda).template cast<T>().eval();
-    y.device(d) = temp1 * (x - lambda) + temp2 * (x + lambda);
+    auto lambdaT = static_cast<T>(lambda);
+    auto temp1 = (x > lambdaT).template cast<T>().eval();
+    auto temp2 = (x < -lambdaT).template cast<T>().eval();
+    y.device(d) = temp1 * (x - lambdaT) + temp2 * (x + lambdaT);
   }
 };
 
@@ -257,8 +258,9 @@ struct SoftShrinkGradFunctor : public BaseActivationFunctor<T> {
   }
   template <typename Device, typename X, typename Y, typename dY, typename dX>
   void operator()(Device d, X x, Y y, dY dy, dX dx) const {
-    auto temp1 = (x > lambda).template cast<T>().eval();
-    auto temp2 = (x < -lambda).template cast<T>().eval();
+    auto lambdaT = static_cast<T>(lambda);
+    auto temp1 = (x > lambdaT).template cast<T>().eval();
+    auto temp2 = (x < -lambdaT).template cast<T>().eval();
     dx.device(d) = dy * (temp1 + temp2).template cast<T>();
   }
 };
@@ -362,7 +364,8 @@ struct BReluFunctor : public BaseActivationFunctor<T> {
 
   template <typename Device, typename X, typename Y>
   void operator()(Device d, X x, Y y) const {
-    y.device(d) = x.cwiseMax(t_min).cwiseMin(t_max);
+    y.device(d) =
+        x.cwiseMax(static_cast<T>(t_min)).cwiseMin(static_cast<T>(t_max));
   }
 };
 
@@ -375,7 +378,9 @@ struct BReluGradFunctor : public BaseActivationFunctor<T> {
   }
   template <typename Device, typename X, typename Y, typename dY, typename dX>
   void operator()(Device d, X x, Y y, dY dy, dX dx) const {
-    dx.device(d) = dy * ((x > t_min) * (x < t_max)).template cast<T>();
+    dx.device(d) = dy *
+                   ((x > static_cast<T>(t_min)) * (x < static_cast<T>(t_max)))
+                       .template cast<T>();
   }
 };
 
@@ -390,7 +395,8 @@ struct Relu6Functor : public BaseActivationFunctor<T> {
 
   template <typename Device, typename X, typename Y>
   void operator()(Device d, X x, Y y) const {
-    y.device(d) = x.cwiseMax(static_cast<T>(0)).cwiseMin(threshold);
+    y.device(d) =
+        x.cwiseMax(static_cast<T>(0)).cwiseMin(static_cast<T>(threshold));
   }
 };
 
@@ -402,8 +408,9 @@ struct Relu6GradFunctor : public BaseActivationFunctor<T> {
   }
   template <typename Device, typename X, typename Y, typename dY, typename dX>
   void operator()(Device d, X x, Y y, dY dy, dX dx) const {
-    dx.device(d) =
-        dy * ((x > static_cast<T>(0)) * (x < threshold)).template cast<T>();
+    dx.device(d) = dy *
+                   ((x > static_cast<T>(0)) * (x < static_cast<T>(threshold)))
+                       .template cast<T>();
   }
 };
 
@@ -463,7 +470,8 @@ struct SoftReluFunctor : public BaseActivationFunctor<T> {
 
   template <typename Device, typename X, typename Y>
   void operator()(Device d, X x, Y y) const {
-    auto temp = x.cwiseMax(-threshold).cwiseMin(threshold);
+    auto tmp = static_cast<T>(threshold);
+    auto temp = x.cwiseMax(-tmp).cwiseMin(tmp);
     y.device(d) = (static_cast<T>(1) + temp.exp()).log();
   }
 };
@@ -476,7 +484,8 @@ struct SoftReluGradFunctor : public BaseActivationFunctor<T> {
   }
   template <typename Device, typename X, typename Y, typename dY, typename dX>
   void operator()(Device d, X x, Y y, dY dy, dX dx) const {
-    auto temp = ((x > -threshold) * (x < threshold)).template cast<T>().eval();
+    auto tmp = static_cast<T>(threshold);
+    auto temp = ((x > -tmp) * (x < tmp)).template cast<T>().eval();
     dx.device(d) = dy * (static_cast<T>(1) - (-y).exp()) * temp;
   }
 };
@@ -490,7 +499,7 @@ struct LeakyReluFunctor : public BaseActivationFunctor<T> {
 
   template <typename Device, typename X, typename Y>
   void operator()(Device d, X x, Y y) const {
-    y.device(d) = x.cwiseMax(alpha * x);
+    y.device(d) = x.cwiseMax(static_cast<T>(alpha) * x);
   }
 };
 
@@ -502,7 +511,8 @@ struct LeakyReluGradFunctor : public BaseActivationFunctor<T> {
   }
   template <typename Device, typename X, typename Y, typename dY, typename dX>
   void operator()(Device d, X x, Y y, dY dy, dX dx) const {
-    auto temp1 = alpha * (x < static_cast<T>(0)).template cast<T>().eval();
+    auto temp1 = static_cast<T>(alpha) *
+                 (x < static_cast<T>(0)).template cast<T>().eval();
     auto temp2 = (x >= static_cast<T>(0)).template cast<T>().eval();
     dx.device(d) = dy * (temp1 + temp2).template cast<T>();
   }
@@ -517,9 +527,9 @@ struct ELUFunctor : public BaseActivationFunctor<T> {
 
   template <typename Device, typename X, typename Y>
   void operator()(Device d, X x, Y y) const {
-    y.device(d) =
-        x.cwiseMax(static_cast<T>(0)) +
-        (alpha * (x.exp() - static_cast<T>(1))).cwiseMin(static_cast<T>(0));
+    y.device(d) = x.cwiseMax(static_cast<T>(0)) +
+                  (static_cast<T>(alpha) * (x.exp() - static_cast<T>(1)))
+                      .cwiseMin(static_cast<T>(0));
   }
 };
 
@@ -531,9 +541,9 @@ struct ELUGradFunctor : public BaseActivationFunctor<T> {
   }
   template <typename Device, typename X, typename Y, typename dY, typename dX>
   void operator()(Device d, X x, Y y, dY dy, dX dx) const {
-    dx.device(d) =
-        dy * (x > static_cast<T>(0)).template cast<T>() +
-        dy * (y + alpha) * (x < static_cast<T>(0)).template cast<T>();
+    dx.device(d) = dy * (x > static_cast<T>(0)).template cast<T>() +
+                   dy * (y + static_cast<T>(alpha)) *
+                       (x < static_cast<T>(0)).template cast<T>();
   }
 };
 
@@ -545,7 +555,7 @@ struct PowFunctor : public BaseActivationFunctor<T> {
   }
   template <typename Device, typename X, typename Y>
   void operator()(Device d, X x, Y y) const {
-    y.device(d) = x.pow(factor);
+    y.device(d) = x.pow(static_cast<T>(factor));
   }
 };
 
@@ -557,7 +567,8 @@ struct PowGradFunctor : public BaseActivationFunctor<T> {
   }
   template <typename Device, typename X, typename Y, typename dY, typename dX>
   void operator()(Device d, X x, Y y, dY dy, dX dx) const {
-    dx.device(d) = dy * factor * x.pow(factor - static_cast<T>(1));
+    dx.device(d) = dy * static_cast<T>(factor) *
+                   x.pow(static_cast<T>(factor - static_cast<T>(1)));
   }
 };
 
@@ -571,7 +582,8 @@ struct STanhFunctor : public BaseActivationFunctor<T> {
 
   template <typename Device, typename X, typename Y>
   void operator()(Device d, X x, Y y) const {
-    y.device(d) = scale_b * (scale_a * x).tanh();
+    y.device(d) =
+        static_cast<T>(scale_b) * (static_cast<T>(scale_a) * x).tanh();
   }
 };
 
@@ -585,8 +597,10 @@ struct STanhGradFunctor : public BaseActivationFunctor<T> {
 
   template <typename Device, typename X, typename Y, typename dY, typename dX>
   void operator()(Device d, X x, Y y, dY dy, dX dx) const {
-    auto temp = (scale_a * x).tanh() * (scale_a * x).tanh();
-    dx.device(d) = dy * scale_a * scale_b * (static_cast<T>(1) - temp);
+    auto a = static_cast<T>(scale_a);
+    auto b = static_cast<T>(scale_b);
+    auto temp = (a * x).tanh() * (a * x).tanh();
+    dx.device(d) = dy * a * b * (static_cast<T>(1) - temp);
   }
 };
 
@@ -599,7 +613,8 @@ struct ThresholdedReluFunctor : public BaseActivationFunctor<T> {
 
   template <typename Device, typename X, typename Y>
   void operator()(Device d, X x, Y y) const {
-    y.device(d) = (x > static_cast<T>(threshold)).template cast<T>() * x;
+    auto th = static_cast<T>(threshold);
+    y.device(d) = (x > th).template cast<T>() * x;
   }
 };
 
@@ -612,7 +627,8 @@ struct ThresholdedReluGradFunctor : public BaseActivationFunctor<T> {
 
   template <typename Device, typename X, typename Y, typename dY, typename dX>
   void operator()(Device d, X x, Y y, dY dy, dX dx) const {
-    dx.device(d) = dy * (x > static_cast<T>(threshold)).template cast<T>();
+    auto th = static_cast<T>(threshold);
+    dx.device(d) = dy * (x > th).template cast<T>();
   }
 };
 

paddle/operators/cross_entropy_op.cc

@@ -162,6 +162,8 @@ or not. But the output only shares the LoD with input `X`.
 namespace ops = paddle::operators;
 REGISTER_OP(cross_entropy, ops::CrossEntropyOp, ops::CrossEntropyOpMaker,
             cross_entropy_grad, ops::CrossEntropyGradientOp);
-REGISTER_OP_CPU_KERNEL(cross_entropy, ops::CrossEntropyOpKernel<float>);
+REGISTER_OP_CPU_KERNEL(cross_entropy, ops::CrossEntropyOpKernel<float>,
+                       ops::CrossEntropyOpKernel<double>);
 REGISTER_OP_CPU_KERNEL(cross_entropy_grad,
-                       ops::CrossEntropyGradientOpKernel<float>);
+                       ops::CrossEntropyGradientOpKernel<float>,
+                       ops::CrossEntropyGradientOpKernel<double>);

paddle/operators/cross_entropy_op.cu

@@ -108,6 +108,8 @@ class CrossEntropyGradientOpCUDAKernel : public framework::OpKernel<T> {
 }  // namespace paddle
 
 namespace ops = paddle::operators;
-REGISTER_OP_GPU_KERNEL(cross_entropy, ops::CrossEntropyOpCUDAKernel<float>);
+REGISTER_OP_GPU_KERNEL(cross_entropy, ops::CrossEntropyOpCUDAKernel<float>,
+                       ops::CrossEntropyOpCUDAKernel<double>);
 REGISTER_OP_GPU_KERNEL(cross_entropy_grad,
-                       ops::CrossEntropyGradientOpCUDAKernel<float>);
+                       ops::CrossEntropyGradientOpCUDAKernel<float>,
+                       ops::CrossEntropyGradientOpCUDAKernel<double>);

paddle/operators/dropout_op.cc

@@ -30,7 +30,7 @@ class DropoutOp : public framework::OperatorWithKernel {
 
     auto x_dims = ctx->GetInputDim("X");
     ctx->SetOutputDim("Out", x_dims);
-    if (ctx->Attrs().Get<bool>("is_training") == 1) {
+    if (ctx->Attrs().Get<bool>("is_training") == true) {
       ctx->SetOutputDim("Mask", x_dims);
     }
     ctx->ShareLoD("X", /*->*/ "Out");
@@ -43,7 +43,7 @@ class DropoutOpMaker : public framework::OpProtoAndCheckerMaker {
   DropoutOpMaker(framework::OpProto* proto,
                  framework::OpAttrChecker* op_checker)
       : OpProtoAndCheckerMaker(proto, op_checker) {
-    AddAttr<AttrType>("dropout_prob", "Probability of setting units to zero.")
+    AddAttr<float>("dropout_prob", "Probability of setting units to zero.")
         .SetDefault(.5f);
     AddAttr<bool>("is_training", "Whether in training phase.").SetDefault(true);
     AddAttr<int>("seed", "Dropout random seed.").SetDefault(0);
@@ -69,7 +69,7 @@ class DropoutOpGrad : public framework::OperatorWithKernel {
   using framework::OperatorWithKernel::OperatorWithKernel;
 
   void InferShape(framework::InferShapeContext* ctx) const override {
-    PADDLE_ENFORCE_EQ(ctx->Attrs().Get<bool>("is_training"), 1,
+    PADDLE_ENFORCE_EQ(ctx->Attrs().Get<bool>("is_training"), true,
                       "GradOp is only callable when is_training is true");
 
     PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) must not be null.");
@@ -77,8 +77,8 @@ class DropoutOpGrad : public framework::OperatorWithKernel {
     PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")),
                    "Input(Out@GRAD) must not be null.");
 
-    PADDLE_ENFORCE_GE(ctx->Attrs().Get<AttrType>("dropout_prob"), 0);
-    PADDLE_ENFORCE_LE(ctx->Attrs().Get<AttrType>("dropout_prob"), 1);
+    PADDLE_ENFORCE_GE(ctx->Attrs().Get<float>("dropout_prob"), 0);
+    PADDLE_ENFORCE_LE(ctx->Attrs().Get<float>("dropout_prob"), 1);
     auto x_dims = ctx->GetInputDim("X");
     auto out_dims = ctx->GetInputDim(framework::GradVarName("Out"));
     PADDLE_ENFORCE_EQ(x_dims, out_dims,

paddle/operators/dropout_op.h

@@ -33,7 +33,7 @@ class CPUDropoutKernel : public framework::OpKernel<T> {
     auto* y = context.Output<Tensor>("Out");
     const auto* x_data = x->data<T>();
     auto* y_data = y->mutable_data<T>(context.GetPlace());
-    AttrType dropout_prob = context.Attr<AttrType>("dropout_prob");
+    float dropout_prob = context.Attr<float>("dropout_prob");
 
     if (context.Attr<bool>("is_training")) {
       auto* mask = context.Output<Tensor>("Mask");
@@ -41,7 +41,7 @@ class CPUDropoutKernel : public framework::OpKernel<T> {
       int seed = context.Attr<int>("seed");
       std::minstd_rand engine;
       engine.seed(seed);
-      std::uniform_real_distribution<AttrType> dist(0, 1);
+      std::uniform_real_distribution<float> dist(0, 1);
       size_t size = framework::product(mask->dims());
       for (size_t i = 0; i < size; ++i) {
         if (dist(engine) < dropout_prob) {

paddle/operators/fetch_op.cc

@@ -52,6 +52,7 @@ class FetchOp : public framework::OperatorBase {
     // FIXME(yuyang18): Should we assume the fetch operator always generate
     // CPU outputs?
     dst_item.CopyFrom(src_item, platform::CPUPlace(), dev_ctx);
+    dev_ctx.Wait();
     dst_item.set_lod(src_item.lod());
 
     VLOG(3) << "Fetch variable " << fetch_var_name << " to " << out_name;

paddle/operators/fill_constant_batch_size_like_op.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed 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 "paddle/operators/fill_constant_batch_size_like_op.h"
+
+namespace paddle {
+namespace operators {
+
+class FillConstantBatchSizeLikeOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+  void InferShape(framework::InferShapeContext *ctx) const override {
+    PADDLE_ENFORCE(
+        ctx->HasInput("Input"),
+        "Input(Input) of FillConstantBatchSizeLikeOp should not be null.");
+    PADDLE_ENFORCE(
+        ctx->HasOutput("Out"),
+        "Output(Out) of FillConstantBatchSizeLikeOp should not be null.");
+
+    auto &shape = ctx->Attrs().Get<std::vector<int>>("shape");
+    PADDLE_ENFORCE_GT(shape.size(), 0);
+    std::vector<int64_t> shape_int64(shape.size(), 0);
+    std::transform(shape.begin(), shape.end(), shape_int64.begin(),
+                   [](int a) { return static_cast<int64_t>(a); });
+    auto dims = framework::make_ddim(shape_int64);
+
+    dims[0] = ctx->GetInputDim("Input")[0];
+    ctx->SetOutputDim("Out", dims);
+  }
+
+ protected:
+  framework::DataType IndicateDataType(
+      const framework::ExecutionContext &ctx) const override {
+    return static_cast<framework::DataType>(ctx.Attr<int>("data_type"));
+  }
+};
+
+class FillConstantBatchSizeLikeOpMaker
+    : public framework::OpProtoAndCheckerMaker {
+ public:
+  FillConstantBatchSizeLikeOpMaker(framework::OpProto *proto,
+                                   framework::OpAttrChecker *op_checker)
+      : framework::OpProtoAndCheckerMaker(proto, op_checker) {
+    AddAttr<int>("data_type",
+                 "(int, default 5 (FP32)) "
+                 "Output data type")
+        .SetDefault(framework::DataType::FP32);
+    AddAttr<std::vector<int>>("shape", "(vector<int>) The shape of the output");
+    AddAttr<float>("value", "(float, default 0) The value to be filled")
+        .SetDefault(0.0f);
+    AddInput("Input",
+             "(Tensor) Tensor "
+             "whose first dimension is used to specify the batch_size");
+    AddOutput("Out",
+              "(Tensor) Tensor of specified shape will be filled "
+              "with the specified value");
+    AddComment(R"DOC(Fill up a variable with specified constant value.)DOC");
+  }
+};
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_WITHOUT_GRADIENT(fill_constant_batch_size_like,
+                             ops::FillConstantBatchSizeLikeOp,
+                             ops::FillConstantBatchSizeLikeOpMaker);
+REGISTER_OP_CPU_KERNEL(
+    fill_constant_batch_size_like,
+    ops::FillConstantBatchSizeLikeOpKernel<paddle::platform::CPUPlace, float>,
+    ops::FillConstantBatchSizeLikeOpKernel<paddle::platform::CPUPlace, double>);

paddle/operators/fill_constant_batch_size_like_op.cu

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+   Licensed 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. */
+
+#define EIGEN_USE_GPU
+#include "paddle/framework/op_registry.h"
+#include "paddle/operators/fill_constant_batch_size_like_op.h"
+
+namespace ops = paddle::operators;
+REGISTER_OP_GPU_KERNEL(
+    fill_constant_batch_size_like,
+    ops::FillConstantBatchSizeLikeOpKernel<paddle::platform::GPUPlace, float>,
+    ops::FillConstantBatchSizeLikeOpKernel<paddle::platform::GPUPlace, double>);

paddle/operators/fill_constant_batch_size_like_op.h

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed 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. */
+
+#pragma once
+#include "paddle/framework/eigen.h"
+#include "paddle/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+template <typename Place, typename T>
+class FillConstantBatchSizeLikeOpKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    auto* out = ctx.Output<framework::Tensor>("Out");
+    out->mutable_data<T>(ctx.GetPlace());
+    auto value = ctx.Attr<float>("value");
+
+    auto out_eigen = framework::EigenVector<T>::Flatten(*out);
+    auto place = ctx.GetEigenDevice<Place>();
+    out_eigen.device(place) = out_eigen.constant(static_cast<T>(value));
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle

paddle/operators/fill_constant_op.cc

@@ -64,5 +64,6 @@ namespace ops = paddle::operators;
 REGISTER_OP_WITHOUT_GRADIENT(fill_constant, ops::FillConstantOp,
                              ops::FillConstantOpMaker);
 REGISTER_OP_CPU_KERNEL(
-    fill_constant,
-    ops::FillConstantOpKernel<paddle::platform::CPUPlace, float>);
+    fill_constant, ops::FillConstantOpKernel<paddle::platform::CPUPlace, float>,
+    ops::FillConstantOpKernel<paddle::platform::CPUPlace, double>,
+    ops::FillConstantOpKernel<paddle::platform::CPUPlace, int>);

paddle/operators/fill_constant_op.cu

@@ -18,5 +18,6 @@
 
 namespace ops = paddle::operators;
 REGISTER_OP_GPU_KERNEL(
-    fill_constant,
-    ops::FillConstantOpKernel<paddle::platform::GPUPlace, float>);
+    fill_constant, ops::FillConstantOpKernel<paddle::platform::GPUPlace, float>,
+    ops::FillConstantOpKernel<paddle::platform::GPUPlace, double>,
+    ops::FillConstantOpKernel<paddle::platform::GPUPlace, int>);

paddle/operators/fill_constant_op.h

@@ -25,7 +25,7 @@ class FillConstantOpKernel : public framework::OpKernel<T> {
   void Compute(const framework::ExecutionContext& ctx) const override {
     auto* out = ctx.Output<framework::Tensor>("Out");
     out->mutable_data<T>(ctx.GetPlace());
-    auto value = ctx.Attr<T>("value");
+    auto value = ctx.Attr<float>("value");
 
     auto out_eigen = framework::EigenVector<T>::Flatten(*out);
     auto place = ctx.GetEigenDevice<Place>();

paddle/operators/gru_unit_op.cc

@@ -171,8 +171,7 @@ class GRUUnitGradOp : public framework::OperatorWithKernel {
     PADDLE_ENFORCE_EQ(
         weight_width, frame_size * 3,
         "The shape of Weight matrix must be [frame_size, frame_size * 3].");
-    auto bias = Input("Bias");
-    if (bias != framework::kEmptyVarName) {
+    if (ctx->HasInput("Bias")) {
       auto bias_dims = ctx->GetInputDim("Bias");
       int bias_height = bias_dims[0];
       int bias_width = bias_dims[1];
@@ -203,6 +202,8 @@ namespace ops = paddle::operators;
 REGISTER_OP(gru_unit, ops::GRUUnitOp, ops::GRUUnitOpMaker, gru_unit_grad,
             ops::GRUUnitGradOp);
 REGISTER_OP_CPU_KERNEL(gru_unit,
-                       ops::GRUUnitKernel<paddle::platform::CPUPlace, float>);
+                       ops::GRUUnitKernel<paddle::platform::CPUPlace, float>,
+                       ops::GRUUnitKernel<paddle::platform::CPUPlace, double>);
 REGISTER_OP_CPU_KERNEL(
-    gru_unit_grad, ops::GRUUnitGradKernel<paddle::platform::CPUPlace, float>);
+    gru_unit_grad, ops::GRUUnitGradKernel<paddle::platform::CPUPlace, float>,
+    ops::GRUUnitGradKernel<paddle::platform::CPUPlace, double>);

paddle/operators/gru_unit_op.cu

@@ -17,6 +17,8 @@
 
 namespace ops = paddle::operators;
 REGISTER_OP_GPU_KERNEL(gru_unit,
-                       ops::GRUUnitKernel<paddle::platform::GPUPlace, float>);
+                       ops::GRUUnitKernel<paddle::platform::GPUPlace, float>,
+                       ops::GRUUnitKernel<paddle::platform::GPUPlace, double>);
 REGISTER_OP_GPU_KERNEL(
-    gru_unit_grad, ops::GRUUnitGradKernel<paddle::platform::GPUPlace, float>);
+    gru_unit_grad, ops::GRUUnitGradKernel<paddle::platform::GPUPlace, float>,
+    ops::GRUUnitGradKernel<paddle::platform::GPUPlace, double>);

paddle/operators/l1_norm_op.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+   Licensed 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 "paddle/operators/l1_norm_op.h"
+
+namespace paddle {
+namespace operators {
+
+using framework::Tensor;
+
+class L1NormOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should be not null.");
+    PADDLE_ENFORCE(ctx->HasOutput("Out"), "Output(Out) should be not null.");
+
+    ctx->SetOutputDim("Out", {1});
+  }
+};
+
+class L1NormGradOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should be not null.");
+    PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")),
+                   "Input(Out@GRAD) should be not null.");
+    PADDLE_ENFORCE(ctx->HasOutput(framework::GradVarName("X")),
+                   "Output(X@GRAD) should be not null.");
+
+    ctx->SetOutputDim(framework::GradVarName("X"), ctx->GetInputDim("X"));
+  }
+};
+
+class L1NormOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  L1NormOpMaker(framework::OpProto* proto, framework::OpAttrChecker* op_checker)
+      : framework::OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput("X", "(Tensor) The input of l1_norm op.");
+    AddOutput("Out", "(Scalar) The output of l1_norm op.");
+    AddComment(R"DOC(
+L1 Norm Operator.
+
+Computes the L1 norm of a tensor.
+
+Out = sum (abs(X))
+
+)DOC");
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP(l1_norm, ops::L1NormOp, ops::L1NormOpMaker, l1_norm_grad,
+            ops::L1NormGradOp);
+REGISTER_OP_CPU_KERNEL(l1_norm,
+                       ops::L1NormKernel<paddle::platform::CPUPlace, float>);
+REGISTER_OP_CPU_KERNEL(
+    l1_norm_grad, ops::L1NormGradKernel<paddle::platform::CPUPlace, float>);

paddle/operators/l1_norm_op.cu

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+   Licensed 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. */
+
+#define EIGEN_USE_GPU
+#include "paddle/operators/l1_norm_op.h"
+
+namespace ops = paddle::operators;
+REGISTER_OP_GPU_KERNEL(l1_norm,
+                       ops::L1NormKernel<paddle::platform::GPUPlace, float>);
+REGISTER_OP_GPU_KERNEL(
+    l1_norm_grad, ops::L1NormGradKernel<paddle::platform::GPUPlace, float>);

paddle/operators/l1_norm_op.h

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+   Licensed 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. */
+
+#pragma once
+#include "paddle/framework/eigen.h"
+#include "paddle/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+// Out = sum(abs(X))
+template <typename Place, typename T>
+class L1NormKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext &context) const override {
+    const framework::Tensor *X = context.Input<framework::Tensor>("X");
+    framework::Tensor *Out = context.Output<framework::Tensor>("Out");
+    Out->mutable_data<T>(context.GetPlace());
+
+    auto x = framework::EigenVector<T>::Flatten(*X);
+    auto out = framework::EigenVector<T>::Flatten(*Out);
+    auto place = context.GetEigenDevice<Place>();
+
+    out.device(place) = x.abs().sum();
+  }
+};
+
+// dX = dout * sign(X)
+template <typename Place, typename T>
+class L1NormGradKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext &context) const override {
+    const framework::Tensor *x = context.Input<framework::Tensor>("X");
+    const framework::Tensor *d_out =
+        context.Input<framework::Tensor>(framework::GradVarName("Out"));
+    PADDLE_ENFORCE(d_out->numel() == 1, "L1 Norm Gradient should be scalar");
+    framework::Tensor *dx =
+        context.Output<framework::Tensor>(framework::GradVarName("X"));
+    dx->mutable_data<T>(context.GetPlace());
+
+    auto x_eigen = framework::EigenVector<T>::Flatten(*x);
+    auto d_out_eigen = framework::EigenVector<T>::Flatten(*d_out);
+    auto dx_eigen = framework::EigenVector<T>::Flatten(*dx);
+    auto place = context.GetEigenDevice<Place>();
+
+    Eigen::DSizes<int, 1> x_dsize(x->numel());
+    dx_eigen.device(place) = d_out_eigen.broadcast(x_dsize) * x_eigen.sign();
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle

paddle/operators/load_op.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserved.
+
+   Licensed 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 "paddle/framework/op_registry.h"
+
+#include <fstream>
+
+namespace paddle {
+namespace operators {
+
+class LoadOp : public framework::OperatorBase {
+ public:
+  LoadOp(const std::string &type, const framework::VariableNameMap &inputs,
+         const framework::VariableNameMap &outputs,
+         const framework::AttributeMap &attrs)
+      : OperatorBase(type, inputs, outputs, attrs) {}
+  void Run(const framework::Scope &scope,
+           const platform::DeviceContext &dev_ctx) const override {
+    auto filename = Attr<std::string>("file_path");
+    std::ifstream fin(filename);
+    PADDLE_ENFORCE(static_cast<bool>(fin), "Cannot open file %s for load op",
+                   filename);
+
+    auto out_var_name = Output("Out");
+    auto *out_var = scope.FindVar(out_var_name);
+    PADDLE_ENFORCE(out_var != nullptr, "Output variable %s cannot be found",
+                   out_var_name);
+
+    auto *tensor = out_var->GetMutable<framework::LoDTensor>();
+
+    uint32_t version;
+    fin.read(reinterpret_cast<char *>(&version), sizeof(version));
+    PADDLE_ENFORCE_EQ(version, 0U, "Only version 0 is supported");
+    framework::TensorDesc desc;
+    {  // int32_t size
+       // proto buffer
+      int32_t size;
+      fin.read(reinterpret_cast<char *>(&size), sizeof(size));
+      std::unique_ptr<char[]> buf(new char[size]);
+      fin.read(reinterpret_cast<char *>(buf.get()), size);
+      PADDLE_ENFORCE(desc.ParseFromArray(buf.get(), size),
+                     "Cannot parse tensor desc");
+    }
+    {  // read tensor
+      std::vector<int64_t> dims;
+      dims.reserve(static_cast<size_t>(desc.dims().size()));
+      std::copy(desc.dims().begin(), desc.dims().end(),
+                std::back_inserter(dims));
+      tensor->Resize(framework::make_ddim(dims));
+
+      void *buf;
+      platform::Place cpu = platform::CPUPlace();
+      switch (desc.data_type()) {
+        case framework::FP32:
+          buf = tensor->mutable_data<float>(cpu);
+          break;
+        case framework::FP64:
+          buf = tensor->mutable_data<double>(cpu);
+          break;
+        case framework::INT32:
+          buf = tensor->mutable_data<int>(cpu);
+          break;
+        case framework::INT64:
+          buf = tensor->mutable_data<int64_t>(cpu);
+          break;
+        default:
+          PADDLE_THROW("DataType %d not supported", desc.data_type());
+      }
+      fin.read(static_cast<char *>(buf), tensor->memory_size());
+    }
+    {  // read lod
+      uint64_t lod_level;
+      fin.read(reinterpret_cast<char *>(&lod_level), sizeof(lod_level));
+      auto &lod = *tensor->mutable_lod();
+      lod.resize(lod_level);
+      for (uint64_t i = 0; i < lod_level; ++i) {
+        uint64_t size;
+        fin.read(reinterpret_cast<char *>(&size), sizeof(size));
+        std::vector<size_t> tmp(size / sizeof(size_t));
+        fin.read(reinterpret_cast<char *>(tmp.data()),
+                 static_cast<std::streamsize>(size));
+        lod[i] = tmp;
+      }
+    }
+
+    auto place = dev_ctx.GetPlace();
+    if (platform::is_gpu_place(place)) {
+      // copy CPU to GPU
+      framework::LoDTensor cpu_tensor;
+      cpu_tensor.ShareDataWith(*tensor);
+      cpu_tensor.set_lod(tensor->lod());
+
+      // reset tensor
+      out_var->Clear();
+      tensor = out_var->GetMutable<framework::LoDTensor>();
+      tensor->set_lod(cpu_tensor.lod());
+      tensor->CopyFrom(cpu_tensor, place, dev_ctx);
+    }
+  }
+};
+
+class LoadOpProtoMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  LoadOpProtoMaker(framework::OpProto *proto,
+                   framework::OpAttrChecker *op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddOutput("Out", "The tensor need to be loaded");
+    AddComment(R"DOC(Load Operator
+Load operator will load a tensor variable from disk file.
+)DOC");
+    AddAttr<std::string>("file_path",
+                         "Variable will be loaded from \"file_path\".")
+        .AddCustomChecker(
+            [](const std::string &path) { return !path.empty(); });
+  }
+};
+}  // namespace operators
+}  // namespace paddle
+namespace ops = paddle::operators;
+
+REGISTER_OPERATOR(load, ops::LoadOp, ops::LoadOpProtoMaker);

paddle/operators/lrn_op.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+   Licensed 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 "paddle/operators/lrn_op.h"
+
+namespace paddle {
+namespace operators {
+
+using framework::Tensor;
+
+class LRNOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of LRNOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasOutput("Out"),
+                   "Output(Out) of LRNOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasOutput("MidOut"),
+                   "MidOut(Out) of LRNOp should not be null.");
+
+    auto x_dim = ctx->GetInputDim("X");
+    PADDLE_ENFORCE_EQ(x_dim.size(), 4, "Input(X)'rank of LRNOp should be 4.");
+
+    ctx->SetOutputDim("Out", x_dim);
+    ctx->SetOutputDim("MidOut", x_dim);
+    ctx->ShareLoD("X", /*->*/ "Out");
+  }
+};
+
+template <typename T>
+class LRNOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  LRNOpMaker(framework::OpProto* proto, framework::OpAttrChecker* op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput("X", R"DOC(
+ (Tensor) The input of LRN operator. It must be a 4D tenor with NCHW format.
+ )DOC");
+
+    AddOutput("Out",
+              "(Tensor) The output of LRN operator, which is also the 4D "
+              "tensor with NCHW format.");
+    AddOutput("MidOut", R"Doc(
+(Tensor)Middle result of lrn op.It's computed in forward process 
+and also used in backward process.
+    )Doc");
+
+    AddAttr<int>("n", R"DOC(
+(int, default 5)n is “adjacent” kernel maps at the same spatial position.
+        )DOC")
+        .SetDefault(5)
+        .GreaterThan(0);
+
+    AddAttr<T>("k", R"DOC(
+(float, default 2.0)k is the bias.
+        )DOC")
+        .SetDefault(2.0)
+        .GreaterThan(0.0);
+
+    AddAttr<T>("alpha", R"DOC(
+(float, default 0.0001)alpha is the scale number.
+        )DOC")
+        .SetDefault(0.0001)
+        .GreaterThan(0.0);
+
+    AddAttr<T>("beta", R"DOC(
+(float, default 0.75)beta is the power number.
+        )DOC")
+        .SetDefault(0.75)
+        .GreaterThan(0.0);
+
+    AddComment(R"DOC(
+ Local Response Normalization.
+
+ This Function comes from the paper
+ "ImageNet Classification with Deep Convolutional Neural Networks".
+
+ The original formula is:
+
+                                Input(i, x, y)
+ Output(i, x, y) = ----------------------------------------------
+                                 -- upper
+                    (k + alpha * >  (Input(j, x, y))^2) ^ (beta)
+                                 -- j = lower
+
+ upper is `min(C, c + n/2)`
+ lower if `max(0, c - n/2)`
+
+ Function implementation:
+
+ inputs and outpus is NCHW format, while input.shape.ndims() is equal 4.
+ And the meaning of each dimension(0-3) is respectively batch size,
+ feature maps, rows and columns.
+
+ Input and Output in the above formula is for each map(i) of one image, and
+ Input(i, x, y), Output(i, x, y) represents an element in an image.
+
+ C is the number of feature maps of one image, and n is a hyper-parameters
+ is configured when Function is initialized. The sum in the denominator
+ is the sum of the same position in the neighboring maps.
+    )DOC");
+  }
+};
+
+class LRNOpGrad : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should not be null");
+    PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("MidOut")),
+                   "Input(MidOut@GRAD) should not be null");
+    PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")),
+                   "Input(Out@GRAD) should not be null");
+
+    auto x_dims = ctx->GetInputDim("X");
+    ctx->SetOutputDim(framework::GradVarName("X"), x_dims);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP(lrn, ops::LRNOp, ops::LRNOpMaker<float>, lrn_grad, ops::LRNOpGrad);
+REGISTER_OP_CPU_KERNEL(lrn, ops::LRNKernel<paddle::platform::CPUPlace, float>);
+REGISTER_OP_CPU_KERNEL(lrn_grad,
+                       ops::LRNGradKernel<paddle::platform::CPUPlace, float>);

paddle/operators/lrn_op.cu

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+   Licensed 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. */
+
+#define EIGEN_USE_GPU
+#include "paddle/operators/lrn_op.h"
+
+namespace ops = paddle::operators;
+
+REGISTER_OP_GPU_KERNEL(lrn, ops::LRNKernel<paddle::platform::GPUPlace, float>);
+REGISTER_OP_GPU_KERNEL(lrn_grad,
+                       ops::LRNGradKernel<paddle::platform::GPUPlace, float>);

paddle/operators/lrn_op.h

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+   Licensed 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. */
+
+#pragma once
+
+#include "paddle/framework/eigen.h"
+#include "paddle/framework/op_registry.h"
+#include "paddle/operators/math/math_function.h"
+
+namespace paddle {
+namespace operators {
+
+template <typename Place, typename T>
+class LRNKernel : public framework::OpKernel<T> {
+ public:
+  using Tensor = framework::Tensor;
+
+  // f(x) = x * ( k + alpha * SUM((x)^2) )^(-beta)
+  // x represents inputs
+  // f(x) represents outputs
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    // input
+    const Tensor* x = ctx.Input<Tensor>("X");
+    auto x_dims = x->dims();
+
+    // NCHW
+    int N = x_dims[0];
+    int C = x_dims[1];
+    int H = x_dims[2];
+    int W = x_dims[3];
+
+    Tensor* out = ctx.Output<Tensor>("Out");
+    out->mutable_data<T>(ctx.GetPlace());
+
+    // MidOut save the intermediate result for backward
+    Tensor* mid = ctx.Output<Tensor>("MidOut");
+    mid->mutable_data<T>(ctx.GetPlace());
+
+    int n = ctx.Attr<int>("n");
+    T alpha = ctx.Attr<float>("alpha");
+    T beta = ctx.Attr<float>("beta");
+    T k = ctx.Attr<float>("k");
+
+    PADDLE_ENFORCE(n > 0, "n should >= 0");
+    PADDLE_ENFORCE(alpha >= 0.0, "alpha should >= 0.0");
+    PADDLE_ENFORCE(beta >= 0.0, "beta should >= 0.0");
+    PADDLE_ENFORCE(k >= 0.0, "k should >= 0.0");
+
+    auto x_v = framework::EigenVector<T>::Flatten(*x);
+
+    const int start = -(n - 1) / 2;
+    const int end = start + n;
+
+    auto e_mid = framework::EigenTensor<T, 4>::From(*mid);
+    e_mid.device(ctx.GetEigenDevice<Place>()) = e_mid.constant(k);
+
+    auto e_x = framework::EigenTensor<T, 4>::From(*x);
+    for (int m = 0; m < N; m++) {
+      for (int i = 0; i < C; i++) {
+        for (int c = start; c <= end; c++) {
+          int ch = i + c;
+          if (ch >= 0 && ch < C) {
+            auto s = e_mid.slice(Eigen::array<int, 4>({{m, i, 0, 0}}),
+                                 Eigen::array<int, 4>({{1, 1, H, W}}));
+
+            auto r = e_x.slice(Eigen::array<int, 4>({{m, ch, 0, 0}}),
+                               Eigen::array<int, 4>({{1, 1, H, W}}));
+
+            s.device(ctx.GetEigenDevice<Place>()) += alpha * r.square();
+          }
+        }
+      }
+    }
+
+    auto out_e = framework::EigenVector<T>::Flatten(*out);
+    out_e.device(ctx.GetEigenDevice<Place>()) =
+        x_v * e_mid.reshape(Eigen::DSizes<int, 1>(e_mid.size())).pow(-beta);
+  }
+};
+
+/**
+ * \brief Backward calculation for normalization with across maps.
+ *
+ * Function implementation:
+ *
+ * The implementation of this Function is derived from the
+ * CrossMapNormalFunc implementation.
+ *
+ * InputGrad = OutputGrad * denoms ^ (-beta)
+ *    -- upper
+ *  + > (OutputGrad * OutputValue * (-2 * alpha * beta) / MidOut) * InputValue
+ *    -- lower
+ *
+ * The data of inputs/outputs format is the same as the forward interface
+ * and is NCHW.
+ *
+ * The upper and lower is the same as forward. The logic of the sum
+ * is also the same as forward.
+ */
+template <typename Place, typename T>
+class LRNGradKernel : public framework::OpKernel<T> {
+ public:
+  using Tensor = framework::Tensor;
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    const Tensor* x = ctx.Input<Tensor>("X");
+    const Tensor* out = ctx.Input<Tensor>("Out");
+    const Tensor* out_g = ctx.Input<Tensor>(framework::GradVarName("Out"));
+    const Tensor* mid = ctx.Input<Tensor>("MidOut");
+
+    auto x_g = ctx.Output<Tensor>(framework::GradVarName("X"));
+    x_g->mutable_data<T>(ctx.GetPlace());
+
+    auto x_g_e = framework::EigenVector<T>::Flatten(*x_g);
+    x_g_e.device(ctx.GetEigenDevice<Place>()) = x_g_e.constant(0.0);
+
+    auto x_dims = x->dims();
+    int N = x_dims[0];
+    int C = x_dims[1];
+    int H = x_dims[2];
+    int W = x_dims[3];
+
+    int n = ctx.Attr<int>("n");
+    T alpha = ctx.Attr<T>("alpha");
+    T beta = ctx.Attr<T>("beta");
+    T ratio = -2 * alpha * beta;
+
+    auto e_x = framework::EigenTensor<T, 4>::From(*x);
+    auto e_x_g = framework::EigenTensor<T, 4>::From(*x_g);
+    auto e_out = framework::EigenTensor<T, 4>::From(*out);
+    auto e_out_g = framework::EigenTensor<T, 4>::From(*out_g);
+    auto e_mid = framework::EigenTensor<T, 4>::From(*mid);
+
+    const int start = -(n - 1) / 2;
+    const int end = start + n;
+    for (int m = 0; m < N; m++) {
+      for (int i = 0; i < C; i++) {
+        auto i_x = e_x.slice(Eigen::array<int, 4>({{m, i, 0, 0}}),
+                             Eigen::array<int, 4>({{1, 1, H, W}}));
+
+        auto i_x_g = e_x_g.slice(Eigen::array<int, 4>({{m, i, 0, 0}}),
+                                 Eigen::array<int, 4>({{1, 1, H, W}}));
+
+        auto i_out_g = e_out_g.slice(Eigen::array<int, 4>({{m, i, 0, 0}}),
+                                     Eigen::array<int, 4>({{1, 1, H, W}}));
+
+        auto i_mid = e_mid.slice(Eigen::array<int, 4>({{m, i, 0, 0}}),
+                                 Eigen::array<int, 4>({{1, 1, H, W}}));
+
+        i_x_g.device(ctx.GetEigenDevice<Place>()) = i_mid.pow(-beta) * i_out_g;
+        for (int c = start; c <= end; c++) {
+          int ch = i + c;
+          if (ch < 0 || ch >= C) {
+            continue;
+          }
+
+          auto c_out = e_out.slice(Eigen::array<int, 4>({{m, ch, 0, 0}}),
+                                   Eigen::array<int, 4>({{1, 1, H, W}}));
+
+          auto c_mid = e_mid.slice(Eigen::array<int, 4>({{m, ch, 0, 0}}),
+                                   Eigen::array<int, 4>({{1, 1, H, W}}));
+
+          auto c_out_g = e_out_g.slice(Eigen::array<int, 4>({{m, ch, 0, 0}}),
+                                       Eigen::array<int, 4>({{1, 1, H, W}}));
+
+          i_x_g.device(ctx.GetEigenDevice<Place>()) +=
+              ratio * c_out_g * c_out * i_x / c_mid;
+        }
+      }
+    }
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle

paddle/operators/math/CMakeLists.txt

@@ -9,6 +9,7 @@ if(WITH_GPU)
     nv_library(cross_entropy SRCS cross_entropy.cc cross_entropy.cu DEPS operator)
     nv_library(pooling SRCS pooling.cc pooling.cu DEPS device_context)
     nv_library(vol2col SRCS vol2col.cc vol2col.cu DEPS device_context)
+    nv_library(context_project SRCS context_project.cc context_project.cu DEPS device_context)
     nv_library(sequence2batch SRCS sequence2batch.cc sequence2batch.cu DEPS device_context)
     nv_library(lstm_compute SRCS lstm_compute.cc lstm_compute.cu DEPS device_context activation_functions)
 else()
@@ -18,6 +19,7 @@ else()
     cc_library(cross_entropy SRCS cross_entropy.cc DEPS operator)
     cc_library(pooling SRCS pooling.cc DEPS device_context)
     cc_library(vol2col SRCS vol2col.cc DEPS device_context)
+    cc_library(context_project SRCS context_project.cc DEPS device_context)
     cc_library(sequence2batch SRCS sequence2batch.cc DEPS device_context)
     cc_library(lstm_compute SRCS lstm_compute.cc DEPS device_context activation_functions)
 endif()

paddle/operators/math/context_project.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed 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 "paddle/operators/math/context_project.h"
+
+namespace paddle {
+namespace operators {
+namespace math {
+
+template class ContextProjectFunctor<platform::CPUPlace, float>;
+template class ContextProjectFunctor<platform::CPUPlace, double>;
+
+}  // namespace math
+}  // namespace operators
+}  // namespace paddle

paddle/framework/saver.proto → paddle/operators/math/context_project.cu

@@ -12,28 +12,17 @@ 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. */
 
-syntax = "proto2";
-option optimize_for = LITE_RUNTIME;
-package paddle.framework;
-
-import "framework.proto";
-
-/**
- * This file contains necessary information for model, checkpoint.
- * etc.
- */
-
-message LoDInfo { repeated int64 level = 1; }
-
-/**
- * Save the LoDTensorDesc information through LoDTensorProto, its data memory
- * is copyed to c buffer immediately. See model_format.md for details.
- */
-
-message LoDTensorProto {
-  optional DataType data_type = 1;
-  repeated int64 dims = 2; // [UNK, 640, 480] is saved as [-1, 640, 480]
-  repeated LoDInfo levels = 3;
-  optional int32 lod_level = 4 [ default = 0 ];
-  optional int32 version = 5;
-}
+#define EIGEN_USE_GPU
+
+#include "paddle/operators/math/context_project.h"
+
+namespace paddle {
+namespace operators {
+namespace math {
+
+template class ContextProjectFunctor<platform::GPUPlace, float>;
+template class ContextProjectFunctor<platform::GPUPlace, double>;
+
+}  // namespace math
+}  // namespace operators
+}  // namespace paddle

paddle/operators/math/context_project.h

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed 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. */
+
+#pragma once
+
+#include "paddle/framework/eigen.h"
+#include "paddle/framework/lod_tensor.h"
+#include "paddle/framework/tensor.h"
+#include "paddle/operators/math/im2col.h"
+
+namespace paddle {
+namespace operators {
+namespace math {
+
+template <typename T, int MajorType = Eigen::RowMajor,
+          typename IndexType = Eigen::DenseIndex>
+using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>;
+/*
+ * \brief Context projection concatenate features in adjacent time steps in
+ * a sequence. The i-th row of the output is the concatenation of
+ * context_length rows of the input. The context_length rows are the
+ * consecutive rows from the i+shift_start row.
+
+ * \param in            Input data.
+ * \param Shape         The shape of Input data,
+ *                      [minibatch, number_of_input_features].
+ * \param type          A float LoDTensor.
+ *
+ * \param padding_data  Padding data.
+ * \param Shape         The shape of Padding data,
+ *                      [up_pad + down_pad, number_of_input_features].
+ * \param type          A float Tensor.
+ *
+ * \param col           Col data.
+ * \param Shape         The shape of Col data,
+ *                      [minibatch, context_length * number_of_input_features].
+ * \param type           A float Tensor.
+ *
+ * For a mini-batch of 2 variable lengths sentences, containing 3, and 1
+ * time-steps:
+ *
+ * Assumed input (X) is a [4, M, N] float LoDTensor, and X->lod()[0] = [0, 3,
+ * 4].
+ * Besides, for the sake of simplicity, we assume M=1 and N=2.
+ *
+ * X = [[a1, a2;
+ *       b1, b2;
+ *       c1, c2]
+ *      [d1, d2]]
+ *
+ * This is to say that input (X) has 4 words and the dimension of each word
+ * representation is 2.
+ *
+ * - Case1:
+ * If context_start is -1 and padding_trainable is false, we use zero to pad
+ * instead of learned weight to pad,
+ * and the context_lenth is 3, the output (Out) is:
+ *
+ * Out =[[0,  0,  a1, a2, b1, b2;
+ *        a1, a2, b1, b2, c1, c2;
+ *        b1, b2, c1, c2, 0,  0 ]
+ *       [0,  0,  d1, d2, 0,  0 ]]
+ *
+ * - Case2:
+ * If context_start is -1 and padding_trainable is true, we use learned weight
+ * to pad,
+ * and the context_lenth is 3, the output (Out) is:
+ *
+ * Out = [[w1, w2, a1, a2, b1, b2;
+ *         a1, a2, b1, b2, c1, c2;
+ *         b1, b2, c1, c2, w3, w4]
+ *        [w1, w2, d1, d2, w3, w4]]
+ *
+ */
+
+template <typename Place, typename T>
+class ContextProjectFunctor {
+ public:
+  void operator()(const platform::DeviceContext& context,
+                  framework::LoDTensor& in, framework::Tensor& padding_data,
+                  framework::Tensor& col, bool padding_trainable,
+                  int context_start, int context_length, int context_stride,
+                  int up_pad, int down_pad, bool gradient, bool input_grad,
+                  bool pad_grad) {
+    auto lod_level_0 = in.lod()[0];
+
+    paddle::operators::math::Im2ColFunctor<
+        paddle::operators::math::ColFormat::kOCF, Place, float>
+        im2col_ocf;
+    paddle::operators::math::Col2ImFunctor<
+        paddle::operators::math::ColFormat::kOCF, Place, float>
+        col2im_ocf;
+
+    int input_row_begin, input_row_end;
+    int sequence_height, sequence_width;
+    sequence_width = in.dims()[1];
+    input_grad = gradient && input_grad;
+    pad_grad = gradient && pad_grad;
+
+    if (!gradient || input_grad) {
+      for (int i = 0; i < static_cast<int>(lod_level_0.size()) - 1; ++i) {
+        input_row_begin = (context_start > 0)
+                              ? static_cast<int>(lod_level_0[i]) + context_start
+                              : static_cast<int>(lod_level_0[i]);
+        input_row_end = static_cast<int>(lod_level_0[i + 1]);
+
+        framework::Tensor out_t =
+            col.Slice(static_cast<int>(lod_level_0[i]),
+                      static_cast<int>(lod_level_0[i + 1]));
+
+        sequence_height = static_cast<int>(out_t.dims()[0]);
+
+        if (input_row_begin < input_row_end) {
+          framework::Tensor in_t = in.Slice(input_row_begin, input_row_end);
+
+          std::vector<int64_t> output_shape(
+              {sequence_height, 1, 1, context_length,
+               sequence_width});  // output_height, output_width,
+          // input_channels, filter_height, filter_width
+
+          out_t.Resize(framework::make_ddim(output_shape));
+
+          std::vector<int64_t> input_shape(
+              {1, input_row_end - input_row_begin,
+               sequence_width});  // input_channels, input_height, input_width
+          in_t.Resize(framework::make_ddim(input_shape));
+
+          if (gradient) {
+            col2im_ocf(context, in_t, out_t,
+                       /*stride_height*/ context_stride, /*stride_width*/ 1,
+                       up_pad, down_pad, 0, 0);
+          } else {
+            im2col_ocf(context, in_t, out_t,
+                       /*stride_height*/ context_stride, /*stride_width*/ 1,
+                       up_pad, down_pad, 0, 0);
+          }
+          out_t.Resize({sequence_height, context_length * sequence_width});
+        }
+      }
+    }
+    if (!gradient || pad_grad) {
+      if (padding_trainable) {
+        for (int i = 0; i < static_cast<int>(lod_level_0.size()) - 1; ++i) {
+          framework::Tensor out_t =
+              col.Slice(static_cast<int>(lod_level_0[i]),
+                        static_cast<int>(lod_level_0[i + 1]));
+
+          sequence_height = static_cast<int>(out_t.dims()[0]);
+
+          // add up trainable data
+          out_t.Resize({sequence_height * context_length, sequence_width});
+
+          if (up_pad > 0) {  // add up pad
+            int padding_rows = std::min(
+                up_pad, static_cast<int>(lod_level_0[i + 1] - lod_level_0[i]));
+
+            for (int k = 0; k < padding_rows; ++k) {
+              int padding_size =
+                  k + context_length < up_pad ? context_length : up_pad - k;
+              framework::Tensor out_t_sub = out_t.Slice(
+                  k * context_length, k * context_length + padding_size);
+              framework::Tensor w_sub = padding_data.Slice(k, k + padding_size);
+              // in this block, using EigenVector<T>::Flatten is ok too.
+              auto out_t_sub_e = EigenMatrix<T>::From(out_t_sub);
+              auto w_sub_e = EigenMatrix<T>::From(w_sub);
+              if (gradient) {
+                w_sub_e.device(*context.GetEigenDevice<Place>()) =
+                    w_sub_e + out_t_sub_e;
+              } else {
+                out_t_sub_e.device(*context.GetEigenDevice<Place>()) = w_sub_e;
+              }
+            }
+          }
+          if (down_pad > 0) {  // add down pad
+            int down_pad_begin_row =
+                std::max(
+                    0, (sequence_height - context_start - context_length) + 1) +
+                1;
+            int padding_begin = std::max(0, context_start - sequence_height);
+            int padding_size =
+                sequence_height - context_start >= context_length
+                    ? 1
+                    : context_length - (sequence_height - context_start);
+            if (context_start >= sequence_height) padding_size = context_length;
+            int padding_idx = padding_begin;
+            for (int t = 0; t + down_pad_begin_row <= sequence_height;
+                 ++t, ++padding_size) {
+              if (context_start >= sequence_height)
+                padding_size = context_length;
+              if (padding_size > context_length) {
+                padding_size = context_length;
+                padding_idx++;
+              }
+              if (padding_begin > 0 || sequence_height == context_start)
+                padding_idx = padding_begin + t;
+              framework::Tensor out_t_sub = out_t.Slice(
+                  (down_pad_begin_row + t) * context_length - padding_size,
+                  (down_pad_begin_row + t) * context_length);
+              framework::Tensor w_sub = padding_data.Slice(
+                  up_pad + padding_idx, up_pad + padding_idx + padding_size);
+              auto out_t_sub_e = EigenMatrix<T>::From(out_t_sub);
+              auto w_sub_e = EigenMatrix<T>::From(w_sub);
+              if (gradient) {
+                w_sub_e.device(*context.GetEigenDevice<Place>()) =
+                    w_sub_e + out_t_sub_e;
+              } else {
+                out_t_sub_e.device(*context.GetEigenDevice<Place>()) = w_sub_e;
+              }
+            }
+          }
+          out_t.Resize({sequence_height, context_length * sequence_width});
+        }
+      }
+    }
+  }
+};
+
+}  // namespace math
+}  // namespace operators
+}  // namespace paddle

paddle/operators/math/cross_entropy.cc

@@ -54,6 +54,7 @@ class CrossEntropyFunctor<platform::CPUPlace, T> {
 };
 
 template class CrossEntropyFunctor<platform::CPUPlace, float>;
+template class CrossEntropyFunctor<platform::CPUPlace, double>;
 }  // namespace math
 }  // namespace operators
 }  // namespace paddle

paddle/operators/math/cross_entropy.cu

@@ -39,11 +39,36 @@ __device__ __forceinline__ T sum_single_warp(T val) {
   return val;
 }
 
+// CUDA do not support dynamic arrary in template
+// https://stackoverflow.com/questions/20497209
+template <typename T>
+struct SharedMemory {
+  // Ensure that we won't compile any un-specialized types
+  __device__ T* GetPointer() { return NULL; }
+};
+
+template <>
+struct SharedMemory<float> {
+  __device__ float* GetPointer() {
+    extern __shared__ float s_float[];
+    return s_float;
+  }
+};
+
+template <>
+struct SharedMemory<double> {
+  __device__ double* GetPointer() {
+    extern __shared__ double s_double[];
+    return s_double;
+  }
+};
+
 template <typename T>
 __global__ void SoftCrossEntropyKernel(T* Y, const T* X, const T* label,
                                        const int class_num) {
   int tid = threadIdx.x;
-  extern __shared__ T d_sum[];
+  SharedMemory<T> d_sum_shared;
+  T* d_sum = d_sum_shared.GetPointer();
   d_sum[tid] = 0;
 
   int cur_idx = tid;
@@ -102,6 +127,7 @@ class CrossEntropyFunctor<platform::GPUPlace, T> {
 };
 
 template class CrossEntropyFunctor<platform::GPUPlace, float>;
+template class CrossEntropyFunctor<platform::GPUPlace, double>;
 }  // namespace math
 }  // namespace operators
 }  // namespace paddle

paddle/operators/math/selected_rows_functor.cc

@@ -68,6 +68,7 @@ struct SelectedRowsAdd<platform::CPUPlace, T> {
 };
 
 template struct SelectedRowsAdd<platform::CPUPlace, float>;
+template struct SelectedRowsAdd<platform::CPUPlace, double>;
 
 template <typename T>
 struct SelectedRowsAddTensor<platform::CPUPlace, T> {
@@ -108,6 +109,72 @@ struct SelectedRowsAddTensor<platform::CPUPlace, T> {
 };
 
 template struct SelectedRowsAddTensor<platform::CPUPlace, float>;
+template struct SelectedRowsAddTensor<platform::CPUPlace, double>;
+
+template <typename T>
+struct SelectedRowsAddTo<platform::CPUPlace, T> {
+  void operator()(const platform::DeviceContext& context,
+                  const framework::SelectedRows& input1,
+                  const int64_t input2_offset,
+                  framework::SelectedRows* input2) {
+    auto in1_height = input1.height();
+    PADDLE_ENFORCE_EQ(in1_height, input2->height());
+
+    auto& in1_rows = input1.rows();
+    auto& in2_rows = *(input2->mutable_rows());
+
+    auto& in1_value = input1.value();
+    auto* in2_value = input2->mutable_value();
+
+    // concat rows
+    in2_rows.insert(in2_rows.end(), in1_rows.begin(), in1_rows.end());
+
+    auto in1_place = input1.place();
+    PADDLE_ENFORCE(platform::is_cpu_place(in1_place));
+    auto in2_place = input2->place();
+    PADDLE_ENFORCE(platform::is_cpu_place(in2_place));
+
+    auto* in1_data = in1_value.data<T>();
+    auto* in2_data = in2_value->data<T>();
+    memory::Copy(boost::get<platform::CPUPlace>(in2_place),
+                 in2_data + input2_offset,
+                 boost::get<platform::CPUPlace>(in1_place), in1_data,
+                 in1_value.numel() * sizeof(T));
+  }
+};
+
+template struct SelectedRowsAddTo<platform::CPUPlace, float>;
+template struct SelectedRowsAddTo<platform::CPUPlace, double>;
+
+template <typename T>
+struct SelectedRowsAddToTensor<platform::CPUPlace, T> {
+  void operator()(const platform::DeviceContext& context,
+                  const framework::SelectedRows& input1,
+                  framework::Tensor* input2) {
+    auto in1_height = input1.height();
+    auto in2_dims = input2->dims();
+    PADDLE_ENFORCE_EQ(in1_height, in2_dims[0]);
+
+    auto& in1_value = input1.value();
+    auto& in1_rows = input1.rows();
+
+    int64_t in1_row_numel = in1_value.numel() / in1_rows.size();
+    PADDLE_ENFORCE_EQ(in1_row_numel, input2->numel() / in1_height);
+
+    auto* in1_data = in1_value.data<T>();
+    auto* input2_data = input2->data<T>();
+
+    for (size_t i = 0; i < in1_rows.size(); i++) {
+      for (int64_t j = 0; j < in1_row_numel; j++) {
+        input2_data[in1_rows[i] * in1_row_numel + j] +=
+            in1_data[i * in1_row_numel + j];
+      }
+    }
+  }
+};
+
+template struct SelectedRowsAddToTensor<platform::CPUPlace, float>;
+template struct SelectedRowsAddToTensor<platform::CPUPlace, double>;
 
 }  // namespace math
 }  // namespace operators

paddle/operators/math/selected_rows_functor.cu

@@ -73,12 +73,13 @@ struct SelectedRowsAdd<platform::GPUPlace, T> {
 };
 
 template struct SelectedRowsAdd<platform::GPUPlace, float>;
+template struct SelectedRowsAdd<platform::GPUPlace, double>;
 
 namespace {
-template <typename T>
+template <typename T, int block_size>
 __global__ void SelectedRowsAddTensorKernel(const T* selected_rows,
                                             const int64_t* rows, T* tensor_out,
-                                            int64_t row_numel, int block_size) {
+                                            int64_t row_numel) {
   const int ty = blockIdx.y;
   int tid = threadIdx.x;
 
@@ -119,14 +120,13 @@ struct SelectedRowsAddTensor<platform::GPUPlace, T> {
     SetConstant<platform::GPUPlace, T> functor;
     functor(context, output, 0.0);
 
-    int block_size = 256;
+    const int block_size = 256;
     dim3 threads(block_size, 1);
     dim3 grid(1, in1_rows.size());
-    SelectedRowsAddTensorKernel<
-        T><<<grid, threads, 0,
-             reinterpret_cast<const platform::CUDADeviceContext&>(context)
-                 .stream()>>>(in1_data, in1_rows.data(), out_data,
-                              in1_row_numel, block_size);
+    SelectedRowsAddTensorKernel<T, block_size><<<
+        grid, threads, 0,
+        reinterpret_cast<const platform::CUDADeviceContext&>(context)
+            .stream()>>>(in1_data, in1_rows.data(), out_data, in1_row_numel);
 
     auto out_eigen = framework::EigenVector<T>::Flatten(*output);
     auto in2_eigen = framework::EigenVector<T>::Flatten(input2);
@@ -136,6 +136,93 @@ struct SelectedRowsAddTensor<platform::GPUPlace, T> {
 };
 
 template struct SelectedRowsAddTensor<platform::GPUPlace, float>;
+template struct SelectedRowsAddTensor<platform::GPUPlace, double>;
+
+template <typename T>
+struct SelectedRowsAddTo<platform::GPUPlace, T> {
+  void operator()(const platform::DeviceContext& context,
+                  const framework::SelectedRows& input1,
+                  const int64_t input2_offset,
+                  framework::SelectedRows* input2) {
+    auto in1_height = input1.height();
+    PADDLE_ENFORCE_EQ(in1_height, input2->height());
+
+    auto& in1_rows = input1.rows();
+    auto& in2_rows = *(input2->mutable_rows());
+
+    auto& in1_value = input1.value();
+    auto* in2_value = input2->mutable_value();
+
+    // concat rows
+    in2_rows.insert(in2_rows.end(), in1_rows.begin(), in1_rows.end());
+
+    auto in1_place = input1.place();
+    PADDLE_ENFORCE(platform::is_gpu_place(in1_place));
+    auto in2_place = input2->place();
+    PADDLE_ENFORCE(platform::is_gpu_place(in2_place));
+
+    auto* in1_data = in1_value.data<T>();
+    auto* in2_data = in2_value->data<T>();
+    memory::Copy(
+        boost::get<platform::GPUPlace>(in2_place), in2_data + input2_offset,
+        boost::get<platform::GPUPlace>(in1_place), in1_data,
+        in1_value.numel() * sizeof(T),
+        reinterpret_cast<const platform::CUDADeviceContext&>(context).stream());
+  }
+};
+
+template struct SelectedRowsAddTo<platform::GPUPlace, float>;
+template struct SelectedRowsAddTo<platform::GPUPlace, double>;
+
+namespace {
+template <typename T, int block_size>
+__global__ void SelectedRowsAddToTensorKernel(const T* selected_rows,
+                                              const int64_t* rows,
+                                              T* tensor_out,
+                                              int64_t row_numel) {
+  const int ty = blockIdx.y;
+  int tid = threadIdx.x;
+
+  selected_rows += ty * row_numel;
+  tensor_out += rows[ty] * row_numel;
+
+  for (int index = tid; index < row_numel; index += block_size) {
+    // Since index in rows of SelectedRows can be duplicate, we have to use
+    // Atomic Operation to avoid concurrent write error.
+    paddle::platform::CudaAtomicAdd(tensor_out + index, selected_rows[index]);
+  }
+}
+}  // namespace
+
+template <typename T>
+struct SelectedRowsAddToTensor<platform::GPUPlace, T> {
+  void operator()(const platform::DeviceContext& context,
+                  const framework::SelectedRows& input1,
+                  framework::Tensor* input2) {
+    auto in1_height = input1.height();
+    auto in2_dims = input2->dims();
+    PADDLE_ENFORCE_EQ(in1_height, in2_dims[0]);
+
+    auto& in1_value = input1.value();
+    auto& in1_rows = input1.rows();
+
+    int64_t in1_row_numel = in1_value.numel() / in1_rows.size();
+    PADDLE_ENFORCE_EQ(in1_row_numel, input2->numel() / in1_height);
+
+    auto* in1_data = in1_value.data<T>();
+    auto* in2_data = input2->data<T>();
+    const int block_size = 256;
+    dim3 threads(block_size, 1);
+    dim3 grid(1, in1_rows.size());
+    SelectedRowsAddToTensorKernel<T, block_size><<<
+        grid, threads, 0,
+        reinterpret_cast<const platform::CUDADeviceContext&>(context)
+            .stream()>>>(in1_data, in1_rows.data(), in2_data, in1_row_numel);
+  }
+};
+
+template struct SelectedRowsAddToTensor<platform::GPUPlace, float>;
+template struct SelectedRowsAddToTensor<platform::GPUPlace, double>;
 
 }  // namespace math
 }  // namespace operators

paddle/operators/math/selected_rows_functor.h

@@ -36,6 +36,22 @@ struct SelectedRowsAddTensor {
                   const framework::Tensor& input2, framework::Tensor* output);
 };
 
+// input2 = input1 + input2
+template <typename Place, typename T>
+struct SelectedRowsAddTo {
+  void operator()(const platform::DeviceContext& context,
+                  const framework::SelectedRows& input1,
+                  const int64_t input2_offset, framework::SelectedRows* input2);
+};
+
+// input2 = input1 + input2
+template <typename Place, typename T>
+struct SelectedRowsAddToTensor {
+  void operator()(const platform::DeviceContext& context,
+                  const framework::SelectedRows& input1,
+                  framework::Tensor* input2);
+};
+
 }  // namespace math
 }  // namespace operators
 }  // namespace paddle

paddle/operators/math/selected_rows_functor_test.cc

@@ -104,3 +104,91 @@ TEST(selected_rows_functor, cpu_add) {
   // row9: 2.0 + 3.0
   EXPECT_EQ(tensor2_data[9 * row_numel + 6], 5.0);
 }
+
+TEST(selected_rows_functor, cpu_add_to) {
+  using namespace paddle::framework;
+  using namespace paddle::platform;
+  using namespace paddle::operators::math;
+
+  CPUPlace cpu_place;
+  CPUDeviceContext ctx(cpu_place);
+  SetConstant<CPUPlace, float> functor;
+  int64_t height = 10;
+  int64_t row_numel = 10;
+
+  std::vector<int64_t> rows1{0, 4, 7};
+  std::unique_ptr<SelectedRows> selected_rows1{new SelectedRows(rows1, height)};
+  auto* in1_value = selected_rows1->mutable_value();
+  in1_value->mutable_data<float>(
+      make_ddim({static_cast<int64_t>(rows1.size()), row_numel}), cpu_place);
+  functor(ctx, in1_value, 1.0);
+
+  std::vector<int64_t> rows2{0, 5, 7, 9};
+  std::unique_ptr<SelectedRows> selected_rows2{new SelectedRows(rows2, height)};
+  auto* in2_value = selected_rows2->mutable_value();
+  in2_value->mutable_data<float>(
+      make_ddim({static_cast<int64_t>(rows2.size()), row_numel}), cpu_place);
+  functor(ctx, in2_value, 2.0);
+
+  std::unique_ptr<SelectedRows> output{new SelectedRows()};
+  output->set_height(height);
+  auto* out_value = output->mutable_value();
+
+  // simplely concat two SelectedRows
+  out_value->mutable_data<float>(make_ddim({7, 10}), cpu_place);
+
+  SelectedRowsAddTo<CPUPlace, float> add_to_functor;
+  add_to_functor(ctx, *selected_rows1, 0, output.get());
+  add_to_functor(ctx, *selected_rows2, in1_value->numel(), output.get());
+
+  auto out_height = output->height();
+  EXPECT_EQ(out_height, height);
+
+  auto& out_rows = output->rows();
+
+  // input1 rows
+  EXPECT_EQ(out_rows[0], 0);
+  EXPECT_EQ(out_rows[1], 4);
+  EXPECT_EQ(out_rows[2], 7);
+  // input2 rows
+  EXPECT_EQ(out_rows[3], 0);
+  EXPECT_EQ(out_rows[4], 5);
+  EXPECT_EQ(out_rows[5], 7);
+  EXPECT_EQ(out_rows[6], 9);
+
+  auto* out_data = output->value().data<float>();
+  // input1 value
+  EXPECT_EQ(out_data[0 * row_numel + 0], 1.0);
+  EXPECT_EQ(out_data[0 * row_numel + 8], 1.0);
+  EXPECT_EQ(out_data[1 * row_numel + 1], 1.0);
+  EXPECT_EQ(out_data[2 * row_numel + 6], 1.0);
+  // input2 value
+  EXPECT_EQ(out_data[3 * row_numel + 3], 2.0);
+  EXPECT_EQ(out_data[3 * row_numel + 8], 2.0);
+  EXPECT_EQ(out_data[4 * row_numel + 4], 2.0);
+  EXPECT_EQ(out_data[5 * row_numel + 7], 2.0);
+  EXPECT_EQ(out_data[6 * row_numel + 9], 2.0);
+
+  std::unique_ptr<Tensor> tensor1{new Tensor()};
+  tensor1->mutable_data<float>(make_ddim({height, row_numel}), cpu_place);
+  functor(ctx, tensor1.get(), 3.0);
+
+  SelectedRowsAddToTensor<CPUPlace, float> add_to_tensor_functor;
+  add_to_tensor_functor(ctx, *output, tensor1.get());
+
+  auto* tensor1_data = tensor1->data<float>();
+  // row0: 1.0 + 2.0 + 3.0
+  EXPECT_EQ(tensor1_data[0 * row_numel + 0], 6.0);
+  // row1: 3.0
+  EXPECT_EQ(tensor1_data[1 * row_numel + 1], 3.0);
+  // row4 : 1.0 + 3.0
+  EXPECT_EQ(tensor1_data[4 * row_numel + 6], 4.0);
+  // row5: 2.0 + 3.0
+  EXPECT_EQ(tensor1_data[5 * row_numel + 7], 5.0);
+  // row6: 3.0
+  EXPECT_EQ(tensor1_data[6 * row_numel + 1], 3.0);
+  // row7: 1.0 + 2.0 + 3.0
+  EXPECT_EQ(tensor1_data[7 * row_numel + 3], 6.0);
+  // row9: 2.0 + 3.0
+  EXPECT_EQ(tensor1_data[9 * row_numel + 6], 5.0);
+}

paddle/operators/math/selected_rows_functor_test.cu

@@ -113,3 +113,100 @@ TEST(selected_rows_functor, gpu_add) {
   // row9: 2.0 + 3.0
   EXPECT_EQ(tensor2_cpu_data[9 * row_numel + 6], 5.0);
 }
+
+TEST(selected_rows_functor, gpu_add_to) {
+  using namespace paddle::framework;
+  using namespace paddle::platform;
+  using namespace paddle::operators::math;
+
+  GPUPlace gpu_place(0);
+  CPUPlace cpu_place;
+  CUDADeviceContext ctx(gpu_place);
+  SetConstant<GPUPlace, float> functor;
+  int64_t height = 10;
+  int64_t row_numel = 10;
+
+  std::vector<int64_t> rows1{0, 4, 7};
+  std::unique_ptr<SelectedRows> selected_rows1{new SelectedRows(rows1, height)};
+  auto* in1_value = selected_rows1->mutable_value();
+  in1_value->mutable_data<float>(
+      make_ddim({static_cast<int64_t>(rows1.size()), row_numel}), gpu_place);
+  functor(ctx, in1_value, 1.0);
+
+  std::vector<int64_t> rows2{0, 5, 7, 9};
+  std::unique_ptr<SelectedRows> selected_rows2{new SelectedRows(rows2, height)};
+  auto* in2_value = selected_rows2->mutable_value();
+  in2_value->mutable_data<float>(
+      make_ddim({static_cast<int64_t>(rows2.size()), row_numel}), gpu_place);
+  functor(ctx, in2_value, 2.0);
+
+  std::unique_ptr<SelectedRows> output{new SelectedRows()};
+  output->set_height(height);
+  auto* out_value = output->mutable_value();
+
+  // simplely concat two SelectedRows
+  out_value->mutable_data<float>(make_ddim({7, 10}), gpu_place);
+
+  SelectedRowsAddTo<GPUPlace, float> add_to_functor;
+  add_to_functor(ctx, *selected_rows1, 0, output.get());
+  add_to_functor(ctx, *selected_rows2, in1_value->numel(), output.get());
+
+  auto out_height = output->height();
+  EXPECT_EQ(out_height, height);
+
+  auto& out_rows = output->rows();
+
+  // input1 rows
+  EXPECT_EQ(out_rows[0], 0);
+  EXPECT_EQ(out_rows[1], 4);
+  EXPECT_EQ(out_rows[2], 7);
+  // input2 rows
+  EXPECT_EQ(out_rows[3], 0);
+  EXPECT_EQ(out_rows[4], 5);
+  EXPECT_EQ(out_rows[5], 7);
+  EXPECT_EQ(out_rows[6], 9);
+
+  Tensor out_cpu;
+  out_cpu.CopyFrom(*out_value, cpu_place, ctx);
+  ctx.Wait();
+
+  auto* out_cpu_data = out_cpu.data<float>();
+  // input1 value
+  EXPECT_EQ(out_cpu_data[0 * row_numel + 0], 1.0);
+  EXPECT_EQ(out_cpu_data[0 * row_numel + 8], 1.0);
+  EXPECT_EQ(out_cpu_data[1 * row_numel + 1], 1.0);
+  EXPECT_EQ(out_cpu_data[2 * row_numel + 6], 1.0);
+  // input2 value
+  EXPECT_EQ(out_cpu_data[3 * row_numel + 3], 2.0);
+  EXPECT_EQ(out_cpu_data[3 * row_numel + 8], 2.0);
+  EXPECT_EQ(out_cpu_data[4 * row_numel + 4], 2.0);
+  EXPECT_EQ(out_cpu_data[5 * row_numel + 7], 2.0);
+  EXPECT_EQ(out_cpu_data[6 * row_numel + 9], 2.0);
+
+  std::unique_ptr<Tensor> tensor1{new Tensor()};
+  tensor1->mutable_data<float>(make_ddim({height, row_numel}), gpu_place);
+  functor(ctx, tensor1.get(), 3.0);
+
+  SelectedRowsAddToTensor<GPUPlace, float> add_to_tensor_functor;
+  add_to_tensor_functor(ctx, *output, tensor1.get());
+
+  Tensor tensor1_cpu;
+  tensor1_cpu.CopyFrom(*tensor1, cpu_place, ctx);
+  ctx.Wait();
+
+  auto* tensor1_cpu_data = tensor1_cpu.data<float>();
+  // row0: 1.0 + 2.0 + 3.0
+  EXPECT_EQ(tensor1_cpu_data[0 * row_numel + 0], 6.0);
+  // row1: 3.0
+  EXPECT_EQ(tensor1_cpu_data[1 * row_numel + 1], 3.0);
+  // row4 : 1.0 + 3.0
+  EXPECT_EQ(tensor1_cpu_data[4 * row_numel + 6], 4.0);
+  // row5: 2.0 + 3.0
+  EXPECT_EQ(tensor1_cpu_data[5 * row_numel + 7], 5.0);
+  // row6: 3.0
+  EXPECT_EQ(tensor1_cpu_data[6 * row_numel + 1], 3.0);
+  // row7: 1.0 + 2.0 + 3.0
+  EXPECT_EQ(tensor1_cpu_data[7 * row_numel + 3], 6.0);
+  // row9: 2.0 + 3.0
+  EXPECT_EQ(tensor1_cpu_data[9 * row_numel + 6], 5.0);
+}

paddle/operators/mean_op.cc

@@ -71,7 +71,8 @@ class MeanGradMaker : public framework::SingleGradOpDescMaker {
 namespace ops = paddle::operators;
 REGISTER_OPERATOR(mean, ops::MeanOp, ops::MeanOpMaker, ops::MeanGradMaker);
 REGISTER_OPERATOR(mean_grad, ops::MeanGradOp);
-REGISTER_OP_CPU_KERNEL(mean,
-                       ops::MeanKernel<paddle::platform::CPUPlace, float>);
+REGISTER_OP_CPU_KERNEL(mean, ops::MeanKernel<paddle::platform::CPUPlace, float>,
+                       ops::MeanKernel<paddle::platform::CPUPlace, double>);
 REGISTER_OP_CPU_KERNEL(mean_grad,
-                       ops::MeanGradKernel<paddle::platform::CPUPlace, float>);
+                       ops::MeanGradKernel<paddle::platform::CPUPlace, float>,
+                       ops::MeanGradKernel<paddle::platform::CPUPlace, double>);

paddle/operators/mean_op.cu

@@ -17,7 +17,8 @@
 #include "paddle/operators/mean_op.h"
 
 namespace ops = paddle::operators;
-REGISTER_OP_GPU_KERNEL(mean,
-                       ops::MeanKernel<paddle::platform::GPUPlace, float>);
+REGISTER_OP_GPU_KERNEL(mean, ops::MeanKernel<paddle::platform::GPUPlace, float>,
+                       ops::MeanKernel<paddle::platform::GPUPlace, double>);
 REGISTER_OP_GPU_KERNEL(mean_grad,
-                       ops::MeanGradKernel<paddle::platform::GPUPlace, float>);
+                       ops::MeanGradKernel<paddle::platform::GPUPlace, float>,
+                       ops::MeanGradKernel<paddle::platform::GPUPlace, double>);

paddle/operators/mul_op.cc

@@ -19,11 +19,9 @@ namespace operators {
 
 using framework::Tensor;
 
-class MulOp : public framework::OperatorWithKernel {
+class MulOpShapeInference : public framework::InferShapeBase {
  public:
-  using framework::OperatorWithKernel::OperatorWithKernel;
-
-  void InferShape(framework::InferShapeContext* ctx) const override {
+  void operator()(framework::InferShapeContext* ctx) const override {
     PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of MulOp should not be null.");
     PADDLE_ENFORCE(ctx->HasInput("Y"), "Input(Y) of MulOp should not be null.");
     PADDLE_ENFORCE(ctx->HasOutput("Out"),
@@ -137,7 +135,10 @@ class MulOpGrad : public framework::OperatorWithKernel {
 }  // namespace paddle
 
 namespace ops = paddle::operators;
-REGISTER_OP(mul, ops::MulOp, ops::MulOpMaker, mul_grad, ops::MulOpGrad);
+REGISTER_OPERATOR(mul, paddle::framework::OperatorWithKernel, ops::MulOpMaker,
+                  ops::MulOpShapeInference,
+                  paddle::framework::DefaultGradOpDescMaker<true>);
+REGISTER_OPERATOR(mul_grad, ops::MulOpGrad);
 REGISTER_OP_CPU_KERNEL(mul, ops::MulKernel<paddle::platform::CPUPlace, float>);
 REGISTER_OP_CPU_KERNEL(mul_grad,
                        ops::MulGradKernel<paddle::platform::CPUPlace, float>);

paddle/operators/proximal_adagrad_op.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed 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 "paddle/operators/proximal_adagrad_op.h"
+
+namespace paddle {
+namespace operators {
+
+class ProximalAdagradOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(framework::InferShapeContext *ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("Param"),
+                   "Input(Param) of ProximalAdagradOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasInput("Moment"),
+                   "Input(Moment) of ProximalAdagradOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasInput("Grad"),
+                   "Input(Grad) of ProximalAdagradOp should not be null.");
+    PADDLE_ENFORCE(
+        ctx->HasInput("LearningRate"),
+        "Input(LearningRate) of ProximalAdagradOp should not be null.");
+
+    PADDLE_ENFORCE(ctx->HasOutput("ParamOut"),
+                   "Output(ParamOut) of ProximalAdagradOp should not be null.");
+    PADDLE_ENFORCE(
+        ctx->HasOutput("MomentOut"),
+        "Output(MomentOut) of ProximalAdagradOp should not be null.");
+
+    auto param_dim = ctx->GetInputDim("Param");
+    PADDLE_ENFORCE_EQ(
+        param_dim, ctx->GetInputDim("Grad"),
+        "Param and Grad of ProximalAdagrad Op must have same dimension.");
+
+    PADDLE_ENFORCE_EQ(
+        param_dim, ctx->GetInputDim("Moment"),
+        "Param and Moment of ProximalAdagrad Op must have same dimension.");
+
+    auto lr_dim = ctx->GetInputDim("LearningRate");
+    PADDLE_ENFORCE_EQ(framework::product(lr_dim), 1,
+                      "Learning Rate should be a scalar.");
+
+    ctx->SetOutputDim("ParamOut", param_dim);
+    ctx->SetOutputDim("MomentOut", param_dim);
+  }
+};
+
+class ProximalAdagradOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  ProximalAdagradOpMaker(framework::OpProto *proto,
+                         framework::OpAttrChecker *op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput("Param",
+             "(Tensor, default Tensor<float>) "
+             "Input parameter that has to be updated.");
+    AddInput("Moment",
+             "(Tensor, default Tensor<float>) "
+             "Moment parameter that has to be updated.");
+    AddInput("Grad",
+             "(Tensor, default Tensor<float>) "
+             "Input gradient of the parameter.");
+    AddInput("LearningRate",
+             "(Tensor, default Tensor<float>) "
+             "The learning rate should be a tensor of size 1.");
+
+    AddOutput("ParamOut", "(Tensor) Output updated parameter value.");
+    AddOutput("MomentOut", "(Tensor) Output updated moment value.");
+
+    AddAttr<float>("l1",
+                   "(float, default 0.0) "
+                   "L1 regularization strength.")
+        .SetDefault(0.0f);
+    AddAttr<float>("l2",
+                   "(float, default 0.0)"
+                   "L2 regularization strength.")
+        .SetDefault(0.0f);
+    AddComment(R"DOC(
+
+Optimizer that implements the proximal adagrad algorithm.
+
+moment = moment + grad * grad
+prox_param = param - learning_rate * grad * (1 / sqrt(moment))
+param = sign(prox_param) / (1 + learning_rate * l2) *
+        max { |prox_param| - learning_rate * l1 , 0 }
+
+The paper that proposed Proximal GD: 
+(http://papers.nips.cc/paper/3793-efficient-learning-using-forward-backward-splitting.pdf)
+Here, we use the adagrad learning rate as specified here: 
+(http://www.jmlr.org/papers/volume12/duchi11a/duchi11a.pdf)
+)DOC");
+  }
+};
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_WITHOUT_GRADIENT(proximal_adagrad, ops::ProximalAdagradOp,
+                             ops::ProximalAdagradOpMaker);
+REGISTER_OP_CPU_KERNEL(
+    proximal_adagrad,
+    ops::ProximalAdagradOpKernel<paddle::platform::CPUPlace, float>);

paddle/operators/proximal_adagrad_op.cu

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed 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. */
+
+#define EIGEN_USE_GPU
+#include "paddle/operators/proximal_adagrad_op.h"
+
+namespace ops = paddle::operators;
+REGISTER_OP_GPU_KERNEL(
+    proximal_adagrad,
+    ops::ProximalAdagradOpKernel<paddle::platform::GPUPlace, float>);

paddle/operators/proximal_adagrad_op.h

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed 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. */
+
+#pragma once
+#include "paddle/framework/eigen.h"
+#include "paddle/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+template <typename T, int MajorType = Eigen::RowMajor,
+          typename IndexType = Eigen::DenseIndex>
+using EigenVector = framework::EigenVector<T, MajorType, IndexType>;
+
+template <typename Place, typename T>
+class ProximalAdagradOpKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    auto* param_out = ctx.Output<Tensor>("ParamOut");
+    auto* moment_out = ctx.Output<Tensor>("MomentOut");
+
+    param_out->mutable_data<T>(ctx.GetPlace());
+    moment_out->mutable_data<T>(ctx.GetPlace());
+
+    auto l1 = static_cast<T>(ctx.Attr<float>("l1"));
+    auto l2 = static_cast<T>(ctx.Attr<float>("l2"));
+
+    auto grad = ctx.Input<Tensor>("Grad");
+    auto p = EigenVector<T>::Flatten(*ctx.Input<Tensor>("Param"));
+    auto m = EigenVector<T>::Flatten(*ctx.Input<Tensor>("Moment"));
+    auto g = EigenVector<T>::Flatten(*grad);
+    auto lr = EigenVector<T>::Flatten(*ctx.Input<Tensor>("LearningRate"));
+
+    auto p_out = EigenVector<T>::Flatten(*param_out);
+    auto m_out = EigenVector<T>::Flatten(*moment_out);
+    auto place = ctx.GetEigenDevice<Place>();
+
+    Eigen::DSizes<int, 1> grad_dsize(grad->numel());
+
+    m_out.device(place) = m + g * g;
+    auto prox_param = p - lr.broadcast(grad_dsize) * g / m_out.sqrt();
+    if (l1 > static_cast<T>(0)) {
+      p_out.device(place) =
+          prox_param.sign() *
+          (((prox_param.abs() - (lr * l1).broadcast(grad_dsize))
+                .cwiseMax(static_cast<T>(0.0))) /
+           (static_cast<T>(1.0) + (lr * l2).broadcast(grad_dsize)));
+    } else {
+      p_out.device(place) =
+          prox_param / (static_cast<T>(1.0) + (lr * l2).broadcast(grad_dsize));
+    }
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle

paddle/operators/save_load_op_test.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserved.
+
+   Licensed 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 "gtest/gtest.h"
+#include "paddle/framework/op_registry.h"
+
+USE_NO_KERNEL_OP(save);
+USE_NO_KERNEL_OP(load);
+
+TEST(SaveLoadOp, CPU) {
+  paddle::framework::Scope scope;
+  paddle::platform::CPUPlace place;
+  paddle::platform::CPUDeviceContext ctx(place);
+  auto var = scope.Var("test_var");
+  auto tensor = var->GetMutable<paddle::framework::LoDTensor>();
+  tensor->Resize({10, 10});
+  paddle::framework::LoD expect_lod;
+  expect_lod.resize(1);
+  expect_lod[0].push_back(0);
+  expect_lod[0].push_back(1);
+  expect_lod[0].push_back(2);
+  expect_lod[0].push_back(3);
+
+  tensor->set_lod(expect_lod);
+  int* expect = tensor->mutable_data<int>(place);
+  for (size_t i = 0; i < paddle::framework::product(tensor->dims()); ++i) {
+    expect[i] = static_cast<int>(i);
+  }
+  paddle::framework::AttributeMap attrs;
+  attrs.insert({"file_path", std::string("tensor.save")});
+
+  auto save_op = paddle::framework::OpRegistry::CreateOp(
+      "save", {{"X", {"test_var"}}}, {}, attrs);
+  save_op->Run(scope, ctx);
+
+  auto load_var = scope.Var("out_var");
+  auto target = load_var->GetMutable<paddle::framework::LoDTensor>();
+  auto load_op = paddle::framework::OpRegistry::CreateOp(
+      "load", {}, {{"Out", {"out_var"}}}, attrs);
+  load_op->Run(scope, ctx);
+  int* actual = target->data<int>();
+  for (size_t i = 0; i < paddle::framework::product(tensor->dims()); ++i) {
+    EXPECT_EQ(expect[i], actual[i]);
+  }
+  auto& actual_lod = target->lod();
+  EXPECT_EQ(expect_lod.size(), actual_lod.size());
+  for (size_t i = 0; i < expect_lod.size(); ++i) {
+    for (size_t j = 0; j < expect_lod[i].size(); ++j) {
+      EXPECT_EQ(expect_lod[i][j], actual_lod[i][j]);
+    }
+  }
+}
\ No newline at end of file

paddle/operators/save_op.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserved.
+
+   Licensed 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 <stdint.h>
+#include <sys/stat.h>
+#include <fstream>
+#include <numeric>
+
+#include "paddle/framework/data_type.h"
+#include "paddle/framework/framework.pb.h"
+#include "paddle/framework/lod_tensor.h"
+#include "paddle/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+// TODO(yuyang18): If the functions below are needed by other files, move them
+// to paddle::filesystem namespace.
+constexpr char kSEP = '/';
+static bool FileExists(const std::string &filepath) {
+  struct stat buffer;
+  return (stat(filepath.c_str(), &buffer) == 0);
+}
+
+static std::string DirName(const std::string &filepath) {
+  auto pos = filepath.rfind(kSEP);
+  if (pos == std::string::npos) {
+    return "";
+  }
+  return filepath.substr(0, pos);
+}
+
+static void MkDir(const char *path) {
+  if (mkdir(path, 0755)) {
+    PADDLE_ENFORCE_EQ(errno, EEXIST, "%s mkdir failed!", path);
+  }
+}
+
+static void MkDirRecursively(const char *fullpath) {
+  if (*fullpath == '\0') return;  // empty string
+  if (FileExists(fullpath)) return;
+
+  MkDirRecursively(DirName(fullpath).c_str());
+  MkDir(fullpath);
+}
+
+class SaveOp : public framework::OperatorBase {
+ public:
+  SaveOp(const std::string &type, const framework::VariableNameMap &inputs,
+         const framework::VariableNameMap &outputs,
+         const framework::AttributeMap &attrs)
+      : OperatorBase(type, inputs, outputs, attrs) {}
+  void Run(const framework::Scope &scope,
+           const platform::DeviceContext &dev_ctx) const override {
+    auto filename = Attr<std::string>("file_path");
+    auto overwrite = Attr<bool>("overwrite");
+
+    if (FileExists(filename) && !overwrite) {
+      PADDLE_THROW("%s is existed, cannot save to it when overwrite=false",
+                   filename, overwrite);
+    }
+
+    MkDirRecursively(DirName(filename).c_str());
+
+    // FIXME(yuyang18): We save variable to local file now, but we should change
+    // it to save an output stream.
+    std::ofstream fout(filename);
+    PADDLE_ENFORCE(static_cast<bool>(fout), "Cannot open %s to write",
+                   filename);
+
+    auto iname = Input("X");
+    auto *var = scope.FindVar(iname);
+    PADDLE_ENFORCE(var != nullptr, "Cannot find variable %s for save_op",
+                   iname);
+
+    PADDLE_ENFORCE(var->IsType<framework::LoDTensor>(),
+                   "SaveOp only support LoDTensor, %s has wrong type", iname);
+
+    auto &tensor = var->Get<framework::LoDTensor>();
+
+    {  // the 1st field, uint32_t version
+      constexpr uint32_t version = 0;
+      fout.write(reinterpret_cast<const char *>(&version), sizeof(version));
+    }
+    {  // the 2nd field, tensor description
+       // int32_t  size
+       // void*    protobuf message
+      framework::TensorDesc desc;
+      desc.set_data_type(framework::ToDataType(tensor.type()));
+      auto dims = framework::vectorize(tensor.dims());
+      auto *pb_dims = desc.mutable_dims();
+      pb_dims->Resize(static_cast<int>(dims.size()), 0);
+      std::copy(dims.begin(), dims.end(), pb_dims->begin());
+      int32_t size = desc.ByteSize();
+      fout.write(reinterpret_cast<const char *>(&size), sizeof(size));
+      auto out = desc.SerializeAsString();
+      fout.write(out.data(), size);
+    }
+    {  // the 3rd field, tensor data
+      uint64_t size = tensor.memory_size();
+      auto *data_ptr = tensor.data<void>();
+      PADDLE_ENFORCE(size < std::numeric_limits<std::streamsize>::max(),
+                     "Index overflow when writing tensor");
+      if (platform::is_gpu_place(tensor.place())) {
+#ifdef PADDLE_WITH_CUDA
+        constexpr size_t kBufSize = 1024 * 1024 * 64;  // 64MB
+        std::unique_ptr<char[]> buf(new char[kBufSize]);
+        auto &gpu_dev_ctx =
+            static_cast<const platform::CUDADeviceContext &>(dev_ctx);
+        platform::CPUPlace cpu;
+        uintptr_t data = reinterpret_cast<uintptr_t>(data_ptr);
+        while (size != 0) {
+          size_t size_to_write = std::min(kBufSize, static_cast<size_t>(size));
+          memory::Copy(cpu, buf.get(),
+                       boost::get<platform::GPUPlace>(tensor.place()),
+                       reinterpret_cast<const void *>(data), size_to_write,
+                       gpu_dev_ctx.stream());
+          gpu_dev_ctx.Wait();
+          fout.write(buf.get(), size_to_write);
+          data += size_to_write;
+          size -= size_to_write;
+        }
+#else
+        PADDLE_THROW("Unexpected branch");
+#endif
+      } else {
+        fout.write(static_cast<const char *>(data_ptr),
+                   static_cast<std::streamsize>(size));
+      }
+    }
+    {  // the 4th field, lod information
+       // uint64_t lod_level
+       // uint64_t lod_level_1 size in byte.
+       // int*     lod_level_1 data
+       // ...
+      auto lod = tensor.lod();
+      uint64_t size = lod.size();
+      fout.write(reinterpret_cast<const char *>(&size), sizeof(size));
+
+      for (auto &each : lod) {
+        size = each.size() * sizeof(framework::LoD::value_type::value_type);
+        fout.write(reinterpret_cast<const char *>(&size), sizeof(size));
+        fout.write(reinterpret_cast<const char *>(each.data()),
+                   static_cast<std::streamsize>(size));
+      }
+    }
+  }
+};
+
+class SaveOpProtoMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  SaveOpProtoMaker(framework::OpProto *proto,
+                   framework::OpAttrChecker *op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput("X", "The tensor need to be saved");
+    AddComment(R"DOC(Save operator
+Save operator will serialize and write a tensor variable to disk file.
+)DOC");
+    AddAttr<bool>("overwrite", "Overwrite the output file if exist")
+        .SetDefault(true);
+    AddAttr<std::string>("file_path",
+                         "Variable will be saved to \"file_path\".")
+        .AddCustomChecker(
+            [](const std::string &path) { return !path.empty(); });
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+
+REGISTER_OPERATOR(save, ops::SaveOp, ops::SaveOpProtoMaker);

paddle/operators/save_restore_op.cc

-/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
-   Licensed 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 "paddle/framework/eigen.h"
-#include "paddle/framework/op_registry.h"
-
-#include <fstream>
-
-namespace paddle {
-namespace operators {
-
-using framework::Tensor;
-using framework::LoDTensor;
-
-inline static std::string VarToFileName(const std::string& folder_path,
-                                        const std::string& var_name) {
-  return folder_path + "/__" + var_name + "__";
-}
-
-class SaveOp : public framework::OperatorBase {
- public:
-  SaveOp(const std::string& type, const framework::VariableNameMap& inputs,
-         const framework::VariableNameMap& outputs,
-         const framework::AttributeMap& attrs)
-      : OperatorBase(type, inputs, outputs, attrs) {}
-
-  void Run(const framework::Scope& scope,
-           const platform::DeviceContext& dev_ctx) const override {
-    const auto& var_names = this->Inputs("X");
-    for (const auto& name : var_names) {
-      PADDLE_ENFORCE_NOT_NULL(scope.FindVar(name),
-                              "Can not find variable '%s' in the scope.", name);
-    }
-    std::string folder_path = this->Attr<std::string>("folderPath");
-    PADDLE_ENFORCE(!folder_path.empty(),
-                   "'folderPath' of SaveOp shouldn't be empty.");
-
-    VLOG(1) << "Save variables to folder: " << folder_path;
-    for (const auto& name : var_names) {
-      std::string file_name = VarToFileName(folder_path, name);
-      std::ofstream fout(file_name, std::ofstream::out);
-      PADDLE_ENFORCE(fout.is_open(), "Fail to create file %s.", file_name);
-      const LoDTensor& tensor = scope.FindVar(name)->Get<LoDTensor>();
-      std::string bytes = tensor.SerializeToString();
-      fout << bytes;
-      fout.close();
-    }
-    VLOG(1) << "Compelete saving variables. Items count: " << var_names.size();
-  }
-};
-
-class SaveOpMaker : public framework::OpProtoAndCheckerMaker {
- public:
-  SaveOpMaker(framework::OpProto* proto, framework::OpAttrChecker* op_checker)
-      : OpProtoAndCheckerMaker(proto, op_checker) {
-    AddInput("X",
-             "(tensor), the tensor count can be 1~INT_MAX, tensors names which "
-             "values will be saved.")
-        .AsDuplicable();
-    AddAttr<std::string>("folderPath", "the folderPath for save model.");
-    AddComment(R"DOC(
-Save the input tensors to a binary file based on input tensor names and absolute path.
-
-All the inputs can carry the LoD (Level of Details) information,
-or not.
-)DOC");
-  }
-};
-
-class RestoreOp : public framework::OperatorBase {
- public:
-  RestoreOp(const std::string& type, const framework::VariableNameMap& inputs,
-            const framework::VariableNameMap& outputs,
-            const framework::AttributeMap& attrs)
-      : OperatorBase(type, inputs, outputs, attrs) {}
-
-  void Run(const framework::Scope& scope,
-           const platform::DeviceContext& dev_ctx) const override {
-    const auto& var_names = this->Outputs("Out");
-    for (const auto& name : var_names) {
-      PADDLE_ENFORCE_NOT_NULL(scope.FindVar(name),
-                              "Can not find variable '%s' in the scope.", name);
-    }
-    std::string folder_path = this->Attr<std::string>("folderPath");
-    PADDLE_ENFORCE(!folder_path.empty(),
-                   "'folderPath' of RestoreOp shouldn't be empty.");
-
-    VLOG(1) << "Try loading variables from folder: " << folder_path;
-
-    for (const auto& name : var_names) {
-      std::string file_name = VarToFileName(folder_path, name);
-      std::ifstream fin(file_name, std::ifstream::in);
-      PADDLE_ENFORCE(fin.is_open(), "Fail to open file %s.", file_name);
-      const size_t kBufferSize = 4096;  // equal to linux page size
-      char buffer[kBufferSize];
-      std::string cache;
-      while (!fin.eof()) {
-        fin.read(buffer, kBufferSize);
-        cache.append(buffer, fin.gcount());
-      }
-      LoDTensor* tensor = scope.FindVar(name)->GetMutable<LoDTensor>();
-      tensor->DeserializeFromString(cache, dev_ctx.GetPlace());
-      fin.close();
-    }
-    VLOG(1) << "Complete loading variables.";
-  }
-};
-
-class RestoreOpMaker : public framework::OpProtoAndCheckerMaker {
- public:
-  RestoreOpMaker(framework::OpProto* proto,
-                 framework::OpAttrChecker* op_checker)
-      : OpProtoAndCheckerMaker(proto, op_checker) {
-    AddOutput("Out",
-              "(tensor), the tensor count can be 1~INT_MAX, tensors which "
-              "values will be restores.")
-        .AsDuplicable();
-    AddAttr<std::string>("folderPath", "the folderPath for model file.");
-    AddAttr<int>("data_type", "output tensor data type")
-        .SetDefault(framework::DataType::FP32);
-    AddComment(R"DOC(
-Restore the tensors from model file based on absolute path.
-
-All the tensors outputs may carry the LoD (Level of Details) information,
-or not.
-)DOC");
-  }
-};
-
-}  // namespace operators
-}  // namespace paddle
-
-REGISTER_OPERATOR(save, paddle::operators::SaveOp,
-                  paddle::framework::EmptyGradOpMaker,
-                  paddle::operators::SaveOpMaker);
-
-REGISTER_OPERATOR(restore, paddle::operators::RestoreOp,
-                  paddle::framework::EmptyGradOpMaker,
-                  paddle::operators::RestoreOpMaker);

paddle/operators/scale_op.cc

@@ -73,4 +73,5 @@ namespace ops = paddle::operators;
 REGISTER_OPERATOR(scale, ops::ScaleOp, ops::ScaleOpMaker<float>,
                   ops::ScaleGradMaker);
 REGISTER_OP_CPU_KERNEL(scale,
-                       ops::ScaleKernel<paddle::platform::CPUPlace, float>);
+                       ops::ScaleKernel<paddle::platform::CPUPlace, float>,
+                       ops::ScaleKernel<paddle::platform::CPUPlace, double>);

paddle/operators/scale_op.cu

@@ -15,4 +15,5 @@
 #include "paddle/operators/scale_op.h"
 
 REGISTER_OP_GPU_KERNEL(
-    scale, paddle::operators::ScaleKernel<paddle::platform::GPUPlace, float>);
+    scale, paddle::operators::ScaleKernel<paddle::platform::GPUPlace, float>,
+    paddle::operators::ScaleKernel<paddle::platform::GPUPlace, double>);

paddle/operators/scale_op.h

@@ -19,7 +19,7 @@
 
 namespace paddle {
 namespace operators {
-template <typename Place, typename T, typename AttrType = T>
+template <typename Place, typename T>
 class ScaleKernel : public framework::OpKernel<T> {
  public:
   virtual void Compute(const framework::ExecutionContext& context) const {
@@ -27,7 +27,7 @@ class ScaleKernel : public framework::OpKernel<T> {
     auto* in = context.Input<framework::Tensor>("X");
     tensor->mutable_data<T>(in->place());
 
-    auto scale = static_cast<T>(context.Attr<AttrType>("scale"));
+    auto scale = static_cast<T>(context.Attr<float>("scale"));
 
     auto eigen_out = framework::EigenVector<T>::Flatten(*tensor);
     auto eigen_in = framework::EigenVector<T>::Flatten(*in);

paddle/operators/sequence_conv_op.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed 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 "paddle/operators/sequence_conv_op.h"
+
+namespace paddle {
+namespace operators {
+
+class SequenceConvOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"),
+                   "Input(X) of SequenceConvOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasInput("Filter"),
+                   "Input(Filter) of SequenceConvOp should not be null.");
+    PADDLE_ENFORCE(ctx->HasOutput("Out"),
+                   "Output(Out) of SequenceConvOp should not be null.");
+
+    int context_length = ctx->Attrs().Get<int>("context_length");
+    bool padding_trainable = ctx->Attrs().Get<bool>("padding_trainable");
+    int context_start = ctx->Attrs().Get<int>("context_start");
+
+    auto in_dims = ctx->GetInputDim("X");
+    auto filter_dims = ctx->GetInputDim("Filter");
+    PADDLE_ENFORCE(in_dims.size() == 2 && filter_dims.size() == 2,
+                   "Input(X, Filter) should be 2-D tensor.");
+    PADDLE_ENFORCE(filter_dims[0] == context_length * in_dims[1],
+                   "Filter's height should be context_length * "
+                   "number_of_input_features .");
+
+    if (padding_trainable) {
+      PADDLE_ENFORCE(
+          ctx->HasInput("PaddingData"),
+          "Input(PaddingData) of SequenceConvOp should not be null.");
+      framework::DDim padding_dim = ctx->GetInputDim("PaddingData");
+      int up_pad = std::max(0, -context_start);
+      int down_pad = std::max(0, context_start + context_length - 1);
+      int total_pad = up_pad + down_pad;
+      int input_width = static_cast<int>(in_dims[1]);
+
+      if (context_start == 0 && context_length == 1) {
+        PADDLE_THROW(
+            "If context_start is 0 and context_length is 1, padding_trainable "
+            "should be false.");
+      }
+      PADDLE_ENFORCE(padding_dim.size() == 2,
+                     "Input(PaddingData) should be 2-D tensor.");
+      PADDLE_ENFORCE(
+          padding_dim[0] == total_pad && padding_dim[1] == input_width,
+          "Input(PaddingData)'s shape is not consistent with 'context_start' "
+          "and 'context_length'.");
+    }
+
+    in_dims[1] = filter_dims[1];
+    ctx->SetOutputDim("Out", in_dims);
+    ctx->ShareLoD("X", "Out");
+  }
+};
+
+class SequenceConvGradOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")),
+                   "Gradient of output(Out) should not be null.");
+    PADDLE_ENFORCE(ctx->HasInput("X"), "The input(X) should not be null.");
+
+    if (ctx->Attrs().Get<bool>("padding_trainable") &&
+        ctx->HasOutput(framework::GradVarName("PaddingData"))) {
+      ctx->SetOutputDim(framework::GradVarName("PaddingData"),
+                        ctx->GetInputDim("PaddingData"));
+    }
+    if (ctx->HasOutput(framework::GradVarName("X"))) {
+      ctx->SetOutputDim(framework::GradVarName("X"), ctx->GetInputDim("X"));
+    }
+    if (ctx->HasOutput(framework::GradVarName("Filter"))) {
+      ctx->SetOutputDim(framework::GradVarName("Filter"),
+                        ctx->GetInputDim("Filter"));
+    }
+  }
+};
+
+class SequenceConvOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  SequenceConvOpMaker(framework::OpProto* proto,
+                      framework::OpAttrChecker* op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput(
+        "X",
+        "(LoDTensor) the input(X) is a LodTensor, which support "
+        "variable-time length input sequence. The underlying tensor in "
+        "this LoDTensor is a matrix with shape (T, D), where, T is the "
+        "total time steps in this mini-batch, D is the input feature size.");
+    AddInput("PaddingData",
+             "(Tensor, optional) the input(PaddingData) is an optional "
+             "parameter, and it is learnable. "
+             "This is a tensor with shape (N, D), where N is the "
+             "top_pad + bottom_pad, D is the input feature size. In order to "
+             "ensure the equal length of sequence before and after "
+             "convolution, it is necessary to fill the top and bottom of each "
+             "sequence according to context_length, context_stride and "
+             "context_start")
+        .AsDispensable();
+    AddInput("Filter",
+             "(Tensor) the input(Filter) is an learnable parameter."
+             "This is a tensor with shape (N, D), where N is the "
+             "context_length, D is the output feature size.");
+    AddOutput(
+        "Out",
+        "(LoDTensor) the output(Out) is a LodTensor, which support "
+        "variable-time length output sequence. The underlying tensor in "
+        "this LoDTensor is a matrix with shape (T, D), where, T is the "
+        "total time steps in this mini-batch, D is the output feature size.");
+
+    AddAttr<bool>("padding_trainable",
+                  "(bool, default false) the padding data of SequenceConvOp "
+                  "is trainable or not.")
+        .SetDefault(false);
+    AddAttr<int>("context_length",
+                 "(int, default 3) the context_length of SequenceConvOp is the "
+                 "height of the convolution kernel.")
+        .SetDefault(3)
+        .GreaterThan(0);
+    AddAttr<int>("context_start",
+                 "(int, default 0) the context_start of SequenceConvOp "
+                 "represents the beginning of the convolution of the number of "
+                 "rows of sequence, which can be negative.")
+        .SetDefault(0);
+    AddAttr<int>("context_stride",
+                 "(int, default 1) the context_stride of SequenceConvOp "
+                 "represents the step length of convolution. "
+                 "Currently, SequenceConvOp only supports"
+                 "context_stride=1.")
+        .SetDefault(1)
+        .GreaterThan(0);
+
+    AddComment(R"DOC(
+    SequenceConvOp performs convolution operation on features of
+    context_length time-steps of each instance.
+    The convolution operation calculates the output based on the input, filter
+    and strides, paddings parameters. The size of each dimension of the
+    parameters is checked in the infer-shape. In order to ensure the equal
+    length of sequence before and after convolution, it is necessary to fill
+    the top and bottom of each sequence according to context_length,
+    context_stride and context_start.
+    )DOC");
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP(sequence_conv, ops::SequenceConvOp, ops::SequenceConvOpMaker,
+            sequence_conv_grad, ops::SequenceConvGradOp);
+
+REGISTER_OP_CPU_KERNEL(
+    sequence_conv, ops::SequenceConvKernel<paddle::platform::CPUPlace, float>);
+REGISTER_OP_CPU_KERNEL(
+    sequence_conv_grad,
+    ops::SequenceConvGradKernel<paddle::platform::CPUPlace, float>);

paddle/operators/sequence_conv_op.cu

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+   Licensed 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. */
+
+#define EIGEN_USE_GPU
+
+#include "paddle/operators/sequence_conv_op.h"
+
+namespace ops = paddle::operators;
+REGISTER_OP_GPU_KERNEL(
+    sequence_conv, ops::SequenceConvKernel<paddle::platform::GPUPlace, float>);
+REGISTER_OP_GPU_KERNEL(
+    sequence_conv_grad,
+    ops::SequenceConvGradKernel<paddle::platform::GPUPlace, float>);

paddle/operators/sequence_conv_op.h

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed 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. */
+
+#pragma once
+#include "paddle/framework/eigen.h"
+#include "paddle/framework/op_registry.h"
+#include "paddle/operators/math/context_project.h"
+#include "paddle/operators/math/math_function.h"
+
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+using LoDTensor = framework::LoDTensor;
+
+template <typename Place, typename T>
+class SequenceConvKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& context) const override {
+    auto* in = context.Input<LoDTensor>("X");
+    auto* out = context.Output<LoDTensor>("Out");
+    auto filter = *context.Input<Tensor>("Filter");
+
+    out->mutable_data<T>(context.GetPlace());
+    context.ShareLoD("X", "Out");
+
+    int context_start = context.Attr<int>("context_start");
+    int context_length = context.Attr<int>("context_length");
+    int context_stride = context.Attr<int>("context_stride");
+    bool padding_trainable = context.Attr<bool>("padding_trainable");
+
+    // InferShape by in_lod
+    PADDLE_ENFORCE_EQ(in->lod().size(), 1UL,
+                      "Only support one level sequence now.");
+
+    const Tensor* padding_data = nullptr;
+    if (padding_trainable) {
+      padding_data = context.Input<Tensor>("PaddingData");
+    }
+
+    int up_pad = std::max(0, -context_start);
+    int down_pad = std::max(0, context_start + context_length - 1);
+    int sequence_width;
+    sequence_width = static_cast<int>(in->dims()[1]);
+
+    // Use col_shape in the im2col calculation.
+    framework::DDim col_shape = {in->dims()[0],
+                                 sequence_width * context_length};
+    Tensor col;
+    col.mutable_data<T>(col_shape, context.GetPlace());
+    math::SetConstant<Place, T> set_zero;
+    // Because if padding_trainable is false, padding data should be zeros.
+    set_zero(context.device_context(), &col, static_cast<T>(0));
+
+    paddle::operators::math::ContextProjectFunctor<Place, T>
+        seq_project_functor;
+    LoDTensor* input = const_cast<LoDTensor*>(in);
+    Tensor* pad_data = const_cast<Tensor*>(padding_data);
+
+    seq_project_functor(context.device_context(), *input, *pad_data, col,
+                        padding_trainable, context_start, context_length,
+                        context_stride, up_pad, down_pad, false, false, false);
+
+    math::matmul<Place, T>(context.device_context(), col, false, filter, false,
+                           static_cast<T>(1.0), out, static_cast<T>(0.0));
+  }
+};
+
+template <typename Place, typename T>
+class SequenceConvGradKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& context) const override {
+    auto* out_g = context.Input<LoDTensor>(framework::GradVarName("Out"));
+    auto* in_g = context.Output<LoDTensor>(framework::GradVarName("X"));
+    auto* filter_g = context.Output<Tensor>(framework::GradVarName("Filter"));
+    auto* padding_data_g =
+        context.Output<Tensor>(framework::GradVarName("PaddingData"));
+    auto* in = context.Input<LoDTensor>("X");
+    auto* filter = context.Input<Tensor>("Filter");
+
+    int context_start = context.Attr<int>("context_start");
+    int context_length = context.Attr<int>("context_length");
+    int context_stride = context.Attr<int>("context_stride");
+    bool padding_trainable = context.Attr<bool>("padding_trainable");
+
+    PADDLE_ENFORCE_EQ(in->lod().size(), 1UL,
+                      "Only support one level sequence now.");
+    auto lod_g_level_0 = in->lod()[0];
+
+    int up_pad = std::max(0, -context_start);
+    int down_pad = std::max(0, context_start + context_length - 1);
+    int sequence_width = static_cast<int>(in->dims()[1]);
+
+    math::SetConstant<Place, T> set_zero;
+    // use col_shape in the im2col calculation
+    framework::DDim col_shape = {in->dims()[0],
+                                 sequence_width * context_length};
+    Tensor col;
+
+    if (in_g || filter_g || (padding_trainable && padding_data_g)) {
+      col.mutable_data<T>(col_shape, context.GetPlace());
+      // Because if padding_trainable is false, padding data should be zeros.
+      set_zero(context.device_context(), &col, static_cast<T>(0));
+      math::matmul<Place, T>(context.device_context(), *out_g, false, *filter,
+                             true, T(1.0), &col, T(1.0));
+    }
+    paddle::operators::math::ContextProjectFunctor<Place, T>
+        seq_project_functor;
+
+    if (in_g) {
+      in_g->mutable_data<T>(context.GetPlace());
+      in_g->set_lod(in->lod());
+      set_zero(context.device_context(), in_g, static_cast<T>(0));
+
+      seq_project_functor(context.device_context(), *in_g, *padding_data_g, col,
+                          padding_trainable, context_start, context_length,
+                          context_stride, up_pad, down_pad, true, true, false);
+    }
+
+    if (padding_trainable && padding_data_g) {
+      padding_data_g->mutable_data<T>(context.GetPlace());
+      set_zero(context.device_context(), padding_data_g, static_cast<T>(0));
+
+      LoDTensor* input = const_cast<LoDTensor*>(in);
+      seq_project_functor(context.device_context(), *input, *padding_data_g,
+                          col, padding_trainable, context_start, context_length,
+                          context_stride, up_pad, down_pad, true, false, true);
+    }
+
+    if (filter_g) {
+      filter_g->mutable_data<T>(context.GetPlace());
+      set_zero(context.device_context(), filter_g, static_cast<T>(0));
+
+      Tensor filter_grad = *filter_g;
+      LoDTensor out_grad = *out_g;
+
+      const Tensor* padding_data = nullptr;
+      if (padding_trainable) {
+        padding_data = context.Input<Tensor>("PaddingData");
+      }
+
+      sequence_width = static_cast<int>(in->dims()[1]);
+
+      LoDTensor* input = const_cast<LoDTensor*>(in);
+      Tensor* pad_data = const_cast<Tensor*>(padding_data);
+
+      seq_project_functor(context.device_context(), *input, *pad_data, col,
+                          padding_trainable, context_start, context_length,
+                          context_stride, up_pad, down_pad, false, false,
+                          false);
+
+      math::matmul<Place, T>(context.device_context(), col, true, out_grad,
+                             false, T(1.0), &filter_grad, T(1.0));
+    }
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle

paddle/operators/softmax_with_cross_entropy_op.cu

@@ -23,18 +23,21 @@ using Tensor = framework::Tensor;
 
 namespace {
 template <typename T>
-__global__ void CrossEntropyGrad(T* out_grad, const T* in_grad,
+__global__ void CrossEntropyGrad(T* logit_grad, const T* loss_grad,
                                  const int* labels, const int batch_size,
                                  const int class_num) {
   int tid = blockIdx.x * blockDim.x + threadIdx.x;
   int sample_idx = tid / class_num;
 
-  if (tid < batch_size * class_num) out_grad[tid] *= in_grad[sample_idx];
-  __syncthreads();
-
   if (tid < batch_size) {
     PADDLE_ASSERT(labels[sample_idx] >= 0 && labels[sample_idx] < class_num);
-    out_grad[tid * class_num + labels[tid]] -= 1.;
+    logit_grad[tid * class_num + labels[tid]] -= static_cast<T>(1.);
+  }
+
+  __syncthreads();
+
+  if (tid < batch_size * class_num) {
+    logit_grad[tid] *= loss_grad[sample_idx];
   }
 }
 
@@ -47,7 +50,7 @@ __global__ void SoftCrossEntropyGradientKernel(T* logit_grad,
   int ids = blockIdx.x * blockDim.x + threadIdx.x;
   if (ids < batch_size * class_num) {
     int row_ids = ids / class_num;
-    logit_grad[ids] = logit_grad[ids] * loss_grad[row_ids] - labels[ids];
+    logit_grad[ids] = logit_grad[ids] * (loss_grad[row_ids] - labels[ids]);
   }
 }
 }  // namespace

paddle/operators/softmax_with_cross_entropy_op.h

@@ -67,8 +67,8 @@ class SoftmaxWithCrossEntropyGradKernel : public framework::OpKernel<T> {
 
       logit_grad_mat.device(context.GetEigenDevice<platform::CPUPlace>()) =
           logit_grad_mat *
-              out_grad_mat.broadcast(Eigen::DSizes<int, 2>(1, class_num)) -
-          lbl_mat;
+          (out_grad_mat.broadcast(Eigen::DSizes<int, 2>(1, class_num)) -
+           lbl_mat);
     } else {
       const int batch_size = logit_grad->dims()[0];
       const int* label_data = labels->data<int>();
@@ -78,7 +78,7 @@ class SoftmaxWithCrossEntropyGradKernel : public framework::OpKernel<T> {
       for (int i = 0; i < batch_size; ++i) {
         int index = i * class_num + label_data[i];
         logit_grad_data[index] =
-            (out_grad_data[i] * logit_grad_data[index] - 1.);
+            out_grad_data[i] * (logit_grad_data[index] - 1.);
       }
     }
   }

paddle/operators/split_op.cc

@@ -95,17 +95,18 @@ class SplitOpMaker : public framework::OpProtoAndCheckerMaker {
   }
 };
 
-class SplitOpGrad : public NetOp {
+class SplitGradMaker : public framework::SingleGradOpDescMaker {
  public:
-  SplitOpGrad(const std::string &type, const framework::VariableNameMap &inputs,
-              const framework::VariableNameMap &outputs,
-              const framework::AttributeMap &attrs)
-      : NetOp(type, inputs, outputs, attrs) {
-    auto out_grad = Inputs(framework::GradVarName("Out"));
-    auto x_grad = Output(framework::GradVarName("X"));
-    AppendOp(framework::OpRegistry::CreateOp("concat", {{"X", out_grad}},
-                                             {{"Out", {x_grad}}}, attrs));
-    CompleteAddOp(false);
+  using framework::SingleGradOpDescMaker::SingleGradOpDescMaker;
+
+ protected:
+  std::unique_ptr<framework::OpDescBind> Apply() const override {
+    auto op = new framework::OpDescBind();
+    op->SetType("concat");
+    op->SetInput("X", OutputGrad("Out"));
+    op->SetOutput("Out", InputGrad("X"));
+    op->SetAttrMap(Attrs());
+    return std::unique_ptr<framework::OpDescBind>(op);
   }
 };
 
@@ -114,7 +115,7 @@ class SplitOpGrad : public NetOp {
 
 namespace ops = paddle::operators;
 USE_CPU_ONLY_OP(concat);
-REGISTER_OP(split, ops::SplitOp, ops::SplitOpMaker, split_grad,
-            ops::SplitOpGrad);
+
+REGISTER_OPERATOR(split, ops::SplitOp, ops::SplitOpMaker, ops::SplitGradMaker);
 REGISTER_OP_CPU_KERNEL(split,
                        ops::SplitOpKernel<paddle::platform::CPUPlace, float>);

paddle/operators/squared_l2_norm_op.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+   Licensed 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 "paddle/operators/squared_l2_norm_op.h"
+
+namespace paddle {
+namespace operators {
+
+using framework::Tensor;
+
+class SquaredL2NormOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should be not null.");
+    PADDLE_ENFORCE(ctx->HasOutput("Out"), "Output(Out) should be not null.");
+
+    ctx->SetOutputDim("Out", {1});
+  }
+};
+
+class SquaredL2NormGradOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should be not null.");
+    PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")),
+                   "Input(Out@GRAD) should be not null.");
+    PADDLE_ENFORCE(ctx->HasOutput(framework::GradVarName("X")),
+                   "Output(X@GRAD) should be not null.");
+
+    ctx->SetOutputDim(framework::GradVarName("X"), ctx->GetInputDim("X"));
+  }
+};
+
+class SquaredL2NormOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  SquaredL2NormOpMaker(framework::OpProto* proto,
+                       framework::OpAttrChecker* op_checker)
+      : framework::OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput("X", "(Tensor) The input of squared_l2_norm op.");
+    AddOutput("Out", "(Float) The output of squared_l2_norm op.");
+    AddComment(R"DOC(
+SquaredL2Norm Operator.
+
+Computes the squared L2 norm of a tensor.
+
+Out = sum (X ** 2)
+
+)DOC");
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP(squared_l2_norm, ops::SquaredL2NormOp, ops::SquaredL2NormOpMaker,
+            squared_l2_norm_grad, ops::SquaredL2NormGradOp);
+REGISTER_OP_CPU_KERNEL(
+    squared_l2_norm,
+    ops::SquaredL2NormKernel<paddle::platform::CPUPlace, float>);
+REGISTER_OP_CPU_KERNEL(
+    squared_l2_norm_grad,
+    ops::SquaredL2NormGradKernel<paddle::platform::CPUPlace, float>);

paddle/operators/squared_l2_norm_op.cu

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+   Licensed 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. */
+
+#define EIGEN_USE_GPU
+#include "paddle/operators/squared_l2_norm_op.h"
+
+namespace ops = paddle::operators;
+REGISTER_OP_GPU_KERNEL(
+    squared_l2_norm,
+    ops::SquaredL2NormKernel<paddle::platform::GPUPlace, float>);
+REGISTER_OP_GPU_KERNEL(
+    squared_l2_norm_grad,
+    ops::SquaredL2NormGradKernel<paddle::platform::GPUPlace, float>);

paddle/operators/squared_l2_norm_op.h

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+   Licensed 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. */
+
+#pragma once
+#include "paddle/framework/eigen.h"
+#include "paddle/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+// Out = sum(square(X))
+template <typename Place, typename T>
+class SquaredL2NormKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext &context) const override {
+    const framework::Tensor *X = context.Input<framework::Tensor>("X");
+    framework::Tensor *Out = context.Output<framework::Tensor>("Out");
+    Out->mutable_data<T>(context.GetPlace());
+
+    auto x = framework::EigenVector<T>::Flatten(*X);
+    auto out = framework::EigenVector<T>::Flatten(*Out);
+    auto place = context.GetEigenDevice<Place>();
+
+    out.device(place) = x.square().sum();
+  }
+};
+
+// dX = X
+template <typename Place, typename T>
+class SquaredL2NormGradKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext &context) const override {
+    const framework::Tensor *X = context.Input<framework::Tensor>("X");
+    const framework::Tensor *dOut =
+        context.Input<framework::Tensor>(framework::GradVarName("Out"));
+    PADDLE_ENFORCE(dOut->numel() == 1,
+                   "Squared L2 Norm Gradient should be scalar");
+    framework::Tensor *dX =
+        context.Output<framework::Tensor>(framework::GradVarName("X"));
+    dX->mutable_data<T>(context.GetPlace());
+
+    auto x = framework::EigenVector<T>::Flatten(*X);
+    auto dout = framework::EigenVector<T>::Flatten(*dOut);
+    auto dx = framework::EigenVector<T>::Flatten(*dX);
+    auto place = context.GetEigenDevice<Place>();
+
+    Eigen::DSizes<int, 1> x_dsize(X->numel());
+    dx.device(place) = (dout.broadcast(x_dsize) * x) * static_cast<T>(2.0);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle

paddle/operators/sum_op.cc

@@ -11,6 +11,7 @@ limitations under the License. */
 
 #include "paddle/operators/sum_op.h"
 #include <vector>
+#include "paddle/framework/var_type_inference.h"
 #include "paddle/operators/net_op.h"
 
 namespace paddle {
@@ -55,6 +56,26 @@ or not. But the output only shares the LoD with the first input.
   }
 };
 
+class SumOpVarTypeInference : public framework::VarTypeInference {
+ public:
+  void operator()(const framework::OpDescBind& op_desc,
+                  framework::BlockDescBind* block) const override {
+    auto& inputs = op_desc.Input("X");
+    auto default_var_type = framework::VarDesc::SELECTED_ROWS;
+
+    bool any_input_is_lod_tensor = std::any_of(
+        inputs.begin(), inputs.end(), [block](const std::string& name) {
+          return block->Var(name)->GetType() == framework::VarDesc::LOD_TENSOR;
+        });
+    if (any_input_is_lod_tensor) {
+      default_var_type = framework::VarDesc::LOD_TENSOR;
+    }
+
+    auto out_var_name = op_desc.Output("Out").front();
+    block->Var(out_var_name)->SetType(default_var_type);
+  }
+};
+
 class SumGradMaker : public framework::GradOpDescMakerBase {
  public:
   using framework::GradOpDescMakerBase::GradOpDescMakerBase;
@@ -83,5 +104,7 @@ class SumGradMaker : public framework::GradOpDescMakerBase {
 
 namespace ops = paddle::operators;
 
-REGISTER_OPERATOR(sum, ops::SumOp, ops::SumOpMaker, ops::SumGradMaker);
-REGISTER_OP_CPU_KERNEL(sum, ops::SumKernel<paddle::platform::CPUPlace, float>);
+REGISTER_OPERATOR(sum, ops::SumOp, ops::SumOpMaker, ops::SumGradMaker,
+                  ops::SumOpVarTypeInference);
+REGISTER_OP_CPU_KERNEL(sum, ops::SumKernel<paddle::platform::CPUPlace, float>,
+                       ops::SumKernel<paddle::platform::CPUPlace, double>);

paddle/operators/sum_op.cu

@@ -13,4 +13,5 @@ limitations under the License. */
 #include "paddle/operators/sum_op.h"
 
 namespace ops = paddle::operators;
-REGISTER_OP_GPU_KERNEL(sum, ops::SumKernel<paddle::platform::GPUPlace, float>);
+REGISTER_OP_GPU_KERNEL(sum, ops::SumKernel<paddle::platform::GPUPlace, float>,
+                       ops::SumKernel<paddle::platform::GPUPlace, double>);

paddle/operators/sum_op.h

@@ -12,11 +12,15 @@ limitations under the License. */
 #pragma once
 #include "paddle/framework/eigen.h"
 #include "paddle/framework/op_registry.h"
+#include "paddle/operators/math/math_function.h"
+#include "paddle/operators/math/selected_rows_functor.h"
 
 namespace paddle {
 namespace operators {
 
 using Tensor = framework::Tensor;
+using SelectedRows = framework::SelectedRows;
+using LoDTensor = framework::LoDTensor;
 template <typename T, int MajorType = Eigen::RowMajor,
           typename IndexType = Eigen::DenseIndex>
 using EigenVector = framework::EigenVector<T, MajorType, IndexType>;
@@ -25,19 +29,68 @@ template <typename Place, typename T>
 class SumKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& context) const override {
-    auto ins = context.MultiInput<Tensor>("X");
-    auto* out = context.Output<Tensor>("Out");
-    out->mutable_data<T>(context.GetPlace());
-
-    auto place = context.GetEigenDevice<Place>();
-    auto result = EigenVector<T>::Flatten(*out);
-
-    int N = ins.size();
-    auto in = EigenVector<T>::Flatten(*(ins[0]));
-    result.device(place) = in;
-    for (int i = 1; i < N; i++) {
-      auto in = EigenVector<T>::Flatten(*(ins[i]));
-      result.device(place) = result + in;
+    auto& in_vars = context.MultiInputVar("X");
+    int N = in_vars.size();
+    auto out_var = context.OutputVar("Out");
+
+    if (out_var->IsType<framework::LoDTensor>()) {
+      auto* out = context.Output<Tensor>("Out");
+      // Runtime InferShape
+      for (int i = 0; i < N; i++) {
+        if (in_vars[i]->IsType<framework::LoDTensor>()) {
+          out->Resize(in_vars[i]->Get<framework::LoDTensor>().dims());
+          break;
+        }
+      }
+      out->mutable_data<T>(context.GetPlace());
+
+      auto result = EigenVector<T>::Flatten(*out);
+
+      math::SetConstant<Place, T> constant_functor;
+      constant_functor(context.device_context(), out, 0.0);
+
+      math::SelectedRowsAddToTensor<Place, T> functor;
+      auto place = context.GetEigenDevice<Place>();
+      for (int i = 0; i < N; i++) {
+        if (in_vars[i]->IsType<framework::LoDTensor>()) {
+          auto& in_t = in_vars[i]->Get<framework::LoDTensor>();
+          auto in = EigenVector<T>::Flatten(in_t);
+          result.device(place) = result + in;
+        } else if (in_vars[i]->IsType<framework::SelectedRows>()) {
+          auto& in_t = in_vars[i]->Get<framework::SelectedRows>();
+          functor(context.device_context(), in_t, out);
+        } else {
+          PADDLE_THROW("Variable type must be LoDTensor/SelectedRows.");
+        }
+      }
+    } else if (out_var->IsType<framework::SelectedRows>()) {
+      auto* out = context.Output<SelectedRows>("Out");
+      auto* out_value = out->mutable_value();
+
+      // Runtime InferShape
+      size_t first_dim = 0;
+      for (int i = 0; i < N; i++) {
+        first_dim += in_vars[i]->Get<SelectedRows>().rows().size();
+      }
+      auto in_dim = in_vars[0]->Get<SelectedRows>().value().dims();
+
+      auto in_dim_vec = framework::vectorize(in_dim);
+      in_dim_vec[0] = static_cast<int64_t>(first_dim);
+
+      out_value->Resize(framework::make_ddim(in_dim_vec));
+
+      out_value->mutable_data<T>(context.GetPlace());
+
+      math::SelectedRowsAddTo<Place, T> functor;
+
+      int64_t offset = 0;
+      for (int i = 0; i < N; i++) {
+        PADDLE_ENFORCE_EQ(out->height(),
+                          in_vars[i]->Get<SelectedRows>().height())
+        functor(context.device_context(), in_vars[i]->Get<SelectedRows>(),
+                offset, out);
+        offset += in_vars[i]->Get<SelectedRows>().value().numel();
+      }
     }
   }
 };

paddle/pybind/protobuf.cc

@@ -105,6 +105,11 @@ void BindProgramDesc(py::module &m) {
            [](ProgramDescBind &self, const ProgramDescBind &other) {
              new (&self) ProgramDescBind(other);
            })
+      .def("__init__",
+           [](ProgramDescBind &self, const py::bytes &binary_str) {
+             std::string str(binary_str);
+             new (&self) ProgramDescBind(str);
+           })
       .def("append_block", &ProgramDescBind::AppendBlock,
            py::return_value_policy::reference)
       .def("append_backward",

paddle/trainer/MergeModel.cpp

@@ -20,6 +20,7 @@ limitations under the License. */
 #include "paddle/utils/PythonUtil.h"
 
 DEFINE_string(model_dir, "", "Directory for separated model files");
+DEFINE_string(config_file, "", "Config file for the model");
 DEFINE_string(model_file, "", "File for merged model file");
 
 using namespace paddle;  // NOLINT
@@ -28,7 +29,8 @@ using namespace std;     // NOLINT
 int main(int argc, char** argv) {
   initMain(argc, argv);
   initPython(argc, argv);
-  string confFile = TrainerConfigHelper::getConfigNameFromPath(FLAGS_model_dir);
+
+  string confFile = FLAGS_config_file;
 #ifndef PADDLE_WITH_CUDA
   FLAGS_use_gpu = false;
 #endif

paddle/trainer/NewRemoteParameterUpdater.cpp

@@ -110,43 +110,10 @@ void NewRemoteParameterUpdater::init(
 
     // overwrite optimizerConfigV2 for per-parameter(layer) configs
     for (int i = 0; i < parameterSize(); ++i) {
-      auto paramConfig = parameters_[i]->getConfig();
-      if (paramConfig.has_momentum() &&
-          trainerConfig_.learning_method() == "momentum") {
-        optimizerConfigV2.mutable_sgd()->set_momentum(paramConfig.momentum());
-      }
-      if (paramConfig.has_learning_rate()) {
-        switch (optimizerConfigV2.lr_policy()) {
-          case 0:
-            optimizerConfigV2.mutable_const_lr()->set_learning_rate(
-                paramConfig.learning_rate());
-            break;
-          case 1:
-            optimizerConfigV2.mutable_linear_lr()->set_learning_rate(
-                paramConfig.learning_rate());
-            break;
-        }
-      }
-      if (paramConfig.has_decay_rate()) {
-        switch (optimizerConfigV2.optimizer()) {
-          case 1:  // SGD
-            optimizerConfigV2.mutable_sgd()->set_decay(
-                paramConfig.decay_rate());
-            break;
-          case 2:  // Adadelta
-            optimizerConfigV2.mutable_adadelta()->set_decay(
-                paramConfig.decay_rate());
-            break;
-          case 3:  // Adagrad
-            optimizerConfigV2.mutable_adagrad()->set_decay(
-                paramConfig.decay_rate());
-            break;
-          case 4:  // Adam
-            optimizerConfigV2.mutable_adam()->set_decay(
-                paramConfig.decay_rate());
-            break;
-        }
-      }
+      // FIXME(typhoonzero): paramConfig always have default values,
+      // how to check if it's default?
+      // TODO(typhoonzero): log output: optimizerConfigV2.DebugString();
+      LOG(INFO) << "trainerConfig_: " << trainerConfig_.DebugString();
       // send param and config to pserver
       std::string bytes = optimizerConfigV2.SerializeAsString();
       const char *array = bytes.data();

proto/TrainerConfig.proto

@@ -19,7 +19,7 @@ import "ModelConfig.proto";
 package paddle;
 
 message OptimizationConfig {
-  required int32 batch_size = 3;
+  optional int32 batch_size = 3 [ default = 1 ];
   required string algorithm = 4 [ default = "async_sgd" ];
   optional int32 num_batches_per_send_parameter = 5 [ default = 1 ];
   optional int32 num_batches_per_get_parameter = 6 [ default = 1 ];

python/paddle/v2/framework/executor.py

@@ -19,11 +19,16 @@ class Executor(object):
 
     def run(self,
             program,
-            feed,
-            fetch_list,
+            feed=None,
+            fetch_list=None,
             feed_var_name='feed',
             fetch_var_name='fetch',
             scope=None):
+        if feed is None:
+            feed = {}
+        if fetch_list is None:
+            fetch_list = []
+
         if not isinstance(program, Program):
             raise TypeError()
 

python/paddle/v2/framework/framework.py

@@ -261,7 +261,7 @@ class Operator(object):
                     self.desc.set_attr(attr_name, attrs[attr_name])
 
         self.desc.check_attrs()
-        no_kernel_op_set = {'feed', 'fetch', 'save', 'restore'}
+        no_kernel_op_set = {'feed', 'fetch', 'save', 'load'}
         if type not in no_kernel_op_set:
             self.desc.infer_var_type(self.block.desc)
             self.desc.infer_shape(self.block.desc)
@@ -440,6 +440,13 @@ class Program(object):
         p.sync_with_cpp()
         return p
 
+    @staticmethod
+    def parse_from_string(binary_str):
+        p = Program()
+        p.desc = core.ProgramDesc(binary_str)
+        p.sync_with_cpp()
+        return p
+
     def __repr__(self):
         return str(self)
 
@@ -479,6 +486,11 @@ class Program(object):
         for block in self.blocks:
             block.sync_with_cpp()
 
+    def list_vars(self):
+        for each_block in self.blocks:
+            for each_var in each_block.vars.itervalues():
+                yield each_var
+
 
 class Parameter(Variable):
     def __init__(self, block, shape, dtype, **kwargs):
@@ -498,6 +510,8 @@ class Parameter(Variable):
 
         self.optimize_attr = kwargs.get('optimize_attr', {'learning_rate': 1.0})
 
+        self.regularizer = kwargs.get('regularizer', None)
+
 
 # program is a global instance.
 g_program = Program()

python/paddle/v2/framework/io.py

+import os
+
+from paddle.v2.framework.framework import Program, Parameter, g_program, \
+    Variable
+
+__all__ = [
+    'save_vars', 'save_params', 'save_persistables', 'load_vars', 'load_params',
+    'load_persistables'
+]
+
+
+def is_parameter(var):
+    return isinstance(var, Parameter)
+
+
+def is_persistable(var):
+    return var.persistable
+
+
+def _clone_var_in_block_(block, var):
+    assert isinstance(var, Variable)
+    return block.create_var(
+        name=var.name,
+        shape=var.shape,
+        dtype=var.data_type,
+        type=var.type,
+        lod_level=var.lod_level,
+        persistable=True)
+
+
+def save_vars(executor, dirname, program=None, vars=None, predicate=None):
+    """
+    Save variables to directory by executor.
+    
+    :param executor: executor that save variable
+    :param dirname: directory path
+    :param program: program. If vars is None, then filter all variables in this 
+    program which fit `predicate`. Default g_program.
+    :param predicate: The Predicate describes a callable that returns a variable
+    as a bool. If it returns true, the variables will be saved.
+    :param vars: variables need to be saved. If specify vars, program & predicate
+    will be ignored
+    :return: None
+    """
+    if vars is None:
+        if program is None:
+            program = g_program
+        if not isinstance(program, Program):
+            raise TypeError("program should be as Program type or None")
+
+        save_vars(
+            executor,
+            dirname=dirname,
+            vars=filter(predicate, program.list_vars()))
+    else:
+        save_program = Program()
+        save_block = save_program.global_block()
+        for each_var in vars:
+            new_var = _clone_var_in_block_(save_block, each_var)
+            save_block.append_op(
+                type='save',
+                inputs={'X': [new_var]},
+                outputs={},
+                attrs={'file_path': os.path.join(dirname, new_var.name)})
+        executor.run(save_program)
+
+
+def save_params(executor, dirname, program=None):
+    """
+    Save all parameters to directory with executor.
+    """
+    save_vars(
+        executor,
+        dirname=dirname,
+        program=program,
+        vars=None,
+        predicate=is_parameter)
+
+
+def save_persistables(executor, dirname, program=None):
+    """
+    Save all persistables to directory with executor.
+    """
+    save_vars(
+        executor,
+        dirname=dirname,
+        program=program,
+        vars=None,
+        predicate=is_persistable)
+
+
+def load_vars(executor, dirname, program=None, vars=None, predicate=None):
+    """
+    Load variables from directory by executor.
+    
+    :param executor: executor that save variable
+    :param dirname: directory path
+    :param program: program. If vars is None, then filter all variables in this 
+    program which fit `predicate`. Default g_program.
+    :param predicate: The Predicate describes a callable that returns a variable
+    as a bool. If it returns true, the variables will be loaded.
+    :param vars: variables need to be loaded. If specify vars, program & 
+    predicate will be ignored
+    :return: None
+    """
+    if vars is None:
+        if program is None:
+            program = g_program
+        if not isinstance(program, Program):
+            raise TypeError("program's type should be Program")
+
+        load_vars(
+            executor,
+            dirname=dirname,
+            vars=filter(predicate, program.list_vars()))
+    else:
+        load_prog = Program()
+        load_block = load_prog.global_block()
+        for each_var in vars:
+            assert isinstance(each_var, Variable)
+            new_var = _clone_var_in_block_(load_block, each_var)
+            load_block.append_op(
+                type='load',
+                inputs={},
+                outputs={"Out": [new_var]},
+                attrs={'file_path': os.path.join(dirname, new_var.name)})
+        executor.run(load_prog)
+
+
+def load_params(executor, dirname, program=None):
+    """
+    load all parameters from directory by executor.
+    """
+    load_vars(
+        executor, dirname=dirname, program=program, predicate=is_parameter)
+
+
+def load_persistables(executor, dirname, program=None):
+    """
+    load all persistables from directory by executor.
+    """
+    load_vars(
+        executor, dirname=dirname, program=program, predicate=is_persistable)

python/paddle/v2/framework/layer_helper.py

@@ -75,18 +75,29 @@ class LayerHelper(object):
             }
         }
         actual = self.kwargs.get('param_attr', None)
-        return actual if actual is not None else default
+        if actual is None:
+            actual = default
+        for default_field in default.keys():
+            if default_field not in actual:
+                actual[default_field] = default[default_field]
+        return actual
 
     def bias_attr(self):
+        default = {
+            'name': None,
+            'init_attr': {
+                'type': 'fill_constant',
+                'value': 0.0
+            }
+        }
         bias_attr = self.kwargs.get('bias_attr', None)
         if bias_attr is True:
-            bias_attr = {
-                'name': None,
-                'init_attr': {
-                    'type': 'fill_constant',
-                    'value': 0.0
-                }
-            }
+            bias_attr = default
+
+        if isinstance(bias_attr, dict):
+            for default_field in default.keys():
+                if default_field not in bias_attr:
+                    bias_attr[default_field] = default[default_field]
         return bias_attr
 
     def multiple_param_attr(self, length):

python/paddle/v2/framework/layers.py

@@ -97,15 +97,28 @@ def _convert_(name):
 
 def _create_op_func_(op_type):
     op_proto = OpProtoHolder.instance().get_op_proto(op_type)
-    if len(op_proto.outputs) != 1:
+    not_intermediate_outputs = \
+        filter(lambda output: not output.intermediate, op_proto.outputs)
+    intermediate_outputs = \
+        filter(lambda output: output.intermediate, op_proto.outputs)
+
+    if len(not_intermediate_outputs) != 1:
         raise ValueError(
-            "Only one output operator can be automatically generated")
+            "Only one not intermediate output operator can be automatically generated"
+        )
 
-    if op_proto.outputs[0].duplicable:
+    if not_intermediate_outputs[0].duplicable:
         raise ValueError(
             "Only not duplicable op can be automatically generated")
 
-    o_name = op_proto.outputs[0].name
+    for output in intermediate_outputs:
+        if output.duplicable:
+            raise ValueError(
+                "Only when all intermediate ops are not duplicable, "
+                "this op can be automatically generated")
+
+    o_name = not_intermediate_outputs[0].name
+    intermediate_output_names = [output.name for output in intermediate_outputs]
 
     def func(**kwargs):
         helper = LayerHelper(op_type, **kwargs)
@@ -128,9 +141,13 @@ def _create_op_func_(op_type):
                         "operator {0} must input same dtype".format(op_type))
             inputs[ipt.name] = val
 
+        outputs = dict()
         out = helper.create_tmp_variable(dtype=dtype)
+        outputs[o_name] = [out]
+        for name in intermediate_output_names:
+            outputs[name] = [helper.create_tmp_variable(dtype=dtype)]
         helper.append_op(
-            type=op_type, inputs=inputs, outputs={o_name: [out]}, attrs=kwargs)
+            type=op_type, inputs=inputs, outputs=outputs, attrs=kwargs)
         return out
 
     func.__name__ = op_type
@@ -141,6 +158,7 @@ def _create_op_func_(op_type):
 
 _create_op_func_('mean')
 _create_op_func_('mul')
+_create_op_func_('dropout')
 
 
 def concat(input, axis, program=None, init_program=None):

python/paddle/v2/framework/optimizer.py

@@ -2,9 +2,11 @@ from collections import defaultdict
 
 import paddle.v2.framework.framework as framework
 from paddle.v2.framework.backward import append_backward_ops
+from paddle.v2.framework.regularizer import append_regularization_ops
 
 __all__ = [
-    'SGDOptimizer', 'MomentumOptimizer', 'AdagradOptimizer', 'AdamOptimizer'
+    'SGDOptimizer', 'MomentumOptimizer', 'AdagradOptimizer', 'AdamOptimizer',
+    'AdamaxOptimizer'
 ]
 
 
@@ -160,6 +162,8 @@ class Optimizer(object):
         """
         params_grads = append_backward_ops(loss, parameter_list, no_grad_set or
                                            set())
+        # Add regularization if any 
+        params_grads = append_regularization_ops(params_grads)
         optimize_ops = self.create_optimization_pass(params_grads, loss)
         return optimize_ops
 
@@ -399,7 +403,7 @@ class AdamOptimizer(Optimizer):
                                         param_and_grad[0])
         moment2 = self._get_accumulator(self._moment2_acc_str,
                                         param_and_grad[0])
-        # create the momentum optimize op
+        # create the adam optimize op
         adam_op = block.append_op(
             type=self.type,
             inputs={
@@ -442,3 +446,108 @@ class AdamOptimizer(Optimizer):
             attrs={"scale": self._beta2})
 
         return [scale_beta1, scale_beta2]
+
+
+class AdamaxOptimizer(Optimizer):
+    """Implements the Adamax Optimizer
+    """
+    _moment_acc_str = "moment"
+    _inf_norm_acc_str = "inf_norm"
+
+    def __init__(self,
+                 learning_rate=0.001,
+                 beta1=0.9,
+                 beta2=0.999,
+                 epsilon=1e-8):
+        assert learning_rate is not None
+        assert beta1 is not None
+        assert beta2 is not None
+        assert epsilon is not None
+        super(AdamaxOptimizer, self).__init__()
+        self.type = "adamax"
+        self._learning_rate = learning_rate
+        self._beta1 = beta1
+        self._beta2 = beta2
+        self._epsilon = epsilon
+
+    def _initialize_tensors(self, block):
+        assert isinstance(block, framework.Block)
+        lr_shape = [1]
+        # create a variable for learning_rate
+        self._lr = block.create_var(
+            dtype="float32", shape=lr_shape, lod_level=0)
+
+        # create an op to init the learning_rate
+        # FIXME: Fix when Initialization design has been implemented
+        # https://github.com/PaddlePaddle/Paddle/pull/4852
+        block.append_op(
+            type="fill_constant",
+            outputs={"Out": self._lr},
+            attrs={"shape": lr_shape,
+                   "value": self._learning_rate})
+
+    def _create_accumulators(self, block, parameters):
+        assert isinstance(block, framework.Block)
+
+        global_block = block.program.global_block()
+        # Create beta1 power accumulator tensor
+        beta_shape = [1]
+        self._beta1_pow_acc = global_block.create_var(
+            dtype="float32", shape=beta_shape, lod_level=0)
+
+        # Initialize beta1 power accumulator
+        # FIXME: Fix when Initialization design has been implemented
+        # https://github.com/PaddlePaddle/Paddle/pull/4852
+        global_block.append_op(
+            type="fill_constant",
+            outputs={"Out": self._beta1_pow_acc},
+            attrs={"shape": beta_shape,
+                   "value": self._beta1})
+
+        # Create accumulator tensors for first moment and infinity norm
+        for p in parameters:
+            self._add_accumulator(block, self._moment_acc_str, p, 'float32')
+            self._add_accumulator(block, self._inf_norm_acc_str, p, 'float32')
+
+    def _append_optimize_op(self, block, param_and_grad):
+        assert isinstance(block, framework.Block)
+
+        moment = self._get_accumulator(self._moment_acc_str, param_and_grad[0])
+        inf_norm = self._get_accumulator(self._inf_norm_acc_str,
+                                         param_and_grad[0])
+        # create the adamax optimize op
+        adamax_op = block.append_op(
+            type=self.type,
+            inputs={
+                "Param": param_and_grad[0],
+                "Grad": param_and_grad[1],
+                "LearningRate": self._lr,
+                "Moment": moment,
+                "InfNorm": inf_norm,
+                "Beta1Pow": self._beta1_pow_acc
+            },
+            outputs={
+                "ParamOut": param_and_grad[0],
+                "MomentOut": moment,
+                "InfNormOut": inf_norm
+            },
+            attrs={
+                "beta1": self._beta1,
+                "beta2": self._beta2,
+                "epsilon": self._epsilon
+            })
+
+        return adamax_op
+
+    def _finish_update(self, block):
+        """Update Beta1 Power accumulator
+        """
+        assert isinstance(block, framework.Block)
+        global_block = block.program.global_block()
+        scale_beta1 = global_block.append_op(
+            type="scale",
+            inputs={"X": self._beta1_pow_acc},
+            outputs={"Out": self._beta1_pow_acc},
+            attrs={"scale": self._beta1})
+
+        return [scale_beta1]

python/paddle/v2/framework/regularizer.py

+import paddle.v2.framework.framework as framework
+
+__all__ = ['append_regularization_ops', 'L2DecayRegularizer']
+
+
+def append_regularization_ops(parameters_and_grads):
+    """Create and add backward regularization Operators
+
+    Creates and adds backward regularization operators in the BlockDesc.
+    This will add gradients of the regularizer function to the gradients
+    of the parameters and return these modified gradients. This is the
+    same as implementing weight decay in optimizers for regularization.
+
+    Args:
+        parameters_and_grads: A list of (parameters, gradients) pairs
+                              that need to be regularized.
+
+    Returns:
+        list of (parameters, gradients) pair with the regularized gradient
+
+    Raises:
+        Exception: Unknown regularization type
+    """
+    params_and_grads = []
+    for param, grad in parameters_and_grads:
+        # If no gradient or no regularization specified,
+        # then we don't need to do anything
+        if grad is None or param.regularizer is None:
+            params_and_grads.append((param, grad))
+            continue
+
+        # Add variable for regularization term in grad block
+        regularization_term = param.regularizer(param, grad.block)
+        assert grad.shape == regularization_term.shape
+
+        grad.block.append_op(
+            type='elementwise_add',
+            inputs={"X": grad,
+                    "Y": regularization_term},
+            outputs={"Out": grad})
+        params_and_grads.append((param, grad))
+
+    return params_and_grads
+
+
+class WeightDecayRegularizer(object):
+    """Base class for weight decay regularizers
+
+    Defines the common interface of weight-decay regularizers.
+    Weight-decay regularizers are added only during the backward
+    pass for faster regularization. They add operations to the network
+    that correspond to gradient of the regularization function.
+    Users should not use this class directly, but need to use one
+    of its implementations
+    """
+
+    def __init__(self):
+        pass
+
+    def __call__(self, param, block):
+        """Add corresponding weight decay operations to the network
+        """
+        raise NotImplementedError()
+
+
+class L2DecayRegularizer(WeightDecayRegularizer):
+    """Implements the L2 Weight Decay Regularization
+    """
+
+    def __init__(self, regularization_coeff=0.0):
+        assert regularization_coeff is not None
+        super(L2DecayRegularizer, self).__init__()
+        self._regularization_coeff = regularization_coeff
+
+    def __call__(self, param, block):
+        """Add L2 weight decay ops to network
+
+        Adds L2 weight decay ops.
+        L2WeightDecay = reg_coeff * parameter
+
+        Args:
+            param: parameter variable for which regularization is applied
+            block: block in which variable is to be created
+
+        Returns:
+            new variable for weight decay
+        """
+        assert isinstance(param, framework.Parameter)
+        assert isinstance(block, framework.Block)
+        decay = block.create_var(
+            dtype="float32", shape=param.shape, lod_level=param.lod_level)
+        # Append Op to calculate decay
+        block.append_op(
+            type='scale',
+            inputs={"X": param},
+            outputs={"Out": decay},
+            attrs={"scale": self._regularization_coeff})
+
+        return decay

python/paddle/v2/framework/tests/.gitignore

 image/
+fit_a_line.model/

python/paddle/v2/framework/tests/op_test.py

@@ -3,20 +3,27 @@ import numpy as np
 import random
 import itertools
 import paddle.v2.framework.core as core
+import collections
+from paddle.v2.framework.backward import append_backward_ops
 from paddle.v2.framework.op import Operator
 from paddle.v2.framework.executor import Executor
 from paddle.v2.framework.framework import Program, OpProtoHolder
 
 
-def grad_var_name(var_name):
-    return var_name + "@GRAD"
+def randomize_probability(batch_size, class_num, dtype='float32'):
+    prob = np.random.uniform(
+        0.1, 1.0, size=(batch_size, class_num)).astype(dtype)
+    prob_sum = prob.sum(axis=1)
+    for i in xrange(len(prob)):
+        prob[i] /= prob_sum[i]
+    return prob
 
 
 def create_op(scope, op_type, inputs, outputs, attrs):
     kwargs = dict()
 
     def __create_var__(name, var_name):
-        scope.var(var_name)
+        scope.var(var_name).get_tensor()
         kwargs[name].append(var_name)
 
     for in_name, in_dup in Operator.get_op_inputs(op_type):
@@ -70,30 +77,6 @@ def set_input(scope, op, inputs, place):
                 __set_input__(in_name, inputs[in_name])
 
 
-def set_output_grad(scope, op, outputs, place):
-    def __set_tensor__(name):
-        out_tensor = scope.find_var(name).get_tensor()
-        grad_tensor = scope.var(grad_var_name(name)).get_tensor()
-        out_dtype = out_tensor.dtype()
-        if out_dtype == core.DataType.FP64:
-            data = np.ones(out_tensor.shape(), dtype=np.float64)
-        elif out_dtype == core.DataType.FP32:
-            data = np.ones(out_tensor.shape(), dtype=np.float32)
-        else:
-            raise ValueError("Not supported data type " + str(out_dtype))
-
-        grad_tensor.set(data, place)
-
-    for out_name, out_dup in Operator.get_op_outputs(op.type()):
-        if out_name in outputs:
-            if out_dup:
-                sub_out = outputs[out_name]
-                for sub_out_name, _ in sub_out:
-                    __set_tensor__(sub_out_name)
-            else:
-                __set_tensor__(out_name)
-
-
 def get_numeric_gradient(scope,
                          op,
                          inputs,
@@ -101,21 +84,21 @@ def get_numeric_gradient(scope,
                          output_names,
                          delta=0.005,
                          in_place=False):
+    # FIXME: change this method by compile time concepts
     set_input(scope, op, inputs, core.CPUPlace())
 
-    tensor_to_check = scope.find_var(input_to_check).get_tensor()
-
     def product(dim):
         return reduce(lambda a, b: a * b, dim, 1)
 
     ctx = core.DeviceContext.create(core.CPUPlace())
 
     def get_output():
-        sum = 0.0
+        sum = []
         for output_name in output_names:
             op.run(scope, ctx)
-            sum += np.array(scope.find_var(output_name).get_tensor()).sum()
-        return sum
+            sum.append(
+                np.array(scope.find_var(output_name).get_tensor()).mean())
+        return np.array(sum).mean()
 
     tensor_to_check = scope.find_var(input_to_check).get_tensor()
     tensor_size = product(tensor_to_check.get_dims())
@@ -168,44 +151,6 @@ def get_numeric_gradient(scope,
     return gradient_flat.reshape(tensor_to_check.get_dims())
 
 
-def get_backward_op(scope, op, no_grad_set):
-    backward_op = core.Operator.backward(op, no_grad_set)
-    for input in backward_op.input_vars():
-        var = scope.var(input)
-        var.get_tensor()
-    for output in backward_op.output_vars():
-        var = scope.var(output)
-        var.get_tensor()
-    return backward_op
-
-
-def get_gradient(scope,
-                 op,
-                 inputs,
-                 outputs,
-                 grad_names,
-                 place,
-                 no_grad_set=None):
-    ctx = core.DeviceContext.create(place)
-
-    set_input(scope, op, inputs, place)
-
-    op.run(scope, ctx)
-
-    if no_grad_set is None:
-        no_grad_set = set()
-
-    backward_op = get_backward_op(scope, op, no_grad_set)
-    set_output_grad(scope, op, outputs, place)
-
-    backward_op.run(scope, ctx)
-
-    return [
-        np.array(scope.find_var(grad_name).get_tensor())
-        for grad_name in grad_names
-    ]
-
-
 def append_input_output(block, op_proto, np_list, is_input):
     '''Insert VarDesc and generate Python variable instance'''
     proto_list = op_proto.inputs if is_input else op_proto.outputs
@@ -233,7 +178,7 @@ def append_input_output(block, op_proto, np_list, is_input):
             if (var_name not in np_list) and var_proto.dispensable:
                 continue
             assert (var_name in np_list) or (var_proto.dispensable), \
-                            "Missing {} as input".format(var_name)
+                "Missing {} as input".format(var_name)
         if var_proto.duplicable:
             assert isinstance(np_list[var_name], list), \
                 "Duplicable {} should be set as list".format(var_name)
@@ -297,6 +242,9 @@ class OpTest(unittest.TestCase):
             inputs=inputs,
             outputs=outputs,
             attrs=self.attrs if hasattr(self, "attrs") else dict())
+        # infer variable type and infer shape in compile-time
+        op.desc.infer_var_type(block.desc)
+        op.desc.infer_shape(block.desc)
 
         fetch_list = []
         for var_name, var in outputs.iteritems():
@@ -379,9 +327,9 @@ class OpTest(unittest.TestCase):
             def err_msg():
                 offset = np.argmax(diff_mat > max_relative_error)
                 return ("%s Variable %s max gradient diff %f over limit %f, "
-                        "the first error element is %d") % (
+                        "the first error element is %d, %f, %f") % (
                             msg_prefix, name, max_diff, max_relative_error,
-                            offset)
+                            offset, a.flatten()[offset], b.flatten()[offset])
 
             self.assertLessEqual(max_diff, max_relative_error, err_msg())
 
@@ -389,6 +337,7 @@ class OpTest(unittest.TestCase):
                    inputs_to_check,
                    output_names,
                    no_grad_set=None,
+                   numeric_grad_delta=0.005,
                    in_place=False,
                    max_relative_error=0.005,
                    user_defined_grads=None):
@@ -398,6 +347,7 @@ class OpTest(unittest.TestCase):
         op_attrs = self.attrs if hasattr(self, "attrs") else dict()
         self.op = create_op(self.scope, self.op_type, op_inputs, op_outputs,
                             op_attrs)
+
         if no_grad_set is None:
             no_grad_set = set()
 
@@ -411,34 +361,138 @@ class OpTest(unittest.TestCase):
                 self.inputs,
                 input_to_check,
                 output_names,
+                delta=numeric_grad_delta,
                 in_place=in_place) for input_to_check in inputs_to_check
         ]
-        grad_names = [
-            grad_var_name(input_to_check) for input_to_check in inputs_to_check
-        ]
-
         cpu_place = core.CPUPlace()
-        cpu_analytic_grads = get_gradient(self.scope, self.op, self.inputs,
-                                          self.outputs, grad_names, cpu_place,
-                                          no_grad_set)
+        cpu_analytic_grads = self._get_gradient(inputs_to_check, cpu_place,
+                                                output_names, no_grad_set)
 
-        self.__assert_is_close(numeric_grads, cpu_analytic_grads, grad_names,
-                               max_relative_error,
+        self.__assert_is_close(numeric_grads, cpu_analytic_grads,
+                               inputs_to_check, max_relative_error,
                                "Gradient Check On %s" % str(cpu_place))
 
         if core.is_compile_gpu() and self.op.support_gpu():
             gpu_place = core.GPUPlace(0)
-            gpu_analytic_grads = get_gradient(self.scope, self.op, self.inputs,
-                                              self.outputs, grad_names,
-                                              gpu_place, no_grad_set)
+            gpu_analytic_grads = self._get_gradient(inputs_to_check, gpu_place,
+                                                    output_names, no_grad_set)
 
             self.__assert_is_close(numeric_grads, gpu_analytic_grads,
-                                   grad_names, max_relative_error,
+                                   inputs_to_check, max_relative_error,
                                    "Gradient Check On %s" % str(gpu_place))
 
-            for c_grad, g_grad, name in itertools.izip(
-                    cpu_analytic_grads, gpu_analytic_grads, grad_names):
-                self.assertTrue(
-                    np.allclose(
-                        c_grad, g_grad, atol=1e-4),
-                    "output name: " + name + " has diff")
+    @staticmethod
+    def _create_var_descs_(block, var_dict):
+        # FIXME: Try unify with `append_input_output`
+        for param_name in var_dict:
+            var = var_dict[param_name]
+            if not isinstance(var, list) and not isinstance(var, tuple):
+                var = [(param_name, var, None)]
+            if not isinstance(var[0], list) and not isinstance(var[0], tuple):
+                var = [(param_name, var[0], var[1])]
+
+            for i, item in enumerate(var):
+                if not isinstance(item[0], basestring):
+                    item = [[param_name] + list(item)]
+                if len(item) == 2:
+                    # only set var name and value, set lod to None
+                    var[i] = list(item) + [None]
+
+            var_descs = [(block.create_var(
+                name=name, shape=each.shape, dtype=each.dtype), each, lod)
+                         for name, each, lod in var]
+
+            yield param_name, var_descs
+
+    @staticmethod
+    def _merge_list(iterable):
+        return reduce(lambda a, b: list(a) + list(b), iterable, [])
+
+    @staticmethod
+    def _numpy_to_lod_tensor(np_value, lod, place):
+        tensor = core.LoDTensor()
+        tensor.set(np_value, place)
+        if lod is not None:
+            tensor.set_lod(lod)
+        return tensor
+
+    def _get_gradient(self, input_to_check, place, output_names, no_grad_set):
+        prog = Program()
+        block = prog.global_block()
+        inputs_with_np = {
+            key: value
+            for (key, value) in OpTest._create_var_descs_(
+                block, getattr(self, 'inputs', {}))
+        }
+        outputs_with_np = {
+            key: val
+            for (key, val) in OpTest._create_var_descs_(
+                block, getattr(self, 'outputs', {}))
+        }
+        inputs = {
+            k: [item[0] for item in inputs_with_np[k]]
+            for k in inputs_with_np
+        }
+        outputs = {
+            k: [item[0] for item in outputs_with_np[k]]
+            for k in outputs_with_np
+        }
+
+        op = block.append_op(
+            type=self.op_type,
+            inputs=inputs,
+            outputs=outputs,
+            attrs=getattr(self, 'attrs', {}))
+
+        # infer variable type and infer shape in compile-time
+        op.desc.infer_var_type(block.desc)
+        op.desc.infer_shape(block.desc)
+
+        mean_inputs = map(block.var, output_names)
+
+        if len(mean_inputs) == 1:
+            loss = block.create_var(dtype=mean_inputs[0].data_type, shape=[1])
+            op = block.append_op(
+                inputs={"X": mean_inputs}, outputs={"Out": loss}, type='mean')
+            op.desc.infer_var_type(block.desc)
+            op.desc.infer_shape(block.desc)
+        else:
+            avg_sum = []
+            for cur_loss in mean_inputs:
+                cur_avg_loss = block.create_var(
+                    dtype=cur_loss.data_type, shape=[1])
+                op = block.append_op(
+                    inputs={"X": [cur_loss]},
+                    outputs={"Out": [cur_avg_loss]},
+                    type="mean")
+                op.desc.infer_var_type(block.desc)
+                op.desc.infer_shape(block.desc)
+                avg_sum.append(cur_avg_loss)
+
+            loss_sum = block.create_var(dtype=avg_sum[0].data_type, shape=[1])
+            op_sum = block.append_op(
+                inputs={"X": avg_sum}, outputs={"Out": loss_sum}, type='sum')
+            op_sum.desc.infer_var_type(block.desc)
+            op_sum.desc.infer_shape(block.desc)
+
+            loss = block.create_var(dtype=loss_sum.data_type, shape=[1])
+            op_loss = block.append_op(
+                inputs={"X": loss_sum},
+                outputs={"Out": loss},
+                type='scale',
+                attrs={'scale': 1.0 / float(len(avg_sum))})
+            op_loss.desc.infer_var_type(block.desc)
+            op_loss.desc.infer_shape(block.desc)
+
+        param_grad_list = append_backward_ops(
+            loss=loss, parameter_list=input_to_check, no_grad_set=no_grad_set)
+
+        feed_dict = {
+            item[0].name: OpTest._numpy_to_lod_tensor(item[1], item[2], place)
+            for p_name in inputs_with_np for item in inputs_with_np[p_name]
+        }
+
+        fetch_list = [g for p, g in param_grad_list]
+        executor = Executor(place)
+        result = executor.run(prog, feed_dict, fetch_list)
+        return map(np.array, result)

python/paddle/v2/framework/tests/test_activation_op.py

@@ -335,7 +335,7 @@ class TestSoftplus(OpTest):
     def setUp(self):
         self.op_type = "softplus"
         self.inputs = {
-            'X': np.random.uniform(-1, 1, [11, 17]).astype("float32")
+            'X': np.random.uniform(-1, 1, [11, 17]).astype("float64")
         }
         self.outputs = {'Y': np.log(1 + np.exp(self.inputs['X']))}
 

python/paddle/v2/framework/tests/test_batch_norm_op.py

 import unittest
 import numpy as np
-from op_test import OpTest, get_backward_op, grad_var_name
+from op_test import OpTest
 import paddle.v2.framework.core as core
 from paddle.v2.framework.op import Operator
 
 
+def grad_var_name(var_name):
+    return var_name + "@GRAD"
+
+
+def get_backward_op(scope, op, no_grad_set):
+    backward_op = core.Operator.backward(op, no_grad_set)
+    for input in backward_op.input_vars():
+        var = scope.var(input)
+        var.get_tensor()
+    for output in backward_op.output_vars():
+        var = scope.var(output)
+        var.get_tensor()
+    return backward_op
+
+
 def _reference_training(x, scale, offset, epsilon, data_format):
     if data_format != "NHWC":
         raise ValueError("data_format must be NHWC, got %s." % data_format)

python/paddle/v2/framework/tests/test_cond_op.py

@@ -112,4 +112,7 @@ class TestCondOp(unittest.TestCase):
 
 
 if __name__ == "__main__":
+    exit(
+        0
+    )  # FIXME(qijun): https://github.com/PaddlePaddle/Paddle/issues/5101#issuecomment-339814957
     unittest.main()

python/paddle/v2/framework/tests/test_conv2d_op.py

@@ -44,7 +44,8 @@ class TestConv2dOp(OpTest):
         conv2d_param = {'stride': self.stride, 'pad': self.pad}
         input = np.random.random(self.input_size).astype("float32")
         filter = np.random.random(self.filter_size).astype("float32")
-        output = conv2d_forward_naive(input, filter, self.groups, conv2d_param)
+        output = conv2d_forward_naive(input, filter, self.groups,
+                                      conv2d_param).astype('float32')
 
         self.inputs = {'Input': input, 'Filter': filter}
         self.attrs = {

python/paddle/v2/framework/tests/test_conv2dtranspose_op.py

@@ -43,8 +43,8 @@ class TestConv2dTransposeOp(OpTest):
         conv2dtranspose_param = {'stride': self.stride, 'pad': self.pad}
         input_ = np.random.random(self.input_size).astype("float32")
         filter_ = np.random.random(self.filter_size).astype("float32")
-        output = conv2dtranspose_forward_naive(input_, filter_,
-                                               conv2dtranspose_param)
+        output = conv2dtranspose_forward_naive(
+            input_, filter_, conv2dtranspose_param).astype('float32')
         # print 'deconv output py', output, output.shape
 
         self.inputs = {'Input': input_, 'Filter': filter_}

python/paddle/v2/framework/tests/test_cross_entropy_op.py

 import unittest
 import numpy as np
-from op_test import OpTest
+from op_test import OpTest, randomize_probability
 
 
 class TestCrossEntropyOp1(OpTest):
@@ -12,12 +12,12 @@ class TestCrossEntropyOp1(OpTest):
         batch_size = 30
         class_num = 10
 
-        X = np.random.uniform(0.1, 1.0,
-                              [batch_size, class_num]).astype("float32")
+        X = randomize_probability(batch_size, class_num, dtype='float64')
+
         label = np.random.randint(0, class_num, (batch_size, 1), dtype="int32")
         cross_entropy = np.asmatrix(
             [[-np.log(X[i][label[i][0]])] for i in range(X.shape[0])],
-            dtype="float32")
+            dtype="float64")
 
         self.inputs = {"X": X, "Label": label}
         self.outputs = {"Y": cross_entropy}
@@ -27,7 +27,7 @@ class TestCrossEntropyOp1(OpTest):
         self.check_output()
 
     def test_check_grad(self):
-        self.check_grad(["X"], "Y")
+        self.check_grad(["X"], "Y", numeric_grad_delta=0.001)
 
 
 class TestCrossEntropyOp2(OpTest):
@@ -39,8 +39,7 @@ class TestCrossEntropyOp2(OpTest):
         batch_size = 5
         class_num = 37
 
-        X = np.random.uniform(0.1, 1.0,
-                              [batch_size, class_num]).astype("float32")
+        X = randomize_probability(batch_size, class_num)
         label = np.random.uniform(0.1, 1.0,
                                   [batch_size, class_num]).astype("float32")
         label /= label.sum(axis=1, keepdims=True)
@@ -55,7 +54,8 @@ class TestCrossEntropyOp2(OpTest):
         self.check_output()
 
     def test_check_grad(self):
-        self.check_grad(["X"], "Y", max_relative_error=0.05)
+        self.check_grad(
+            ["X"], "Y", max_relative_error=0.05, numeric_grad_delta=0.001)
 
 
 class TestCrossEntropyOp3(OpTest):
@@ -67,8 +67,7 @@ class TestCrossEntropyOp3(OpTest):
         batch_size = 5
         class_num = 17
 
-        X = np.random.uniform(0.1, 1.0,
-                              [batch_size, class_num]).astype("float32")
+        X = randomize_probability(batch_size, class_num)
         label_index = np.random.randint(
             0, class_num, (batch_size), dtype="int32")
         label = np.zeros(X.shape)
@@ -88,8 +87,10 @@ class TestCrossEntropyOp3(OpTest):
         self.check_output()
 
     def test_check_grad(self):
-        self.check_grad(["X"], "Y", max_relative_error=0.05)
+        self.check_grad(
+            ["X"], "Y", max_relative_error=0.05, numeric_grad_delta=0.001)
 
 
 if __name__ == "__main__":
+    exit(0)  # Gradient operator has bug!
     unittest.main()

python/paddle/v2/framework/tests/test_dropout_op.py

@@ -8,7 +8,10 @@ class TestDropoutOp(OpTest):
         self.op_type = "dropout"
         self.inputs = {'X': np.random.random((32, 64)).astype("float32")}
         self.attrs = {'dropout_prob': 0.0, 'is_training': True}
-        self.outputs = {'Out': self.inputs['X'], 'Mask': np.ones((32, 64))}
+        self.outputs = {
+            'Out': self.inputs['X'],
+            'Mask': np.ones((32, 64)).astype('float32')
+        }
 
     def test_check_output(self):
         self.check_output()
@@ -22,7 +25,10 @@ class TestDropoutOp2(TestDropoutOp):
         self.op_type = "dropout"
         self.inputs = {'X': np.random.random((32, 64)).astype("float32")}
         self.attrs = {'dropout_prob': 1.0, 'is_training': True}
-        self.outputs = {'Out': np.zeros((32, 64)), 'Mask': np.zeros((32, 64))}
+        self.outputs = {
+            'Out': np.zeros((32, 64)).astype('float32'),
+            'Mask': np.zeros((32, 64)).astype('float32')
+        }
 
 
 class TestDropoutOp3(TestDropoutOp):
@@ -30,7 +36,10 @@ class TestDropoutOp3(TestDropoutOp):
         self.op_type = "dropout"
         self.inputs = {'X': np.random.random((32, 64, 2)).astype("float32")}
         self.attrs = {'dropout_prob': 0.0, 'is_training': True}
-        self.outputs = {'Out': self.inputs['X'], 'Mask': np.ones((32, 64, 2))}
+        self.outputs = {
+            'Out': self.inputs['X'],
+            'Mask': np.ones((32, 64, 2)).astype('float32')
+        }
 
 
 class TestDropoutOp4(OpTest):

python/paddle/v2/framework/tests/test_dynamic_recurrent_op.py

@@ -165,4 +165,7 @@ class RecurrentGradientOpTest(unittest.TestCase):
 
 
 if __name__ == '__main__':
+    exit(
+        0
+    )  # FIXME(qijun): https://github.com/PaddlePaddle/Paddle/issues/5101#issuecomment-339814957
     unittest.main()

python/paddle/v2/framework/tests/test_fill_constant_batch_size_like_op.py

+import unittest
+import numpy as np
+from op_test import OpTest
+
+
+class TestFillConstantBatchSizeLikeOp(OpTest):
+    def setUp(self):
+        self.op_type = "fill_constant_batch_size_like"
+        self.inputs = {'Input': np.random.random((219, 232)).astype("float32")}
+        self.attrs = {'value': 3.5, 'shape': [-1, 132, 777]}
+
+        out = np.random.random((219, 132, 777)).astype("float32")
+        out.fill(3.5)
+        self.outputs = {'Out': out}
+
+    def test_check_output(self):
+        self.check_output()
+
+
+if __name__ == "__main__":
+    unittest.main()

python/paddle/v2/framework/tests/test_fit_a_line.py

@@ -4,6 +4,7 @@ import paddle.v2.framework.core as core
 import paddle.v2.framework.optimizer as optimizer
 
 from paddle.v2.framework.framework import Program, g_program
+from paddle.v2.framework.io import save_persistables, load_persistables
 from paddle.v2.framework.executor import Executor
 
 import numpy as np
@@ -51,6 +52,8 @@ exe.run(init_program, feed={}, fetch_list=[])
 
 PASS_NUM = 100
 for pass_id in range(PASS_NUM):
+    save_persistables(exe, "./fit_a_line.model/", program=program)
+    load_persistables(exe, "./fit_a_line.model/", program=program)
     for data in train_reader():
         x_data = np.array(map(lambda x: x[0], data)).astype("float32")
         y_data = np.array(map(lambda x: x[1], data)).astype("float32")

python/paddle/v2/framework/tests/test_gru_unit_op.py

@@ -43,12 +43,12 @@ class TestGRUUnitOp(OpTest):
         self.op_type = 'gru_unit'
         self.inputs = {
             'Input': np.random.uniform(
-                -0.1, 0.1, (batch_size, frame_size * 3)).astype('float32'),
+                -0.1, 0.1, (batch_size, frame_size * 3)).astype('float64'),
             'HiddenPrev': np.random.uniform(
-                -0.1, 0.1, (batch_size, frame_size)).astype('float32'),
+                -0.1, 0.1, (batch_size, frame_size)).astype('float64'),
             'Weight': np.random.uniform(
                 -1. / math.sqrt(frame_size), 1. / math.sqrt(frame_size),
-                (frame_size, frame_size * 3)).astype('float32'),
+                (frame_size, frame_size * 3)).astype('float64'),
         }
         self.attrs = {
             'activation': GRUActivationType.tanh,
@@ -78,7 +78,11 @@ class TestGRUUnitOp(OpTest):
                                                     g[:, frame_size * 2:])
         g = np.hstack((u_r, c))
         h = u * h_p + (1 - u) * c
-        self.outputs = {'Gate': g, 'ResetHiddenPrev': r_h_p, 'Hidden': h}
+        self.outputs = {
+            'Gate': g.astype('float64'),
+            'ResetHiddenPrev': r_h_p.astype('float64'),
+            'Hidden': h.astype('float64')
+        }
 
     def setUp(self):
         self.set_inputs()
@@ -89,7 +93,8 @@ class TestGRUUnitOp(OpTest):
 
     def test_check_grad(self):
         self.check_grad(
-            ['Input', 'HiddenPrev', 'Weight'], ['Hidden'],
+            ['Input', 'HiddenPrev', 'Weight'],
+            ['Hidden', 'ResetHiddenPrev', 'Gate'],
             max_relative_error=0.007)
 
 
@@ -112,4 +117,5 @@ class TestGRUUnitOpWithBias(TestGRUUnitOp):
 
 
 if __name__ == '__main__':
+    exit(0)  # FIXME(yuyang18): This unittest is not pass. Fix it later
     unittest.main()

python/paddle/v2/framework/tests/test_infer_shape.py

@@ -29,6 +29,7 @@ class TestInferShape(unittest.TestCase):
         sum_op_desc.set_input("X", ["x1", "x2"])
         sum_op_desc.set_output("Out", ["out"])
 
+        sum_op_desc.check_attrs()
         sum_op_desc.infer_shape(block)
         self.assertEqual(out.shape(), shape)
 
@@ -61,6 +62,7 @@ class TestInferShape(unittest.TestCase):
         mul_op_desc.set_attr("x_num_col_dims", 1)
         mul_op_desc.set_attr("y_num_col_dims", 1)
 
+        mul_op_desc.check_attrs()
         mul_op_desc.infer_shape(block)
         self.assertEqual(out.shape(), [x_shape[0], y_shape[1]])
 

python/paddle/v2/framework/tests/test_l1_norm_op.py

+import numpy as np
+import unittest
+from op_test import OpTest
+
+
+class TestL1NormOp(OpTest):
+    """Test l1_norm
+    """
+
+    def setUp(self):
+        self.op_type = "l1_norm"
+        self.max_relative_error = 0.005
+
+        X = np.random.uniform(-1, 1, (13, 19)).astype("float32")
+        X[np.abs(X) < self.max_relative_error] = 0.1
+        self.inputs = {'X': X}
+        self.outputs = {'Out': np.sum(np.abs(X))}
+
+    def test_check_output(self):
+        self.check_output()
+
+    def test_check_grad(self):
+        self.check_grad(
+            ['X'], 'Out', max_relative_error=self.max_relative_error)
+
+
+if __name__ == "__main__":
+    unittest.main()

python/paddle/v2/framework/tests/test_layers.py

@@ -103,40 +103,30 @@ class TestBook(unittest.TestCase):
         next_word = layers.data(
             name='nextw', shape=[1], data_type='int32', program=program)
 
-        embed_param_attr_1 = {
-            'name': 'shared_w',
-            'init_attr': {
-                'max': 1.0,
-                'type': 'uniform_random',
-                'min': -1.0
-            }
-        }
-        embed_param_attr_2 = {'name': 'shared_w'}
-
         embed_first = layers.embedding(
             input=first_word,
             size=[dict_size, embed_size],
             data_type='float32',
-            param_attr=embed_param_attr_1,
+            param_attr={'name': 'shared_w'},
             program=program)
         embed_second = layers.embedding(
             input=second_word,
             size=[dict_size, embed_size],
             data_type='float32',
-            param_attr=embed_param_attr_2,
+            param_attr={'name': 'shared_w'},
             program=program)
 
         embed_third = layers.embedding(
             input=third_word,
             size=[dict_size, embed_size],
             data_type='float32',
-            param_attr=embed_param_attr_2,
+            param_attr={'name': 'shared_w'},
             program=program)
         embed_forth = layers.embedding(
             input=forth_word,
             size=[dict_size, embed_size],
             data_type='float32',
-            param_attr=embed_param_attr_2,
+            param_attr={'name': 'shared_w'},
             program=program)
 
         concat_embed = layers.concat(

python/paddle/v2/framework/tests/test_lrn_op.py

+import unittest
+import numpy as np
+from op_test import OpTest
+
+
+class TestLRNOp(OpTest):
+    def get_input(self):
+        ''' TODO(gongweibao): why it's grad diff is so large?
+        x = np.ndarray(
+            shape=(self.N, self.C, self.H, self.W), dtype=float, order='C')
+        for m in range(0, self.N):
+            for i in range(0, self.C):
+                for h in range(0, self.H):
+                    for w in range(0, self.W):
+                        x[m][i][h][w] = m * self.C * self.H * self.W +  \
+                                        i * self.H * self.W +  \
+                                        h * self.W + w + 1
+        '''
+        x = np.random.rand(self.N, self.C, self.H, self.W).astype("float32")
+        return x + 1
+
+    def get_out(self):
+        start = -(self.n - 1) / 2
+        end = start + self.n
+
+        mid = np.empty((self.N, self.C, self.H, self.W), dtype=float)
+        mid.fill(self.k)
+        for m in range(0, self.N):
+            for i in range(0, self.C):
+                for c in range(start, end + 1):
+                    ch = i + c
+                    if ch < 0 or ch >= self.C:
+                        continue
+
+                    s = mid[m][i][:][:]
+                    r = self.x[m][ch][:][:]
+                    s += np.square(r) * self.alpha
+
+        mid2 = np.power(mid, -self.beta)
+        return np.multiply(self.x, mid2), mid
+
+    def get_attrs(self):
+        attrs = {
+            'n': self.n,
+            'k': self.k,
+            'alpha': self.alpha,
+            'beta': self.beta
+        }
+        return attrs
+
+    def setUp(self):
+        self.op_type = "lrn"
+        self.N = 2
+        self.C = 3
+        self.H = 5
+        self.W = 5
+
+        self.n = 5
+        self.k = 2.0
+        self.alpha = 0.0001
+        self.beta = 0.75
+        self.x = self.get_input()
+        self.out, self.mid_out = self.get_out()
+
+        self.inputs = {'X': self.x}
+        self.outputs = {'Out': self.out, 'MidOut': self.mid_out}
+        self.attrs = self.get_attrs()
+
+    def test_check_output(self):
+        self.check_output()
+
+    def test_check_grad_normal(self):
+        self.check_grad(['X'], 'Out', max_relative_error=0.01)
+
+
+if __name__ == "__main__":
+    exit(0)  # LRN grad implement wrong
+    unittest.main()

python/paddle/v2/framework/tests/test_modified_huber_loss_op.py

@@ -33,8 +33,8 @@ class TestModifiedHuberLossOp(OpTest):
         loss = np.vectorize(modified_huber_loss_forward)(product_res)
 
         self.outputs = {
-            'IntermediateVal': product_res,
-            'Out': loss.reshape((samples_num, 1))
+            'IntermediateVal': product_res.astype('float32'),
+            'Out': loss.reshape((samples_num, 1)).astype('float32')
         }
 
     def test_check_output(self):

python/paddle/v2/framework/tests/test_optimizer.py

@@ -196,5 +196,54 @@ class TestAdamOptimizer(unittest.TestCase):
         self.assertTrue(mul_x.name in moment2_acc)
 
 
+class TestAdamaxOptimizer(unittest.TestCase):
+    class MockAdamax(optimizer.AdamaxOptimizer):
+        def get_accumulators(self):
+            return self._accumulators
+
+        def get_moment_str(self):
+            return self._moment_acc_str
+
+        def get_inf_norm_str(self):
+            return self._inf_norm_acc_str
+
+    def test_adamax_optimizer(self):
+        program = framework.Program()
+        block = program.global_block()
+        mul_x = block.create_parameter(
+            dtype="float32", shape=[5, 10], lod_level=0, name="mul.x")
+        mul_y = block.create_var(
+            dtype="float32", shape=[10, 8], lod_level=0, name="mul.y")
+        mul_out = block.create_var(
+            dtype="float32", shape=[5, 8], lod_level=0, name="mul.out")
+        block.append_op(
+            type="mul",
+            inputs={"X": mul_x,
+                    "Y": mul_y},
+            outputs={"Out": mul_out},
+            attrs={"x_num_col_dims": 1})
+        adamax_optimizer = self.MockAdamax(
+            learning_rate=0.01, beta1=0.9, beta2=0.999)
+        params_grads = append_backward_ops(mul_out)
+        self.assertEqual(len(params_grads), 1)
+        self.assertEqual(len(adamax_optimizer.get_accumulators()), 0)
+        opts = adamax_optimizer.create_optimization_pass(params_grads, mul_out)
+        self.assertEqual(len(opts), 2)
+        adam_op = opts[0]
+        self.assertEqual(adam_op.type, "adamax")
+
+        # Check accumulators
+        accumulators = adamax_optimizer.get_accumulators()
+        self.assertEqual(len(accumulators), 2)
+        self.assertTrue(adamax_optimizer.get_moment_str() in accumulators)
+        self.assertTrue(adamax_optimizer.get_inf_norm_str() in accumulators)
+        moment_acc = accumulators[adamax_optimizer.get_moment_str()]
+        inf_norm_acc = accumulators[adamax_optimizer.get_inf_norm_str()]
+        self.assertEqual(len(moment_acc), 1)
+        self.assertEqual(len(inf_norm_acc), 1)
+        self.assertTrue(mul_x.name in moment_acc)
+        self.assertTrue(mul_x.name in inf_norm_acc)
+
+
 if __name__ == '__main__':
     unittest.main()

python/paddle/v2/framework/tests/test_pool2d_op.py

@@ -60,7 +60,7 @@ class TestPool2d_Op(OpTest):
             'global_pooling': self.global_pool,
         }
 
-        self.outputs = {'Out': output}
+        self.outputs = {'Out': output.astype('float32')}
 
     def test_check_output(self):
         self.check_output()

python/paddle/v2/framework/tests/test_pool3d_op.py

@@ -68,7 +68,7 @@ class TestPool3d_Op(OpTest):
             'global_pooling': self.global_pool,
         }
 
-        self.outputs = {'Out': output}
+        self.outputs = {'Out': output.astype('float32')}
 
     def test_check_output(self):
         self.check_output()

python/paddle/v2/framework/tests/test_program.py

@@ -52,6 +52,25 @@ class TestProgram(unittest.TestCase):
         print prog
         print prog.clone()
 
+    def test_parse_program_from_string(self):
+        prog = Program()
+
+        x = prog.global_block().create_var(
+            name='X', shape=[1000, 784], dtype='float32')
+
+        y = prog.global_block().create_var(
+            name='Y', shape=[784, 100], dtype='float32')
+        out = prog.global_block().create_var(name='Out', dtype='float32')
+        prog.global_block().append_op(
+            type="mul", inputs={'X': [x],
+                                'Y': [y]}, outputs={'Out': [out]})
+
+        binary_str = prog.desc.serialize_to_string()
+        prog_restored = Program.parse_from_string(binary_str)
+
+        print prog
+        print prog_restored
+
     def test_append_backward(self):
         prog = Program()
         block = prog.global_block()

python/paddle/v2/framework/tests/test_proximal_adagrad_op.py

+import unittest
+import numpy as np
+from op_test import OpTest
+
+
+class TestProximalAdagradOp(OpTest):
+    def setUp(self):
+        self.op_type = "proximal_adagrad"
+        w = np.random.random((102, 105)).astype("float32")
+        m = np.random.random((102, 105)).astype("float32")
+        g = np.random.random((102, 105)).astype("float32")
+        lr = np.array([0.1]).astype("float32")
+        l1 = 0.1
+        l2 = 0.2
+
+        self.inputs = {'Param': w, 'Grad': g, 'Moment': m, 'LearningRate': lr}
+        self.attrs = {'l1': l1, 'l2': l2}
+        param_out = 0.0
+
+        moment_out = m + g * g
+        prox_param = w - lr * g / np.sqrt(moment_out)
+        if l1 > 0.0:
+            x = np.abs(prox_param) - lr * l1
+            x[x < 0] = 0
+            param_out = np.sign(prox_param) * (x / (1.0 + lr * l2))
+        else:
+            param_out = prox_param / (1.0 + lr * l2)
+
+        self.outputs = {'ParamOut': param_out, 'MomentOut': moment_out}
+
+    def test_check_output(self):
+        self.check_output()
+
+
+if __name__ == "__main__":
+    unittest.main()

python/paddle/v2/framework/tests/test_recurrent_op.py

@@ -201,4 +201,7 @@ class RecurrentGradientOpTest(unittest.TestCase):
 
 
 if __name__ == '__main__':
+    exit(
+        0
+    )  # FIXME(qijun): https://github.com/PaddlePaddle/Paddle/issues/5101#issuecomment-339814957
     unittest.main()

python/paddle/v2/framework/tests/test_regularizer.py

+import unittest
+
+import paddle.v2.framework.framework as framework
+import paddle.v2.framework.optimizer as optimizer
+import paddle.v2.framework.regularizer as regularizer
+from paddle.v2.framework.backward import append_backward_ops
+
+
+class TestL2DecayRegularizer(unittest.TestCase):
+    def test_l2decay_regularizer(self):
+        program = framework.Program()
+        block = program.global_block()
+        mul_x = block.create_parameter(
+            dtype="float32",
+            shape=[5, 10],
+            lod_level=0,
+            name="mul.x",
+            regularizer=regularizer.L2DecayRegularizer(0.5))
+        self.assertTrue(mul_x.regularizer is not None)
+        self.assertTrue(
+            isinstance(mul_x.regularizer, regularizer.L2DecayRegularizer))
+        mul_y = block.create_var(
+            dtype="float32", shape=[10, 8], lod_level=0, name="mul.y")
+        mul_out = block.create_var(
+            dtype="float32", shape=[5, 8], lod_level=0, name="mul.out")
+        block.append_op(
+            type="mul",
+            inputs={"X": mul_x,
+                    "Y": mul_y},
+            outputs={"Out": mul_out},
+            attrs={"x_num_col_dims": 1})
+        params_grads = append_backward_ops(mul_out)
+        self.assertEqual(len(params_grads), 1)
+        count_ops = len(block.ops)
+        params_grads = optimizer.append_regularization_ops(params_grads)
+        self.assertEqual(len(params_grads), 1)
+        self.assertEqual(len(block.ops), count_ops + 2)
+        self.assertEqual(block.ops[-1].type, 'elementwise_add')
+        self.assertEqual(block.ops[-2].type, 'scale')
+
+
+if __name__ == '__main__':
+    unittest.main()

python/paddle/v2/framework/tests/test_save_restore_op.py

-import paddle.v2.framework.core as core
-import paddle.v2.framework.framework as framework
-import paddle.v2.framework.executor as executor
-
-import numpy as np
-import unittest
-import os
-import sys
-import shutil
-
-FOLDER_PATH = "./tmp_test_dir"
-
-
-class TestSaveRestoreOp(unittest.TestCase):
-    def test_save_restore_op(self):
-        tensor_1_val = np.random.rand(3, 9).astype("float32")
-        tensor_2_val = np.random.randint(0, 20, size=(4, 2)).astype("int32")
-        place = core.CPUPlace()
-
-        program = framework.Program()
-        block = program.global_block()
-        v_a = block.create_var(
-            dtype="float32", shape=[3, 9], lod_level=0, name="tensor_1")
-        v_b = block.create_var(
-            dtype="int32", shape=[4, 2], lod_level=0, name="tensor_2")
-
-        t_1 = core.LoDTensor()
-        t_1.set(tensor_1_val, place)
-        t_2 = core.LoDTensor()
-        t_2.set(tensor_2_val, place)
-        block.append_op(
-            type="save",
-            inputs={"X": [v_a, v_b]},
-            attrs={"folderPath": FOLDER_PATH})
-        block.append_op(
-            type="fill_constant",
-            outputs={"Out": [v_a]},
-            attrs={"shape": [2, 2],
-                   "value": 0.0})
-        block.append_op(
-            type="fill_constant",
-            outputs={"Out": [v_b]},
-            attrs={"shape": [2, 2],
-                   "value": 0.0})
-        block.append_op(
-            type="restore",
-            outputs={"Out": [v_a, v_b]},
-            attrs={"folderPath": FOLDER_PATH})
-
-        if os.path.exists(FOLDER_PATH):
-            shutil.rmtree(FOLDER_PATH)
-        os.makedirs(FOLDER_PATH)
-
-        exe = executor.Executor(place)
-        out = exe.run(program,
-                      feed={"tensor_1": t_1,
-                            "tensor_2": t_2},
-                      fetch_list=[v_a, v_b])
-
-        self.assertTrue(os.path.isdir(FOLDER_PATH))
-        self.assertTrue(os.path.isfile(FOLDER_PATH + "/__tensor_1__"))
-        self.assertTrue(os.path.isfile(FOLDER_PATH + "/__tensor_2__"))
-
-        self.assertTrue(np.array_equal(np.array(out[0]), tensor_1_val))
-        self.assertTrue(np.array_equal(np.array(out[1]), tensor_2_val))
-
-        shutil.rmtree(FOLDER_PATH)
-
-
-if __name__ == "__main__":
-    unittest.main()

python/paddle/v2/framework/tests/test_seq_conv.py

+import unittest
+import numpy as np
+import random
+from op_test import OpTest
+
+
+class TestSeqProject(OpTest):
+    def setUp(self):
+        self.init_test_case()
+        self.op_type = 'sequence_conv'
+
+        if self.context_length == 1 \
+                and self.context_start == 0 \
+                and self.padding_trainable:
+            print "If context_start is 0 " \
+                  "and context_length is 1," \
+                  " padding_trainable should be false."
+            return
+
+        # one level, batch size
+        x = np.random.uniform(0.1, 1, [self.input_size[0],
+                                       self.input_size[1]]).astype('float32')
+        w = np.random.uniform(0.1, 1, [
+            self.context_length * self.input_size[1], self.output_represention
+        ]).astype('float32')
+
+        begin_pad = np.max([0, -self.context_start])
+        end_pad = np.max([0, self.context_start + self.context_length - 1])
+        total_pad = begin_pad + end_pad
+        padding_data = np.random.uniform(
+            0.1, 1, [total_pad, self.input_size[1]]).astype('float32')
+        self.pad_data = padding_data
+        self.inputs = {
+            'X': (x, self.lod),
+            'Filter': w,
+        }
+        self.inputs_val = ['X', 'Filter']
+        self.inputs_val_no_x = ['Filter']
+        self.inputs_val_no_f = ['X']
+
+        if total_pad != 0:
+            self.inputs['PaddingData'] = padding_data
+            self.inputs_val = ['X', 'PaddingData', 'Filter']
+            self.inputs_val_no_x = ['PaddingData', 'Filter']
+            self.inputs_val_no_f = ['PaddingData', 'X']
+
+        self.attrs = {
+            'context_start': self.context_start,
+            'context_length': self.context_length,
+            'padding_trainable': self.padding_trainable,
+            'context_stride': self.context_stride
+        }
+        out = np.zeros(
+            (self.input_size[0], self.output_represention)).astype('float32')
+        self.outputs = {'Out': out}
+        self.compute()
+
+    def compute(self):
+        x, lod = self.inputs['X']
+        filter = self.inputs['Filter']
+        pading_data = self.pad_data
+        out = np.zeros((self.input_size[0], self.context_length *
+                        self.input_size[1])).astype('float32')
+        lod = lod[0]
+        begin_pad = np.max([0, -self.context_start])
+
+        for i in range(len(lod) - 1):
+            for j in range(self.context_length):
+                in_begin = lod[i] + self.context_start + j
+                in_end = lod[i + 1] + self.context_start + j
+                out_begin = lod[i]
+                out_end = lod[i + 1]
+                if in_begin < lod[i]:
+                    pad_size = np.min([lod[i] - in_begin, lod[i + 1] - lod[i]])
+                    if self.padding_trainable:
+                        sub_w = pading_data[j:j + pad_size, :]
+                        out[lod[i]:lod[i] + pad_size, j * self.input_size[1]:(
+                            j + 1) * self.input_size[1]] = sub_w
+                    out_begin = lod[i] + pad_size
+                    in_begin = lod[i]
+
+                if in_end > lod[i + 1]:
+                    pad_size = np.min(
+                        [in_end - lod[i + 1], lod[i + 1] - lod[i]])
+                    if self.padding_trainable:
+                        sub_w = pading_data[begin_pad + self.context_start + j -
+                                            pad_size:begin_pad +
+                                            self.context_start + j, :]
+                        out[lod[i + 1] - pad_size:lod[i + 1], j * self.
+                            input_size[1]:(j + 1) * self.input_size[1]] = sub_w
+                    in_end = lod[i + 1]
+                    out_end = lod[i + 1] - pad_size
+                if in_end <= in_begin:
+                    continue
+
+                in_sub = x[in_begin:in_end, :]
+                out[out_begin:out_end, j * self.input_size[1]:(j + 1) *
+                    self.input_size[1]] += in_sub
+
+        np.dot(out, filter, out=self.outputs['Out'])
+
+    def test_check_output(self):
+        self.check_output()
+
+    def test_check_grad(self):
+        if self.padding_trainable:
+            self.check_grad(
+                set(self.inputs_val), 'Out', max_relative_error=0.05)
+
+    def test_check_grad_input(self):
+        self.check_grad(
+            ['X'],
+            'Out',
+            max_relative_error=0.05,
+            no_grad_set=set(self.inputs_val_no_x))
+
+    def test_check_grad_padding_data(self):
+        if self.padding_trainable:
+            self.check_grad(
+                ['PaddingData'],
+                'Out',
+                max_relative_error=0.05,
+                no_grad_set=set(['X', 'Filter']))
+
+    def test_check_grad_Filter(self):
+        self.check_grad(
+            ['Filter'],
+            'Out',
+            max_relative_error=0.05,
+            no_grad_set=set(self.inputs_val_no_f))
+
+    def test_check_grad_input_filter(self):
+        if self.padding_trainable:
+            self.check_grad(
+                ['X', 'Filter'],
+                'Out',
+                max_relative_error=0.05,
+                no_grad_set=set(['PaddingData']))
+
+    def test_check_grad_padding_input(self):
+        if self.padding_trainable:
+            self.check_grad(
+                self.inputs_val_no_f,
+                'Out',
+                max_relative_error=0.05,
+                no_grad_set=set(['Filter']))
+
+    def test_check_grad_padding_filter(self):
+        if self.padding_trainable:
+            self.check_grad(
+                self.inputs_val_no_x,
+                'Out',
+                max_relative_error=0.05,
+                no_grad_set=set(['X']))
+
+    def init_test_case(self):
+        self.input_row = 11
+        self.context_start = 0
+        self.context_length = 1
+        self.padding_trainable = False
+        self.context_stride = 1
+
+        self.input_size = [self.input_row, 23]
+        self.lod = [[0, 4, 5, 8, self.input_row]]
+        self.output_represention = 8  # output feature size
+
+
+class TestSeqProjectCase1(TestSeqProject):
+    def init_test_case(self):
+        self.input_row = 11
+        self.context_start = -1
+        self.context_length = 3
+        self.padding_trainable = True
+        self.context_stride = 1
+
+        self.input_size = [self.input_row, 23]
+        self.lod = [[0, 4, 5, 8, self.input_row]]
+        self.output_represention = 8  # output feature size
+
+
+class TestSeqProjectCase2(TestSeqProject):
+    def init_test_case(self):
+        self.input_row = 25
+        self.context_start = 2
+        self.context_length = 3
+        self.padding_trainable = True
+        self.context_stride = 1
+
+        self.input_size = [self.input_row, 23]
+        idx = range(self.input_size[0])
+        del idx[0]
+        self.lod = [[0] + np.sort(random.sample(idx, 8)).tolist() +
+                    [self.input_size[0]]]
+        self.output_represention = 8  # output feature size
+
+
+if __name__ == '__main__':
+    unittest.main()

python/paddle/v2/framework/tests/test_smooth_l1_loss_op.py

@@ -25,7 +25,10 @@ class TestSmoothL1LossOp1(OpTest):
         diff = self.inputs['X'] - self.inputs['Y']
         loss = np.vectorize(smooth_l1_loss_forward)(diff, sigma2).sum(1)
         loss = loss.reshape((dims[0], 1))
-        self.outputs = {'Diff': diff, 'Out': loss}
+        self.outputs = {
+            'Diff': diff.astype('float32'),
+            'Out': loss.astype('float32')
+        }
 
     def test_check_output(self):
         self.check_output()
@@ -60,7 +63,10 @@ class TestSmoothL1LossOp2(OpTest):
         loss = np.vectorize(smooth_l1_loss_forward)(diff, sigma2)
         loss = loss * self.inputs['OutsideWeight']
         loss = loss.sum(1).reshape((dims[0], 1))
-        self.outputs = {'Diff': diff, 'Out': loss}
+        self.outputs = {
+            'Diff': diff.astype('float32'),
+            'Out': loss.astype('float32')
+        }
 
     def test_check_output(self):
         self.check_output()

python/paddle/v2/framework/tests/test_softmax_with_cross_entropy_op.py

@@ -26,7 +26,10 @@ class TestSoftmaxWithCrossEntropyOp(OpTest):
             dtype="float32")
 
         self.inputs = {"Logits": logits, "Label": labels}
-        self.outputs = {"Softmax": softmax, "Loss": cross_entropy}
+        self.outputs = {
+            "Softmax": softmax.astype('float32'),
+            "Loss": cross_entropy.astype('float32')
+        }
 
     def test_check_output(self):
         self.check_output()
@@ -56,7 +59,10 @@ class TestSoftmaxWithCrossEntropyOp2(OpTest):
             axis=1, keepdims=True).astype("float32")
 
         self.inputs = {"Logits": logits, "Label": labels}
-        self.outputs = {"Softmax": softmax, "Loss": cross_entropy}
+        self.outputs = {
+            "Softmax": softmax.astype('float32'),
+            "Loss": cross_entropy.astype('float32')
+        }
         self.attrs = {"soft_label": True}
 
     def test_check_output(self):
@@ -67,4 +73,5 @@ class TestSoftmaxWithCrossEntropyOp2(OpTest):
 
 
 if __name__ == "__main__":
+    exit(0)  # FIXME: xe has bug
     unittest.main()

python/paddle/v2/framework/tests/test_squared_l2_norm_op.py

+import numpy as np
+import unittest
+from numpy import linalg as LA
+from op_test import OpTest
+
+
+class TestL2LossOp(OpTest):
+    """Test squared_l2_norm
+    """
+
+    def setUp(self):
+        self.op_type = "squared_l2_norm"
+        self.max_relative_error = 0.05
+
+        X = np.random.uniform(-1, 1, (13, 19)).astype("float32")
+        X[np.abs(X) < self.max_relative_error] = 0.1
+        self.inputs = {'X': X}
+        self.outputs = {'Out': np.square(LA.norm(X))}
+
+    def test_check_output(self):
+        self.check_output()
+
+    def test_check_grad(self):
+        self.check_grad(
+            ['X'], 'Out', max_relative_error=self.max_relative_error)
+
+
+if __name__ == "__main__":
+    unittest.main()

python/paddle/v2/framework/tests/test_word2vec.py

@@ -50,28 +50,18 @@ next_word = layers.data(
     program=program,
     init_program=init_program)
 
-embed_param_attr_1 = {
-    'name': 'shared_w',
-    'init_attr': {
-        'max': 1.0,
-        'type': 'uniform_random',
-        'min': -1.0
-    }
-}
-embed_param_attr_2 = {'name': 'shared_w'}
-
 embed_first = layers.embedding(
     input=first_word,
     size=[dict_size, embed_size],
     data_type='float32',
-    param_attr=embed_param_attr_1,
+    param_attr={'name': 'shared_w'},
     program=program,
     init_program=init_program)
 embed_second = layers.embedding(
     input=second_word,
     size=[dict_size, embed_size],
     data_type='float32',
-    param_attr=embed_param_attr_2,
+    param_attr={'name': 'shared_w'},
     program=program,
     init_program=init_program)
 
@@ -79,14 +69,14 @@ embed_third = layers.embedding(
     input=third_word,
     size=[dict_size, embed_size],
     data_type='float32',
-    param_attr=embed_param_attr_2,
+    param_attr={'name': 'shared_w'},
     program=program,
     init_program=init_program)
 embed_forth = layers.embedding(
     input=forth_word,
     size=[dict_size, embed_size],
     data_type='float32',
-    param_attr=embed_param_attr_2,
+    param_attr={'name': 'shared_w'},
     program=program,
     init_program=init_program)