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

Dockerfile

@@ -53,10 +53,14 @@ RUN localedef -i en_US -f UTF-8 en_US.UTF-8
 
 # FIXME: due to temporary ipykernel dependency issue, specify ipykernel jupyter
 # version util jupyter fixes this issue.
+
+# specify sphinx version as 1.5.6 and remove -U option for [pip install -U
+# sphinx-rtd-theme] since -U option will cause sphinx being updated to newest
+# version(1.7.1 for now), which causes building documentation failed.
 RUN pip install --upgrade pip && \
     pip install -U wheel && \
-    pip install -U docopt PyYAML sphinx && \
-    pip install -U sphinx-rtd-theme==0.1.9 recommonmark
+    pip install -U docopt PyYAML sphinx==1.5.6 && \
+    pip install sphinx-rtd-theme==0.1.9 recommonmark
 
 RUN pip install pre-commit 'ipython==5.3.0' && \
     pip install 'ipykernel==4.6.0' 'jupyter==1.0.0' && \

doc/design/cpp_data_feeding.md

-# C++ Data Feeding
-
-While using Paddle V2 API for Training, data feeding completely depends on the Python code. To get rid of the Python environment and achieve the goal of "wrapping the whole training by a while loop op" in Paddle Fluid, a C++ data feeding mechanism is required. 
-
-In this document we show the fundamental design of a C++ data feeding process, which includes data reading, shuffling and batching.
-
-## Reader
-
-In order to handle the above mentioned problem, a new concept called 'Reader' is introduced. `Reader` is a series of inherited classes which can be held by our `Variable` and they are used to read or process file data.
-
-
-### `ReaderBase`
-
-`ReaderBase` is the abstract base class for all readers. It defines the interface for all readers.
-
-```cpp
-class ReaderBase {
- public:
-  explicit ReaderBase(const std::vector<DDim>& shapes) : shapes_(shapes) {
-    PADDLE_ENFORCE(!shapes_.empty());
-  }
-  // Read the next batch of data. (A 'batch' can be only one instance)
-  // If the next batch doesn't exist, '*out' will be an empty std::vector.
-  virtual void ReadNext(std::vector<LoDTensor>* out) = 0;
-  
-  // Reinitialize the reader and read the file from the beginning.
-  virtual void ReInit() = 0;
-  
-  // Get a certain read in data's shape.
-  DDim shape(size_t idx) const;
-  // Get shapes of all read in data.
-  std::vector<DDim> shapes() const { return shapes_; }
-  // Set shapes of read in data.
-  void set_shapes(const std::vector<DDim>& shapes) { shapes_ = shapes; }
-
-  virtual ~ReaderBase() {}
-
- protected:
-  std::vector<DDim> shapes_;
-};
-```
-
-### `FileReader` and `DecoratedReader`
-
-These two classes are derived from the `ReaderBase` and will further be derived by more specific readers. Thus, in our design, there are two kinds of readers: file readers and decorated readers. A file reader reads from a file of some specific format, and yield only one instance of data at a time. For example, RecordIO reader, jpg reader, .... A decorated reader takes another reader(both file reader and decorated reader are OK) as its 'underlying reader'. It gets data from its underlying reader, does some processing on them(shuffling, or batching), then yields processed data. The output data of a decorated reader can be a single instance or a batch. `ShuffleReader` and `BatchReader` are both decorated readers.
-
-All the readers share exactly the same interface as defined in `ReaderBase`. So they can be decorated for more than one time: We can **shuffle** a reader's outputs and then **batch** the shuffle outputs. The interface consistency also allows related ops use readers without knowing what they are exactly.
-
-
-### `ReaderHolder`
-
-Different readers belong to different class types. This leads to a problem: How can we drop them into `Variable`s and fetch them out by a unified method? For example, if a Variable holds a `BatchReader`, we can not get it by the following code:
-
-```cpp
-var->Get<ReaderBase>("batch_reader");
-```
-
-We would have to write:
-
-```cpp
-var->Get<BatchReader>("batch_reader");
-```
-
-This requires that in order to get a reader from a variable, every time, we must know the reader's type exactly. This is nearly impossible.
-
-To solve this problem, we introduce `ReaderHolder` as a wrapper. It acts as an empty decorator of `ReaderBase`, which hides reader's type. With `ReaderHolder` we are able to fetch all types of readers by `var->Get<ReaderHolder>("...")` and regard the obtained object as a reader.
-
-## Related Operators
-
-To create and invoke readers, some new ops are introduced:
-
-### `CreateReaderOp`
-
-Each reader has its creation op. File readers' creation ops have no input and yield the created file reader as its output. Decorated readers' creation ops take the underlying readers as inputs and then yield new decorated readers.
-
-### `ReadOp`
-
-A reader is only a Variable. It cannot trigger the reading process by itself. So we add the `ReadOp` to execute it. A `ReadOp` takes a reader Variable as its input. Each time it runs, it invokes the reader‘s `ReadNext()` function and gets a new batch of data(or only one instance of data, if we use file reader directly). The output data of a reader are in the form of `std::vector<LoDTenosr>`, so the `ReadOp` also needs to split the vector and move LoDTensors to their respective output Variables.

doc/fluid/design/concepts/cpp_data_feeding.md

+# C++ Data Feeding
+
+While using Paddle V2 API for training, data feeding completely depends on the Python code. To get rid of the Python environment and achieve the goal of "wrapping the whole training by a while loop op" in Paddle Fluid, a C++ data feeding mechanism is required.
+
+In this document, we show the fundamental design of a C++ data feeding process, which includes data reading, shuffling and batching.
+
+## Overview
+
+![](images/readers.png)
+
+## Reader
+
+In order to handle the above-mentioned problem, a new concept called 'Reader' is introduced. `Reader` is a series of inherited classes which can be held by our `Variable` and they are used to read or process file data.
+
+
+### ReaderBase
+
+`ReaderBase` is the abstract base class for all readers. It defines the interface for all readers.
+
+```cpp
+class ReaderBase {
+ public:
+  // Reads the next batch of data. (A 'batch' can be only one instance)
+  // If the next batch doesn't exist, it throws an exception
+  virtual void ReadNext(std::vector<LoDTensor>* out) = 0;
+  
+  // Checks whether the next instance exists.
+  virtual bool HasNext() = 0;
+  
+  // Reinitializes the reader and read the file from the beginning.
+  virtual void ReInit() = 0;
+
+  virtual ~ReaderBase();
+};
+```
+
+### FileReader
+
+`FileReader` is derived from the `ReaderBase`. It is still an abstract class and will further be derived by Readers of respective specific format.
+
+```cpp
+class FileReader : public ReaderBase {
+ public:
+  explicit FileReader(const std::vector<DDim>& dims);
+
+  void ReadNext(std::vector<LoDTensor>* out) override;
+
+ protected:
+  virtual void ReadNextImpl(std::vector<LoDTensor>* out) = 0;
+
+ private:
+  std::vector<DDim> dims_;
+};
+```
+
+A file reader binds with a single file and reads one data instance at a time. Each type of file reader shall implement its own `ReadNextImpl()`, `HasNext()` and `ReInit()`.
+
+The `ReadNextImpl()` is invoked by `ReadNext()`. Besides invoking `ReadNextImpl()`, `ReadNext()` is also responsible for checking the output, making sure that each shape of `LoDTensor` in `*out` is consistent with the one in `dims_`.  
+
+### DecoratedReader
+
+A decorated reader takes another reader(both file reader and decorated reader are OK) as its 'underlying reader'. It gets data from its underlying reader, does some processing on them(shuffling,  batching or something else), then yields processed data. The output data of a decorated reader can be a single instance or a batch. `ShuffleReader` and `BatchReader` are both decorated readers.
+
+```cpp
+class DecoratedReader : public ReaderBase {
+ public:
+  explicit DecoratedReader(ReaderBase* reader) : ReaderBase(), reader_(reader) {
+    PADDLE_ENFORCE_NOT_NULL(reader_);
+  }
+
+  void ReInit() override { reader_->ReInit(); }
+
+  bool HasNext() const override { return reader_->HasNext(); }
+
+ protected:
+  ReaderBase* reader_;
+};
+```
+
+Both the `FileReader` and `DecoratedReader` share exactly the same interface as defined in `ReaderBase`. So they can be decorated for multiple times: We can **shuffle** a reader's outputs and then **batch** the shuffled outputs. The interface consistency also allows related ops use readers without knowing their underlying type.
+
+### MultipleReader
+
+All `FileReader` binds with a single file and are single-threaded. However, sometimes we need to read data from more than one file. In this case, it's not enough to only have `FileReader` and `DecoratedReader`.
+
+So `MultipleReader` is introduced. It is also derived from `ReaderBase`. A `MultipleReader` holds several prefetching `FileReaders` and these readers run concurrently. Another pivotal part of a `MultipleReader` is a buffer channel. The channel collects data yield by all prefetching readers and makes subsequent OPs or decorated readers be able to fetch data without concerning about multiple readers scheduling.
+
+![](images/multiple_reader.png)
+
+This graph shows how a `MultipleReader` works with three prefetching file readers and two GPUs. There is a queue of files which are going to be read. Each time when a prefetching file reader is free(complete reading from one file), it fetches a new file from the queue. Each prefetching file reader runs in a separated prefetch thread and dumps their outputs to the same channel.
+
+To the subsequent two decorated readers, the `MultipleReader` is **a single reader**. They don't need to concern about how prefetch readers are scheduled. They only need to invoke `MultipleReader::ReadNext()` to get the next data from the buffer channel. 
+
+### ReaderHolder
+
+Different readers belong to different class types. This leads to a problem: How can we drop them into `Variable`s and fetch them out by a unified method? For example, if a Variable holds a `BatchReader`, we can not get it by the following code:
+
+```cpp
+var->Get<ReaderBase>("batch_reader");
+```
+
+We would have to write:
+
+```cpp
+var->Get<BatchReader>("batch_reader");
+```
+
+This requires that in order to get a reader from a variable, every time, we must know the reader's type exactly. This is nearly impossible.
+
+To solve this problem, we introduce `ReaderHolder` as a wrapper. It acts as an empty decorator of `ReaderBase`, which hides reader's type. With `ReaderHolder` we are able to fetch all types of readers by `var->Get<ReaderHolder>("...")` and regard the obtained object as a reader.
+
+## Related Operators
+
+To create and invoke readers, some new ops are introduced:
+
+### CreateReaderOp
+
+Each reader has its creation op. File readers' creation ops have no input and yield the created file reader as its output. Decorated readers' creation ops take the underlying readers as inputs and then yield new decorated readers.
+
+However, direct usage of file readers' creation ops is not recommended because a file reader can only read one file via a single thread. Using `OpenFilesOp` is a better choice.
+
+### OpenFilesOp
+
+The `OpenFilesOp` is the creation op of `MultipleReader`. It takes no input but requires a list of file names as one of its attributes. The newly created `MultipleReader` then creates its own prefetching readers according to given file names.
+
+To make sure that created prefetching readers match file formats, we need a name prefix rule to append file format tags to file names, as well as a file reader registry mechanism to map file format tags to their corresponding file readers' constructors.
+
+### HasNextOp
+
+`HasNextOp` is used to check whether the next data batch exists via the reader's `HasNext()` interface.
+
+### ResetOp
+
+`ResetOp` is used to reset a reader via its `ReInit()` interface.
+
+### ReadOp
+
+A reader is only a Variable. It cannot trigger the reading process by itself. So we add the `ReadOp` to execute it. A `ReadOp` takes a reader Variable as its input. Each time it runs, it invokes the reader‘s `ReadNext()` function and gets a new batch of data(or only one instance of data, if we use file reader directly). The output data of a reader are in the form of `std::vector<LoDTenosr>`, so the `ReadOp` also needs to split the vector and move LoDTensors to their respective output Variables.
+
+## Program with Readers
+
+A `Program` holds readers as its persistable variables. These variables are created by `CreateReaderOp` or `OpenFilesOp`. These ops shall run only once. So they shall be settled in the `startup_program`. `HasNextOp`, `ResetOp` and `ReadOp` are required by training loop, so they shall be in the `main_program`.
+
+The ops of a `startup_program` with readers would be like this:
+
+```
+multiple_reader = open_files_op(...)
+batch_reader = create_batch_reader_op(multiple_reader)
+double_buffer_reader = create_double_buffer_op(batch_reader)
+... (other initializers)
+```
+
+The forwarding ops of the corresponding `main_program` would be like this:
+
+```
+while_op {
+    has_next = has_next_op(double_buffer_reader)
+    if_else_op(has_next) {
+        batch_data = read_op(double_buffer_reader)
+        ... (subsequent training ops)
+    } else {
+        reset_op(double_buffer_reader)
+    }
+}
+```
+
+Two important considerations for these programs are as follows:
+
+1. The multiple\_reader is the batch\_reader's underlying reader, and the batch\_reader is the double\_buffer\_reader's underlying reader. `read_op`, `has_next_op` and other reader related ops will only invoke the top-most reader. In this case, it's the double\_buffer\_reader.
+
+2. All readers exist in both `startup_program` and `main_program`. And they are persistable.

doc/design/speech/deep_speech_2.md → doc/fluid/design/network/deep_speech_2.md

@@ -94,7 +94,7 @@ The classical DS2 network contains 15 layers (from bottom to top):
 - **One** CTC-loss layer
 
 <div align="center">
-<img src="image/ds2_network.png" width=350><br/>
+<img src="images/ds2_network.png" width=350><br/>
 Figure 1. Archetecture of Deep Speech 2 Network.
 </div>
 
@@ -141,7 +141,7 @@ TODO by Assignees
 ### Beam Search with CTC and LM
 
 <div align="center">
-<img src="image/beam_search.png" width=600><br/>
+<img src="images/beam_search.png" width=600><br/>
 Figure 2. Algorithm for CTC Beam Search Decoder.
 </div>
 

doc/v2/howto/cluster/index_en.rst

 Distributed Training
 ====================
 
-In this section, we'll explain how to run distributed training jobs with PaddlePaddle on different types of clusters. The diagram below shows the main architecture of a distributed trainning job:
-
+The effectiveness of the deep learning model is often directly related to the scale of the data: it can generally achieve better results after increasing the size of the dataset on the same model. However, it can not fit in one single computer when the amount of data increases to a certain extent. At this point, using multiple computers for distributed training is a natural solution. In distributed training, the training data is divided into multiple copies (sharding), and multiple machines participating in the training read their own data for training and collaboratively update the parameters of the overall model.
 .. image:: src/ps_en.png
    :width: 500
 
@@ -10,13 +9,27 @@ In this section, we'll explain how to run distributed training jobs with PaddleP
 - Trainer: each trainer reads the data shard, and train the neural network. Then the trainer will upload calculated "gradients" to parameter servers, and wait for parameters to be optimized on the parameter server side. When that finishes, the trainer download optimized parameters and continues its training.
 - Parameter server: every parameter server stores part of the whole neural network model data. They will do optimization calculations when gradients are uploaded from trainers, and then send updated parameters to trainers.
 
-PaddlePaddle can support both synchronize stochastic gradient descent (SGD) and asynchronous SGD.
+The training of synchronous random gradient descent for neural network can be achieved by cooperation of trainers and parameter servers.
+
+PaddlePaddle supports both synchronize stochastic gradient descent (SGD) and asynchronous SGD.
 
-When training with synchronize SGD, PaddlePaddle uses an internal "synchronize barrier" which makes gradients update and parameter download in strict order. On the other hand, asynchronous SGD won't wait for all trainers to finish upload at a single step, this will increase the parallelism of distributed training: parameter servers do not depend on each other, they'll do parameter optimization concurrently. Parameter servers will not wait for trainers, so trainers will also do their work concurrently. But asynchronous SGD will introduce more randomness and noises in the gradient.
+Before starting the cluster training, you need to prepare the cluster configuration, PaddlePaddle installation, and other preparations. To understand how to configure the basic environment for distributed training, check the link below:
 
 ..  toctree::
   :maxdepth: 1
 
   preparations_en.md
+
+Cluster training has a large number of configurable parameters, such as the number of machines used, communication ports, etc. To learn how to configure the distributed training process by setting startup these parameters, check the link below:
+
+..  toctree::
+  :maxdepth: 1
+
   cmd_argument_en.md
+
+PaddlePaddle is compatible with a variety of different clusters. Each cluster has its own advantages, To learn how to run PaddlePaddle in different types of them, check the link below:
+
+..  toctree::
+  :maxdepth: 1
+
   multi_cluster/index_en.rst

paddle/capi/CMakeLists.txt

@@ -36,7 +36,7 @@ target_include_directories(paddle_capi PUBLIC ${CMAKE_CURRENT_BINARY_DIR})
 add_style_check_target(paddle_capi ${CAPI_SOURCES} ${CAPI_HEADER}
   ${CAPI_PRIVATE_HEADER})
 
-add_dependencies(paddle_capi paddle_proto)
+add_dependencies(paddle_capi paddle_proto paddle_gserver)
 
 # TODO: paddle_capi_whole will be removed.
 set(PADDLE_CAPI_LAYERS_LIBS

paddle/fluid/framework/CMakeLists.txt

@@ -103,4 +103,5 @@ cc_test(cow_ptr_tests SRCS details/cow_ptr_test.cc)
 cc_test(channel_test SRCS channel_test.cc)
 cc_test(tuple_test SRCS tuple_test.cc )
 cc_test(concurrency_test SRCS concurrency_test.cc DEPS go_op channel_close_op channel_create_op
-        channel_send_op channel_recv_op sum_op elementwise_add_op executor proto_desc)
+        channel_send_op channel_recv_op sum_op select_op elementwise_add_op compare_op
+        conditional_block_op while_op assign_op print_op executor proto_desc)

paddle/fluid/framework/channel.h

@@ -162,24 +162,12 @@ class ChannelHolder {
     }
   }
 
-  template <typename T>
   void RemoveFromSendQ(const void* referrer) {
-    if (IsInitialized()) {
-      Channel<T>* channel = static_cast<Channel<T>*>(holder_->Ptr());
-      if (channel != nullptr) {
-        channel->RemoveFromSendQ(referrer);
-      }
-    }
+    if (IsInitialized()) holder_->RemoveFromSendQ(referrer);
   }
 
-  template <typename T>
   void RemoveFromReceiveQ(const void* referrer) {
-    if (IsInitialized()) {
-      Channel<T>* channel = static_cast<Channel<T>*>(holder_->Ptr());
-      if (channel != nullptr) {
-        channel->RemoveFromReceiveQ(referrer);
-      }
-    }
+    if (IsInitialized()) holder_->RemoveFromReceiveQ(referrer);
   }
 
   inline bool IsInitialized() const { return holder_ != nullptr; }
@@ -201,6 +189,8 @@ class ChannelHolder {
     virtual bool IsClosed() = 0;
     virtual bool CanSend() = 0;
     virtual bool CanReceive() = 0;
+    virtual void RemoveFromSendQ(const void* referrer) = 0;
+    virtual void RemoveFromReceiveQ(const void* referrer) = 0;
     virtual void Close() = 0;
     virtual void Lock() = 0;
     virtual void Unlock() = 0;
@@ -238,6 +228,18 @@ class ChannelHolder {
       return false;
     }
 
+    virtual void RemoveFromSendQ(const void* referrer) {
+      if (channel_) {
+        channel_->RemoveFromSendQ(referrer);
+      }
+    }
+
+    virtual void RemoveFromReceiveQ(const void* referrer) {
+      if (channel_) {
+        channel_->RemoveFromReceiveQ(referrer);
+      }
+    }
+
     virtual void Close() {
       if (channel_) channel_->Close();
     }

paddle/fluid/framework/channel_impl.h

@@ -151,7 +151,7 @@ bool ChannelImpl<T>::Send(T *item) {
       // We do not care about notifying other
       // because they would have been notified
       // by the executed select case.
-      return Send(item);
+      return send_return(Send(item));
 
     // Wake up the blocked process and unlock
     m->Notify();
@@ -214,7 +214,7 @@ bool ChannelImpl<T>::Receive(T *item) {
       // We do not care about notifying other
       // because they would have been notified
       // by the executed select case.
-      return Receive(item);
+      return recv_return(Receive(item));
 
     // Wake up the blocked process and unlock
     m->Notify();
@@ -331,7 +331,6 @@ void ChannelImpl<T>::RemoveFromSendQ(const void *referrer) {
 
     if (sendMsg->referrer == referrer) {
       it = sendq.erase(it);
-      send_ctr--;
     } else {
       ++it;
     }
@@ -347,7 +346,6 @@ void ChannelImpl<T>::RemoveFromReceiveQ(const void *referrer) {
 
     if (recvMsg->referrer == referrer) {
       it = recvq.erase(it);
-      recv_ctr--;
     } else {
       ++it;
     }

paddle/fluid/framework/concurrency_test.cc

@@ -19,7 +19,6 @@ limitations under the License. */
 #include "paddle/fluid/framework/channel.h"
 #include "paddle/fluid/framework/executor.h"
 #include "paddle/fluid/framework/op_registry.h"
-#include "paddle/fluid/framework/program_desc.h"
 
 USE_NO_KERNEL_OP(go);
 USE_NO_KERNEL_OP(channel_close);
@@ -27,6 +26,12 @@ USE_NO_KERNEL_OP(channel_create);
 USE_NO_KERNEL_OP(channel_recv);
 USE_NO_KERNEL_OP(channel_send);
 USE_NO_KERNEL_OP(elementwise_add);
+USE_NO_KERNEL_OP(select);
+USE_NO_KERNEL_OP(conditional_block);
+USE_NO_KERNEL_OP(equal);
+USE_NO_KERNEL_OP(assign);
+USE_NO_KERNEL_OP(while);
+USE_NO_KERNEL_OP(print);
 
 namespace f = paddle::framework;
 namespace p = paddle::platform;
@@ -35,27 +40,15 @@ namespace paddle {
 namespace framework {
 
 template <typename T>
-void CreateIntVariable(Scope &scope, p::CPUPlace &place, std::string name,
+LoDTensor *CreateVariable(Scope &scope, p::CPUPlace &place, std::string name,
                           T value) {
-  // Create LoDTensor<int> of dim [1,1]
+  // Create LoDTensor<int> of dim [1]
   auto var = scope.Var(name);
   auto tensor = var->GetMutable<LoDTensor>();
-  tensor->Resize({1, 1});
+  tensor->Resize({1});
   T *expect = tensor->mutable_data<T>(place);
   expect[0] = value;
-}
-
-void InitTensorsInScope(Scope &scope, p::CPUPlace &place) {
-  p::CPUDeviceContext ctx(place);
-
-  // Create channel variable
-  scope.Var("Channel");
-
-  // Create Variables, x0 will be put into channel,
-  // result will be pulled from channel
-  CreateIntVariable(scope, place, "Status", false);
-  CreateIntVariable(scope, place, "x0", 99);
-  CreateIntVariable(scope, place, "result", 0);
+  return tensor;
 }
 
 void AddOp(const std::string &type, const VariableNameMap &inputs,
@@ -73,12 +66,116 @@ void AddOp(const std::string &type, const VariableNameMap &inputs,
   op->SetAttrMap(attrs);
 }
 
+void AddCase(ProgramDesc *program, Scope *scope, p::CPUPlace *place,
+             BlockDesc *casesBlock, int caseId, int caseType,
+             std::string caseChannel, std::string caseVarName,
+             std::function<void(BlockDesc *, Scope *)> func) {
+  std::string caseCondName = std::string("caseCond") + std::to_string(caseId);
+  std::string caseCondXVarName =
+      std::string("caseCondX") + std::to_string(caseId);
+
+  BlockDesc *caseBlock = program->AppendBlock(*casesBlock);
+  func(caseBlock, scope);
+
+  CreateVariable(*scope, *place, caseCondName, false);
+  CreateVariable(*scope, *place, caseCondXVarName, caseId);
+  CreateVariable(*scope, *place, caseVarName, caseId);
+
+  scope->Var("step_scope");
+
+  AddOp("equal", {{"X", {caseCondXVarName}}, {"Y", {"caseToExecute"}}},
+        {{"Out", {caseCondName}}}, {}, casesBlock);
+
+  AddOp("conditional_block", {{"X", {caseCondName}}, {"Params", {}}},
+        {{"Out", {}}, {"Scope", {"step_scope"}}},
+        {{"sub_block", caseBlock}, {"is_scalar_condition", true}}, casesBlock);
+}
+
+void AddFibonacciSelect(Scope *scope, p::CPUPlace *place, ProgramDesc *program,
+                        BlockDesc *parentBlock, std::string dataChanName,
+                        std::string quitChanName) {
+  BlockDesc *whileBlock = program->AppendBlock(*parentBlock);
+
+  CreateVariable(*scope, *place, "whileExitCond", true);
+  CreateVariable(*scope, *place, "caseToExecute", -1);
+  CreateVariable(*scope, *place, "case1var", 0);
+
+  CreateVariable(*scope, *place, "xtemp", 0);
+
+  // TODO(thuan): Need to create fibXToSend, since channel send moves the actual
+  // data,
+  // which causes the data to be no longer accessible to do the fib calculation
+  // TODO(abhinav): Change channel send to do a copy instead of a move!
+  CreateVariable(*scope, *place, "fibXToSend", 0);
+
+  CreateVariable(*scope, *place, "fibX", 0);
+  CreateVariable(*scope, *place, "fibY", 1);
+  CreateVariable(*scope, *place, "quitVar", 0);
+
+  BlockDesc *casesBlock = program->AppendBlock(*whileBlock);
+  std::function<void(BlockDesc * caseBlock)> f = [](BlockDesc *caseBlock) {};
+
+  // TODO(thuan): Remove this once we change channel send to do a copy instead
+  // of move
+  AddOp("assign", {{"X", {"fibX"}}}, {{"Out", {"fibXToSend"}}}, {}, whileBlock);
+
+  // Case 0: Send to dataChanName
+  std::function<void(BlockDesc * caseBlock, Scope * scope)> case0Func = [&](
+      BlockDesc *caseBlock, Scope *scope) {
+    AddOp("assign", {{"X", {"fibX"}}}, {{"Out", {"xtemp"}}}, {}, caseBlock);
+    AddOp("assign", {{"X", {"fibY"}}}, {{"Out", {"fibX"}}}, {}, caseBlock);
+    AddOp("elementwise_add", {{"X", {"xtemp"}}, {"Y", {"fibY"}}},
+          {{"Out", {"fibY"}}}, {}, caseBlock);
+  };
+  AddCase(program, scope, place, casesBlock, 0, 1, dataChanName, "fibXToSend",
+          case0Func);
+  std::string case0Config =
+      std::string("0,1,") + dataChanName + std::string(",fibXToSend");
+
+  // Case 1: Receive from quitChanName
+  std::function<void(BlockDesc * caseBlock, Scope * scope)> case2Func = [&](
+      BlockDesc *caseBlock, Scope *scope) {
+    // Exit the while loop after we receive from quit channel.
+    // We assign a false to "whileExitCond" variable, which will
+    // break out of while_op loop
+    CreateVariable(*scope, *place, "whileFalse", false);
+    AddOp("assign", {{"X", {"whileFalse"}}}, {{"Out", {"whileExitCond"}}}, {},
+          caseBlock);
+  };
+  AddCase(program, scope, place, casesBlock, 1, 2, quitChanName, "quitVar",
+          case2Func);
+  std::string case1Config =
+      std::string("1,2,") + quitChanName + std::string(",quitVar");
+
+  // Select block
+  AddOp("select", {{"X", {dataChanName, quitChanName}},
+                   {"case_to_execute", {"caseToExecute"}}},
+        {}, {{"sub_block", casesBlock},
+             {"cases", std::vector<std::string>{case0Config, case1Config}}},
+        whileBlock);
+
+  scope->Var("stepScopes");
+  AddOp("while",
+        {{"X", {dataChanName, quitChanName}}, {"Condition", {"whileExitCond"}}},
+        {{"Out", {}}, {"StepScopes", {"stepScopes"}}},
+        {{"sub_block", whileBlock}}, parentBlock);
+}
+
 TEST(Concurrency, Go_Op) {
   Scope scope;
   p::CPUPlace place;
 
   // Initialize scope variables
-  InitTensorsInScope(scope, place);
+  p::CPUDeviceContext ctx(place);
+
+  // Create channel variable
+  scope.Var("Channel");
+
+  // Create Variables, x0 will be put into channel,
+  // result will be pulled from channel
+  CreateVariable(scope, place, "Status", false);
+  CreateVariable(scope, place, "x0", 99);
+  CreateVariable(scope, place, "result", 0);
 
   framework::Executor executor(place);
   ProgramDesc program;
@@ -118,5 +215,78 @@ TEST(Concurrency, Go_Op) {
   auto *finalData = tensor.data<int>();
   EXPECT_EQ(finalData[0], 99);
 }
+
+/**
+ * This test implements the fibonacci function using go_op and select_op
+ */
+TEST(Concurrency, Select) {
+  Scope scope;
+  p::CPUPlace place;
+
+  // Initialize scope variables
+  p::CPUDeviceContext ctx(place);
+
+  CreateVariable(scope, place, "Status", false);
+  CreateVariable(scope, place, "result", 0);
+  CreateVariable(scope, place, "currentXFib", 0);
+
+  framework::Executor executor(place);
+  ProgramDesc program;
+  BlockDesc *block = program.MutableBlock(0);
+
+  // Create channel OP
+  std::string dataChanName = "Channel";
+  scope.Var(dataChanName);
+  AddOp("channel_create", {}, {{"Out", {dataChanName}}},
+        {{"capacity", 0}, {"data_type", f::proto::VarType::LOD_TENSOR}}, block);
+
+  std::string quitChanName = "Quit";
+  scope.Var(quitChanName);
+  AddOp("channel_create", {}, {{"Out", {quitChanName}}},
+        {{"capacity", 0}, {"data_type", f::proto::VarType::LOD_TENSOR}}, block);
+
+  // Create Go Op routine, which loops 10 times over fibonacci sequence
+  CreateVariable(scope, place, "xReceiveVar", 0);
+
+  BlockDesc *goOpBlock = program.AppendBlock(program.Block(0));
+  for (int i = 0; i < 10; ++i) {
+    AddOp("channel_recv", {{"Channel", {dataChanName}}},
+          {{"Status", {"Status"}}, {"Out", {"currentXFib"}}}, {}, goOpBlock);
+    AddOp("print", {{"In", {"currentXFib"}}}, {{"Out", {"currentXFib"}}},
+          {{"first_n", 100},
+           {"summarize", -1},
+           {"print_tensor_name", false},
+           {"print_tensor_type", true},
+           {"print_tensor_shape", false},
+           {"print_tensor_lod", false},
+           {"print_phase", std::string("FORWARD")},
+           {"message", std::string("X: ")}},
+          goOpBlock);
+  }
+
+  CreateVariable(scope, place, "quitSignal", 0);
+  AddOp("channel_send", {{"Channel", {quitChanName}}, {"X", {"quitSignal"}}},
+        {{"Status", {"Status"}}}, {}, goOpBlock);
+
+  // Create Go Op
+  AddOp("go", {{"X", {dataChanName, quitChanName}}}, {},
+        {{"sub_block", goOpBlock}}, block);
+
+  AddFibonacciSelect(&scope, &place, &program, block, dataChanName,
+                     quitChanName);
+
+  // Create Channel Close Op
+  AddOp("channel_close", {{"Channel", {dataChanName}}}, {}, {}, block);
+  AddOp("channel_close", {{"Channel", {quitChanName}}}, {}, {}, block);
+
+  executor.Run(program, &scope, 0, true, true);
+
+  // After we call executor.run, "result" variable should be equal to 34
+  // (which is 10 loops through fibonacci sequence)
+  const LoDTensor &tensor = (scope.FindVar("currentXFib"))->Get<LoDTensor>();
+  auto *finalData = tensor.data<int>();
+  EXPECT_EQ(finalData[0], 34);
+}
+
 }  // namespace framework
 }  // namespace paddle

paddle/fluid/framework/executor.cc

@@ -25,6 +25,7 @@ limitations under the License. */
 #include "paddle/fluid/framework/op_registry.h"
 #include "paddle/fluid/framework/reader.h"
 #include "paddle/fluid/platform/place.h"
+#include "paddle/fluid/platform/profiler.h"
 
 DECLARE_bool(benchmark);
 DEFINE_bool(check_nan_inf, false,
@@ -33,12 +34,17 @@ DEFINE_bool(check_nan_inf, false,
 
 namespace paddle {
 namespace framework {
+namespace {
+// block id starts from 0. This id is used to represent the codeblock
+// wrapping the first block 0.
+int kProgramId = -1;
+}  // namespace
 
 struct ExecutorPrepareContext {
   ExecutorPrepareContext(const framework::ProgramDesc& prog, size_t block_id)
       : prog_(prog), block_id_(block_id) {}
 
-  framework::ProgramDesc prog_;
+  const framework::ProgramDesc& prog_;
   size_t block_id_;
   std::vector<std::unique_ptr<OperatorBase>> ops_;
 };
@@ -94,6 +100,7 @@ static void CheckTensorNANOrInf(const std::string& name,
 
 void Executor::Run(const ProgramDesc& pdesc, Scope* scope, int block_id,
                    bool create_local_scope, bool create_vars) {
+  platform::RecordBlock b(block_id);
   auto* ctx = Prepare(pdesc, block_id);
   RunPreparedContext(ctx, scope, create_local_scope, create_vars);
   delete ctx;
@@ -184,6 +191,7 @@ void Executor::Run(const ProgramDesc& program, Scope* scope,
                    std::map<std::string, LoDTensor*>& fetch_targets,
                    const std::string& feed_holder_name,
                    const std::string& fetch_holder_name) {
+  platform::RecordBlock b(kProgramId);
   auto* copy_program = new ProgramDesc(program);
   auto* global_block = copy_program->MutableBlock(0);
 

paddle/fluid/operators/CMakeLists.txt

@@ -165,7 +165,6 @@ op_library(cond_op DEPS framework_proto tensor net_op)
 op_library(cross_entropy_op DEPS cross_entropy)
 op_library(softmax_with_cross_entropy_op DEPS cross_entropy softmax)
 op_library(softmax_op DEPS softmax)
-op_library(detection_output_op DEPS softmax)
 op_library(sequence_softmax_op DEPS softmax)
 op_library(sum_op DEPS selected_rows_functor)
 op_library(sgd_op DEPS selected_rows_functor)
@@ -203,6 +202,11 @@ op_library(save_combine_op DEPS lod_tensor)
 op_library(load_combine_op DEPS lod_tensor)
 op_library(concat_op DEPS concat)
 
+# FIXME(thuan): Move CSP operators to paddle/fluid/framework/operators/concurrency
+add_subdirectory(concurrency)
+op_library(channel_send_op DEPS concurrency)
+op_library(channel_recv_op DEPS concurrency)
+
 list(REMOVE_ITEM GENERAL_OPS ${DEPS_OPS})
 foreach(src ${GENERAL_OPS})
     op_library(${src})

paddle/fluid/operators/cast_op.cc

@@ -14,6 +14,7 @@ limitations under the License. */
 
 #include "paddle/fluid/operators/cast_op.h"
 #include "paddle/fluid/framework/op_registry.h"
+#include "paddle/fluid/platform/float16.h"
 
 namespace paddle {
 namespace operators {
@@ -88,4 +89,5 @@ REGISTER_OP_CPU_KERNEL(cast, ops::CastOpKernel<CPU, float>,
                        ops::CastOpKernel<CPU, double>,
                        ops::CastOpKernel<CPU, int>,
                        ops::CastOpKernel<CPU, int64_t>,
-                       ops::CastOpKernel<CPU, bool>);
+                       ops::CastOpKernel<CPU, bool>,
+                       ops::CastOpKernel<CPU, paddle::platform::float16>);

paddle/fluid/operators/cast_op.cu

@@ -13,6 +13,7 @@ See the License for the specific language governing permissions and
 limitations under the License. */
 
 #include "paddle/fluid/operators/cast_op.h"
+#include "paddle/fluid/platform/float16.h"
 
 template <typename T>
 using CastOpKernel =
@@ -20,4 +21,5 @@ using CastOpKernel =
 
 REGISTER_OP_CUDA_KERNEL(cast, CastOpKernel<float>, CastOpKernel<double>,
                         CastOpKernel<int>, CastOpKernel<int64_t>,
-                        CastOpKernel<bool>);
+                        CastOpKernel<bool>,
+                        CastOpKernel<paddle::platform::float16>);

paddle/fluid/operators/channel_recv_op.cc

@@ -18,6 +18,7 @@ limitations under the License. */
 #include <paddle/fluid/framework/reader.h>
 #include "paddle/fluid/framework/op_registry.h"
 #include "paddle/fluid/framework/var_type.h"
+#include "paddle/fluid/operators/concurrency/channel_util.h"
 #include "paddle/fluid/operators/math/math_function.h"
 
 static constexpr char Channel[] = "Channel";
@@ -36,25 +37,6 @@ void SetReceiveStatus(const platform::Place &dev_place,
   status_tensor[0] = status;
 }
 
-bool ChannelReceive(framework::ChannelHolder *ch, framework::Variable *var) {
-  // Get type of channel and use that to call mutable data for Variable
-  auto type = framework::ToVarType(ch->Type());
-  if (type == framework::proto::VarType_Type_LOD_TENSOR)
-    return ch->Receive(var->GetMutable<framework::LoDTensor>());
-  else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE)
-    return ch->Receive(var->GetMutable<framework::LoDRankTable>());
-  else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY)
-    return ch->Receive(var->GetMutable<framework::LoDTensorArray>());
-  else if (type == framework::proto::VarType_Type_SELECTED_ROWS)
-    return ch->Receive(var->GetMutable<framework::SelectedRows>());
-  else if (type == framework::proto::VarType_Type_READER)
-    return ch->Receive(var->GetMutable<framework::ReaderHolder>());
-  else if (type == framework::proto::VarType_Type_CHANNEL)
-    return ch->Receive(var->GetMutable<framework::ChannelHolder>());
-  else
-    PADDLE_THROW("ChannelReceive:Unsupported type");
-}
-
 class ChannelRecvOp : public framework::OperatorBase {
  public:
   ChannelRecvOp(const std::string &type,
@@ -81,7 +63,7 @@ class ChannelRecvOp : public framework::OperatorBase {
         scope.FindVar(Input(Channel))->GetMutable<framework::ChannelHolder>();
     auto output_var = scope.FindVar(Output(Out));
     // Receive the data from the channel.
-    bool ok = ChannelReceive(ch, output_var);
+    bool ok = concurrency::ChannelReceive(ch, output_var);
 
     // Set the status output of the `ChannelReceive` call.
     SetReceiveStatus(dev_place, *scope.FindVar(Output(Status)), ok);

paddle/fluid/operators/channel_send_op.cc

@@ -18,6 +18,7 @@ limitations under the License. */
 #include <paddle/fluid/framework/reader.h>
 #include "paddle/fluid/framework/op_registry.h"
 #include "paddle/fluid/framework/var_type.h"
+#include "paddle/fluid/operators/concurrency/channel_util.h"
 #include "paddle/fluid/operators/math/math_function.h"
 
 static constexpr char Channel[] = "Channel";
@@ -37,24 +38,6 @@ void SetSendStatus(const platform::Place &dev_place,
   status_tensor[0] = status;
 }
 
-bool ChannelSend(framework::ChannelHolder *ch, framework::Variable *var) {
-  auto type = framework::ToVarType(var->Type());
-  if (type == framework::proto::VarType_Type_LOD_TENSOR)
-    return ch->Send(var->GetMutable<framework::LoDTensor>());
-  else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE)
-    return ch->Send(var->GetMutable<framework::LoDRankTable>());
-  else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY)
-    return ch->Send(var->GetMutable<framework::LoDTensorArray>());
-  else if (type == framework::proto::VarType_Type_SELECTED_ROWS)
-    return ch->Send(var->GetMutable<framework::SelectedRows>());
-  else if (type == framework::proto::VarType_Type_READER)
-    return ch->Send(var->GetMutable<framework::ReaderHolder>());
-  else if (type == framework::proto::VarType_Type_CHANNEL)
-    return ch->Send(var->GetMutable<framework::ChannelHolder>());
-  else
-    PADDLE_THROW("ChannelSend:Unsupported type");
-}
-
 class ChannelSendOp : public framework::OperatorBase {
  public:
   ChannelSendOp(const std::string &type,
@@ -82,7 +65,7 @@ class ChannelSendOp : public framework::OperatorBase {
     auto input_var = scope.FindVar(Input(X));
 
     // Send the input data through the channel.
-    bool ok = ChannelSend(ch, input_var);
+    bool ok = concurrency::ChannelSend(ch, input_var);
 
     // Set the status output of the `ChannelSend` call.
     SetSendStatus(dev_place, *scope.FindVar(Output(Status)), ok);

paddle/fluid/operators/concurrency/CMakeLists.txt

+cc_library(concurrency SRCS channel_util.cc DEPS device_context framework_proto boost eigen3)

paddle/fluid/operators/concurrency/channel_util.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 "channel_util.h"
+#include "paddle/fluid/framework/var_type.h"
+
+namespace poc = paddle::operators::concurrency;
+
+bool poc::ChannelSend(framework::ChannelHolder *ch, framework::Variable *var) {
+  auto type = framework::ToVarType(var->Type());
+  if (type == framework::proto::VarType_Type_LOD_TENSOR)
+    return ch->Send(var->GetMutable<framework::LoDTensor>());
+  else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE)
+    return ch->Send(var->GetMutable<framework::LoDRankTable>());
+  else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY)
+    return ch->Send(var->GetMutable<framework::LoDTensorArray>());
+  else if (type == framework::proto::VarType_Type_SELECTED_ROWS)
+    return ch->Send(var->GetMutable<framework::SelectedRows>());
+  else if (type == framework::proto::VarType_Type_READER)
+    return ch->Send(var->GetMutable<framework::ReaderHolder>());
+  else if (type == framework::proto::VarType_Type_CHANNEL)
+    return ch->Send(var->GetMutable<framework::ChannelHolder>());
+  else
+    PADDLE_THROW("ChannelSend:Unsupported type");
+}
+
+bool poc::ChannelReceive(framework::ChannelHolder *ch,
+                         framework::Variable *var) {
+  // Get type of channel and use that to call mutable data for Variable
+  auto type = framework::ToVarType(ch->Type());
+  if (type == framework::proto::VarType_Type_LOD_TENSOR)
+    return ch->Receive(var->GetMutable<framework::LoDTensor>());
+  else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE)
+    return ch->Receive(var->GetMutable<framework::LoDRankTable>());
+  else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY)
+    return ch->Receive(var->GetMutable<framework::LoDTensorArray>());
+  else if (type == framework::proto::VarType_Type_SELECTED_ROWS)
+    return ch->Receive(var->GetMutable<framework::SelectedRows>());
+  else if (type == framework::proto::VarType_Type_READER)
+    return ch->Receive(var->GetMutable<framework::ReaderHolder>());
+  else if (type == framework::proto::VarType_Type_CHANNEL)
+    return ch->Receive(var->GetMutable<framework::ChannelHolder>());
+  else
+    PADDLE_THROW("ChannelReceive:Unsupported type");
+}
+
+void poc::ChannelAddToSendQ(framework::ChannelHolder *ch, const void *referrer,
+                            framework::Variable *var,
+                            std::shared_ptr<std::condition_variable_any> cond,
+                            std::function<bool(framework::ChannelAction)> cb) {
+  auto type = framework::ToVarType(var->Type());
+  if (type == framework::proto::VarType_Type_LOD_TENSOR) {
+    ch->AddToSendQ(referrer, var->GetMutable<framework::LoDTensor>(), cond, cb);
+  } else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE) {
+    ch->AddToSendQ(referrer, var->GetMutable<framework::LoDRankTable>(), cond,
+                   cb);
+  } else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY) {
+    ch->AddToSendQ(referrer, var->GetMutable<framework::LoDTensorArray>(), cond,
+                   cb);
+  } else if (type == framework::proto::VarType_Type_SELECTED_ROWS) {
+    ch->AddToSendQ(referrer, var->GetMutable<framework::SelectedRows>(), cond,
+                   cb);
+  } else if (type == framework::proto::VarType_Type_READER) {
+    ch->AddToSendQ(referrer, var->GetMutable<framework::ReaderHolder>(), cond,
+                   cb);
+  } else if (type == framework::proto::VarType_Type_CHANNEL) {
+    ch->AddToSendQ(referrer, var->GetMutable<framework::ChannelHolder>(), cond,
+                   cb);
+  } else {
+    PADDLE_THROW("ChannelAddToSendQ:Unsupported type");
+  }
+}
+
+void poc::ChannelAddToReceiveQ(
+    framework::ChannelHolder *ch, const void *referrer,
+    framework::Variable *var, std::shared_ptr<std::condition_variable_any> cond,
+    std::function<bool(framework::ChannelAction)> cb) {
+  auto type = framework::ToVarType(var->Type());
+  if (type == framework::proto::VarType_Type_LOD_TENSOR) {
+    ch->AddToReceiveQ(referrer, var->GetMutable<framework::LoDTensor>(), cond,
+                      cb);
+  } else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE) {
+    ch->AddToReceiveQ(referrer, var->GetMutable<framework::LoDRankTable>(),
+                      cond, cb);
+  } else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY) {
+    ch->AddToReceiveQ(referrer, var->GetMutable<framework::LoDTensorArray>(),
+                      cond, cb);
+  } else if (type == framework::proto::VarType_Type_SELECTED_ROWS) {
+    ch->AddToReceiveQ(referrer, var->GetMutable<framework::SelectedRows>(),
+                      cond, cb);
+  } else if (type == framework::proto::VarType_Type_READER) {
+    ch->AddToReceiveQ(referrer, var->GetMutable<framework::ReaderHolder>(),
+                      cond, cb);
+  } else if (type == framework::proto::VarType_Type_CHANNEL) {
+    ch->AddToReceiveQ(referrer, var->GetMutable<framework::ChannelHolder>(),
+                      cond, cb);
+  } else {
+    PADDLE_THROW("ChannelAddToReceiveQ:Unsupported type");
+  }
+}

paddle/fluid/operators/detection_output_op.cu.cc → paddle/fluid/operators/concurrency/channel_util.h

@@ -2,7 +2,7 @@
 
 Licensed under the Apache License, Version 2.0 (the "License");
 you may not use this file except in compliance with the License.
-Indicesou may obtain a copy of the License at
+You may obtain a copy of the License at
 
     http://www.apache.org/licenses/LICENSE-2.0
 
@@ -12,10 +12,27 @@ 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/fluid/operators/detection_output_op.h"
+#pragma once
 
-namespace ops = paddle::operators;
-REGISTER_OP_CUDA_KERNEL(
-    detection_output,
-    ops::DetectionOutputKernel<paddle::platform::CUDADeviceContext, float>,
-    ops::DetectionOutputKernel<paddle::platform::CUDADeviceContext, double>);
+#include "paddle/fluid/framework/channel.h"
+#include "paddle/fluid/framework/variable.h"
+
+namespace paddle {
+namespace operators {
+namespace concurrency {
+
+bool ChannelSend(framework::ChannelHolder *ch, framework::Variable *var);
+bool ChannelReceive(framework::ChannelHolder *ch, framework::Variable *var);
+
+void ChannelAddToSendQ(framework::ChannelHolder *ch, const void *referrer,
+                       framework::Variable *var,
+                       std::shared_ptr<std::condition_variable_any> cond,
+                       std::function<bool(framework::ChannelAction)> cb);
+void ChannelAddToReceiveQ(framework::ChannelHolder *ch, const void *referrer,
+                          framework::Variable *var,
+                          std::shared_ptr<std::condition_variable_any> cond,
+                          std::function<bool(framework::ChannelAction)> cb);
+
+}  // namespace concurrency
+}  // namespace operators
+}  // namespace paddle

paddle/fluid/operators/conv_cudnn_op.cu.cc

@@ -18,6 +18,7 @@ limitations under the License. */
 #include "paddle/fluid/operators/conv_op.h"
 #include "paddle/fluid/platform/assert.h"
 #include "paddle/fluid/platform/cudnn_helper.h"
+#include "paddle/fluid/platform/float16.h"
 
 namespace paddle {
 namespace operators {
@@ -27,6 +28,8 @@ using ScopedTensorDescriptor = platform::ScopedTensorDescriptor;
 using ScopedFilterDescriptor = platform::ScopedFilterDescriptor;
 using ScopedConvolutionDescriptor = platform::ScopedConvolutionDescriptor;
 using DataLayout = platform::DataLayout;
+template <typename T>
+using ScalingParamType = typename platform::CudnnDataType<T>::ScalingParamType;
 
 static constexpr size_t kCONV_CUDNN_WORKSPACE_LIMIT_BYTES =
     static_cast<size_t>(1024) * 1024 * 1024;
@@ -133,7 +136,7 @@ class CUDNNConvOpKernel : public framework::OpKernel<T> {
     platform::CUDAPlace gpu = boost::get<platform::CUDAPlace>(ctx.GetPlace());
     cudnn_workspace = paddle::memory::Alloc(gpu, workspace_size_in_bytes);
     // ------------------- cudnn conv forward ---------------------
-    T alpha = 1.0f, beta = 0.0f;
+    ScalingParamType<T> alpha = 1.0f, beta = 0.0f;
     for (int i = 0; i < groups; i++) {
       PADDLE_ENFORCE(platform::dynload::cudnnConvolutionForward(
           handle, &alpha, cudnn_input_desc, input_data + i * group_offset_in,
@@ -280,7 +283,7 @@ class CUDNNConvGradOpKernel : public framework::OpKernel<T> {
     platform::CUDAPlace gpu = boost::get<platform::CUDAPlace>(ctx.GetPlace());
     cudnn_workspace = paddle::memory::Alloc(gpu, workspace_size_in_bytes);
     // ------------------- cudnn conv backward data ---------------------
-    T alpha = 1.0f, beta = 0.0f;
+    ScalingParamType<T> alpha = 1.0f, beta = 0.0f;
     if (input_grad) {
       T* input_grad_data = input_grad->mutable_data<T>(ctx.GetPlace());
       // Because beta is zero, it is unnecessary to reset input_grad.
@@ -315,16 +318,18 @@ class CUDNNConvGradOpKernel : public framework::OpKernel<T> {
 }  // namespace operators
 }  // namespace paddle
 
-REGISTER_OP_KERNEL(conv2d, CUDNN, ::paddle::platform::CUDAPlace,
+namespace plat = paddle::platform;
+REGISTER_OP_KERNEL(conv2d, CUDNN, plat::CUDAPlace,
                    paddle::operators::CUDNNConvOpKernel<float>,
-                   paddle::operators::CUDNNConvOpKernel<double>);
-REGISTER_OP_KERNEL(conv2d_grad, CUDNN, ::paddle::platform::CUDAPlace,
+                   paddle::operators::CUDNNConvOpKernel<double>,
+                   paddle::operators::CUDNNConvOpKernel<plat::float16>);
+REGISTER_OP_KERNEL(conv2d_grad, CUDNN, plat::CUDAPlace,
                    paddle::operators::CUDNNConvGradOpKernel<float>,
                    paddle::operators::CUDNNConvGradOpKernel<double>);
 
-REGISTER_OP_KERNEL(conv3d, CUDNN, ::paddle::platform::CUDAPlace,
+REGISTER_OP_KERNEL(conv3d, CUDNN, plat::CUDAPlace,
                    paddle::operators::CUDNNConvOpKernel<float>,
                    paddle::operators::CUDNNConvOpKernel<double>);
-REGISTER_OP_KERNEL(conv3d_grad, CUDNN, ::paddle::platform::CUDAPlace,
+REGISTER_OP_KERNEL(conv3d_grad, CUDNN, plat::CUDAPlace,
                    paddle::operators::CUDNNConvGradOpKernel<float>,
                    paddle::operators::CUDNNConvGradOpKernel<double>);

paddle/fluid/operators/conv_op.cc

@@ -83,12 +83,23 @@ framework::OpKernelType ConvOp::GetExpectedKernelType(
   }
 #endif
 
+  auto input_data_type =
+      framework::ToDataType(ctx.Input<Tensor>("Input")->type());
+  auto filter_data_type =
+      framework::ToDataType(ctx.Input<Tensor>("Filter")->type());
+  PADDLE_ENFORCE_EQ(input_data_type, filter_data_type,
+                    "input and filter data type should be consistent");
+
+  if (input_data_type == framework::proto::VarType::FP16) {
+    PADDLE_ENFORCE_EQ(library_, framework::LibraryType::kCUDNN,
+                      "float16 can only be used when CUDNN is used");
+  }
+
   std::string data_format = ctx.Attr<std::string>("data_format");
   // TODO(pzelazko-intel): enable MKLDNN layout when it's ready
   framework::DataLayout layout_ = framework::StringToDataLayout(data_format);
-  return framework::OpKernelType(
-      framework::ToDataType(ctx.Input<Tensor>("Input")->type()), ctx.GetPlace(),
-      layout_, library_);
+  return framework::OpKernelType(input_data_type, ctx.GetPlace(), layout_,
+                                 library_);
 }
 
 Conv2DOpMaker::Conv2DOpMaker(OpProto* proto, OpAttrChecker* op_checker)

paddle/fluid/operators/detection_output_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.
-Indicesou 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/fluid/operators/detection_output_op.h"
-namespace paddle {
-namespace operators {
-
-class DetectionOutputOpMaker : public framework::OpProtoAndCheckerMaker {
- public:
-  DetectionOutputOpMaker(OpProto* proto, OpAttrChecker* op_checker)
-      : OpProtoAndCheckerMaker(proto, op_checker) {
-    AddInput("Loc",
-             "(Tensor) The input tensor of detection_output operator."
-             "The input predict locations"
-             "The format of input tensor is kNCHW. Where K is priorbox point "
-             "numbers,"
-             "N is How many boxes are there on each point, "
-             "C is 4, H and W both are 1.");
-    AddInput("Conf",
-             "(Tensor) The input tensor of detection_output operator."
-             "The input priorbox confidence."
-             "The format of input tensor is kNCHW. Where K is priorbox point "
-             "numbers,"
-             "N is How many boxes are there on each point, "
-             "C is the number of classes, H and W both are 1.");
-    AddInput("PriorBox",
-             "(Tensor) The input tensor of detection_output operator."
-             "The format of input tensor is the position and variance "
-             "of the boxes");
-    AddOutput("Out",
-              "(Tensor) The output tensor of detection_output operator.");
-    AddAttr<int>("background_label_id", "(int), The background class index.");
-    AddAttr<int>("num_classes", "(int), The number of the classification.");
-    AddAttr<float>("nms_threshold",
-                   "(float), The Non-maximum suppression threshold.");
-    AddAttr<float>("confidence_threshold",
-                   "(float), The classification confidence threshold.");
-    AddAttr<int>("top_k", "(int), The bbox number kept of the layer’s output.");
-    AddAttr<int>("nms_top_k",
-                 "(int), The bbox number kept of the NMS’s output.");
-    AddComment(R"DOC(
-          detection output for SSD(single shot multibox detector)
-          Apply the NMS to the output of network and compute the predict
-          bounding box location. The output’s shape of this layer could
-          be zero if there is no valid bounding box.
-        )DOC");
-  }
-};
-
-class DetectionOutputOp : public framework::OperatorWithKernel {
- public:
-  using framework::OperatorWithKernel::OperatorWithKernel;
-  void InferShape(framework::InferShapeContext* ctx) const override {
-    PADDLE_ENFORCE(ctx->HasInput("Loc"),
-                   "Input(X) of DetectionOutputOp"
-                   "should not be null.");
-    PADDLE_ENFORCE(ctx->HasInput("Conf"),
-                   "Input(X) of DetectionOutputOp"
-                   "should not be null.");
-    PADDLE_ENFORCE(ctx->HasInput("PriorBox"),
-                   "Input(X) of DetectionOutputOp"
-                   "should not be null.");
-    PADDLE_ENFORCE(ctx->HasOutput("Out"),
-                   "Output(Out) of DetectionOutputOp should not be null.");
-    std::vector<int64_t> output_shape({1, 7});
-    ctx->SetOutputDim("Out", framework::make_ddim(output_shape));
-  }
-};
-}  // namespace operators
-}  // namespace paddle
-
-namespace ops = paddle::operators;
-REGISTER_OP_WITHOUT_GRADIENT(detection_output, ops::DetectionOutputOp,
-                             ops::DetectionOutputOpMaker);
-REGISTER_OP_CPU_KERNEL(
-    detection_output,
-    ops::DetectionOutputKernel<paddle::platform::CPUDeviceContext, float>,
-    ops::DetectionOutputKernel<paddle::platform::CPUDeviceContext, double>);

paddle/fluid/operators/detection_output_op.h

-/* 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.
-   Indicesou 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/fluid/framework/op_registry.h"
-#include "paddle/fluid/framework/tensor.h"
-#include "paddle/fluid/operators/math/detection_util.h"
-#include "paddle/fluid/operators/math/math_function.h"
-#include "paddle/fluid/operators/math/softmax.h"
-#include "paddle/fluid/operators/strided_memcpy.h"
-namespace paddle {
-namespace operators {
-template <typename DeviceContext, typename T>
-inline void transpose_fun(const framework::ExecutionContext& context,
-                          const framework::Tensor& src,
-                          framework::Tensor* dst) {
-  int input_nums = src.dims()[0];
-  int offset = 0;
-  for (int j = 0; j < input_nums; ++j) {
-    framework::Tensor in_p_tensor = src.Slice(j, j + 1);
-    std::vector<int64_t> shape_vec(
-        {in_p_tensor.dims()[0], in_p_tensor.dims()[1], in_p_tensor.dims()[3],
-         in_p_tensor.dims()[4], in_p_tensor.dims()[2]});
-    framework::DDim shape(framework::make_ddim(shape_vec));
-    framework::Tensor in_p_tensor_transpose;
-    in_p_tensor_transpose.mutable_data<T>(shape, context.GetPlace());
-    std::vector<int> shape_axis({0, 1, 3, 4, 2});
-    math::Transpose<DeviceContext, T, 5> trans5;
-    trans5(context.template device_context<DeviceContext>(), in_p_tensor,
-           &in_p_tensor_transpose, shape_axis);
-    auto dst_stride = framework::stride(dst->dims());
-    auto src_stride = framework::stride(in_p_tensor_transpose.dims());
-    StridedMemcpy<T>(context.device_context(), in_p_tensor_transpose.data<T>(),
-                     src_stride, in_p_tensor_transpose.dims(), dst_stride,
-                     dst->data<T>() + offset);
-    offset += in_p_tensor_transpose.dims()[4] * src_stride[4];
-  }
-}
-template <typename DeviceContext, typename T>
-class DetectionOutputKernel : public framework::OpKernel<T> {
- public:
-  void Compute(const framework::ExecutionContext& context) const override {
-    const framework::Tensor* in_loc = context.Input<framework::Tensor>("Loc");
-    const framework::Tensor* in_conf = context.Input<framework::Tensor>("Conf");
-    const framework::Tensor* in_priorbox =
-        context.Input<framework::Tensor>("PriorBox");
-    auto* out = context.Output<framework::Tensor>("Out");
-    int num_classes = context.template Attr<int>("num_classes");
-    int top_k = context.template Attr<int>("top_k");
-    int nms_top_k = context.template Attr<int>("nms_top_k");
-    int background_label_id = context.template Attr<int>("background_label_id");
-    float nms_threshold = context.template Attr<float>("nms_threshold");
-    float confidence_threshold =
-        context.template Attr<float>("confidence_threshold");
-    size_t batch_size = in_conf->dims()[1];
-    int conf_sum_size = in_conf->numel();
-    // for softmax
-    std::vector<int64_t> conf_shape_softmax_vec(
-        {conf_sum_size / num_classes, num_classes});
-    framework::DDim conf_shape_softmax(
-        framework::make_ddim(conf_shape_softmax_vec));
-    // for knchw => nhwc
-    std::vector<int64_t> loc_shape_vec({1, in_loc->dims()[1], in_loc->dims()[3],
-                                        in_loc->dims()[4],
-                                        in_loc->dims()[2] * in_loc->dims()[0]});
-    std::vector<int64_t> conf_shape_vec(
-        {1, in_conf->dims()[1], in_conf->dims()[3], in_conf->dims()[4],
-         in_conf->dims()[2] * in_conf->dims()[0]});
-    framework::DDim loc_shape(framework::make_ddim(loc_shape_vec));
-    framework::DDim conf_shape(framework::make_ddim(conf_shape_vec));
-    framework::Tensor loc_tensor;
-    framework::Tensor conf_tensor;
-    loc_tensor.mutable_data<T>(loc_shape, context.GetPlace());
-    conf_tensor.mutable_data<T>(conf_shape, context.GetPlace());
-    // for cpu
-    framework::Tensor loc_cpu;
-    framework::Tensor conf_cpu;
-    framework::Tensor priorbox_cpu;
-    const T* priorbox_data = in_priorbox->data<T>();
-    transpose_fun<DeviceContext, T>(context, *in_loc, &loc_tensor);
-    transpose_fun<DeviceContext, T>(context, *in_conf, &conf_tensor);
-    conf_tensor.Resize(conf_shape_softmax);
-    math::SoftmaxFunctor<DeviceContext, T>()(
-        context.template device_context<DeviceContext>(), &conf_tensor,
-        &conf_tensor);
-    T* loc_data = loc_tensor.data<T>();
-    T* conf_data = conf_tensor.data<T>();
-    if (platform::is_gpu_place(context.GetPlace())) {
-      loc_cpu.mutable_data<T>(loc_tensor.dims(), platform::CPUPlace());
-      framework::TensorCopy(loc_tensor, platform::CPUPlace(),
-                            context.device_context(), &loc_cpu);
-      loc_data = loc_cpu.data<T>();
-      conf_cpu.mutable_data<T>(conf_tensor.dims(), platform::CPUPlace());
-      framework::TensorCopy(conf_tensor, platform::CPUPlace(),
-                            context.device_context(), &conf_cpu);
-      conf_data = conf_cpu.data<T>();
-      priorbox_cpu.mutable_data<T>(in_priorbox->dims(), platform::CPUPlace());
-      framework::TensorCopy(*in_priorbox, platform::CPUPlace(),
-                            context.device_context(), &priorbox_cpu);
-      priorbox_data = priorbox_cpu.data<T>();
-    }
-    // get decode bboxes
-    size_t num_priors = in_priorbox->numel() / 8;
-    std::vector<std::vector<operators::math::BBox<T>>> all_decoded_bboxes;
-    for (size_t n = 0; n < batch_size; ++n) {
-      std::vector<operators::math::BBox<T>> decoded_bboxes;
-      for (size_t i = 0; i < num_priors; ++i) {
-        size_t prior_offset = i * 8;
-        size_t loc_pred_offset = n * num_priors * 4 + i * 4;
-        std::vector<math::BBox<T>> prior_bbox_vec;
-        math::GetBBoxFromPriorData<T>(priorbox_data + prior_offset, 1,
-                                      prior_bbox_vec);
-        std::vector<std::vector<T>> prior_bbox_var;
-        math::GetBBoxVarFromPriorData<T>(priorbox_data + prior_offset, 1,
-                                         prior_bbox_var);
-        std::vector<T> loc_pred_data;
-        for (size_t j = 0; j < 4; ++j)
-          loc_pred_data.push_back(*(loc_data + loc_pred_offset + j));
-        math::BBox<T> bbox = math::DecodeBBoxWithVar<T>(
-            prior_bbox_vec[0], prior_bbox_var[0], loc_pred_data);
-        decoded_bboxes.push_back(bbox);
-      }
-      all_decoded_bboxes.push_back(decoded_bboxes);
-    }
-    std::vector<std::map<size_t, std::vector<size_t>>> all_indices;
-    int num_kept = math::GetDetectionIndices<T>(
-        conf_data, num_priors, num_classes, background_label_id, batch_size,
-        confidence_threshold, nms_top_k, nms_threshold, top_k,
-        all_decoded_bboxes, &all_indices);
-
-    if (num_kept <= 0) {
-      std::vector<int64_t> out_shape_vec({0, 0});
-      framework::DDim out_shape(framework::make_ddim(out_shape_vec));
-      out->Resize(out_shape);
-      return;
-    }
-    std::vector<int64_t> out_shape_vec({num_kept, 7});
-    framework::DDim out_shape(framework::make_ddim(out_shape_vec));
-    out->mutable_data<T>(out_shape, context.GetPlace());
-    framework::Tensor out_cpu;
-    T* out_data = out->data<T>();
-    if (platform::is_gpu_place(context.GetPlace())) {
-      out_cpu.mutable_data<T>(out->dims(), platform::CPUPlace());
-      out_data = out_cpu.data<T>();
-    }
-    math::GetDetectionOutput<T>(conf_data, num_kept, num_priors, num_classes,
-                                batch_size, all_indices, all_decoded_bboxes,
-                                out_data);
-    if (platform::is_gpu_place(context.GetPlace())) {
-      framework::TensorCopy(out_cpu, platform::CUDAPlace(),
-                            context.device_context(), out);
-    }
-  }
-};
-}  // namespace operators
-}  // namespace paddle

paddle/fluid/operators/lrn_op.cc

@@ -36,7 +36,7 @@ struct LRNFunctor<platform::CPUDeviceContext, T> {
     auto e_x = framework::EigenTensor<T, 4>::From(input);
     for (int m = 0; m < N; m++) {
       for (int i = 0; i < C; i++) {
-        for (int c = start; c <= end; c++) {
+        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}}),
@@ -92,7 +92,7 @@ struct LRNGradFunctor<platform::CPUDeviceContext, T> {
                                  Eigen::array<int, 4>({{1, 1, H, W}}));
 
         i_x_g = i_mid.pow(-beta) * i_out_g;
-        for (int c = start; c <= end; c++) {
+        for (int c = start; c < end; c++) {
           int ch = i + c;
           if (ch < 0 || ch >= C) {
             continue;

paddle/fluid/operators/math/CMakeLists.txt

@@ -38,7 +38,7 @@ math_library(lstm_compute DEPS activation_functions)
 math_library(math_function DEPS cblas)
 math_library(maxouting)
 math_library(pooling)
-math_library(selected_rows_functor DEPS selected_rows)
+math_library(selected_rows_functor DEPS selected_rows math_function)
 math_library(sequence2batch)
 math_library(sequence_padding)
 math_library(sequence_pooling DEPS math_function)

paddle/fluid/operators/math/detection_util.h

-/* 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. */
-#pragma once
-#include <map>
-#include "paddle/fluid/framework/selected_rows.h"
-#include "paddle/fluid/platform/device_context.h"
-
-namespace paddle {
-namespace operators {
-namespace math {
-template <typename T>
-struct BBox {
-  BBox(T x_min, T y_min, T x_max, T y_max)
-      : x_min(x_min),
-        y_min(y_min),
-        x_max(x_max),
-        y_max(y_max),
-        is_difficult(false) {}
-
-  BBox() {}
-
-  T get_width() const { return x_max - x_min; }
-
-  T get_height() const { return y_max - y_min; }
-
-  T get_center_x() const { return (x_min + x_max) / 2; }
-
-  T get_center_y() const { return (y_min + y_max) / 2; }
-
-  T get_area() const { return get_width() * get_height(); }
-
-  // coordinate of bounding box
-  T x_min;
-  T y_min;
-  T x_max;
-  T y_max;
-  // whether difficult object (e.g. object with heavy occlusion is difficult)
-  bool is_difficult;
-};
-// KNCHW ==> NHWC
-// template <typename T>
-template <typename T>
-void GetBBoxFromPriorData(const T* prior_data, const size_t num_bboxes,
-                          std::vector<BBox<T>>& bbox_vec);
-template <typename T>
-void GetBBoxVarFromPriorData(const T* prior_data, const size_t num,
-                             std::vector<std::vector<T>>& var_vec);
-template <typename T>
-BBox<T> DecodeBBoxWithVar(BBox<T>& prior_bbox,
-                          const std::vector<T>& prior_bbox_var,
-                          const std::vector<T>& loc_pred_data);
-template <typename T1, typename T2>
-bool SortScorePairDescend(const std::pair<T1, T2>& pair1,
-                          const std::pair<T1, T2>& pair2);
-template <typename T>
-bool SortScorePairDescend(const std::pair<T, BBox<T>>& pair1,
-                          const std::pair<T, BBox<T>>& pair2);
-template <typename T>
-T jaccard_overlap(const BBox<T>& bbox1, const BBox<T>& bbox2);
-
-template <typename T>
-void ApplyNmsFast(const std::vector<BBox<T>>& bboxes, const T* conf_score_data,
-                  size_t class_idx, size_t top_k, T conf_threshold,
-                  T nms_threshold, size_t num_priors, size_t num_classes,
-                  std::vector<size_t>* indices);
-template <typename T>
-int GetDetectionIndices(
-    const T* conf_data, const size_t num_priors, const size_t num_classes,
-    const size_t background_label_id, const size_t batch_size,
-    const T conf_threshold, const size_t nms_top_k, const T nms_threshold,
-    const size_t top_k,
-    const std::vector<std::vector<BBox<T>>>& all_decoded_bboxes,
-    std::vector<std::map<size_t, std::vector<size_t>>>* all_detection_indices);
-template <typename T>
-BBox<T> ClipBBox(const BBox<T>& bbox);
-template <typename T>
-void GetDetectionOutput(
-    const T* conf_data, const size_t num_kept, const size_t num_priors,
-    const size_t num_classes, const size_t batch_size,
-    const std::vector<std::map<size_t, std::vector<size_t>>>& all_indices,
-    const std::vector<std::vector<BBox<T>>>& all_decoded_bboxes, T* out_data);
-template <typename T>
-void GetBBoxFromPriorData(const T* prior_data, const size_t num_bboxes,
-                          std::vector<BBox<T>>& bbox_vec) {
-  size_t out_offset = bbox_vec.size();
-  bbox_vec.resize(bbox_vec.size() + num_bboxes);
-  for (size_t i = 0; i < num_bboxes; ++i) {
-    BBox<T> bbox;
-    bbox.x_min = *(prior_data + i * 8);
-    bbox.y_min = *(prior_data + i * 8 + 1);
-    bbox.x_max = *(prior_data + i * 8 + 2);
-    bbox.y_max = *(prior_data + i * 8 + 3);
-    bbox_vec[out_offset + i] = bbox;
-  }
-}
-template <typename T>
-void GetBBoxVarFromPriorData(const T* prior_data, const size_t num,
-                             std::vector<std::vector<T>>& var_vec) {
-  size_t out_offset = var_vec.size();
-  var_vec.resize(var_vec.size() + num);
-  for (size_t i = 0; i < num; ++i) {
-    std::vector<T> var;
-    var.push_back(*(prior_data + i * 8 + 4));
-    var.push_back(*(prior_data + i * 8 + 5));
-    var.push_back(*(prior_data + i * 8 + 6));
-    var.push_back(*(prior_data + i * 8 + 7));
-    var_vec[out_offset + i] = var;
-  }
-}
-template <typename T>
-BBox<T> DecodeBBoxWithVar(BBox<T>& prior_bbox,
-                          const std::vector<T>& prior_bbox_var,
-                          const std::vector<T>& loc_pred_data) {
-  T prior_bbox_width = prior_bbox.get_width();
-  T prior_bbox_height = prior_bbox.get_height();
-  T prior_bbox_center_x = prior_bbox.get_center_x();
-  T prior_bbox_center_y = prior_bbox.get_center_y();
-
-  T decoded_bbox_center_x =
-      prior_bbox_var[0] * loc_pred_data[0] * prior_bbox_width +
-      prior_bbox_center_x;
-  T decoded_bbox_center_y =
-      prior_bbox_var[1] * loc_pred_data[1] * prior_bbox_height +
-      prior_bbox_center_y;
-  T decoded_bbox_width =
-      std::exp(prior_bbox_var[2] * loc_pred_data[2]) * prior_bbox_width;
-  T decoded_bbox_height =
-      std::exp(prior_bbox_var[3] * loc_pred_data[3]) * prior_bbox_height;
-
-  BBox<T> decoded_bbox;
-  decoded_bbox.x_min = decoded_bbox_center_x - decoded_bbox_width / 2;
-  decoded_bbox.y_min = decoded_bbox_center_y - decoded_bbox_height / 2;
-  decoded_bbox.x_max = decoded_bbox_center_x + decoded_bbox_width / 2;
-  decoded_bbox.y_max = decoded_bbox_center_y + decoded_bbox_height / 2;
-
-  return decoded_bbox;
-}
-template <typename T1, typename T2>
-bool SortScorePairDescend(const std::pair<T1, T2>& pair1,
-                          const std::pair<T1, T2>& pair2) {
-  return pair1.first > pair2.first;
-}
-template <typename T>
-T jaccard_overlap(const BBox<T>& bbox1, const BBox<T>& bbox2) {
-  if (bbox2.x_min > bbox1.x_max || bbox2.x_max < bbox1.x_min ||
-      bbox2.y_min > bbox1.y_max || bbox2.y_max < bbox1.y_min) {
-    return 0.0;
-  } else {
-    T inter_x_min = std::max(bbox1.x_min, bbox2.x_min);
-    T inter_y_min = std::max(bbox1.y_min, bbox2.y_min);
-    T interX_max = std::min(bbox1.x_max, bbox2.x_max);
-    T interY_max = std::min(bbox1.y_max, bbox2.y_max);
-
-    T inter_width = interX_max - inter_x_min;
-    T inter_height = interY_max - inter_y_min;
-    T inter_area = inter_width * inter_height;
-
-    T bbox_area1 = bbox1.get_area();
-    T bbox_area2 = bbox2.get_area();
-
-    return inter_area / (bbox_area1 + bbox_area2 - inter_area);
-  }
-}
-
-template <typename T>
-void ApplyNmsFast(const std::vector<BBox<T>>& bboxes, const T* conf_score_data,
-                  size_t class_idx, size_t top_k, T conf_threshold,
-                  T nms_threshold, size_t num_priors, size_t num_classes,
-                  std::vector<size_t>* indices) {
-  std::vector<std::pair<T, size_t>> scores;
-  for (size_t i = 0; i < num_priors; ++i) {
-    size_t conf_offset = i * num_classes + class_idx;
-    if (conf_score_data[conf_offset] > conf_threshold)
-      scores.push_back(std::make_pair(conf_score_data[conf_offset], i));
-  }
-  std::stable_sort(scores.begin(), scores.end(),
-                   SortScorePairDescend<T, size_t>);
-  if (top_k > 0 && top_k < scores.size()) scores.resize(top_k);
-  while (scores.size() > 0) {
-    const size_t idx = scores.front().second;
-    bool keep = true;
-    for (size_t i = 0; i < indices->size(); ++i) {
-      if (keep) {
-        const size_t saved_idx = (*indices)[i];
-        T overlap = jaccard_overlap<T>(bboxes[idx], bboxes[saved_idx]);
-        keep = overlap <= nms_threshold;
-      } else {
-        break;
-      }
-    }
-    if (keep) indices->push_back(idx);
-    scores.erase(scores.begin());
-  }
-}
-template <typename T>
-int GetDetectionIndices(
-    const T* conf_data, const size_t num_priors, const size_t num_classes,
-    const size_t background_label_id, const size_t batch_size,
-    const T conf_threshold, const size_t nms_top_k, const T nms_threshold,
-    const size_t top_k,
-    const std::vector<std::vector<BBox<T>>>& all_decoded_bboxes,
-    std::vector<std::map<size_t, std::vector<size_t>>>* all_detection_indices) {
-  int total_keep_num = 0;
-  for (size_t n = 0; n < batch_size; ++n) {
-    const std::vector<BBox<T>>& decoded_bboxes = all_decoded_bboxes[n];
-    size_t num_detected = 0;
-    std::map<size_t, std::vector<size_t>> indices;
-    size_t conf_offset = n * num_priors * num_classes;
-    for (size_t c = 0; c < num_classes; ++c) {
-      if (c == background_label_id) continue;
-      ApplyNmsFast<T>(decoded_bboxes, conf_data + conf_offset, c, nms_top_k,
-                      conf_threshold, nms_threshold, num_priors, num_classes,
-                      &(indices[c]));
-      num_detected += indices[c].size();
-    }
-    if (top_k > 0 && num_detected > top_k) {
-      // std::vector<pair<T,T>> score_index_pairs;
-      std::vector<std::pair<T, std::pair<size_t, size_t>>> score_index_pairs;
-      for (size_t c = 0; c < num_classes; ++c) {
-        const std::vector<size_t>& label_indices = indices[c];
-        for (size_t i = 0; i < label_indices.size(); ++i) {
-          size_t idx = label_indices[i];
-          score_index_pairs.push_back(
-              std::make_pair((conf_data + conf_offset)[idx * num_classes + c],
-                             std::make_pair(c, idx)));
-        }
-      }
-      std::sort(score_index_pairs.begin(), score_index_pairs.end(),
-                SortScorePairDescend<T, std::pair<size_t, size_t>>);
-      score_index_pairs.resize(top_k);
-      std::map<size_t, std::vector<size_t>> new_indices;
-      for (size_t i = 0; i < score_index_pairs.size(); ++i) {
-        size_t label = score_index_pairs[i].second.first;
-        size_t idx = score_index_pairs[i].second.second;
-        new_indices[label].push_back(idx);
-      }
-      all_detection_indices->push_back(new_indices);
-      total_keep_num += top_k;
-    } else {
-      all_detection_indices->push_back(indices);
-      total_keep_num += num_detected;
-    }
-  }
-  return total_keep_num;
-}
-template <typename T>
-BBox<T> ClipBBox(const BBox<T>& bbox) {
-  T one = static_cast<T>(1.0);
-  T zero = static_cast<T>(0.0);
-  BBox<T> clipped_bbox;
-  clipped_bbox.x_min = std::max(std::min(bbox.x_min, one), zero);
-  clipped_bbox.y_min = std::max(std::min(bbox.y_min, one), zero);
-  clipped_bbox.x_max = std::max(std::min(bbox.x_max, one), zero);
-  clipped_bbox.y_max = std::max(std::min(bbox.y_max, one), zero);
-  return clipped_bbox;
-}
-template <typename T>
-void GetDetectionOutput(
-    const T* conf_data, const size_t num_kept, const size_t num_priors,
-    const size_t num_classes, const size_t batch_size,
-    const std::vector<std::map<size_t, std::vector<size_t>>>& all_indices,
-    const std::vector<std::vector<BBox<T>>>& all_decoded_bboxes, T* out_data) {
-  size_t count = 0;
-  for (size_t n = 0; n < batch_size; ++n) {
-    for (std::map<size_t, std::vector<size_t>>::const_iterator it =
-             all_indices[n].begin();
-         it != all_indices[n].end(); ++it) {
-      size_t label = it->first;
-      const std::vector<size_t>& indices = it->second;
-      const std::vector<BBox<T>>& decoded_bboxes = all_decoded_bboxes[n];
-      for (size_t i = 0; i < indices.size(); ++i) {
-        size_t idx = indices[i];
-        size_t conf_offset = n * num_priors * num_classes + idx * num_classes;
-        out_data[count * 7] = n;
-        out_data[count * 7 + 1] = label;
-        out_data[count * 7 + 2] = (conf_data + conf_offset)[label];
-        BBox<T> clipped_bbox = ClipBBox<T>(decoded_bboxes[idx]);
-        out_data[count * 7 + 3] = clipped_bbox.x_min;
-        out_data[count * 7 + 4] = clipped_bbox.y_min;
-        out_data[count * 7 + 5] = clipped_bbox.x_max;
-        out_data[count * 7 + 6] = clipped_bbox.y_max;
-        ++count;
-      }
-    }
-  }
-}
-}  // namespace math
-}  // namespace operators
-}  // namespace paddle

paddle/fluid/operators/math/softmax.cu

@@ -14,13 +14,86 @@ limitations under the License. */
 
 #define EIGEN_USE_GPU
 
+#include "paddle/fluid/operators/math/math_function.h"
 #include "paddle/fluid/operators/math/softmax.h"
 #include "paddle/fluid/operators/math/softmax_impl.h"
+#include "paddle/fluid/platform/cudnn_helper.h"
 
 namespace paddle {
 namespace operators {
 namespace math {
 
+using Tensor = framework::Tensor;
+using ScopedTensorDescriptor = platform::ScopedTensorDescriptor;
+using DataLayout = platform::DataLayout;
+template <typename T>
+using CudnnDataType = platform::CudnnDataType<T>;
+
+template <typename T>
+void SoftmaxCUDNNFunctor<T>::operator()(
+    const platform::CUDADeviceContext& context, const framework::Tensor* X,
+    framework::Tensor* Y) {
+  // ------------------- cudnn descriptors ---------------------
+  ScopedTensorDescriptor xDesc;
+  ScopedTensorDescriptor yDesc;
+  std::vector<int> cudnn_tensor_dims = framework::vectorize2int(X->dims());
+  DataLayout layout = DataLayout::kNCHW;
+  if (cudnn_tensor_dims.size() == 5) {
+    layout = DataLayout::kNCDHW;
+  }
+  // NOTE(*) : cudnn softmax only support >= 4D Tensor,
+  // fill 1 at unused dims
+  if (cudnn_tensor_dims.size() <= 2) {
+    cudnn_tensor_dims.resize(4, 1);
+  }
+  cudnnTensorDescriptor_t cudnn_x_desc =
+      xDesc.descriptor<T>(layout, cudnn_tensor_dims);
+  cudnnTensorDescriptor_t cudnn_y_desc =
+      xDesc.descriptor<T>(layout, cudnn_tensor_dims);
+  PADDLE_ENFORCE(platform::dynload::cudnnSoftmaxForward(
+      context.cudnn_handle(), CUDNN_SOFTMAX_ACCURATE,
+      CUDNN_SOFTMAX_MODE_INSTANCE, CudnnDataType<T>::kOne(), cudnn_x_desc,
+      X->data<T>(), CudnnDataType<T>::kZero(), cudnn_y_desc,
+      Y->mutable_data<T>(context.GetPlace())));
+}
+
+template <typename T>
+void SoftmaxGradCUDNNFunctor<T>::operator()(
+    const platform::CUDADeviceContext& context, const framework::Tensor* Y,
+    const framework::Tensor* YGrad, framework::Tensor* XGrad) {
+  // ------------------- cudnn descriptors ---------------------
+  ScopedTensorDescriptor yDesc;
+  ScopedTensorDescriptor dyDesc;
+  ScopedTensorDescriptor dxDesc;
+  std::vector<int> cudnn_tensor_dims = framework::vectorize2int(Y->dims());
+  DataLayout layout = DataLayout::kNCHW;
+  if (cudnn_tensor_dims.size() == 5) {
+    layout = DataLayout::kNCDHW;
+  }
+  // NOTE(*) : cudnn softmax only support >= 4D Tensor,
+  // fill 1 at unused dims
+  if (cudnn_tensor_dims.size() <= 2) {
+    cudnn_tensor_dims.resize(4, 1);
+  }
+  cudnnTensorDescriptor_t cudnn_y_desc =
+      yDesc.descriptor<T>(layout, cudnn_tensor_dims);
+  cudnnTensorDescriptor_t cudnn_xgrad_desc =
+      dxDesc.descriptor<T>(layout, cudnn_tensor_dims);
+  cudnnTensorDescriptor_t cudnn_ygrad_desc =
+      dyDesc.descriptor<T>(layout, cudnn_tensor_dims);
+  PADDLE_ENFORCE(platform::dynload::cudnnSoftmaxBackward(
+      context.cudnn_handle(), CUDNN_SOFTMAX_ACCURATE,
+      CUDNN_SOFTMAX_MODE_INSTANCE, CudnnDataType<T>::kOne(), cudnn_y_desc,
+      Y->data<T>(), cudnn_ygrad_desc, YGrad->data<T>(),
+      CudnnDataType<T>::kZero(), cudnn_xgrad_desc,
+      XGrad->mutable_data<T>(context.GetPlace())));
+}
+
+template class SoftmaxCUDNNFunctor<float>;
+template class SoftmaxCUDNNFunctor<double>;
+template class SoftmaxGradCUDNNFunctor<float>;
+template class SoftmaxGradCUDNNFunctor<double>;
+
 template class SoftmaxFunctor<platform::CUDADeviceContext, float>;
 template class SoftmaxFunctor<platform::CUDADeviceContext, double>;
 template class SoftmaxGradFunctor<platform::CUDADeviceContext, float>;

paddle/fluid/operators/math/softmax.h

@@ -33,6 +33,23 @@ class SoftmaxGradFunctor {
                   const framework::Tensor* y_grad, framework::Tensor* x_grad);
 };
 
+#ifdef PADDLE_WITH_CUDA
+template <typename T>
+class SoftmaxCUDNNFunctor {
+ public:
+  void operator()(const platform::CUDADeviceContext& context,
+                  const framework::Tensor* X, framework::Tensor* Y);
+};
+
+template <typename T>
+class SoftmaxGradCUDNNFunctor {
+ public:
+  void operator()(const platform::CUDADeviceContext& context,
+                  const framework::Tensor* Y, const framework::Tensor* y_grad,
+                  framework::Tensor* x_grad);
+};
+#endif
+
 }  // namespace math
 }  // namespace operators
 }  // namespace paddle

paddle/fluid/operators/parallel_do_op.cc

@@ -18,6 +18,7 @@ limitations under the License. */
 #include "paddle/fluid/framework/op_registry.h"
 #include "paddle/fluid/framework/threadpool.h"
 #include "paddle/fluid/operators/detail/safe_ref.h"
+#include "paddle/fluid/platform/profiler.h"
 
 namespace paddle {
 namespace operators {
@@ -158,7 +159,10 @@ class ParallelDoOp : public framework::OperatorBase {
       auto &place = places[place_idx];
       auto *cur_scope = sub_scopes[place_idx];
 
-      workers.emplace_back(framework::Async([program, cur_scope, place, block] {
+      workers.emplace_back(
+          framework::Async([program, cur_scope, place, block, place_idx] {
+            // Give the thread an id to distinguish parallel block with same id.
+            platform::RecordThread rt(static_cast<int>(place_idx) + 1);
             framework::Executor executor(place);
             executor.Run(*program, cur_scope, block->ID(),
                          false /*create_local_scope*/);
@@ -234,7 +238,10 @@ class ParallelDoGradOp : public framework::OperatorBase {
       auto *cur_scope = sub_scopes[i];
 
       // execute
-      workers.emplace_back(framework::Async([program, cur_scope, place, block] {
+      workers.emplace_back(
+          framework::Async([program, cur_scope, place, block, i] {
+            // Give the thread an id to distinguish parallel block with same id.
+            platform::RecordThread rt(static_cast<int>(i) + 1);
             framework::Executor executor(place);
             executor.Run(*program, cur_scope, block->ID(),
                          false /*create_local_scope*/);

paddle/fluid/operators/pool_cudnn_op.cu.cc

@@ -24,6 +24,8 @@ using ScopedTensorDescriptor = platform::ScopedTensorDescriptor;
 using ScopedPoolingDescriptor = platform::ScopedPoolingDescriptor;
 using DataLayout = platform::DataLayout;
 using PoolingMode = platform::PoolingMode;
+template <typename T>
+using ScalingParamType = typename platform::CudnnDataType<T>::ScalingParamType;
 
 template <typename T>
 class PoolCUDNNOpKernel : public framework::OpKernel<T> {
@@ -78,8 +80,7 @@ class PoolCUDNNOpKernel : public framework::OpKernel<T> {
 
     // ------------------- cudnn pool algorithm ---------------------
     auto handle = ctx.cuda_device_context().cudnn_handle();
-    T alpha = 1.0f, beta = 0.0f;
-
+    ScalingParamType<T> alpha = 1.0f, beta = 0.0f;
     PADDLE_ENFORCE(platform::dynload::cudnnPoolingForward(
         handle, cudnn_pool_desc, &alpha, cudnn_input_desc, input_data, &beta,
         cudnn_output_desc, output_data));
@@ -144,8 +145,7 @@ class PoolCUDNNGradOpKernel : public framework::OpKernel<T> {
 
     // ------------------- cudnn pool algorithm ---------------------
     auto handle = ctx.cuda_device_context().cudnn_handle();
-    T alpha = 1.0f, beta = 0.0f;
-
+    ScalingParamType<T> alpha = 1.0f, beta = 0.0f;
     if (input_grad) {
       T *input_grad_data = input_grad->mutable_data<T>(ctx.GetPlace());
       // Because beta is zero, it is unnecessary to reset input_grad.
@@ -162,17 +162,19 @@ class PoolCUDNNGradOpKernel : public framework::OpKernel<T> {
 }  // namespace paddle
 
 namespace ops = paddle::operators;
+namespace plat = paddle::platform;
 
-REGISTER_OP_KERNEL(pool2d, CUDNN, ::paddle::platform::CUDAPlace,
+REGISTER_OP_KERNEL(pool2d, CUDNN, plat::CUDAPlace,
                    ops::PoolCUDNNOpKernel<float>,
-                   ops::PoolCUDNNOpKernel<double>);
-REGISTER_OP_KERNEL(pool2d_grad, CUDNN, ::paddle::platform::CUDAPlace,
+                   ops::PoolCUDNNOpKernel<double>,
+                   ops::PoolCUDNNOpKernel<plat::float16>);
+REGISTER_OP_KERNEL(pool2d_grad, CUDNN, plat::CUDAPlace,
                    ops::PoolCUDNNGradOpKernel<float>,
                    ops::PoolCUDNNGradOpKernel<double>);
 
-REGISTER_OP_KERNEL(pool3d, CUDNN, ::paddle::platform::CUDAPlace,
+REGISTER_OP_KERNEL(pool3d, CUDNN, plat::CUDAPlace,
                    ops::PoolCUDNNOpKernel<float>,
                    ops::PoolCUDNNOpKernel<double>);
-REGISTER_OP_KERNEL(pool3d_grad, CUDNN, ::paddle::platform::CUDAPlace,
+REGISTER_OP_KERNEL(pool3d_grad, CUDNN, plat::CUDAPlace,
                    ops::PoolCUDNNGradOpKernel<float>,
                    ops::PoolCUDNNGradOpKernel<double>);

paddle/fluid/operators/pool_op.cc

@@ -124,11 +124,15 @@ framework::OpKernelType PoolOpGrad::GetExpectedKernelType(
   }
 #endif
 
+  auto input_data_type = framework::ToDataType(ctx.Input<Tensor>("X")->type());
+  if (input_data_type == framework::proto::VarType::FP16) {
+    PADDLE_ENFORCE_EQ(library_, framework::LibraryType::kCUDNN,
+                      "float16 can only be used when CUDNN is used");
+  }
   std::string data_format = ctx.Attr<std::string>("data_format");
   framework::DataLayout layout_ = framework::StringToDataLayout(data_format);
-  return framework::OpKernelType(
-      framework::ToDataType(ctx.Input<Tensor>("X")->type()), ctx.GetPlace(),
-      layout_, library_);
+  return framework::OpKernelType(input_data_type, ctx.GetPlace(), layout_,
+                                 library_);
 }
 
 Pool2dOpMaker::Pool2dOpMaker(OpProto *proto, OpAttrChecker *op_checker)

paddle/fluid/operators/select_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 <boost/tokenizer.hpp>
+#include <memory>
+#include <thread>
+#include <vector>
+#include "paddle/fluid/framework/channel.h"
+#include "paddle/fluid/framework/executor.h"
+#include "paddle/fluid/framework/lod_tensor.h"
+#include "paddle/fluid/framework/op_registry.h"
+#include "paddle/fluid/operators/concurrency/channel_util.h"
+
+namespace paddle {
+namespace operators {
+
+static constexpr char kX[] = "X";
+static constexpr char kCaseToExecute[] = "case_to_execute";
+
+static constexpr char kCases[] = "cases";
+static constexpr char kCasesBlock[] = "sub_block";
+
+class SelectOp : public framework::OperatorBase {
+ public:
+  SelectOp(const std::string &type, const framework::VariableNameMap &inputs,
+           const framework::VariableNameMap &outputs,
+           const framework::AttributeMap &attrs)
+      : framework::OperatorBase(type, inputs, outputs, attrs) {}
+
+ private:
+  enum class SelectOpCaseType {
+    DEFAULT = 0,
+    SEND = 1,
+    RECEIVE = 2,
+  };
+
+  struct SelectOpCase {
+    int caseIndex;
+    SelectOpCaseType caseType;
+    std::string channelName;
+    std::string varName;
+
+    SelectOpCase() {}
+
+    SelectOpCase(int caseIndex, SelectOpCaseType caseType,
+                 std::string channelName, std::string varName)
+        : caseIndex(caseIndex),
+          caseType(caseType),
+          channelName(channelName),
+          varName(varName) {}
+  };
+
+  void RunImpl(const framework::Scope &scope,
+               const platform::Place &dev_place) const override {
+    std::vector<std::string> casesConfigs =
+        Attr<std::vector<std::string>>(kCases);
+
+    framework::BlockDesc *casesBlock =
+        Attr<framework::BlockDesc *>(kCasesBlock);
+
+    framework::Scope &casesBlockScope = scope.NewScope();
+
+    std::string caseToExecuteVarName = Input(kCaseToExecute);
+    framework::Variable *caseToExecuteVar =
+        casesBlockScope.FindVar(caseToExecuteVarName);
+
+    // Construct cases from "conditional_block_op"(s) in the casesBlock
+    std::vector<std::shared_ptr<SelectOpCase>> cases =
+        ParseAndShuffleCases(&casesConfigs);
+
+    // Get all unique channels involved in select
+    std::set<framework::ChannelHolder *> channelsSet;
+    for (auto c : cases) {
+      if (!c->channelName.empty()) {
+        auto channelVar = scope.FindVar(c->channelName);
+        framework::ChannelHolder *ch =
+            channelVar->GetMutable<framework::ChannelHolder>();
+
+        if (channelsSet.find(ch) == channelsSet.end()) {
+          channelsSet.insert(ch);
+        }
+      }
+    }
+
+    // Order all channels by their pointer address
+    std::vector<framework::ChannelHolder *> channels(channelsSet.begin(),
+                                                     channelsSet.end());
+    std::sort(channels.begin(), channels.end());
+
+    // Poll all cases
+    int32_t caseToExecute = pollCases(&scope, &cases, channels);
+
+    // At this point, the case to execute has already been determined,
+    // so we can proceed with executing the cases block
+    framework::LoDTensor *caseToExecuteTensor =
+        caseToExecuteVar->GetMutable<framework::LoDTensor>();
+    caseToExecuteTensor->data<int32_t>()[0] = caseToExecute;
+
+    // Execute the cases block, only one case will be executed since we set the
+    // case_to_execute value to the index of the case we want to execute
+    framework::Executor executor(dev_place);
+    framework::ProgramDesc *program = casesBlock->Program();
+    executor.Run(*program, &casesBlockScope, casesBlock->ID(),
+                 false /*create_local_scope*/);
+  }
+
+  /**
+   * Goes through all operators in the casesConfigs and processes
+   * "conditional_block" operators.  These operators are mapped to our
+   * SelectOpCase objects.  We randomize the case orders, and set the
+   * default case (if any exists) as the last case)
+   * @param casesBlock
+   * @return
+   */
+  std::vector<std::shared_ptr<SelectOpCase>> ParseAndShuffleCases(
+      std::vector<std::string> *casesConfigs) const {
+    std::vector<std::shared_ptr<SelectOpCase>> cases;
+    std::shared_ptr<SelectOpCase> defaultCase;
+
+    if (casesConfigs != nullptr) {
+      boost::char_delimiters_separator<char> sep(false, ",", "");
+      for (std::vector<std::string>::iterator itr = casesConfigs->begin();
+           itr < casesConfigs->end(); ++itr) {
+        std::string caseConfig = *itr;
+        boost::tokenizer<> tokens(caseConfig, sep);
+
+        boost::tokenizer<>::iterator tok_iter = tokens.begin();
+        PADDLE_ENFORCE(tok_iter != tokens.end(), "Cannot get case index");
+        std::string caseIndexString = *tok_iter;
+        int caseIndex = std::stoi(caseIndexString);
+
+        ++tok_iter;
+        PADDLE_ENFORCE(tok_iter != tokens.end(), "Cannot get case type");
+        std::string caseTypeString = *tok_iter;
+        SelectOpCaseType caseType = (SelectOpCaseType)std::stoi(caseTypeString);
+
+        std::string caseChannel;
+        std::string caseChannelVar;
+
+        ++tok_iter;
+        if (caseType != SelectOpCaseType::DEFAULT) {
+          PADDLE_ENFORCE(tok_iter != tokens.end(), "Cannot get case channel");
+          caseChannel = *tok_iter;
+
+          ++tok_iter;
+          PADDLE_ENFORCE(tok_iter != tokens.end(),
+                         "Cannot get case channel variable");
+          caseChannelVar = *tok_iter;
+        }
+
+        auto c = std::make_shared<SelectOpCase>(caseIndex, caseType,
+                                                caseChannel, caseChannelVar);
+
+        if (caseType == SelectOpCaseType::DEFAULT) {
+          PADDLE_ENFORCE(defaultCase == nullptr,
+                         "Select can only contain one default case.");
+          defaultCase = c;
+        } else {
+          cases.push_back(c);
+        }
+      }
+    }
+
+    // Randomly sort cases, with default case being last
+    std::random_shuffle(cases.begin(), cases.end());
+    if (defaultCase != nullptr) {
+      cases.push_back(defaultCase);
+    }
+
+    return cases;
+  }
+
+  /**
+   * This method will recursively poll the cases and determines if any case
+   * condition is true.
+   * If none of the cases conditions are true (and there is no default case),
+   * then block
+   * the thread.  The thread may be woken up by a channel operation, at which
+   * point we
+   * execute the case.
+   * @param scope
+   * @param cases
+   * @param channels
+   * @return
+   */
+  int32_t pollCases(const framework::Scope *scope,
+                    std::vector<std::shared_ptr<SelectOpCase>> *cases,
+                    std::vector<framework::ChannelHolder *> channels) const {
+    // Lock all involved channels
+    lockChannels(channels);
+
+    std::atomic<int> caseToExecute(-1);
+
+    std::vector<std::shared_ptr<SelectOpCase>>::iterator it = cases->begin();
+    while (it != cases->end()) {
+      std::shared_ptr<SelectOpCase> c = *it;
+
+      auto chVar = scope->FindVar(c->channelName);
+      framework::ChannelHolder *ch =
+          chVar->GetMutable<framework::ChannelHolder>();
+
+      switch (c->caseType) {
+        case SelectOpCaseType::SEND:
+          PADDLE_ENFORCE(!ch->IsClosed(), "Cannot send to a closed channel");
+          if (ch->CanSend()) {
+            // We can send to channel directly, send the data to channel
+            // and execute case
+            auto chVar = scope->FindVar(c->varName);
+            concurrency::ChannelSend(ch, chVar);
+            caseToExecute = c->caseIndex;
+          }
+          break;
+        case SelectOpCaseType::RECEIVE:
+          if (ch->CanReceive()) {
+            // We can receive from channel directly, send the data to channel
+            // and execute case
+            auto chVar = scope->FindVar(c->varName);
+            concurrency::ChannelReceive(ch, chVar);
+            caseToExecute = c->caseIndex;
+          }
+          break;
+        case SelectOpCaseType::DEFAULT:
+          caseToExecute = c->caseIndex;
+          break;
+      }
+
+      if (caseToExecute != -1) {
+        // We found a case to execute, stop looking at other case statements
+        break;
+      }
+
+      ++it;
+    }
+
+    if (caseToExecute == -1) {
+      // None of the cases are eligible to execute, enqueue current thread
+      // into all the sending/receiving queue of each involved channel
+      std::atomic<bool> completed(false);
+      std::recursive_mutex mutex;
+      std::unique_lock<std::recursive_mutex> lock{mutex};
+      // std::condition_variable_any selectCond;
+      auto selectCond = std::make_shared<std::condition_variable_any>();
+
+      std::recursive_mutex callbackMutex;
+      pushThreadOnChannelQueues(scope, cases, selectCond, caseToExecute,
+                                completed, callbackMutex);
+
+      // TODO(thuan): Atomically unlock all channels and sleep current thread
+      unlockChannels(channels);
+      selectCond->wait(lock, [&completed]() { return completed.load(); });
+
+      // Select has been woken up by case operation
+      lockChannels(channels);
+      removeThreadOnChannelQueues(scope, cases);
+
+      if (caseToExecute == -1) {
+        // Recursively poll cases, since we were woken up by a channel close
+        // TODO(thuan): Need to test if this is a valid case
+        unlockChannels(channels);
+        return pollCases(scope, cases, channels);
+      }
+    }
+
+    // At this point, caseToExecute != -1, and we can proceed with executing
+    // the case block
+    unlockChannels(channels);
+
+    return caseToExecute;
+  }
+
+  void lockChannels(std::vector<framework::ChannelHolder *> chs) const {
+    std::vector<framework::ChannelHolder *>::iterator it = chs.begin();
+    while (it != chs.end()) {
+      framework::ChannelHolder *ch = *it;
+      ch->Lock();
+      ++it;
+    }
+  }
+
+  void unlockChannels(std::vector<framework::ChannelHolder *> chs) const {
+    std::vector<framework::ChannelHolder *>::reverse_iterator it = chs.rbegin();
+    while (it != chs.rend()) {
+      framework::ChannelHolder *ch = *it;
+      ch->Unlock();
+      ++it;
+    }
+  }
+
+  void pushThreadOnChannelQueues(
+      const framework::Scope *scope,
+      std::vector<std::shared_ptr<SelectOpCase>> *cases,
+      std::shared_ptr<std::condition_variable_any> rCond,
+      std::atomic<int> &caseToExecute, std::atomic<bool> &completed,
+      std::recursive_mutex &callbackMutex) const {
+    std::vector<std::shared_ptr<SelectOpCase>>::iterator it = cases->begin();
+    while (it != cases->end()) {
+      std::shared_ptr<SelectOpCase> c = *it;
+
+      auto chVar = scope->FindVar(c->channelName);
+      framework::ChannelHolder *ch =
+          chVar->GetMutable<framework::ChannelHolder>();
+
+      std::function<bool(framework::ChannelAction channelAction)> cb =
+          [&caseToExecute, &completed, &callbackMutex,
+           c](framework::ChannelAction channelAction) {
+            std::lock_guard<std::recursive_mutex> lock{callbackMutex};
+
+            bool canProcess = false;
+            if (!completed) {
+              // If the channel wasn't closed, we set the caseToExecute index
+              // as this current case
+              if (channelAction != framework::ChannelAction::CLOSE) {
+                caseToExecute = c->caseIndex;
+              }
+              // This will allow our conditional variable to break out of wait
+              completed = true;
+              canProcess = true;
+            }
+
+            return canProcess;
+          };
+
+      switch (c->caseType) {
+        case SelectOpCaseType::SEND: {
+          auto chOutputVar = scope->FindVar(c->varName);
+          concurrency::ChannelAddToSendQ(ch, this, chOutputVar, rCond, cb);
+          break;
+        }
+        case SelectOpCaseType::RECEIVE: {
+          auto chOutputVar = scope->FindVar(c->varName);
+          concurrency::ChannelAddToReceiveQ(ch, this, chOutputVar, rCond, cb);
+          break;
+        }
+        default:
+          break;
+      }
+      ++it;
+    }
+  }
+
+  void removeThreadOnChannelQueues(
+      const framework::Scope *scope,
+      std::vector<std::shared_ptr<SelectOpCase>> *cases) const {
+    std::vector<std::shared_ptr<SelectOpCase>>::iterator it = cases->begin();
+    while (it != cases->end()) {
+      std::shared_ptr<SelectOpCase> c = *it;
+
+      auto chVar = scope->FindVar(c->channelName);
+      framework::ChannelHolder *ch =
+          chVar->GetMutable<framework::ChannelHolder>();
+      switch (c->caseType) {
+        case SelectOpCaseType::SEND: {
+          ch->RemoveFromSendQ(this);
+          break;
+        }
+        case SelectOpCaseType::RECEIVE: {
+          ch->RemoveFromReceiveQ(this);
+          break;
+        }
+        default:
+          break;
+      }
+      ++it;
+    }
+  }
+};
+
+class SelectOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  SelectOpMaker(OpProto *proto, OpAttrChecker *op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput(kX,
+             "A set of variables, which are required by operators inside the "
+             "cases of Select Op")
+        .AsDuplicable();
+    AddInput(kCaseToExecute,
+             "(Int) The variable the sets the index of the case to execute, "
+             "after evaluating the channels being sent to and received from")
+        .AsDuplicable();
+    AddAttr<std::vector<std::string>>(kCases,
+                                      "(String vector) Serialized list of"
+                                      "all cases in the select op. Each"
+                                      "case is serialized as: "
+                                      "'<index>,<type>,<channel>,<value>'"
+                                      "where type is 0 for default, 1 for"
+                                      "send, and 2 for receive"
+                                      "No channel and values are needed for"
+                                      "default cases.");
+    AddAttr<framework::BlockDesc *>(kCasesBlock,
+                                    "The cases block inside select_op");
+    AddComment(R"DOC(
+)DOC");
+  }
+};
+
+// TODO(thuan): Implement Gradient Operator for SELECT_OP
+
+}  // namespace operators
+}  // namespace paddle
+
+REGISTER_OPERATOR(select, paddle::operators::SelectOp,
+                  paddle::framework::EmptyGradOpMaker,
+                  paddle::operators::SelectOpMaker);

paddle/fluid/operators/sequence_softmax_cudnn_op.cu.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/fluid/framework/op_registry.h"
+#include "paddle/fluid/operators/math/math_function.h"
+#include "paddle/fluid/operators/math/softmax.h"
+
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+using LoDTensor = framework::LoDTensor;
+
+template <typename T>
+class SequenceSoftmaxCUDNNKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    auto* x = ctx.Input<LoDTensor>("X");
+    auto* out = ctx.Output<LoDTensor>("Out");
+
+    auto lod = x->lod();
+    auto dims = x->dims();
+
+    const size_t level = lod.size() - 1;
+    PADDLE_ENFORCE_EQ(dims[0], static_cast<int64_t>(lod[level].back()),
+                      "The first dimension of Input(X) should be equal to the "
+                      "sum of all sequences' lengths.");
+    PADDLE_ENFORCE_EQ(dims[0], x->numel(),
+                      "The width of each timestep in Input(X) of "
+                      "SequenceSoftmaxOp should be 1.");
+
+    out->mutable_data<T>(ctx.GetPlace());
+    for (int i = 0; i < static_cast<int>(lod[level].size()) - 1; ++i) {
+      int start_pos = static_cast<int>(lod[level][i]);
+      int end_pos = static_cast<int>(lod[level][i + 1]);
+      Tensor x_i = x->Slice(start_pos, end_pos);
+      Tensor out_i = out->Slice(start_pos, end_pos);
+
+      // Reshape from (end_pos - start_pos) x 1UL to 1UL x (end_pos - start_pos)
+      framework::DDim dims_i =
+          // framework::make_ddim({1UL, end_pos - start_pos, 1UL, 1UL});
+          framework::make_ddim({1UL, end_pos - start_pos});
+      x_i.Resize(dims_i);
+      out_i.Resize(dims_i);
+      math::SoftmaxCUDNNFunctor<T>()(
+          ctx.template device_context<platform::CUDADeviceContext>(), &x_i,
+          &out_i);
+    }
+  }
+};
+
+template <typename T>
+class SequenceSoftmaxGradCUDNNKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    auto* out = ctx.Input<LoDTensor>("Out");
+    auto* out_grad = ctx.Input<LoDTensor>(framework::GradVarName("Out"));
+    auto* x = ctx.Input<LoDTensor>("X");
+    auto* x_grad = ctx.Output<LoDTensor>(framework::GradVarName("X"));
+
+    auto lod = x->lod();
+    const size_t level = lod.size() - 1;
+
+    x_grad->mutable_data<T>(ctx.GetPlace());
+    for (int i = 0; i < static_cast<int>(lod[level].size()) - 1; ++i) {
+      int start_pos = static_cast<int>(lod[level][i]);
+      int end_pos = static_cast<int>(lod[level][i + 1]);
+
+      Tensor out_i = out->Slice(start_pos, end_pos);
+      Tensor out_grad_i = out_grad->Slice(start_pos, end_pos);
+      Tensor x_grad_i = x_grad->Slice(start_pos, end_pos);
+
+      // Reshape from (end_pos - start_pos) x 1UL to 1UL x (end_pos - start_pos)
+      framework::DDim dims_i = framework::make_ddim({1UL, end_pos - start_pos});
+      out_i.Resize(dims_i);
+      out_grad_i.Resize(dims_i);
+      x_grad_i.Resize(dims_i);
+      math::SoftmaxGradCUDNNFunctor<T>()(
+          ctx.template device_context<platform::CUDADeviceContext>(), &out_i,
+          &out_grad_i, &x_grad_i);
+    }
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_KERNEL(sequence_softmax, CUDNN, ::paddle::platform::CUDAPlace,
+                   ops::SequenceSoftmaxCUDNNKernel<float>,
+                   ops::SequenceSoftmaxCUDNNKernel<double>)
+REGISTER_OP_KERNEL(sequence_softmax_grad, CUDNN, ::paddle::platform::CUDAPlace,
+                   ops::SequenceSoftmaxGradCUDNNKernel<float>,
+                   ops::SequenceSoftmaxGradCUDNNKernel<double>)

paddle/fluid/operators/sequence_softmax_op.cc

@@ -29,6 +29,29 @@ class SequenceSoftmaxOp : public framework::OperatorWithKernel {
     ctx->SetOutputDim("Out", ctx->GetInputDim("X"));
     ctx->ShareLoD("X", /*->*/ "Out");
   }
+
+ protected:
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext& ctx) const override {
+    // choose cudnn kernel if the runtime supported.
+    bool use_cudnn = ctx.Attr<bool>("use_cudnn");
+    bool runtime_cudnn_support = false;
+#ifdef PADDLE_WITH_CUDA
+    if (platform::is_gpu_place(ctx.GetPlace())) {
+      auto& dev_ctx =
+          ctx.template device_context<platform::CUDADeviceContext>();
+      runtime_cudnn_support = dev_ctx.cudnn_handle() != nullptr ? true : false;
+    }
+#endif
+    framework::LibraryType library_ = framework::LibraryType::kPlain;
+    if (use_cudnn && runtime_cudnn_support) {
+      library_ = framework::LibraryType::kCUDNN;
+    }
+    std::string data_format = ctx.Attr<std::string>("data_format");
+    return framework::OpKernelType(
+        framework::ToDataType(ctx.Input<Tensor>("X")->type()), ctx.GetPlace(),
+        framework::StringToDataLayout(data_format), library_);
+  }
 };
 
 class SequenceSoftmaxOpMaker : public framework::OpProtoAndCheckerMaker {
@@ -41,6 +64,17 @@ class SequenceSoftmaxOpMaker : public framework::OpProtoAndCheckerMaker {
     AddOutput("Out",
               "(LoDTensor) 1-D or 2-D output LoDTensor with the 2-nd dimension "
               "of length 1.");
+    AddAttr<bool>(
+        "use_cudnn",
+        "(bool, default false) Only used in cudnn kernel, need install cudnn")
+        .SetDefault(false);
+    AddAttr<std::string>(
+        "data_format",
+        "(string, default NCHW) Only used in "
+        "An optional string from: \"NHWC\", \"NCHW\". "
+        "Defaults to \"NHWC\". Specify the data format of the output data, "
+        "the input will be transformed automatically. ")
+        .SetDefault("AnyLayout");
     AddComment(R"DOC(
 Sequence Softmax Operator.
 
@@ -91,6 +125,29 @@ class SequenceSoftmaxGradOp : public framework::OperatorWithKernel {
 
     ctx->SetOutputDim(framework::GradVarName("X"), ctx->GetInputDim("X"));
   }
+
+ protected:
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext& ctx) const override {
+    // choose cudnn kernel if the runtime supported.
+    bool use_cudnn = ctx.Attr<bool>("use_cudnn");
+    bool runtime_cudnn_support = false;
+#ifdef PADDLE_WITH_CUDA
+    if (platform::is_gpu_place(ctx.GetPlace())) {
+      auto& dev_ctx =
+          ctx.template device_context<platform::CUDADeviceContext>();
+      runtime_cudnn_support = dev_ctx.cudnn_handle() != nullptr ? true : false;
+    }
+#endif
+    framework::LibraryType library_ = framework::LibraryType::kPlain;
+    if (use_cudnn && runtime_cudnn_support) {
+      library_ = framework::LibraryType::kCUDNN;
+    }
+    std::string data_format = ctx.Attr<std::string>("data_format");
+    return framework::OpKernelType(
+        framework::ToDataType(ctx.Input<Tensor>("X")->type()), ctx.GetPlace(),
+        framework::StringToDataLayout(data_format), library_);
+  }
 };
 
 }  // namespace operators
@@ -102,7 +159,9 @@ REGISTER_OP(sequence_softmax, ops::SequenceSoftmaxOp,
             ops::SequenceSoftmaxGradOp);
 REGISTER_OP_CPU_KERNEL(
     sequence_softmax,
-    ops::SequenceSoftmaxKernel<paddle::platform::CPUDeviceContext, float>);
+    ops::SequenceSoftmaxKernel<paddle::platform::CPUDeviceContext, float>,
+    ops::SequenceSoftmaxKernel<paddle::platform::CPUDeviceContext, double>);
 REGISTER_OP_CPU_KERNEL(
     sequence_softmax_grad,
-    ops::SequenceSoftmaxGradKernel<paddle::platform::CPUDeviceContext, float>);
+    ops::SequenceSoftmaxGradKernel<paddle::platform::CPUDeviceContext, float>,
+    ops::SequenceSoftmaxGradKernel<paddle::platform::CPUDeviceContext, double>);

paddle/fluid/operators/sequence_softmax_op.cu.cc

@@ -17,7 +17,10 @@ limitations under the License. */
 namespace ops = paddle::operators;
 REGISTER_OP_CUDA_KERNEL(
     sequence_softmax,
-    ops::SequenceSoftmaxKernel<paddle::platform::CUDADeviceContext, float>)
+    ops::SequenceSoftmaxKernel<paddle::platform::CUDADeviceContext, float>,
+    ops::SequenceSoftmaxKernel<paddle::platform::CUDADeviceContext, double>)
 REGISTER_OP_CUDA_KERNEL(
     sequence_softmax_grad,
-    ops::SequenceSoftmaxGradKernel<paddle::platform::CUDADeviceContext, float>);
+    ops::SequenceSoftmaxGradKernel<paddle::platform::CUDADeviceContext, float>,
+    ops::SequenceSoftmaxGradKernel<paddle::platform::CUDADeviceContext,
+                                   double>);

paddle/fluid/operators/softmax_cudnn_op.cu.cc

+/* Copyright (c) 2018 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/fluid/operators/math/softmax.h"
+#include "paddle/fluid/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+
+template <typename T>
+class SoftmaxCUDNNKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& context) const override {
+    auto* X = context.Input<Tensor>("X");
+    auto* Out = context.Output<Tensor>("Out");
+
+    // allocate memory on device.
+    Out->mutable_data<T>(context.GetPlace());
+
+    math::SoftmaxCUDNNFunctor<T>()(
+        context.template device_context<platform::CUDADeviceContext>(), X, Out);
+  }
+};
+
+template <typename T>
+class SoftmaxGradCUDNNKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& context) const override {
+    auto* Out = context.Input<Tensor>("Out");
+    auto* dOut = context.Input<Tensor>(framework::GradVarName("Out"));
+    auto* dX = context.Output<Tensor>(framework::GradVarName("X"));
+
+    // allocate memory on device.
+    dX->mutable_data<T>(context.GetPlace());
+
+    math::SoftmaxGradCUDNNFunctor<T>()(
+        context.template device_context<platform::CUDADeviceContext>(), Out,
+        dOut, dX);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_KERNEL(softmax, CUDNN, ::paddle::platform::CUDAPlace,
+                   ops::SoftmaxCUDNNKernel<float>);
+REGISTER_OP_KERNEL(softmax_grad, CUDNN, ::paddle::platform::CUDAPlace,
+                   ops::SoftmaxGradCUDNNKernel<float>);

paddle/fluid/operators/softmax_op.cc

@@ -33,6 +33,29 @@ class SoftmaxOp : public framework::OperatorWithKernel {
     ctx->SetOutputDim("Out", x_dims);
     ctx->ShareLoD("X", /*->*/ "Out");
   }
+
+ protected:
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext& ctx) const override {
+    // choose cudnn kernel if the runtime supported.
+    bool use_cudnn = ctx.Attr<bool>("use_cudnn");
+    bool runtime_cudnn_support = false;
+#ifdef PADDLE_WITH_CUDA
+    if (platform::is_gpu_place(ctx.GetPlace())) {
+      auto& dev_ctx =
+          ctx.template device_context<platform::CUDADeviceContext>();
+      runtime_cudnn_support = dev_ctx.cudnn_handle() != nullptr ? true : false;
+    }
+#endif
+    framework::LibraryType library_ = framework::LibraryType::kPlain;
+    if (use_cudnn && runtime_cudnn_support) {
+      library_ = framework::LibraryType::kCUDNN;
+    }
+    std::string data_format = ctx.Attr<std::string>("data_format");
+    return framework::OpKernelType(
+        framework::ToDataType(ctx.Input<Tensor>("X")->type()), ctx.GetPlace(),
+        framework::StringToDataLayout(data_format), library_);
+  }
 };
 
 class SoftmaxOpMaker : public framework::OpProtoAndCheckerMaker {
@@ -43,6 +66,17 @@ class SoftmaxOpMaker : public framework::OpProtoAndCheckerMaker {
              "The input tensor of softmax. "
              "2-D with shape [batch_size, input_feature_dimensions].");
     AddOutput("Out", "The normalized values with the same shape as X.");
+    AddAttr<bool>(
+        "use_cudnn",
+        "(bool, default false) Only used in cudnn kernel, need install cudnn")
+        .SetDefault(false);
+    AddAttr<std::string>(
+        "data_format",
+        "(string, default NCHW) Only used in "
+        "An optional string from: \"NHWC\", \"NCHW\". "
+        "Defaults to \"NHWC\". Specify the data format of the output data, "
+        "the input will be transformed automatically. ")
+        .SetDefault("AnyLayout");
     AddComment(R"DOC(
 Softmax Operator.
 
@@ -80,6 +114,29 @@ class SoftmaxOpGrad : public framework::OperatorWithKernel {
 
     ctx->SetOutputDim(framework::GradVarName("X"), ctx->GetInputDim("X"));
   }
+
+ protected:
+  framework::OpKernelType GetExpectedKernelType(
+      const framework::ExecutionContext& ctx) const override {
+    // choose cudnn kernel if the runtime supported.
+    bool use_cudnn = ctx.Attr<bool>("use_cudnn");
+    bool runtime_cudnn_support = false;
+#ifdef PADDLE_WITH_CUDA
+    if (platform::is_gpu_place(ctx.GetPlace())) {
+      auto& dev_ctx =
+          ctx.template device_context<platform::CUDADeviceContext>();
+      runtime_cudnn_support = dev_ctx.cudnn_handle() != nullptr ? true : false;
+    }
+#endif
+    framework::LibraryType library_ = framework::LibraryType::kPlain;
+    if (use_cudnn && runtime_cudnn_support) {
+      library_ = framework::LibraryType::kCUDNN;
+    }
+    std::string data_format = ctx.Attr<std::string>("data_format");
+    return framework::OpKernelType(
+        framework::ToDataType(ctx.Input<Tensor>("X")->type()), ctx.GetPlace(),
+        framework::StringToDataLayout(data_format), library_);
+  }
 };
 
 }  // namespace operators

paddle/fluid/operators/softmax_with_cross_entropy_op.cu

@@ -23,21 +23,21 @@ using Tensor = framework::Tensor;
 
 namespace {
 template <typename T>
-__global__ void CrossEntropyGrad(T* logit_grad, const T* loss_grad,
-                                 const int64_t* 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) {
-    PADDLE_ASSERT(labels[sample_idx] >= 0 && labels[sample_idx] < class_num);
-    logit_grad[tid * class_num + labels[tid]] -= static_cast<T>(1.);
+__global__ void CrossEntropyGrad(T* logit_grad, const int64_t* labels,
+                                 const int batch_size, const int class_num) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < batch_size;
+       i += blockDim.x * gridDim.x) {
+    int idx = i * class_num + labels[i];
+    logit_grad[idx] -= static_cast<T>(1.);
   }
+}
 
-  __syncthreads();
-
-  if (tid < batch_size * class_num) {
-    logit_grad[tid] *= loss_grad[sample_idx];
+template <typename T>
+__global__ void Scale(T* logit_grad, const T* loss_grad, const int num,
+                      const int class_num) {
+  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < num;
+       i += blockDim.x * gridDim.x) {
+    logit_grad[i] *= loss_grad[i / class_num];
   }
 }
 
@@ -94,22 +94,22 @@ class SoftmaxWithCrossEntropyGradCUDAKernel : public framework::OpKernel<T> {
     const int batch_size = logit_grad->dims()[0];
     const int class_num = logit_grad->dims()[1];
     int block = 512;
-    int grid = (batch_size * class_num + block - 1) / block;
+    auto stream = context.cuda_device_context().stream();
 
     if (context.Attr<bool>("soft_label")) {
+      int grid = (batch_size * class_num + block - 1) / block;
       const T* label_data = labels->data<T>();
-      SoftCrossEntropyGradientKernel<
-          T><<<grid, block, 0,
-               context.template device_context<platform::CUDADeviceContext>()
-                   .stream()>>>(logit_grad_data, loss_grad_data, label_data,
-                                batch_size, class_num);
+      SoftCrossEntropyGradientKernel<T><<<grid, block, 0, stream>>>(
+          logit_grad_data, loss_grad_data, label_data, batch_size, class_num);
     } else {
+      int grid = (batch_size + block - 1) / block;
       const int64_t* label_data = labels->data<int64_t>();
-      CrossEntropyGrad<
-          T><<<grid, block, 0,
-               context.template device_context<platform::CUDADeviceContext>()
-                   .stream()>>>(logit_grad_data, loss_grad_data, label_data,
-                                batch_size, class_num);
+      CrossEntropyGrad<T><<<grid, block, 0, stream>>>(
+          logit_grad_data, label_data, batch_size, class_num);
+      int num = batch_size * class_num;
+      grid = (num + block - 1) / block;
+      Scale<T><<<grid, block, 0, stream>>>(logit_grad_data, loss_grad_data, num,
+                                           class_num);
     }
   }
 };

paddle/fluid/platform/cudnn_helper.h

@@ -19,6 +19,7 @@ limitations under the License. */
 #include "paddle/fluid/framework/operator.h"
 #include "paddle/fluid/platform/dynload/cudnn.h"
 #include "paddle/fluid/platform/enforce.h"
+#include "paddle/fluid/platform/float16.h"
 #include "paddle/fluid/platform/macros.h"
 
 namespace paddle {
@@ -80,6 +81,22 @@ enum class PoolingMode {
 template <typename T>
 class CudnnDataType;
 
+template <>
+class CudnnDataType<float16> {
+ public:
+  static const cudnnDataType_t type = CUDNN_DATA_HALF;
+  // The scaling param type is float for HALF and FLOAT tensors
+  typedef const float ScalingParamType;
+  static ScalingParamType* kOne() {
+    static ScalingParamType v = 1.0;
+    return &v;
+  }
+  static ScalingParamType* kZero() {
+    static ScalingParamType v = 0.0;
+    return &v;
+  }
+};
+
 template <>
 class CudnnDataType<float> {
  public:
@@ -289,7 +306,7 @@ inline bool CanCUDNNBeUsed(const framework::ExecutionContext& ctx) {
   use_cudnn &= paddle::platform::is_gpu_place(ctx.GetPlace());
 #ifdef PADDLE_WITH_CUDA
   if (use_cudnn) {
-    auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
+    auto& dev_ctx = ctx.device_context<platform::CUDADeviceContext>();
     use_cudnn &= dev_ctx.cudnn_handle() != nullptr;
   }
 #endif

paddle/fluid/platform/device_tracer.cc

@@ -26,8 +26,14 @@ limitations under the License. */
 namespace paddle {
 namespace platform {
 namespace {
+// Current thread's id. Note, we don't distinguish nested threads
+// for now.
+thread_local int cur_thread_id = 0;
+// Tracking the nested block stacks of each thread.
+thread_local std::deque<int> block_id_stack;
+// Tracking the nested event stacks.
+thread_local std::deque<std::string> annotation_stack;
 
-thread_local const char *cur_annotation = nullptr;
 std::once_flag tracer_once_flag;
 DeviceTracer *tracer = nullptr;
 }  // namespace
@@ -191,19 +197,19 @@ class DeviceTracerImpl : public DeviceTracer {
     correlations_[id] = anno;
   }
 
-  void AddCPURecords(const char *anno, uint64_t start_ns, uint64_t end_ns) {
-    if (!anno) {
-      // TODO(panyx0718): Currently, it doesn't support nested situation
-      // Up-level can be cleared by low-level and therefore get nullptr
-      // here.
+  void AddCPURecords(const std::string &anno, uint64_t start_ns,
+                     uint64_t end_ns, int64_t device_id, int64_t thread_id) {
+    if (anno.empty()) {
+      VLOG(1) << "Empty timeline annotation.";
       return;
     }
     std::lock_guard<std::mutex> l(trace_mu_);
-    cpu_records_.push_back(CPURecord{anno, start_ns, end_ns, 0});
+    cpu_records_.push_back(
+        CPURecord{anno, start_ns, end_ns, device_id, thread_id});
   }
 
   void AddMemRecords(const std::string &name, uint64_t start_ns,
-                     uint64_t end_ns, uint32_t device_id, uint32_t stream_id,
+                     uint64_t end_ns, int64_t device_id, int64_t stream_id,
                      uint32_t correlation_id, uint64_t bytes) {
     // 0 means timestamp information could not be collected for the kernel.
     if (start_ns == 0 || end_ns == 0) {
@@ -215,8 +221,8 @@ class DeviceTracerImpl : public DeviceTracer {
                                      stream_id, correlation_id, bytes});
   }
 
-  void AddKernelRecords(uint64_t start, uint64_t end, uint32_t device_id,
-                        uint32_t stream_id, uint32_t correlation_id) {
+  void AddKernelRecords(uint64_t start, uint64_t end, int64_t device_id,
+                        int64_t stream_id, uint32_t correlation_id) {
     // 0 means timestamp information could not be collected for the kernel.
     if (start == 0 || end == 0) {
       VLOG(3) << correlation_id << " cannot be traced";
@@ -270,27 +276,30 @@ class DeviceTracerImpl : public DeviceTracer {
         continue;
       }
       auto *event = profile_pb.add_events();
+      event->set_type(proto::Event::GPUKernel);
       event->set_name(correlations_.at(r.correlation_id));
       event->set_start_ns(r.start_ns);
       event->set_end_ns(r.end_ns);
-      event->set_stream_id(r.stream_id);
+      event->set_sub_device_id(r.stream_id);
       event->set_device_id(r.device_id);
     }
 
     for (const CPURecord &r : cpu_records_) {
       auto *event = profile_pb.add_events();
+      event->set_type(proto::Event::CPU);
       event->set_name(r.name);
       event->set_start_ns(r.start_ns);
       event->set_end_ns(r.end_ns);
-      event->set_stream_id(r.thread_id);
-      event->set_device_id(-1);
+      event->set_sub_device_id(r.thread_id);
+      event->set_device_id(r.device_id);
     }
     for (const MemRecord &r : mem_records_) {
       auto *event = profile_pb.add_events();
+      event->set_type(proto::Event::GPUKernel);
       event->set_name(r.name);
       event->set_start_ns(r.start_ns);
       event->set_end_ns(r.end_ns);
-      event->set_stream_id(r.stream_id);
+      event->set_sub_device_id(r.stream_id);
       event->set_device_id(r.device_id);
       event->mutable_memcopy()->set_bytes(r.bytes);
     }
@@ -323,8 +332,9 @@ class DeviceTracerImpl : public DeviceTracer {
     if ((domain == CUPTI_CB_DOMAIN_DRIVER_API) &&
         (cbid == CUPTI_DRIVER_TRACE_CBID_cuLaunchKernel)) {
       if (cbInfo->callbackSite == CUPTI_API_ENTER) {
-        const std::string anno =
-            cur_annotation ? cur_annotation : cbInfo->symbolName;
+        const std::string anno = !annotation_stack.empty()
+                                     ? annotation_stack.back()
+                                     : cbInfo->symbolName;
         tracer->AddAnnotation(cbInfo->correlationId, anno);
       }
     } else {
@@ -351,14 +361,15 @@ class DeviceTracerDummy : public DeviceTracer {
 
   void AddAnnotation(uint64_t id, const std::string &anno) {}
 
-  void AddCPURecords(const char *anno, uint64_t start_ns, uint64_t end_ns) {}
+  void AddCPURecords(const std::string &anno, uint64_t start_ns,
+                     uint64_t end_ns, int64_t device_id, int64_t thread_id) {}
 
   void AddMemRecords(const std::string &name, uint64_t start_ns,
-                     uint64_t end_ns, uint32_t device_id, uint32_t stream_id,
+                     uint64_t end_ns, int64_t device_id, int64_t stream_id,
                      uint32_t correlation_id, uint64_t bytes) {}
 
-  void AddKernelRecords(uint64_t start, uint64_t end, uint32_t device_id,
-                        uint32_t stream_id, uint32_t correlation_id) {}
+  void AddKernelRecords(uint64_t start, uint64_t end, int64_t device_id,
+                        int64_t stream_id, uint32_t correlation_id) {}
 
   bool IsEnabled() { return false; }
 
@@ -384,11 +395,28 @@ DeviceTracer *GetDeviceTracer() {
   return tracer;
 }
 
-void SetCurAnnotation(const char *anno) { cur_annotation = anno; }
+void SetCurAnnotation(const std::string &anno) {
+  annotation_stack.push_back(anno);
+}
+
+void ClearCurAnnotation() { annotation_stack.pop_back(); }
+
+std::string CurAnnotation() {
+  if (annotation_stack.empty()) return "";
+  return annotation_stack.back();
+}
+
+void SetCurBlock(int block_id) { block_id_stack.push_back(block_id); }
+
+void ClearCurBlock() { block_id_stack.pop_back(); }
+
+int BlockDepth() { return block_id_stack.size(); }
+
+void SetCurThread(int thread_id) { cur_thread_id = thread_id; }
 
-void ClearCurAnnotation() { cur_annotation = nullptr; }
+void ClearCurThread() { cur_thread_id = 0; }
 
-const char *CurAnnotation() { return cur_annotation; }
+int CurThread() { return cur_thread_id; }
 
 }  // namespace platform
 }  // namespace paddle

paddle/fluid/platform/device_tracer.h

@@ -32,22 +32,23 @@ class DeviceTracer {
   struct KernelRecord {
     uint64_t start_ns;
     uint64_t end_ns;
-    uint32_t device_id;
-    uint32_t stream_id;
+    int64_t device_id;
+    int64_t stream_id;
     uint32_t correlation_id;
   };
   struct CPURecord {
     std::string name;
     uint64_t start_ns;
     uint64_t end_ns;
-    uint64_t thread_id;
+    int64_t device_id;
+    int64_t thread_id;
   };
   struct MemRecord {
     std::string name;
     uint64_t start_ns;
     uint64_t end_ns;
-    uint32_t device_id;
-    uint32_t stream_id;
+    int64_t device_id;
+    int64_t stream_id;
     uint32_t correlation_id;
     uint64_t bytes;
   };
@@ -64,18 +65,18 @@ class DeviceTracer {
   virtual void AddAnnotation(uint64_t id, const std::string& anno) = 0;
 
   virtual void AddMemRecords(const std::string& name, uint64_t start_ns,
-                             uint64_t end_ns, uint32_t device_id,
-                             uint32_t stream_id, uint32_t correlation_id,
+                             uint64_t end_ns, int64_t device_id,
+                             int64_t stream_id, uint32_t correlation_id,
                              uint64_t bytes) = 0;
 
-  virtual void AddCPURecords(const char* anno, uint64_t start_ns,
-                             uint64_t end_ns) = 0;
+  virtual void AddCPURecords(const std::string& anno, uint64_t start_ns,
+                             uint64_t end_ns, int64_t device_id,
+                             int64_t thread_id) = 0;
 
   // Add a cuda kernel stats. `correlation_id` will be mapped to annotation
   // added before for human readability.
-  virtual void AddKernelRecords(uint64_t start, uint64_t end,
-                                uint32_t device_id, uint32_t stream_id,
-                                uint32_t correlation_id) = 0;
+  virtual void AddKernelRecords(uint64_t start, uint64_t end, int64_t device_id,
+                                int64_t stream_id, uint32_t correlation_id) = 0;
 
   // Generate a proto after done (Disabled).
   virtual proto::Profile GenProfile(const std::string& profile_path) = 0;
@@ -87,10 +88,18 @@ class DeviceTracer {
 DeviceTracer* GetDeviceTracer();
 
 // Set a name for the cuda kernel operation being launched by the thread.
-void SetCurAnnotation(const char* anno);
+void SetCurAnnotation(const std::string& anno);
 // Clear the name after the operation is done.
 void ClearCurAnnotation();
 // Current name of the operation being run in the thread.
-const char* CurAnnotation();
+std::string CurAnnotation();
+
+void SetCurBlock(int block_id);
+void ClearCurBlock();
+int BlockDepth();
+
+void SetCurThread(int thread_id);
+void ClearCurThread();
+int CurThread();
 }  // namespace platform
 }  // namespace paddle

paddle/fluid/platform/profiler.cc

@@ -147,19 +147,48 @@ RecordEvent::RecordEvent(const std::string& name, const DeviceContext* dev_ctx)
   name_ = name;
   PushEvent(name_, dev_ctx_);
   // Maybe need the same push/pop behavior.
-  SetCurAnnotation(name_.c_str());
+  SetCurAnnotation(name_);
 }
 
 RecordEvent::~RecordEvent() {
   if (g_state == ProfilerState::kDisabled) return;
   DeviceTracer* tracer = GetDeviceTracer();
   if (tracer) {
-    tracer->AddCPURecords(CurAnnotation(), start_ns_, PosixInNsec());
+    tracer->AddCPURecords(CurAnnotation(), start_ns_, PosixInNsec(),
+                          BlockDepth(), CurThread());
   }
   ClearCurAnnotation();
   PopEvent(name_, dev_ctx_);
 }
 
+RecordBlock::RecordBlock(int block_id) : start_ns_(PosixInNsec()) {
+  if (g_state == ProfilerState::kDisabled) return;
+  SetCurBlock(block_id);
+  name_ = string::Sprintf("block_%d", block_id);
+}
+
+RecordBlock::~RecordBlock() {
+  if (g_state == ProfilerState::kDisabled) return;
+  DeviceTracer* tracer = GetDeviceTracer();
+  if (tracer) {
+    // We try to put all blocks at the same nested depth in the
+    // same timeline lane. and distinguish the using thread_id.
+    tracer->AddCPURecords(name_, start_ns_, PosixInNsec(), BlockDepth(),
+                          CurThread());
+  }
+  ClearCurBlock();
+}
+
+RecordThread::RecordThread(int thread_id) {
+  if (g_state == ProfilerState::kDisabled) return;
+  SetCurThread(thread_id);
+}
+
+RecordThread::~RecordThread() {
+  if (g_state == ProfilerState::kDisabled) return;
+  ClearCurThread();
+}
+
 void EnableProfiler(ProfilerState state) {
   PADDLE_ENFORCE(state != ProfilerState::kDisabled,
                  "Can't enbale profling, since the input state is ",

paddle/fluid/platform/profiler.h

@@ -118,6 +118,20 @@ struct RecordEvent {
   std::string full_name_;
 };
 
+struct RecordBlock {
+  explicit RecordBlock(int block_id);
+  ~RecordBlock();
+
+ private:
+  std::string name_;
+  uint64_t start_ns_;
+};
+
+struct RecordThread {
+  explicit RecordThread(int thread_id);
+  ~RecordThread();
+};
+
 // Return the event list of all threads. Assumed the returned value calls
 // event_lists, event_lists[i][j] represents the j-th Event of i-th thread.
 std::vector<std::vector<Event>> GetAllEvents();

paddle/fluid/platform/profiler.proto

@@ -18,12 +18,17 @@ package paddle.platform.proto;
 message MemCopy { optional uint64 bytes = 1; }
 
 message Event {
+  enum EventType {
+    CPU = 0;
+    GPUKernel = 1;
+  }
+  optional EventType type = 8;
   optional string name = 1;
   optional uint64 start_ns = 2;
   optional uint64 end_ns = 3;
   // When positive, it represents gpu id. When -1, it represents CPU.
   optional int64 device_id = 5;
-  optional uint32 stream_id = 6;
+  optional int64 sub_device_id = 6;
 
   optional MemCopy memcopy = 7;
 }

paddle/fluid/pybind/tensor_py.h

@@ -72,6 +72,7 @@ struct CastToPyBufferImpl<true, I, ARGS...> {
         paddle::platform::GpuMemcpyAsync(
             dst_ptr, src_ptr, sizeof(CUR_TYPE) * tensor.numel(),
             cudaMemcpyDeviceToHost, dev_ctx->stream());
+        dev_ctx->Wait();
 #else
         PADDLE_THROW("'CUDAPlace' is not supported in CPU only device.");
 #endif

python/paddle/fluid/__init__.py

@@ -35,7 +35,7 @@ from core import LoDTensor, CPUPlace, CUDAPlace
 from distribute_transpiler import DistributeTranspiler
 from distribute_transpiler_simple import SimpleDistributeTranspiler
 from concurrency import (Go, make_channel, channel_send, channel_recv,
-                         channel_close)
+                         channel_close, Select)
 import clip
 from memory_optimization_transpiler import memory_optimize, release_memory
 import profiler

python/paddle/fluid/concurrency.py

@@ -12,17 +12,14 @@
 # See the License for the specific language governing permissions and
 # limitations under the License.
 
-from layers.control_flow import BlockGuard
+from layers.control_flow import BlockGuard, Select
 from layer_helper import LayerHelper, unique_name
 from layers import fill_constant
 import core
 
 __all__ = [
-    'Go',
-    'make_channel',
-    'channel_send',
-    'channel_recv',
-    'channel_close',
+    'Go', 'make_channel', 'channel_send', 'channel_recv', 'channel_close',
+    'Select'
 ]
 
 
@@ -198,7 +195,7 @@ def channel_recv(channel, return_value):
 
           ch = fluid.make_channel(dtype='int32', capacity=10)
           with fluid.Go():
-            returned_value = fluid.channel_recv(ch, 'int32')
+            returned_value, return_status = fluid.channel_recv(ch, 'int32')
 
           # Code to send data through the channel.
     """

python/paddle/fluid/framework.py

@@ -487,7 +487,7 @@ class Operator(object):
             'rnn_memory_helper_grad', 'conditional_block', 'while', 'send',
             'recv', 'listen_and_serv', 'parallel_do', 'save_combine',
             'load_combine', 'ncclInit', 'channel_create', 'channel_close',
-            'channel_send', 'channel_recv'
+            'channel_send', 'channel_recv', 'select'
         }
         if type not in no_kernel_op_set:
             self.desc.infer_var_type(self.block.desc)

python/paddle/fluid/layers/control_flow.py

@@ -16,7 +16,7 @@ import contextlib
 from layer_function_generator import autodoc
 from tensor import assign, fill_constant
 from .. import core
-from ..framework import Program, Variable, Operator
+from ..framework import Program, Variable, Operator, Block
 from ..layer_helper import LayerHelper, unique_name
 from ops import logical_and, logical_not, logical_or
 
@@ -29,6 +29,7 @@ __all__ = [
     'WhileGuard',
     'While',
     'Switch',
+    'Select',
     'lod_rank_table',
     'max_sequence_len',
     'topk',
@@ -1211,6 +1212,186 @@ class Switch(object):
         return True
 
 
+class SelectCase(object):
+    DEFAULT = 0
+    SEND = 1
+    RECEIVE = 2
+
+    def __init__(self,
+                 case_idx,
+                 case_to_execute,
+                 channel_action_fn=None,
+                 channel=None,
+                 value=None):
+        self.helper = LayerHelper('conditional_block')
+        self.main_program = self.helper.main_program
+        self.is_scalar_condition = True
+
+        self.case_to_execute = case_to_execute
+        self.idx = case_idx
+
+        # Since we aren't going to use the `channel_send` or `channel_recv`
+        # functions directly, we just need to capture the name.
+        self.action = (self.SEND
+                       if channel_action_fn.__name__ == ('channel_send') else
+                       self.RECEIVE) if channel_action_fn else (self.DEFAULT)
+        self.value = value
+        self.channel = channel
+
+    def __enter__(self):
+        self.block = self.main_program.create_block()
+
+    def construct_op(self):
+        main_program = self.helper.main_program
+        cases_block = main_program.current_block()
+
+        inner_outputs = set()
+        input_set = set()
+        params = set()
+
+        for op in self.block.ops:
+            # Iterate over all operators, get all the inputs
+            # and add as input to the SelectCase operator.
+            for iname in op.input_names:
+                for in_var_name in op.input(iname):
+                    if in_var_name not in inner_outputs:
+                        input_set.add(in_var_name)
+
+            for oname in op.output_names:
+                for out_var_name in op.output(oname):
+                    inner_outputs.add(out_var_name)
+
+        param_list = [
+            cases_block.var(each_name) for each_name in params
+            if each_name not in input_set
+        ]
+
+        # Iterate over all operators, get all the outputs
+        # add to the output list of SelectCase operator only if
+        # they exist in the parent block.
+        out_vars = []
+        for inner_out_name in inner_outputs:
+            if inner_out_name in cases_block.vars:
+                out_vars.append(cases_block.var(inner_out_name))
+
+        # First, create an op that will determine whether or not this is the
+        # conditional variable to execute.
+        should_execute_block = equal(
+            fill_constant(
+                shape=[1], dtype=core.VarDesc.VarType.INT32, value=self.idx),
+            self.case_to_execute)
+
+        step_scope = cases_block.create_var(
+            type=core.VarDesc.VarType.STEP_SCOPES)
+
+        cases_block.append_op(
+            type='conditional_block',
+            inputs={'X': [should_execute_block],
+                    'Params': param_list},
+            outputs={'Out': out_vars,
+                     'Scope': [step_scope]},
+            attrs={
+                'sub_block': self.block,
+                'is_scalar_condition': self.is_scalar_condition
+            })
+
+        return '%s,%s,%s,%s' % (self.idx, self.action, self.channel.name
+                                if self.channel else '', self.value.name
+                                if self.value else '')
+
+    def __exit__(self, exc_type, exc_val, exc_tb):
+        self.main_program.rollback()
+        if exc_type is not None:
+            return False  # re-raise exception
+        return True
+
+
+class Select(BlockGuard):
+    def __init__(self, name=None):
+        self.helper = LayerHelper('select', name=name)
+        self.cases = []
+
+        super(Select, self).__init__(self.helper.main_program)
+        self.case_to_execute = fill_constant(
+            shape=[1], dtype=core.VarDesc.VarType.INT32, value=-1)
+
+    def __enter__(self):
+        super(Select, self).__enter__()
+        return self
+
+    def case(self, channel_action_fn, channel, value):
+        """Create a new block for this condition.
+        """
+        select_case = SelectCase(
+            len(self.cases), self.case_to_execute, channel_action_fn, channel,
+            value)
+
+        self.cases.append(select_case)
+
+        return select_case
+
+    def default(self):
+        """Create a default case block for this condition.
+        """
+        default_case = SelectCase(len(self.cases), self.case_to_execute)
+
+        self.cases.append(default_case)
+
+        return default_case
+
+    def __exit__(self, exc_type, exc_val, exc_tb):
+        if exc_type is not None:
+            return False
+
+        # Create a select op and another block to wrap its
+        # case blocks.
+        select_block = self.helper.main_program.current_block()
+        parent_block = self.helper.main_program.block(select_block.parent_idx)
+
+        # Construct each case op, inside the newly created select block.
+        serialized_cases = []
+        for case in self.cases:
+            serialized_cases.append(case.construct_op())
+
+        intermediate = set()
+        params = set()
+
+        for case_block in select_block.ops:
+            if case_block.attrs and 'sub_block' in case_block.attrs:
+                for each_op in case_block.attrs['sub_block'].ops:
+                    assert isinstance(each_op, Operator)
+                    for iname in each_op.input_names:
+                        for in_var_name in each_op.input(iname):
+                            if in_var_name not in intermediate:
+                                params.add(in_var_name)
+
+                    for oname in each_op.output_names:
+                        for out_var_name in each_op.output(oname):
+                            intermediate.add(out_var_name)
+
+        # TODO(varunarora): Figure out if defining output is needed.
+        out_list = [
+            parent_block.var(var_name) for var_name in parent_block.vars
+            if var_name in intermediate
+        ]
+
+        X = [select_block.var_recursive(x_name) for x_name in params]
+
+        # Needs to be used by `equal` inside the cases block.
+        X.append(self.case_to_execute)
+
+        # Construct the select op.
+        parent_block.append_op(
+            type='select',
+            inputs={'X': X,
+                    'case_to_execute': self.case_to_execute},
+            attrs={'sub_block': select_block,
+                   'cases': serialized_cases},
+            outputs={})
+
+        return super(Select, self).__exit__(exc_type, exc_val, exc_tb)
+
+
 class IfElseBlockGuard(object):
     def __init__(self, is_true, ifelse):
         if not isinstance(ifelse, IfElse):

python/paddle/fluid/layers/detection.py

@@ -132,7 +132,7 @@ def detection_output(loc,
 
     old_shape = scores.shape
     scores = ops.reshape(x=scores, shape=(-1, old_shape[-1]))
-    scores = ops.softmax(x=scores)
+    scores = nn.softmax(input=scores)
     scores = ops.reshape(x=scores, shape=old_shape)
     scores = nn.transpose(scores, perm=[0, 2, 1])
 

python/paddle/fluid/layers/layer_function_generator.py

@@ -16,10 +16,7 @@ import cStringIO
 import functools
 import warnings
 
-from .. import proto
-
-framework_pb2 = proto.framework_pb2
-
+from ..proto import framework_pb2
 from ..framework import OpProtoHolder, Variable
 from ..layer_helper import LayerHelper
 

python/paddle/fluid/layers/nn.py

@@ -39,6 +39,8 @@ __all__ = [
     'sequence_conv',
     'conv2d',
     'sequence_pool',
+    'sequence_softmax',
+    'softmax',
     'pool2d',
     'batch_norm',
     'beam_search_decode',
@@ -1085,6 +1087,30 @@ def sequence_conv(input,
     return helper.append_activation(pre_act)
 
 
+def sequence_softmax(input, param_attr=None, bias_attr=None, use_cudnn=True):
+    helper = LayerHelper('sequence_softmax', **locals())
+    dtype = helper.input_dtype()
+    softmax_out = helper.create_tmp_variable(dtype)
+    helper.append_op(
+        type="sequence_softmax",
+        inputs={"X": input},
+        outputs={"Out": softmax_out},
+        attrs={"use_cudnn": use_cudnn})
+    return softmax_out
+
+
+def softmax(input, param_attr=None, bias_attr=None, use_cudnn=True):
+    helper = LayerHelper('softmax', **locals())
+    dtype = helper.input_dtype()
+    softmax_out = helper.create_tmp_variable(dtype)
+    helper.append_op(
+        type="softmax",
+        inputs={"X": input},
+        outputs={"Out": softmax_out},
+        attrs={"use_cudnn": use_cudnn})
+    return softmax_out
+
+
 def conv2d(input,
            num_filters,
            filter_size,

python/paddle/fluid/layers/ops.py

@@ -45,13 +45,31 @@ __activations__ = [
 ]
 
 __all__ = [
-    'mean', 'mul', 'reshape', 'scale', 'sigmoid_cross_entropy_with_logits',
-    'elementwise_add', 'elementwise_div', 'elementwise_sub', 'elementwise_mul',
-    'elementwise_max', 'elementwise_min', 'elementwise_pow', 'clip',
-    'clip_by_norm', 'softmax', 'sequence_softmax', 'logical_and', 'logical_or',
-    'logical_xor', 'logical_not', 'uniform_random',
-    'uniform_random_batch_size_like', 'gaussian_random',
-    'gaussian_random_batch_size_like', 'cumsum', 'scatter', 'sum'
+    'mean',
+    'mul',
+    'reshape',
+    'scale',
+    'sigmoid_cross_entropy_with_logits',
+    'elementwise_add',
+    'elementwise_div',
+    'elementwise_sub',
+    'elementwise_mul',
+    'elementwise_max',
+    'elementwise_min',
+    'elementwise_pow',
+    'clip',
+    'clip_by_norm',
+    'logical_and',
+    'logical_or',
+    'logical_xor',
+    'logical_not',
+    'uniform_random',
+    'uniform_random_batch_size_like',
+    'gaussian_random',
+    'gaussian_random_batch_size_like',
+    'cumsum',
+    'scatter',
+    'sum',
 ] + __activations__
 
 for _OP in set(__all__):

python/paddle/fluid/optimizer.py

@@ -233,6 +233,8 @@ class Optimizer(object):
         params_grads = append_backward(loss, parameter_list, no_grad_set,
                                        [error_clip_callback])
 
+        params_grads = sorted(params_grads, key=lambda x: x[0].name)
+
         params_grads = append_gradient_clip_ops(params_grads)
 
         # Add regularization if any

python/paddle/fluid/regularizer.py

@@ -44,6 +44,11 @@ def append_regularization_ops(parameters_and_grads, regularization=None):
     """
     params_and_grads = []
     for param, grad in parameters_and_grads:
+        # If no gradient then we don't need to do anything
+        if grad is None:
+            params_and_grads.append((param, grad))
+            continue
+
         regularization_term = None
         if param.regularizer is not None:
             # Add variable for regularization term in grad block
@@ -51,9 +56,8 @@ def append_regularization_ops(parameters_and_grads, regularization=None):
         elif regularization is not None:
             regularization_term = regularization(param, grad, grad.block)
 
-        # If no gradient or no regularization specified,
-        # then we don't need to do anything
-        if grad is None or regularization_term is None:
+        # If no regularization specified, then we don't need to do anything
+        if regularization_term is None:
             params_and_grads.append((param, grad))
             continue
 

python/paddle/fluid/tests/test_concurrency.py

@@ -15,9 +15,9 @@
 import unittest
 import paddle.fluid as fluid
 import paddle.fluid.core as core
-from paddle.fluid import framework, unique_name
+from paddle.fluid import framework, unique_name, layer_helper
 from paddle.fluid.executor import Executor
-from paddle.fluid.layers import fill_constant
+from paddle.fluid.layers import fill_constant, assign, While, elementwise_add, Print
 
 
 class TestRoutineOp(unittest.TestCase):
@@ -86,8 +86,7 @@ class TestRoutineOp(unittest.TestCase):
         self.assertEqual(leftmost_data[0][0], n + 1)
 
     def _create_one_dim_tensor(self, value):
-        one_dim_tensor = fill_constant(
-            shape=[1], dtype=core.VarDesc.VarType.INT64, value=value)
+        one_dim_tensor = fill_constant(shape=[1], dtype='int', value=value)
         one_dim_tensor.stop_gradient = True
         return one_dim_tensor
 
@@ -95,6 +94,180 @@ class TestRoutineOp(unittest.TestCase):
         return framework.default_main_program().current_block().create_var(
             name=unique_name.generate(name), type=type, dtype=dtype)
 
+    def _create_persistable_tensor(self, name, type, dtype):
+        return framework.default_main_program().current_block().create_var(
+            name=unique_name.generate(name),
+            type=type,
+            dtype=dtype,
+            persistable=True)
+
+    def test_select(self):
+        with framework.program_guard(framework.Program()):
+            ch1 = fluid.make_channel(
+                dtype=core.VarDesc.VarType.LOD_TENSOR, capacity=1)
+
+            result1 = self._create_tensor('return_value',
+                                          core.VarDesc.VarType.LOD_TENSOR,
+                                          core.VarDesc.VarType.FP64)
+
+            input_value = fill_constant(
+                shape=[1], dtype=core.VarDesc.VarType.FP64, value=10)
+
+            with fluid.Select() as select:
+                with select.case(fluid.channel_send, ch1, input_value):
+                    # Execute something.
+                    pass
+
+                with select.default():
+                    pass
+
+            # This should not block because we are using a buffered channel.
+            result1, status = fluid.channel_recv(ch1, result1)
+            fluid.channel_close(ch1)
+
+            cpu = core.CPUPlace()
+            exe = Executor(cpu)
+
+            result = exe.run(fetch_list=[result1])
+            self.assertEqual(result[0][0], 10)
+
+    def test_fibonacci(self):
+        """
+        Mimics Fibonacci Go example: https://tour.golang.org/concurrency/5
+        """
+        with framework.program_guard(framework.Program()):
+            quit_ch_input_var = self._create_persistable_tensor(
+                'quit_ch_input', core.VarDesc.VarType.LOD_TENSOR,
+                core.VarDesc.VarType.INT32)
+            quit_ch_input = fill_constant(
+                shape=[1],
+                dtype=core.VarDesc.VarType.INT32,
+                value=0,
+                out=quit_ch_input_var)
+
+            result = self._create_persistable_tensor(
+                'result', core.VarDesc.VarType.LOD_TENSOR,
+                core.VarDesc.VarType.INT32)
+            fill_constant(
+                shape=[1],
+                dtype=core.VarDesc.VarType.INT32,
+                value=0,
+                out=result)
+
+            x = fill_constant(
+                shape=[1], dtype=core.VarDesc.VarType.INT32, value=0)
+            y = fill_constant(
+                shape=[1], dtype=core.VarDesc.VarType.INT32, value=1)
+
+            while_cond = fill_constant(
+                shape=[1], dtype=core.VarDesc.VarType.BOOL, value=True)
+
+            while_false = fill_constant(
+                shape=[1], dtype=core.VarDesc.VarType.BOOL, value=False)
+
+            x_tmp = fill_constant(
+                shape=[1], dtype=core.VarDesc.VarType.INT32, value=0)
+
+            def fibonacci(channel, quit_channel):
+                while_op = While(cond=while_cond)
+                with while_op.block():
+                    result2 = fill_constant(
+                        shape=[1], dtype=core.VarDesc.VarType.INT32, value=0)
+                    x_to_send_tmp = fill_constant(
+                        shape=[1], dtype=core.VarDesc.VarType.INT32, value=0)
+
+                    # TODO(abhinav): Need to perform copy when doing a channel send.
+                    #   Once this is complete, we can remove these lines
+                    assign(input=x, output=x_to_send_tmp)
+
+                    with fluid.Select() as select:
+                        with select.case(fluid.channel_send, channel,
+                                         x_to_send_tmp):
+                            assign(input=x, output=x_tmp)
+                            assign(input=y, output=x)
+                            assign(elementwise_add(x=x_tmp, y=y), output=y)
+
+                        with select.case(fluid.channel_recv, quit_channel,
+                                         result2):
+                            # Quit
+                            helper = layer_helper.LayerHelper('assign')
+                            helper.append_op(
+                                type='assign',
+                                inputs={'X': [while_false]},
+                                outputs={'Out': [while_cond]})
+
+            ch1 = fluid.make_channel(dtype=core.VarDesc.VarType.LOD_TENSOR)
+            quit_ch = fluid.make_channel(dtype=core.VarDesc.VarType.LOD_TENSOR)
+
+            with fluid.Go():
+                for i in xrange(10):
+                    fluid.channel_recv(ch1, result)
+                    Print(result)
+
+                fluid.channel_send(quit_ch, quit_ch_input)
+
+            fibonacci(ch1, quit_ch)
+
+            fluid.channel_close(ch1)
+            fluid.channel_close(quit_ch)
+
+            cpu = core.CPUPlace()
+            exe = Executor(cpu)
+
+            exe_result = exe.run(fetch_list=[result])
+            self.assertEqual(exe_result[0][0], 34)
+
+    def test_ping_pong(self):
+        """
+        Mimics Ping Pong example: https://gobyexample.com/channel-directions
+        """
+        with framework.program_guard(framework.Program()):
+            result = self._create_tensor('return_value',
+                                         core.VarDesc.VarType.LOD_TENSOR,
+                                         core.VarDesc.VarType.FP64)
+
+            ping_result = self._create_tensor('ping_return_value',
+                                              core.VarDesc.VarType.LOD_TENSOR,
+                                              core.VarDesc.VarType.FP64)
+
+            pong_result = self._create_tensor('pong_return_value',
+                                              core.VarDesc.VarType.LOD_TENSOR,
+                                              core.VarDesc.VarType.FP64)
+
+            def ping(ch, message):
+                message_to_send_tmp = fill_constant(
+                    shape=[1], dtype=core.VarDesc.VarType.FP64, value=0)
+
+                assign(input=message, output=message_to_send_tmp)
+                fluid.channel_send(ch, message_to_send_tmp)
+
+            def pong(ch1, ch2):
+                fluid.channel_recv(ch1, ping_result)
+                assign(input=ping_result, output=pong_result)
+                fluid.channel_send(ch2, pong_result)
+
+            pings = fluid.make_channel(
+                dtype=core.VarDesc.VarType.LOD_TENSOR, capacity=1)
+            pongs = fluid.make_channel(
+                dtype=core.VarDesc.VarType.LOD_TENSOR, capacity=1)
+
+            msg = fill_constant(
+                shape=[1], dtype=core.VarDesc.VarType.FP64, value=9)
+
+            ping(pings, msg)
+            pong(pings, pongs)
+
+            fluid.channel_recv(pongs, result)
+
+            fluid.channel_close(pings)
+            fluid.channel_close(pongs)
+
+            cpu = core.CPUPlace()
+            exe = Executor(cpu)
+
+            exe_result = exe.run(fetch_list=[result])
+            self.assertEqual(exe_result[0][0], 9)
+
 
 if __name__ == '__main__':
     unittest.main()

python/paddle/fluid/tests/unittests/CMakeLists.txt

@@ -11,7 +11,6 @@ list(REMOVE_ITEM TEST_OPS test_lstm_unit_op) # # FIXME(qijun) https://github.com
 list(REMOVE_ITEM TEST_OPS test_nce) # IXME(qijun) https://github.com/PaddlePaddle/Paddle/issues/7778
 list(REMOVE_ITEM TEST_OPS test_recurrent_op) # FIXME(qijun) https://github.com/PaddlePaddle/Paddle/issues/6152
 list(REMOVE_ITEM TEST_OPS test_cond_op) # FIXME(qijun): https://github.com/PaddlePaddle/Paddle/issues/5101#issuecomment-339814957
-list(REMOVE_ITEM TEST_OPS test_detection_output_op) # FIXME: detection_output_op will be rewritten. This unittest should be
 
 list(REMOVE_ITEM TEST_OPS op_test) # op_test is a helper python file, not a test
 list(REMOVE_ITEM TEST_OPS decorators) # decorators is a helper python file, not a test

python/paddle/fluid/tests/unittests/op_test.py

@@ -469,6 +469,28 @@ class OpTest(unittest.TestCase):
             tensor.set_lod(lod)
         return tensor
 
+    @staticmethod
+    def np_dtype_to_fluid_dtype(input):
+        """Change the dtype of float16 numpy array
+
+        numpy float16 is binded to paddle::platform::float16 
+        in tensor_py.h via the help of uint16 data type since
+        the internal memory representation of float16 is 
+        uint16_t in paddle and np.uint16 in numpy, which are
+        themselves binded together by pybind.
+
+        Args:
+            input: input numpy array
+
+        Returns:
+            input: The dtype of input will be changed to np.uint16 if 
+                it is originally np.float16, such that the internal memory
+                of input will be reinterpreted as of dtype np.uint16. 
+        """
+        if input.dtype == np.float16:
+            input.dtype = np.uint16
+        return input
+
     def _get_gradient(self, input_to_check, place, output_names, no_grad_set):
         prog = Program()
         block = prog.global_block()

python/paddle/fluid/tests/unittests/test_cast_op.py

@@ -18,7 +18,7 @@ import numpy as np
 import paddle.fluid.core as core
 
 
-class TestCastOp(op_test.OpTest):
+class TestCastOp1(op_test.OpTest):
     def setUp(self):
         ipt = np.random.random(size=[10, 10])
         self.inputs = {'X': ipt.astype('float32')}
@@ -36,5 +36,36 @@ class TestCastOp(op_test.OpTest):
         self.check_grad(['X'], ['Out'])
 
 
+class TestCastOp2(op_test.OpTest):
+    def setUp(self):
+        ipt = np.random.random(size=[10, 10])
+        # numpy float16 is binded to fluid float16 via uint16
+        self.inputs = {'X': ipt.astype('float16').view(np.uint16)}
+        self.outputs = {'Out': ipt.astype('float32')}
+        self.attrs = {
+            'in_dtype': int(core.VarDesc.VarType.FP16),
+            'out_dtype': int(core.VarDesc.VarType.FP32)
+        }
+        self.op_type = 'cast'
+
+    def test_check_output(self):
+        self.check_output(atol=1e-3)
+
+
+class TestCastOp3(op_test.OpTest):
+    def setUp(self):
+        ipt = np.random.random(size=[10, 10])
+        self.inputs = {'X': ipt.astype('float32')}
+        self.outputs = {'Out': ipt.astype('float16')}
+        self.attrs = {
+            'in_dtype': int(core.VarDesc.VarType.FP32),
+            'out_dtype': int(core.VarDesc.VarType.FP16)
+        }
+        self.op_type = 'cast'
+
+    def test_check_output(self):
+        self.check_output(atol=1e-3)
+
+
 if __name__ == '__main__':
     unittest.main()

python/paddle/fluid/tests/unittests/test_conv2d_op.py

@@ -63,9 +63,11 @@ def conv2d_forward_naive(input, filter, group, conv_param):
 
 class TestConv2dOp(OpTest):
     def setUp(self):
+        self.op_type = "conv2d"
         self.use_cudnn = False
         self.use_mkldnn = False
-        self.init_op_type()
+        self.dtype = np.float32
+        self.init_kernel_type()
         self.init_group()
         self.init_dilation()
         self.init_test_case()
@@ -75,12 +77,16 @@ class TestConv2dOp(OpTest):
             'pad': self.pad,
             'dilation': self.dilations
         }
-        input = np.random.random(self.input_size).astype("float32")
-        filter = np.random.random(self.filter_size).astype("float32")
+
+        input = np.random.random(self.input_size).astype(self.dtype)
+        filter = np.random.random(self.filter_size).astype(self.dtype)
         output = conv2d_forward_naive(input, filter, self.groups,
-                                      conv2d_param).astype('float32')
+                                      conv2d_param).astype(self.dtype)
 
-        self.inputs = {'Input': input, 'Filter': filter}
+        self.inputs = {
+            'Input': OpTest.np_dtype_to_fluid_dtype(input),
+            'Filter': OpTest.np_dtype_to_fluid_dtype(filter)
+        }
         self.attrs = {
             'strides': self.stride,
             'paddings': self.pad,
@@ -99,6 +105,8 @@ class TestConv2dOp(OpTest):
             self.check_output()
 
     def test_check_grad(self):
+        if self.dtype == np.float16:
+            return
         if self.use_cudnn:
             place = core.CUDAPlace(0)
             self.check_grad_with_place(
@@ -111,6 +119,8 @@ class TestConv2dOp(OpTest):
                 set(['Input', 'Filter']), 'Output', max_relative_error=0.02)
 
     def test_check_grad_no_filter(self):
+        if self.dtype == np.float16:
+            return
         if self.use_cudnn:
             place = core.CUDAPlace(0)
             self.check_grad_with_place(
@@ -126,6 +136,8 @@ class TestConv2dOp(OpTest):
                 no_grad_set=set(['Filter']))
 
     def test_check_grad_no_input(self):
+        if self.dtype == np.float16:
+            return
         if self.use_cudnn:
             place = core.CUDAPlace(0)
             self.check_grad_with_place(
@@ -154,8 +166,8 @@ class TestConv2dOp(OpTest):
     def init_group(self):
         self.groups = 1
 
-    def init_op_type(self):
-        self.op_type = "conv2d"
+    def init_kernel_type(self):
+        pass
 
 
 class TestWithPad(TestConv2dOp):
@@ -227,39 +239,105 @@ class TestWithInput1x1Filter1x1(TestConv2dOp):
 
 #----------------Conv2dCUDNN----------------
 class TestCUDNN(TestConv2dOp):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "conv2d"
+
+
+class TestFP16CUDNN(TestConv2dOp):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=2e-2)
 
 
 class TestCUDNNWithPad(TestWithPad):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "conv2d"
+
+
+class TestFP16CUDNNWithPad(TestWithPad):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=2e-2)
 
 
 class TestCUDNNWithStride(TestWithStride):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "conv2d"
+
+
+class TestFP16CUDNNWithStride(TestWithStride):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=2e-2)
 
 
 class TestCUDNNWithGroup(TestWithGroup):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "conv2d"
+
+
+class TestFP16CUDNNWithGroup(TestWithGroup):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=2e-2)
 
 
 class TestCUDNNWith1x1(TestWith1x1):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "conv2d"
+
+
+class TestFP16CUDNNWith1x1(TestWith1x1):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=2e-2)
 
 
 class TestCUDNNWithInput1x1Filter1x1(TestWithInput1x1Filter1x1):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "conv2d"
+
+
+class TestFP16CUDNNWithInput1x1Filter1x1(TestWithInput1x1Filter1x1):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=2e-2)
 
 
 class TestDepthwiseConv(TestConv2dOp):
@@ -295,21 +373,18 @@ class TestDepthwiseConv2(TestConv2dOp):
 
 #----------------Conv2dMKLDNN----------------
 class TestMKLDNN(TestConv2dOp):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_mkldnn = True
-        self.op_type = "conv2d"
 
 
 class TestMKLDNNWithPad(TestWithPad):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_mkldnn = True
-        self.op_type = "conv2d"
 
 
 class TestMKLDNNWithStride(TestWithStride):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_mkldnn = True
-        self.op_type = "conv2d"
 
 
 if __name__ == '__main__':

python/paddle/fluid/tests/unittests/test_detection_output_op.py

-#   Copyright (c) 2018 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.
-
-import unittest
-import numpy as np
-from op_test import OpTest
-
-
-class TestUnpoolOp(OpTest):
-    def setUp(self):
-        self.op_type = "detection_output"
-        self.init_test_case()
-
-        #loc.shape ((1, 4, 4, 1, 1))
-        #conf.shape ((1, 4, 2, 1, 1))
-
-        loc = np.array([[[[[0.1]], [[0.1]], [[0.1]], [[0.1]]],
-                         [[[0.1]], [[0.1]], [[0.1]], [[0.1]]],
-                         [[[0.1]], [[0.1]], [[0.1]], [[0.1]]],
-                         [[[0.1]], [[0.1]], [[0.1]], [[0.1]]]]])
-        conf = np.array([[[[[0.1]], [[0.9]]], [[[0.2]], [[0.8]]],
-                          [[[0.3]], [[0.7]]], [[[0.4]], [[0.6]]]]])
-        priorbox = np.array([
-            0.1, 0.1, 0.5, 0.5, 0.1, 0.1, 0.2, 0.2, 0.2, 0.2, 0.6, 0.6, 0.1,
-            0.1, 0.2, 0.2, 0.3, 0.3, 0.7, 0.7, 0.1, 0.1, 0.2, 0.2, 0.4, 0.4,
-            0.8, 0.8, 0.1, 0.1, 0.2, 0.2
-        ])
-
-        output = np.array([
-            0, 1, 0.68997443, 0.099959746, 0.099959746, 0.50804031, 0.50804031
-        ])
-        self.inputs = {
-            'Loc': loc.astype('float32'),
-            'Conf': conf.astype('float32'),
-            'PriorBox': priorbox.astype('float32')
-        }
-        self.attrs = {
-            'num_classes': self.num_classes,
-            'top_k': self.top_k,
-            'nms_top_k': self.nms_top_k,
-            'background_label_id': self.background_label_id,
-            'nms_threshold': self.nms_threshold,
-            'confidence_threshold': self.confidence_threshold,
-        }
-        self.outputs = {'Out': output.astype('float32')}
-
-    def test_check_output(self):
-        self.check_output()
-
-    def init_test_case(self):
-        self.num_classes = 2
-        self.top_k = 10
-        self.nms_top_k = 20
-        self.background_label_id = 0
-        self.nms_threshold = 0.01
-        self.confidence_threshold = 0.01
-
-
-if __name__ == '__main__':
-    unittest.main()

python/paddle/fluid/tests/unittests/test_layers.py

@@ -220,7 +220,7 @@ class TestBook(unittest.TestCase):
             seq_data = layers.data(
                 name='seq_data', shape=[10, 10], dtype='float32', lod_level=1)
             seq = layers.fc(input=seq_data, size=20)
-            self.assertIsNotNone(layers.sequence_softmax(x=seq))
+            self.assertIsNotNone(layers.sequence_softmax(seq))
         print(str(program))
 
     def test_softmax(self):
@@ -228,7 +228,7 @@ class TestBook(unittest.TestCase):
         with program_guard(program):
             data = layers.data(name='data', shape=[10], dtype='float32')
             hid = layers.fc(input=data, size=20)
-            self.assertIsNotNone(layers.softmax(x=hid))
+            self.assertIsNotNone(layers.softmax(hid))
         print(str(program))
 
     def test_get_places(self):

python/paddle/fluid/tests/unittests/test_lrn_op.py

@@ -41,7 +41,7 @@ class TestLRNOp(OpTest):
         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):
+                for c in range(start, end):
                     ch = i + c
                     if ch < 0 or ch >= self.C:
                         continue

python/paddle/fluid/tests/unittests/test_pool2d_op.py

@@ -78,20 +78,22 @@ def avg_pool2D_forward_naive(x,
 
 class TestPool2d_Op(OpTest):
     def setUp(self):
+        self.op_type = "pool2d"
         self.use_cudnn = False
         self.use_mkldnn = False
+        self.dtype = np.float32
         self.init_test_case()
         self.init_global_pool()
-        self.init_op_type()
+        self.init_kernel_type()
         self.init_pool_type()
         self.init_ceil_mode()
         if self.global_pool:
             self.paddings = [0 for _ in range(len(self.paddings))]
-        input = np.random.random(self.shape).astype("float32")
+        input = np.random.random(self.shape).astype(self.dtype)
         output = self.pool2D_forward_naive(input, self.ksize, self.strides,
                                            self.paddings, self.global_pool,
-                                           self.ceil_mode).astype("float32")
-        self.inputs = {'X': input}
+                                           self.ceil_mode).astype(self.dtype)
+        self.inputs = {'X': OpTest.np_dtype_to_fluid_dtype(input)}
 
         self.attrs = {
             'strides': self.strides,
@@ -105,7 +107,7 @@ class TestPool2d_Op(OpTest):
             'data_format': 'AnyLayout'  # TODO(dzhwinter) : should be fix latter
         }
 
-        self.outputs = {'Out': output.astype('float32')}
+        self.outputs = {'Out': output}
 
     def test_check_output(self):
         if self.use_cudnn:
@@ -115,6 +117,8 @@ class TestPool2d_Op(OpTest):
             self.check_output()
 
     def test_check_grad(self):
+        if self.dtype == np.float16:
+            return
         if self.use_cudnn and self.pool_type != "max":
             place = core.CUDAPlace(0)
             self.check_grad_with_place(
@@ -128,8 +132,8 @@ class TestPool2d_Op(OpTest):
         self.strides = [1, 1]
         self.paddings = [0, 0]
 
-    def init_op_type(self):
-        self.op_type = "pool2d"
+    def init_kernel_type(self):
+        pass
 
     def init_pool_type(self):
         self.pool_type = "avg"
@@ -149,9 +153,6 @@ class TestCase1(TestPool2d_Op):
         self.strides = [1, 1]
         self.paddings = [0, 0]
 
-    def init_op_type(self):
-        self.op_type = "pool2d"
-
     def init_pool_type(self):
         self.pool_type = "avg"
         self.pool2D_forward_naive = avg_pool2D_forward_naive
@@ -167,9 +168,6 @@ class TestCase2(TestPool2d_Op):
         self.strides = [1, 1]
         self.paddings = [1, 1]
 
-    def init_op_type(self):
-        self.op_type = "pool2d"
-
     def init_pool_type(self):
         self.pool_type = "avg"
         self.pool2D_forward_naive = avg_pool2D_forward_naive
@@ -179,27 +177,18 @@ class TestCase2(TestPool2d_Op):
 
 
 class TestCase3(TestPool2d_Op):
-    def init_op_type(self):
-        self.op_type = "pool2d"
-
     def init_pool_type(self):
         self.pool_type = "max"
         self.pool2D_forward_naive = max_pool2D_forward_naive
 
 
 class TestCase4(TestCase1):
-    def init_op_type(self):
-        self.op_type = "pool2d"
-
     def init_pool_type(self):
         self.pool_type = "max"
         self.pool2D_forward_naive = max_pool2D_forward_naive
 
 
 class TestCase5(TestCase2):
-    def init_op_type(self):
-        self.op_type = "pool2d"
-
     def init_pool_type(self):
         self.pool_type = "max"
         self.pool2D_forward_naive = max_pool2D_forward_naive
@@ -207,39 +196,105 @@ class TestCase5(TestCase2):
 
 #--------------------test pool2d--------------------
 class TestCUDNNCase1(TestPool2d_Op):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "pool2d"
+
+
+class TestFP16CUDNNCase1(TestPool2d_Op):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=1e-3)
 
 
 class TestCUDNNCase2(TestCase1):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "pool2d"
+
+
+class TestFP16CUDNNCase2(TestCase1):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=1e-3)
 
 
 class TestCUDNNCase3(TestCase2):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "pool2d"
+
+
+class TestFP16CUDNNCase3(TestCase2):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=1e-3)
 
 
 class TestCUDNNCase4(TestCase3):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "pool2d"
+
+
+class TestFP16CUDNNCase4(TestCase3):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=1e-3)
 
 
 class TestCUDNNCase5(TestCase4):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "pool2d"
+
+
+class TestFP16CUDNNCase5(TestCase4):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=1e-3)
 
 
 class TestCUDNNCase6(TestCase5):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_cudnn = True
-        self.op_type = "pool2d"
+
+
+class TestFP16CUDNNCase6(TestCase5):
+    def init_kernel_type(self):
+        self.use_cudnn = True
+        self.dtype = np.float16
+
+    def test_check_output(self):
+        if core.is_compiled_with_cuda():
+            place = core.CUDAPlace(0)
+            if core.is_float16_supported(place):
+                self.check_output_with_place(place, atol=1e-3)
 
 
 class TestCeilModeCase1(TestCUDNNCase1):
@@ -264,39 +319,33 @@ class TestCeilModeCase4(TestCase2):
 
 #--------------------test pool2d MKLDNN--------------------
 class TestMKLDNNCase1(TestPool2d_Op):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_mkldnn = True
-        self.op_type = "pool2d"
 
 
 class TestMKLDNNCase2(TestCase1):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_mkldnn = True
-        self.op_type = "pool2d"
 
 
 class TestMKLDNNCase3(TestCase2):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_mkldnn = True
-        self.op_type = "pool2d"
 
 
 class TestMKLDNNCase4(TestCase3):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_mkldnn = True
-        self.op_type = "pool2d"
 
 
 class TestMKLDNNCase5(TestCase4):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_mkldnn = True
-        self.op_type = "pool2d"
 
 
 class TestMKLDNNCase6(TestCase5):
-    def init_op_type(self):
+    def init_kernel_type(self):
         self.use_mkldnn = True
-        self.op_type = "pool2d"
 
 
 if __name__ == '__main__':

python/paddle/fluid/tests/unittests/test_profiler.py

@@ -31,8 +31,22 @@ class TestProfiler(unittest.TestCase):
 
         with fluid.program_guard(main_program, startup_program):
             image = fluid.layers.data(name='x', shape=[784], dtype='float32')
-            hidden1 = fluid.layers.fc(input=image, size=128, act='relu')
-            hidden2 = fluid.layers.fc(input=hidden1, size=64, act='relu')
+            hidden1 = fluid.layers.fc(input=image, size=64, act='relu')
+            i = layers.zeros(shape=[1], dtype='int64')
+            counter = fluid.layers.zeros(
+                shape=[1], dtype='int64', force_cpu=True)
+            until = layers.fill_constant([1], dtype='int64', value=10)
+            data_arr = layers.array_write(hidden1, i)
+            cond = fluid.layers.less_than(x=counter, y=until)
+            while_op = fluid.layers.While(cond=cond)
+            with while_op.block():
+                hidden_n = fluid.layers.fc(input=hidden1, size=64, act='relu')
+                layers.array_write(hidden_n, i, data_arr)
+                fluid.layers.increment(x=counter, value=1, in_place=True)
+                layers.less_than(x=counter, y=until, cond=cond)
+
+            hidden_n = layers.array_read(data_arr, i)
+            hidden2 = fluid.layers.fc(input=hidden_n, size=64, act='relu')
             predict = fluid.layers.fc(input=hidden2, size=10, act='softmax')
             label = fluid.layers.data(name='y', shape=[1], dtype='int64')
             cost = fluid.layers.cross_entropy(input=predict, label=label)

python/paddle/fluid/tests/unittests/test_sequence_softmax_op.py

@@ -16,11 +16,15 @@ import unittest
 import numpy as np
 from op_test import OpTest
 from test_softmax_op import stable_softmax
+import paddle.fluid.core as core
 
 
 class TestSequenceSoftmaxOp(OpTest):
     def setUp(self):
         self.op_type = "sequence_softmax"
+        self.use_cudnn = False
+        self.init_op_type()
+
         x = np.random.uniform(0.1, 1, (11, 1)).astype("float32")
         lod = [[0, 4, 5, 8, 11]]
 
@@ -34,13 +38,32 @@ class TestSequenceSoftmaxOp(OpTest):
 
         self.inputs = {"X": (x, lod)}
         self.outputs = {"Out": out}
+        self.attrs = {'use_cudnn': self.use_cudnn, }
+
+    def init_op_type(self):
+        pass
 
     def test_check_output(self):
+        if self.use_cudnn:
+            place = core.CUDAPlace(0)
+            self.check_output_with_place(place, atol=1e-5)
+        else:
             self.check_output()
 
     def test_check_grad(self):
+        if self.use_cudnn:
+            place = core.CUDAPlace(0)
+            self.check_grad_with_place(
+                place, ["X"], "Out", max_relative_error=0.01)
+        else:
             self.check_grad(["X"], "Out", max_relative_error=0.01)
 
 
+# ----------------cudnn Sequencesoftmax----------------
+class TestSequenceSoftmaxCUDNNOp(TestSequenceSoftmaxOp):
+    def init_op_type(self):
+        self.use_cudnn = True
+
+
 if __name__ == "__main__":
     unittest.main()

python/paddle/fluid/tests/unittests/test_softmax_op.py

@@ -15,6 +15,7 @@
 import unittest
 import numpy as np
 from op_test import OpTest
+import paddle.fluid.core as core
 
 
 def stable_softmax(x):
@@ -27,18 +28,37 @@ def stable_softmax(x):
 class TestSoftmaxOp(OpTest):
     def setUp(self):
         self.op_type = "softmax"
+        self.use_cudnn = False
         self.inputs = {
             'X': np.random.uniform(0.1, 1, [10, 10]).astype("float32")
         }
         self.outputs = {
             'Out': np.apply_along_axis(stable_softmax, 1, self.inputs['X'])
         }
+        self.attrs = {'use_cudnn': self.use_cudnn, }
+
+    def init_op_type(self):
+        pass
 
     def test_check_output(self):
+        if self.use_cudnn:
+            place = core.CUDAPlace(0)
+            self.check_output_with_place(place, atol=1e-5)
+        else:
             self.check_output()
 
     def test_check_grad(self):
-        self.check_grad(['X'], 'Out')
+        if self.use_cudnn:
+            place = core.CUDAPlace(0)
+            self.check_grad_with_place(
+                place, ["X"], "Out", max_relative_error=0.01)
+        else:
+            self.check_grad(["X"], "Out", max_relative_error=0.01)
+
+
+class TestSoftmaxCUDNNOp(TestSoftmaxOp):
+    def init_op_type(self):
+        self.use_cudnn = True
 
 
 if __name__ == "__main__":

python/paddle/fluid/tests/unittests/test_softmax_with_cross_entropy_op.py

@@ -26,7 +26,7 @@ class TestSoftmaxWithCrossEntropyOp(OpTest):
 
     def setUp(self):
         self.op_type = "softmax_with_cross_entropy"
-        batch_size = 2
+        batch_size = 41
         class_num = 37
 
         logits = np.random.uniform(0.1, 1.0,
@@ -59,7 +59,7 @@ class TestSoftmaxWithCrossEntropyOp2(OpTest):
 
     def setUp(self):
         self.op_type = "softmax_with_cross_entropy"
-        batch_size = 2
+        batch_size = 41
         class_num = 37
 
         logits = np.random.uniform(0.1, 1.0,

tools/timeline.py

@@ -121,27 +121,34 @@ class Timeline(object):
 
     def _allocate_pids(self):
         for event in self._profile_pb.events:
-            if event.device_id not in self._devices:
+            if event.type == profiler_pb2.Event.CPU:
+                if (event.device_id, "CPU") not in self._devices:
                     pid = self._allocate_pid()
-                self._devices[event.device_id] = pid
-                if event.device_id >= 0:
-                    self._chrome_trace.emit_pid("gpu:%s:stream:%d" %
-                                                (pid, event.stream_id), pid)
-                elif event.device_id == -1:
-                    self._chrome_trace.emit_pid("cpu:thread_hash:%d" %
-                                                event.stream_id, pid)
+                    self._devices[(event.device_id, "CPU")] = pid
+                    self._chrome_trace.emit_pid("cpu:block:%d" %
+                                                (event.device_id), pid)
+            elif event.type == profiler_pb2.Event.GPUKernel:
+                if (event.device_id, "GPUKernel") not in self._devices:
+                    pid = self._allocate_pid()
+                    self._devices[(event.device_id, "GPUKernel")] = pid
+                    self._chrome_trace.emit_pid("gpu:%d" % (event.device_id),
+                                                pid)
 
     def _allocate_events(self):
         for event in self._profile_pb.events:
-            pid = self._devices[event.device_id]
+            if event.type == profiler_pb2.Event.CPU:
+                type = "CPU"
+            elif event.type == profiler_pb2.Event.GPUKernel:
+                type = "GPUKernel"
+            pid = self._devices[(event.device_id, type)]
             args = {'name': event.name}
             if event.memcopy.bytes > 0:
                 args = {'mem_bytes': event.memcopy.bytes}
             # TODO(panyx0718): Chrome tracing only handles ms. However, some
             # ops takes micro-seconds. Hence, we keep the ns here.
-            self._chrome_trace.emit_region(event.start_ns,
-                                           (event.end_ns - event.start_ns) /
-                                           1.0, pid, 0, 'Op', event.name, args)
+            self._chrome_trace.emit_region(
+                event.start_ns, (event.end_ns - event.start_ns) / 1.0, pid,
+                event.sub_device_id, 'Op', event.name, args)
 
     def generate_chrome_trace(self):
         self._allocate_pids()