Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into add_FLAGS_use_deterministic_algo

benchmark/fluid/machine_translation.py

@@ -21,7 +21,7 @@ import argparse
 import time
 import distutils.util
 
-import paddle.v2 as paddle
+import paddle
 import paddle.fluid as fluid
 import paddle.fluid.core as core
 import paddle.fluid.framework as framework

benchmark/fluid/mnist.py

@@ -20,7 +20,7 @@ import numpy as np
 import argparse
 import time
 
-import paddle.v2 as paddle
+import paddle
 import paddle.fluid as fluid
 import paddle.fluid.profiler as profiler
 

benchmark/fluid/resnet.py

@@ -23,7 +23,7 @@ import time
 
 import cProfile, pstats, StringIO
 
-import paddle.v2 as paddle
+import paddle
 import paddle.fluid as fluid
 import paddle.fluid.core as core
 import paddle.fluid.profiler as profiler

benchmark/fluid/stacked_dynamic_lstm.py

@@ -23,10 +23,10 @@ import random
 import time
 
 import numpy
-import paddle.v2 as paddle
+import paddle
-import paddle.v2.dataset.imdb as imdb
+import paddle.dataset.imdb as imdb
 import paddle.fluid as fluid
-from paddle.v2 import batch
+import paddle.batch as batch
 import paddle.fluid.profiler as profiler
 
 

benchmark/fluid/vgg.py

@@ -17,7 +17,7 @@ from __future__ import print_function
 import sys
 import time
 import numpy as np
-import paddle.v2 as paddle
+import paddle
 import paddle.fluid as fluid
 import paddle.fluid.core as core
 import argparse

doc/fluid/api/data/data_reader.rst

@@ -56,11 +56,11 @@ DataFeeder
 Reader
 ======
 
-..  automodule:: paddle.v2.reader
+..  automodule:: paddle.reader
     :members:
     :noindex:
 
-..  automodule:: paddle.v2.reader.creator
+..  automodule:: paddle.reader.creator
     :members:
     :noindex:
 

doc/fluid/api/layers.rst

@@ -479,6 +479,13 @@ label_smooth
 ..  autofunction:: paddle.fluid.layers.label_smooth
     :noindex:
 
+roi_pool
+---------
+
+..  autofunction:: paddle.fluid.layers.roi_pool
+    :noindex:
+
+    
 ops
 ===
 
@@ -820,3 +827,5 @@ topk
 
 ..  autofunction:: paddle.fluid.layers.topk
     :noindex:
+
+

doc/fluid/design/data_type/float16.md

@@ -3,7 +3,7 @@
 ## Why float16
 Half precision (float16) is a binary floating-point format that occupies 16 bits in memory. float16 is half the size of traditional 32-bit single precision format (float) and has lower precision and smaller range. 
 
-When high precision computation is not required, using float16 data type could potentially 
+When high precision computation is not required (which is usually the case at least in the deep learning inference stage), using float16 data type could potentially 
 
 - reduce storage space, memory bandwidth, and power usages; 
 - increase the chance of data fitting into a smaller cache of lower latency; 
@@ -12,7 +12,7 @@ When high precision computation is not required, using float16 data type could p
 ## Survey of current float16 support
 A brief survey of float16 support on different compilers, hardwares, and libraries can be found below. Interested readers can refer to [link1](https://github.com/PaddlePaddle/Paddle/issues/4853) and [link2](https://github.com/Xreki/Xreki.github.io/blob/master/multi_data_types_in_dl_framework/ppt/float16_and_quantized_type.md) for more info.
 
-The goal of float16 is to serve as a key for the executor to find and run the correct version of compute method specialized for float16 in operator kernel. It should be compatible with various natively supported float16 implementations including `__half` for cuda, `float16_t` for ARM, and `Eigen::half` for Eigen to make writing customized float16 kernels easier. 
+The goal of float16 is to serve as a key for the executor to find and run the correct version of compute method specialized for float16 in operator kernels. It should be compatible with various natively supported float16 implementations including `__half` for cuda, `float16_t` for ARM, and `Eigen::half` for Eigen to make writing customized float16 kernels easier. 
 
 ### Compiler
 - nvcc supports `__half` data type after CUDA 7.5.
@@ -95,11 +95,89 @@ float half_to_float(float16 h);
 ```
 which provides one-to-one conversion between float32 and float16. These twos functions will do different conversion routines based on the current hardware. CUDA/ARM instrinsics will be used when the corresonding hardware is available. If the hardware or compiler level does not support float32 to float16 conversion, software emulation will be performed to do the conversion.
 
-## To do
+## float16 inference
-After float16 class is available, some of the future items are below:
+In Fluid, a neural network is represented as a protobuf message called [ProgramDesc](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/fluid/design/concepts/program.md), whose Python wrapper is a [Program](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/fluid/design/modules/python_api.md#program). The basic structure of a program is some nested [blocks](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/fluid/design/modules/python_api.md#block), where each block consists of some [variable](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/fluid/design/modules/python_api.md#variable) definitions and a sequence of [operators](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/fluid/design/modules/python_api.md#operator). An [executor](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/fluid/design/concepts/executor.md) will run a given program desc by executing the sequence of operators in the entrance block of the program one by one.  
 
-- Update pybind/tensor_py.h to bind c++ float16 with numpy float16. 
+### Operator level requirement
+Each operator has many kernels for different data types, devices, and library types. The operator will select the appropriate kernel to run based on, among other things, the data type of the input variables. By default, every Fluid operator has a float data type kernel that takes float variables as input and generates float output. 
 
-- Modify `GetKernelType()` method in `framework/operator.h` to make it compatible with float16.
+This means that if we provide float input to the first operator in a program, then each opeartor will use float kernel to compute float output and send it as input to the next operator to trigger the float kernel. Overall, the program will run in float mode and give us a final output of float data type.
 
-- Create a type-casting operator that can convert the data type in tensor between float16 and other types.
+The same principle applies if we want a program to run in float16 mode. We provide input variable of float16 data type to the first operator, and then one by one, each operator in the program will run the float16 kernel (provided that each operator in this program has float16 kernels registered) until we finally obtain a float16 output variable.
+
+So the preliminary requirement for float16 inference is to add float16 kernel to operators that are needed in a specific kind of program. For example, float16 inference on an image classification neural network like Vgg or Resnet, typically requires the following operators to have float16 kernels: convolution, pooling, multiplication, addition, batch norm, dropout, relu, and softmax. Please refer to [new_op_en](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/fluid/dev/new_op_en.md) for details of how to add new kernels to an operator.
+
+### Variable level requirement
+Operators including convolution and multiplication (used in fully-connected layers) takes as input not only the variables generated by the preceding operators but also [parameter](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/fluid/design/modules/python_api.md#parameter) variables, which contains the trained weights to apply to the input data. These weights are obtained in the Fluid training process and are by default of float data type.
+
+When these operators are running in float16 mode, the float16 kernel requires those parameter variables to contain weights of Fluid float16 data type. Thus, we need a convenient way to convert the original float weights to float16 weights. 
+
+In Fluid, we use tensor to hold actual data for a variable on the c++ end. [Pybind](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/fluid/pybind/tensor_py.h) is used to bind c++ tensors of certain data type with numpy array of the correponding numpy data type on the Python end. Each common c++ built-in data type has a corresponding numpy data type of the same name. However, since there is no built-in float16 type in c++, we cannot directly bind numpy float16 data type with the Fluid float16 class. Since both Fluid float16 and numpy float16 use uint16 as the internal data storage type, we use c++ built-in type `uint16_t` and the corresponding numpy uint16 data type to bridge the gap via [Pybind](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/fluid/pybind/tensor_py.h). 
+
+The following code demonstrates how to do the tensor conversion.
+```Python
+# var is the variable of float weights
+# tensor is a numpy array of data copied from the tensor data in var 
+# fp16_var is the variable that will contain float16 weights converted from var  
+tensor = numpy.array(var.get_tensor())
+fp16_tensor = fp16_var.get_tensor()
+
+# After the original tensor data is converted to numpy float16 data type, 
+# view(numpy.uint16) is used so that the internal memory of the numpy array 
+# will be reinterpreted to be of uint16 data type, which is binded to 
+# Fluid float16 class via pybind with the help of uint16_t built-in c++ type
+fp16_tensor.set(tensor.astype(numpy.float16).view(numpy.uint16), GPUPlace)  
+```
+
+### Consistent API requirement
+The basic inference in float16 mode requires users to feed input and obtain output both of float16 data type. However, in this way, the inference APIs are not consistent between float16 mode and float mode, and users may find it confusing and diffcult to use float16 inference since they need to do extra steps to provide float16 input data and convert float16 output data back to float. To have consistent API for different inference modes, we need to transpile the program desc in some way so that we can run float16 inference by feeding and fetching variables of float data type.
+
+This problem can be solved by introducing a type-casting operator which takes an input variable of certain data type, cast it to another specified data type, and put the casted data into the output variable. Insert cast operator where needed can make a program internally run in float16 mode.   
+
+### float16 transpiler
+Put all the above requirements in mind, we designed a float16 inference transpiler that can tranpile a float32 mode inference program desc to a float16 mode one.
+
+Given a float inference program and the corresponding variables of float32 weights in the [scope](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/fluid/design/concepts/scope.md),
+this transpiler mainly does the following modifications:
+
+1. Insert cast operators at the beginning of the program so that the input float data will be converted to float16 data type before feeding to subsequent operators to invoke the float16 kernel. 
+
+2. Insert cast operators at the end of the program so that the output float16 data will be converted back to float data type before users obtain the result.
+
+3. For each parameter variable of float weights, create in the scope a corresponding variable of float16 weights which are converted from the corresponding float weights and add this new float16 variable to the program.
+
+4. Update the operator information in the program so that each relevant operator use the newly created float16 variable instead of its float counterpart.
+
+Below is an example of usage:
+```Python
+# Get the float inference program
+[float_inference_program, feed_target_names,
+ fetch_targets] = fluid.io.load_inference_model(save_dirname, exe)
+
+# Prepare the float input data
+tensor_img = numpy.random.rand(1, 3, 32, 32).astype(numpy.float32)
+
+# Running inference_program in float mode
+float_results = exe.run(float_inference_program,
+                        feed={feed_target_names[0]: tensor_img},
+                        fetch_list=fetch_targets)
+
+# Use float16 transpiler to speedup
+float16_inference_program = float_inference_program.clone()
+t = fluid.InferenceTranspiler()
+t.float16_transpile(float16_inference_program, GPUPlace)
+
+# Running 
+float16_results = exe.run(float16_inference_program,
+                          feed={feed_target_names[0]: tensor_img},
+                          fetch_list=fetch_targets)
+```
+
+As we can see from the example above, users can simply use the `float16_transpile` method provided by the infernece transpiler class on an existing float inference program to run inference in float16 mode.
+
+### Speedup on GPU
+Currently, Fluid inference in float16 mode is only supported on Nvidia GPU device. There is no motivation to support float16 inference on non-ARM CPUs because float16 is not natively supported there and float16 calculation will only be slower than its float counterpart. 
+
+Nvidia started to support its native float16 data type (which has the same internal memory representation as Fluid float16 class) on CUDA 7.5. Moreover, float16 speedups on common computational intensive tasks including GEMM (general matrix-matrix multiplication) and convolution are supported since cublas 7.5 and cuDNN 5.0.
+
+Recently, the introduction of [tensor core](https://devblogs.nvidia.com/programming-tensor-cores-cuda-9/) in volta architecture GPUs and the support of tensor core calculation in CUDA 9.0 and cuDNN 7.0 make float16 truly superior to float in certain deep learning applications. Please refer to this [benchmark report](https://github.com/kexinzhao/Paddle_benchmark/blob/master/float16_benchmark.md) for more details.

doc/v2/api/data/data_reader.rst

@@ -56,11 +56,11 @@ DataFeeder
 Reader
 ======
 
-..  automodule:: paddle.v2.reader
+..  automodule:: paddle.reader
     :members:
     :noindex:
 
-..  automodule:: paddle.v2.reader.creator
+..  automodule:: paddle.reader.creator
     :members:
     :noindex:
 

doc/v2/api/data/dataset.rst

@@ -1,82 +1,82 @@
 Dataset
 =======
 
-..  automodule:: paddle.v2.dataset
+..  automodule:: paddle.dataset
     :members:
     :noindex:
 
 mnist
 +++++
 
-..  automodule:: paddle.v2.dataset.mnist
+..  automodule:: paddle.dataset.mnist
     :members:
     :noindex:
 
 cifar
 +++++
 
-..  automodule:: paddle.v2.dataset.cifar
+..  automodule:: paddle.dataset.cifar
     :members:
     :noindex:
 
 conll05
 +++++++
 
-..  automodule:: paddle.v2.dataset.conll05
+..  automodule:: paddle.dataset.conll05
     :members: get_dict,get_embedding,test
     :noindex:
 
 imdb
 ++++
 
-..  automodule:: paddle.v2.dataset.imdb
+..  automodule:: paddle.dataset.imdb
     :members:
     :noindex:
 
 imikolov
 ++++++++
 
-..  automodule:: paddle.v2.dataset.imikolov
+..  automodule:: paddle.dataset.imikolov
     :members:
     :noindex:
 
 movielens
 +++++++++
 
-..  automodule:: paddle.v2.dataset.movielens
+..  automodule:: paddle.dataset.movielens
     :members:
     :noindex:
 
-..  autoclass:: paddle.v2.dataset.movielens.MovieInfo
+..  autoclass:: paddle.dataset.movielens.MovieInfo
     :noindex:
-    
+
-..  autoclass:: paddle.v2.dataset.movielens.UserInfo
+..  autoclass:: paddle.dataset.movielens.UserInfo
     :noindex:
 
 sentiment
 +++++++++
 
-..  automodule:: paddle.v2.dataset.sentiment
+..  automodule:: paddle.dataset.sentiment
     :members:
     :noindex:
 
 uci_housing
 +++++++++++
 
-..  automodule:: paddle.v2.dataset.uci_housing
+..  automodule:: paddle.dataset.uci_housing
     :members:
     :noindex:
 
 wmt14
 +++++
 
-..  automodule:: paddle.v2.dataset.wmt14
+..  automodule:: paddle.dataset.wmt14
     :members:
     :noindex:
 
 wmt16
 +++++
 
-..  automodule:: paddle.v2.dataset.wmt16
+..  automodule:: paddle.dataset.wmt16
     :members:
     :noindex:

doc/v2/howto/cluster/multi_cluster/index_en.rst

@@ -1,19 +1,35 @@
 Use different clusters
 ======================
 
-PaddlePaddle supports running jobs on several platforms including:
+The user's cluster environment is not the same. To facilitate everyone's deployment, we provide a variety of cluster deployment methods to facilitate the submission of cluster training tasks, which will be introduced as follows:
-- `Kubernetes <http://kubernetes.io>`_ open-source system for automating deployment, scaling, and management of containerized applications from Google.
-- `OpenMPI <https://www.open-mpi.org>`_ Mature high performance parallel computing framework.
-- `Fabric <http://www.fabfile.org>`_ A cluster management tool. Write scripts to submit jobs or manage the cluster.
 
-We'll introduce cluster job management on these platforms. The examples can be found under `cluster_train_v2 <https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/scripts/cluster_train_v2>`_ .
+`Kubernetes <http://kubernetes.io>`_ is a scheduling framework of Google open source container cluster, supporting a complete cluster solution for large-scale cluster production environment. The following guidelines show PaddlePaddle's support for Kubernetes:
 
-These cluster platforms provide API or environment variables for training processes, when the job is dispatched to different nodes. Like node ID, IP or total number of nodes etc.
+..  toctree::
+  :maxdepth: 1
+
+  k8s_cn.md
+  k8s_distributed_cn.md
+
+`OpenMPI <https://www.open-mpi.org>`_ is a mature high-performance parallel computing framework, which is widely used in the field of HPC. The following guide describes how to use OpenMPI to build PaddlePaddle's cluster training task:
 
 ..  toctree::
   :maxdepth: 1
 
-  fabric_en.md
+  openmpi_cn.md
-  openmpi_en.md
+
-  k8s_en.md
+`Fabric <http://www.fabfile.org>`_ is a convenient tool for program deployment and management. We provide a way to deploy and manage with Fabric. If you want to know more about it, please read the following guidelines:
-  k8s_aws_en.md
+
+..  toctree::
+  :maxdepth: 1
+
+  fabric_cn.md
+
+We also support the deployment of PaddlePaddle on AWS. Learn more about:
+
+..  toctree::
+  :maxdepth: 1
+
+  k8s_aws_cn.md
+
+The examples can be found under `cluster_train_v2 <https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/scripts/cluster_train_v2>`_ .

paddle/fluid/framework/details/broadcast_op_handle_test.cc

@@ -139,7 +139,7 @@ struct TestBroadcastOpHandle {
       PADDLE_ENFORCE_EQ(out_tensor.lod(), lod, "lod is not equal.");
 
       f::Tensor result_tensor;
-      f::TensorCopy(out_tensor, cpu_place, *(ctxs_[j]), &result_tensor);
+      f::TensorCopySync(out_tensor, cpu_place, &result_tensor);
       float* ct = result_tensor.mutable_data<float>(cpu_place);
 
       for (int64_t i = 0; i < f::product(kDims); ++i) {
@@ -185,7 +185,7 @@ struct TestBroadcastOpHandle {
       }
 
       f::Tensor result_tensor;
-      f::TensorCopy(rt, cpu_place, *(ctxs_[j]), &result_tensor);
+      f::TensorCopySync(rt, cpu_place, &result_tensor);
       float* ct = result_tensor.data<float>();
 
       for (int64_t i = 0; i < f::product(kDims); ++i) {

paddle/fluid/framework/details/fetch_op_handle.cc

@@ -66,8 +66,7 @@ void FetchOpHandle::RunImpl() {
     auto &t = var->Get<framework::LoDTensor>();
     if (platform::is_gpu_place(t.place())) {
 #ifdef PADDLE_WITH_CUDA
-      TensorCopy(t, cpu, *dev_ctxes_[t.place()], &tensors_[i], true);
+      TensorCopySync(t, cpu, &tensors_[i]);
-      dev_ctxes_.at(t.place())->Wait();
 #endif
     } else {
       tensors_[i].ShareDataWith(t);

paddle/fluid/framework/details/multi_devices_graph_builder.cc

@@ -58,23 +58,20 @@ MultiDevSSAGraphBuilder::MultiDevSSAGraphBuilder(
 
 void MultiDevSSAGraphBuilder::CreateOpHandleIOs(SSAGraph *result,
                                                 const OpDesc &op,
-                                                const platform::Place &p,
+                                                size_t place_id) const {
-                                                const size_t &i) const {
+  auto p = places_[place_id];
   auto *op_handle = result->ops_.back().get();
   op_handle->SetDeviceContext(p,
                               platform::DeviceContextPool::Instance().Get(p));
 
-  auto var_names = op.InputArgumentNames();
+  for (auto &each_var_name : op.InputArgumentNames()) {
-
+    VarHandle *var =
-  for (auto &each_var_name : var_names) {
+        CreateOrGetLatestVarHandle(result, each_var_name, p, place_id);
-    VarHandle *var = CreateOrGetLatestVarHandle(result, each_var_name, p, i);
     op_handle->AddInput(var);
   }
 
-  var_names = op.OutputArgumentNames();
+  for (auto &each_var_name : op.OutputArgumentNames()) {
-
+    CreateOpOutput(result, op_handle, each_var_name, p, place_id);
-  for (auto &each_var_name : var_names) {
-    CreateOpOutput(result, op_handle, each_var_name, p, i);
   }
 }
 
@@ -84,17 +81,18 @@ bool MultiDevSSAGraphBuilder::IsDistTrainOp(const OpDesc &op,
     return false;
   }
 
-  auto checker = [&](const std::vector<std::string> opvars,
+  /**
-                     const std::vector<std::string> sendvars) -> bool {
+   * Check any of opvars contains `.block` and in sendvars
-    bool is_dist_train_op = false;
+   */
+  auto checker = [](const std::vector<std::string> &opvars,
+                    const std::vector<std::string> &sendvars) -> bool {
     for (auto &var : opvars) {
       if (var.find(".block") != std::string::npos &&
           std::find(sendvars.begin(), sendvars.end(), var) != sendvars.end()) {
-        is_dist_train_op = true;
+        return true;
-        break;
       }
     }
-    return is_dist_train_op;
+    return false;
   };
 
   if (op.Type() == "split") {
@@ -117,13 +115,7 @@ std::unique_ptr<SSAGraph> MultiDevSSAGraphBuilder::Build(
       places_.size());
 
   // Find "send" op first for split is in front of send.
-  OpDesc *send_op = nullptr;
+  OpDesc *send_op = GetSendOpDesc(program);
-  for (auto *op : program.Block(0).AllOps()) {
-    if (op->Type() == "send") {
-      send_op = op;
-      break;
-    }
-  }
 
   bool is_forwarding = true;
   for (auto *op : program.Block(0).AllOps()) {
@@ -134,6 +126,7 @@ std::unique_ptr<SSAGraph> MultiDevSSAGraphBuilder::Build(
     } else if (IsDistTrainOp(*op, send_op)) {
       CreateComputationalOps(&result, *op, 1);
     } else if (IsScaleLossOp(*op)) {
+      // user can customize loss@grad if skip_scale_loss_
       if (!skip_scale_loss_) {
         CreateScaleLossGradOp(&result);
       }
@@ -142,10 +135,7 @@ std::unique_ptr<SSAGraph> MultiDevSSAGraphBuilder::Build(
       CreateComputationalOps(&result, *op, places_.size());
       if (!is_forwarding) {
         // Currently, we assume that once gradient is generated, it can be
-        // broadcast, and each gradient is only broadcast once. But there are no
+        // broadcast, and each gradient is only broadcast once.
-        // other cases, for example, we need to adjust the gradient according to
-        // the input when we get the gradient, which is not considered at
-        // present.
         for (auto &og : op->OutputArgumentNames()) {
           if (IsParameterGradientOnce(og, &og_has_been_broadcast)) {
             InsertNCCLAllReduceOp(&result, og);
@@ -175,6 +165,16 @@ std::unique_ptr<SSAGraph> MultiDevSSAGraphBuilder::Build(
   return std::unique_ptr<SSAGraph>(graph);
 }
 
+OpDesc *MultiDevSSAGraphBuilder::GetSendOpDesc(
+    const ProgramDesc &program) const {
+  for (auto *op : program.Block(0).AllOps()) {
+    if (op->Type() == "send") {
+      return op;
+    }
+  }
+  return nullptr;
+}
+
 void MultiDevSSAGraphBuilder::InsertNCCLAllReduceOp(
     SSAGraph *result, const std::string &og) const {
 #ifdef PADDLE_WITH_CUDA
@@ -243,7 +243,7 @@ void MultiDevSSAGraphBuilder::CreateComputationalOps(SSAGraph *result,
     auto p = places_[scope_idx];
     auto s = local_scopes_[scope_idx];
     result->ops_.emplace_back(new ComputationOpHandle(op, s, p));
-    CreateOpHandleIOs(result, op, p, scope_idx);
+    CreateOpHandleIOs(result, op, scope_idx);
   }
 }
 
@@ -255,7 +255,7 @@ void MultiDevSSAGraphBuilder::CreateSendOp(SSAGraph *result,
   result->ops_.emplace_back(new SendOpHandle(op, s, p));
   // Create inputs for output on original place and no ssa output
   // is created for send op.
-  CreateOpHandleIOs(result, op, p, 0);
+  CreateOpHandleIOs(result, op, 0);
 }
 
 bool MultiDevSSAGraphBuilder::IsScaleLossOp(const OpDesc &op) const {

paddle/fluid/framework/details/multi_devices_graph_builder.h

@@ -48,7 +48,7 @@ class MultiDevSSAGraphBuilder : public SSAGraphBuilder {
 
  private:
   void CreateOpHandleIOs(SSAGraph *result, const OpDesc &op,
-                         const platform::Place &p, const size_t &i) const;
+                         size_t place_id) const;
 
  private:
   std::string loss_var_name_;
@@ -65,6 +65,9 @@ class MultiDevSSAGraphBuilder : public SSAGraphBuilder {
 
   void CreateSendOp(SSAGraph *result, const OpDesc &op) const;
 
+  /**
+   * Is this operator as the end-point operator before/after send operator.
+   */
   bool IsDistTrainOp(const OpDesc &op, OpDesc *send_op) const;
 
   void CreateComputationalOps(SSAGraph *result, const OpDesc &op,
@@ -77,6 +80,12 @@ class MultiDevSSAGraphBuilder : public SSAGraphBuilder {
       std::unordered_set<std::string> *og_has_been_broadcast) const;
 
   void InsertNCCLAllReduceOp(SSAGraph *result, const std::string &og) const;
+
+  /**
+   * Get send op in the global block of program.
+   * nullptr if not found.
+   */
+  OpDesc *GetSendOpDesc(const ProgramDesc &program) const;
 };
 }  // namespace details
 }  // namespace framework

paddle/fluid/framework/details/reduce_op_handle_test.cc

@@ -194,7 +194,7 @@ struct TestReduceOpHandle {
     }
 
     f::Tensor result_tensor;
-    f::TensorCopy(rt, cpu_place, *(ctxs_[output_scope_idx]), &result_tensor);
+    f::TensorCopySync(rt, cpu_place, &result_tensor);
     float *ct = result_tensor.data<float>();
 
     for (int64_t j = 0; j < f::product(result_tensor.dims()); ++j) {
@@ -239,7 +239,7 @@ struct TestReduceOpHandle {
     auto &rt = out_var->Get<f::LoDTensor>();
 
     f::Tensor result_tensor;
-    f::TensorCopy(rt, cpu_place, *(ctxs_[output_scope_idx]), &result_tensor);
+    f::TensorCopySync(rt, cpu_place, &result_tensor);
     float *ct = result_tensor.data<float>();
 
     for (int64_t j = 0; j < f::product(result_tensor.dims()); ++j) {

paddle/fluid/framework/details/scale_loss_grad_op_handle.cc

@@ -46,6 +46,7 @@ void ScaleLossGradOpHandle::RunImpl() {
               ->stream();
       memory::Copy(boost::get<platform::CUDAPlace>(place_), tmp,
                    platform::CPUPlace(), &coeff_, sizeof(float), stream);
+      VLOG(1) << place_ << "RUN Scale loss grad op";
     });
 #endif
   }

paddle/fluid/framework/details/ssa_graph.h

@@ -25,12 +25,22 @@ namespace paddle {
 namespace framework {
 namespace details {
 
+// A SSA graph used by parallel executor.
 struct SSAGraph {
+  // all variable in each devices.
+  // The outside vector is the device vector. Each element of this vector is a
+  // map from variable name to variables. The variables, who have the same name,
+  // will have a different version. The offset in the
+  // `std::vector<std::unique_ptr<VarHandle>>` is the version of varaibles.
   std::vector<
       std::unordered_map<std::string, std::vector<std::unique_ptr<VarHandle>>>>
       vars_;
+
   // aux variables to represent dependency. Useful to resolve data hazard.
   std::unordered_set<std::unique_ptr<VarHandleBase>> dep_vars_;
+
+  // all operators. NOTE that even we use a vector here, the operators is
+  // unordered.
   std::vector<std::unique_ptr<OpHandleBase>> ops_;
 };
 

paddle/fluid/framework/details/ssa_graph_builder.h

@@ -48,6 +48,8 @@ class SSAGraphBuilder {
                                                const platform::Place &place,
                                                size_t place_offset);
 
+  // Add an output variable (each_var_name, place, place_offset) to op_handle,
+  // which belongs to graph
   static void CreateOpOutput(SSAGraph *graph, OpHandleBase *op_handle,
                              const std::string &each_var_name,
                              const platform::Place &place, size_t place_offset);

paddle/fluid/framework/executor.cc

@@ -226,15 +226,15 @@ static bool has_fetch_operators(
 }
 
 void Executor::Run(const ProgramDesc& program, Scope* scope,
-                   std::map<std::string, const LoDTensor*>& feed_targets,
+                   std::map<std::string, const LoDTensor*>* feed_targets,
-                   std::map<std::string, LoDTensor*>& fetch_targets,
+                   std::map<std::string, LoDTensor*>* fetch_targets,
                    bool create_vars, const std::string& feed_holder_name,
                    const std::string& fetch_holder_name) {
   platform::RecordBlock b(kProgramId);
   bool has_feed_ops =
-      has_feed_operators(program.Block(0), feed_targets, feed_holder_name);
+      has_feed_operators(program.Block(0), *feed_targets, feed_holder_name);
   bool has_fetch_ops =
-      has_fetch_operators(program.Block(0), fetch_targets, fetch_holder_name);
+      has_fetch_operators(program.Block(0), *fetch_targets, fetch_holder_name);
 
   ProgramDesc* copy_program = const_cast<ProgramDesc*>(&program);
   if (!has_feed_ops || !has_fetch_ops) {
@@ -250,7 +250,7 @@ void Executor::Run(const ProgramDesc& program, Scope* scope,
     feed_holder->SetPersistable(true);
 
     int i = 0;
-    for (auto& feed_target : feed_targets) {
+    for (auto& feed_target : (*feed_targets)) {
       std::string var_name = feed_target.first;
       VLOG(3) << "feed target's name: " << var_name;
 
@@ -273,7 +273,7 @@ void Executor::Run(const ProgramDesc& program, Scope* scope,
     fetch_holder->SetPersistable(true);
 
     int i = 0;
-    for (auto& fetch_target : fetch_targets) {
+    for (auto& fetch_target : (*fetch_targets)) {
       std::string var_name = fetch_target.first;
       VLOG(3) << "fetch target's name: " << var_name;
 
@@ -361,16 +361,16 @@ void Executor::RunPreparedContext(ExecutorPrepareContext* ctx, Scope* scope,
 
 void Executor::RunPreparedContext(
     ExecutorPrepareContext* ctx, Scope* scope,
-    std::map<std::string, const LoDTensor*>& feed_targets,
+    std::map<std::string, const LoDTensor*>* feed_targets,
-    std::map<std::string, LoDTensor*>& fetch_targets, bool create_vars,
+    std::map<std::string, LoDTensor*>* fetch_targets, bool create_vars,
     const std::string& feed_holder_name, const std::string& fetch_holder_name) {
   auto& global_block = ctx->prog_.Block(ctx->block_id_);
 
   PADDLE_ENFORCE(
-      has_feed_operators(global_block, feed_targets, feed_holder_name),
+      has_feed_operators(global_block, *feed_targets, feed_holder_name),
       "Program in ExecutorPrepareContext should has feed_ops.");
   PADDLE_ENFORCE(
-      has_fetch_operators(global_block, fetch_targets, fetch_holder_name),
+      has_fetch_operators(global_block, *fetch_targets, fetch_holder_name),
       "Program in the prepared context should has fetch_ops.");
 
   // map the data of feed_targets to feed_holder
@@ -378,8 +378,8 @@ void Executor::RunPreparedContext(
     if (op->Type() == kFeedOpType) {
       std::string feed_target_name = op->Output("Out")[0];
       int idx = boost::get<int>(op->GetAttr("col"));
-      SetFeedVariable(scope, *feed_targets[feed_target_name], feed_holder_name,
+      SetFeedVariable(scope, *(*feed_targets)[feed_target_name],
-                      idx);
+                      feed_holder_name, idx);
     }
   }
 
@@ -390,7 +390,7 @@ void Executor::RunPreparedContext(
     if (op->Type() == kFetchOpType) {
       std::string fetch_target_name = op->Input("X")[0];
       int idx = boost::get<int>(op->GetAttr("col"));
-      *fetch_targets[fetch_target_name] =
+      *(*fetch_targets)[fetch_target_name] =
           GetFetchVariable(*scope, fetch_holder_name, idx);
     }
   }

paddle/fluid/framework/executor.h

@@ -55,8 +55,8 @@ class Executor {
            bool create_local_scope = true, bool create_vars = true);
 
   void Run(const ProgramDesc& program, Scope* scope,
-           std::map<std::string, const LoDTensor*>& feed_targets,
+           std::map<std::string, const LoDTensor*>* feed_targets,
-           std::map<std::string, LoDTensor*>& fetch_targets,
+           std::map<std::string, LoDTensor*>* fetch_targets,
            bool create_vars = true,
            const std::string& feed_holder_name = "feed",
            const std::string& fetch_holder_name = "fetch");
@@ -74,8 +74,8 @@ class Executor {
                           bool create_vars = true);
 
   void RunPreparedContext(ExecutorPrepareContext* ctx, Scope* scope,
-                          std::map<std::string, const LoDTensor*>& feed_targets,
+                          std::map<std::string, const LoDTensor*>* feed_targets,
-                          std::map<std::string, LoDTensor*>& fetch_targets,
+                          std::map<std::string, LoDTensor*>* fetch_targets,
                           bool create_vars = true,
                           const std::string& feed_holder_name = "feed",
                           const std::string& fetch_holder_name = "fetch");

paddle/fluid/framework/init.cc

@@ -15,7 +15,6 @@ limitations under the License. */
 #include <algorithm>
 #include <stdexcept>
 #include <string>
-#include <vector>
 
 #include "paddle/fluid/framework/init.h"
 #include "paddle/fluid/framework/operator.h"
@@ -31,6 +30,7 @@ std::once_flag p2p_init_flag;
 
 void InitGflags(std::vector<std::string> argv) {
   std::call_once(gflags_init_flag, [&]() {
+    argv.insert(argv.begin(), "dummy");
     int argc = argv.size();
     char **arr = new char *[argv.size()];
     std::string line;
@@ -44,20 +44,23 @@ void InitGflags(std::vector<std::string> argv) {
   });
 }
 
-void InitP2P(int count) {
+void InitP2P(std::vector<int> devices) {
 #ifdef PADDLE_WITH_CUDA
   std::call_once(p2p_init_flag, [&]() {
+    int count = devices.size();
     for (int i = 0; i < count; ++i) {
       for (int j = 0; j < count; ++j) {
-        if (i == j) continue;
+        if (devices[i] == devices[j]) continue;
         int can_acess = -1;
-        PADDLE_ENFORCE(cudaDeviceCanAccessPeer(&can_acess, i, j),
+        PADDLE_ENFORCE(
-                       "Failed to test P2P access.");
+            cudaDeviceCanAccessPeer(&can_acess, devices[i], devices[j]),
+            "Failed to test P2P access.");
         if (can_acess != 1) {
-          LOG(WARNING) << "Cannot enable P2P access from " << i << " to " << j;
+          LOG(WARNING) << "Cannot enable P2P access from " << devices[i]
+                       << " to " << devices[j];
         } else {
-          cudaSetDevice(i);
+          cudaSetDevice(devices[i]);
-          cudaDeviceEnablePeerAccess(j, 0);
+          cudaDeviceEnablePeerAccess(devices[j], 0);
         }
       }
     }
@@ -67,11 +70,26 @@ void InitP2P(int count) {
 
 void InitDevices(bool init_p2p) {
   /*Init all available devices by default */
+  std::vector<int> devices;
+#ifdef PADDLE_WITH_CUDA
+  try {
+    int count = platform::GetCUDADeviceCount();
+    for (int i = 0; i < count; ++i) {
+      devices.push_back(i);
+    }
+  } catch (const std::exception &exp) {
+    LOG(WARNING) << "Compiled with WITH_GPU, but no GPU found in runtime.";
+  }
+#else
+  LOG(WARNING)
+      << "'CUDA' is not supported, Please re-compile with WITH_GPU option";
+#endif
+  InitDevices(init_p2p, devices);
+}
 
+void InitDevices(bool init_p2p, const std::vector<int> devices) {
   std::vector<platform::Place> places;
-  places.emplace_back(platform::CPUPlace());
   int count = 0;
-
 #ifdef PADDLE_WITH_CUDA
   try {
     count = platform::GetCUDADeviceCount();
@@ -83,12 +101,17 @@ void InitDevices(bool init_p2p) {
       << "'CUDA' is not supported, Please re-compile with WITH_GPU option";
 #endif
 
-  for (int i = 0; i < count; ++i) {
+  for (size_t i = 0; i < devices.size(); ++i) {
-    places.emplace_back(platform::CUDAPlace(i));
+    if (devices[i] >= count || devices[i] < 0) {
+      LOG(WARNING) << "Invalid devices id.";
+      continue;
+    }
+    places.emplace_back(platform::CUDAPlace(devices[i]));
   }
   if (init_p2p) {
-    InitP2P(count);
+    InitP2P(devices);
   }
+  places.emplace_back(platform::CPUPlace());
   platform::DeviceContextPool::Init(places);
 }
 

paddle/fluid/framework/init.h

@@ -28,5 +28,7 @@ void InitGLOG(const std::string &prog_name);
 
 void InitDevices(bool init_p2p);
 
+void InitDevices(bool init_p2p, const std::vector<int> devices);
+
 }  // namespace framework
 }  // namespace paddle

paddle/fluid/framework/tensor_util.cc

@@ -20,7 +20,7 @@ namespace paddle {
 namespace framework {
 
 void TensorCopy(const Tensor& src, const platform::Place& dst_place,
-                const platform::DeviceContext& ctx, Tensor* dst, bool sync) {
+                const platform::DeviceContext& ctx, Tensor* dst) {
   VLOG(3) << "TensorCopy " << src.dims() << " from " << src.place() << " to "
           << dst_place;
   src.check_memory_size();
@@ -48,9 +48,7 @@ void TensorCopy(const Tensor& src, const platform::Place& dst_place,
     auto ctx_gpu_place = boost::get<platform::CUDAPlace>(ctx_place);
     PADDLE_ENFORCE_EQ(src_gpu_place, ctx_gpu_place);
     auto stream =
-        sync ? nullptr
+        reinterpret_cast<const platform::CUDADeviceContext&>(ctx).stream();
-             : reinterpret_cast<const platform::CUDADeviceContext&>(ctx)
-                   .stream();
     memory::Copy(dst_cpu_place, dst_ptr, src_gpu_place, src_ptr, size, stream);
   } else if (platform::is_cpu_place(src_place) &&
              platform::is_gpu_place(dst_place)) {
@@ -61,9 +59,7 @@ void TensorCopy(const Tensor& src, const platform::Place& dst_place,
     auto ctx_gpu_place = boost::get<platform::CUDAPlace>(ctx_place);
     PADDLE_ENFORCE_EQ(dst_gpu_place, ctx_gpu_place);
     auto stream =
-        sync ? nullptr
+        reinterpret_cast<const platform::CUDADeviceContext&>(ctx).stream();
-             : reinterpret_cast<const platform::CUDADeviceContext&>(ctx)
-                   .stream();
     memory::Copy(dst_gpu_place, dst_ptr, src_cpu_place, src_ptr, size, stream);
   } else if (platform::is_gpu_place(src_place) &&
              platform::is_gpu_place(dst_place)) {
@@ -72,9 +68,7 @@ void TensorCopy(const Tensor& src, const platform::Place& dst_place,
     auto ctx_place = ctx.GetPlace();
     PADDLE_ENFORCE(platform::is_gpu_place(ctx_place));
     auto stream =
-        sync ? nullptr
+        reinterpret_cast<const platform::CUDADeviceContext&>(ctx).stream();
-             : reinterpret_cast<const platform::CUDADeviceContext&>(ctx)
-                   .stream();
     memory::Copy(dst_gpu_place, dst_ptr, src_gpu_place, src_ptr, size, stream);
   }
 #endif
@@ -92,6 +86,41 @@ void TensorCopy(const Tensor& src, const platform::Place& dst_place,
   TensorCopy(src, dst_place, *dev_ctx, dst);
 }
 
+void TensorCopySync(const Tensor& src, const platform::Place& dst_place,
+                    Tensor* dst) {
+  VLOG(3) << "TensorCopySync " << src.dims() << " from " << src.place()
+          << " to " << dst_place;
+  src.check_memory_size();
+  dst->Resize(src.dims());
+  dst->set_layout(src.layout());
+  auto src_place = src.place();
+  auto src_ptr = src.data<void>();
+  auto dst_ptr = dst->mutable_data(dst_place, src.type());
+  auto size = src.numel() * SizeOfType(src.type());
+  if (platform::is_cpu_place(src_place) && platform::is_cpu_place(dst_place)) {
+    memory::Copy(boost::get<platform::CPUPlace>(dst_place), dst_ptr,
+                 boost::get<platform::CPUPlace>(src_place), src_ptr, size);
+  }
+#ifdef PADDLE_WITH_CUDA
+  else if (platform::is_gpu_place(src_place) &&  // NOLINT
+           platform::is_cpu_place(dst_place)) {
+    auto src_gpu_place = boost::get<platform::CUDAPlace>(src_place);
+    auto dst_cpu_place = boost::get<platform::CPUPlace>(dst_place);
+    memory::Copy(dst_cpu_place, dst_ptr, src_gpu_place, src_ptr, size, nullptr);
+  } else if (platform::is_cpu_place(src_place) &&
+             platform::is_gpu_place(dst_place)) {
+    auto src_cpu_place = boost::get<platform::CPUPlace>(src_place);
+    auto dst_gpu_place = boost::get<platform::CUDAPlace>(dst_place);
+    memory::Copy(dst_gpu_place, dst_ptr, src_cpu_place, src_ptr, size, nullptr);
+  } else if (platform::is_gpu_place(src_place) &&
+             platform::is_gpu_place(dst_place)) {
+    auto src_gpu_place = boost::get<platform::CUDAPlace>(src_place);
+    auto dst_gpu_place = boost::get<platform::CUDAPlace>(dst_place);
+    memory::Copy(dst_gpu_place, dst_ptr, src_gpu_place, src_ptr, size, nullptr);
+  }
+#endif
+}
+
 template <typename Predicate, typename DevCtx>
 struct AnyDTypeVisitor {
   Predicate predicate_;

paddle/fluid/framework/tensor_util.h

@@ -24,10 +24,11 @@ namespace paddle {
 namespace framework {
 
 void TensorCopy(const Tensor& src, const platform::Place& dst_place,
-                const platform::DeviceContext& ctx, Tensor* dst,
+                const platform::DeviceContext& ctx, Tensor* dst);
-                bool sync = false);
 void TensorCopy(const Tensor& src, const platform::Place& dst_place,
                 Tensor* dst);
+void TensorCopySync(const Tensor& src, const platform::Place& dst_place,
+                    Tensor* dst);
 
 template <typename T>
 void TensorFromVector(const std::vector<T>& src,