diff --git a/CMakeLists.txt b/CMakeLists.txt index c86889c05c8cf0d521dce9adbf3e918ba91729a1..0ec65bac84b0b0d89123473a8941f80c90f1b339 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -53,7 +53,7 @@ option(COVERALLS_UPLOAD "Package code coverage data to coveralls" OFF) option(ON_TRAVIS "Exclude special unit test on Travis CI" OFF) option(WITH_C_API "Compile PaddlePaddle with C-API(Prediction)" OFF) # TODO: Only compile PaddlePaddle fluid version by WITH_FLUID option. -option(WITH_FLUID "Compile PaddlePaddle fluid only(TODO)" ON) +option(WITH_FLUID "Compile PaddlePaddle fluid only(TODO)" OFF) option(WITH_GOLANG "Compile PaddlePaddle with GOLANG" OFF) option(GLIDE_INSTALL "Download and install go dependencies " ON) option(USE_NNPACK "Compile PaddlePaddle with NNPACK library" OFF) diff --git a/Dockerfile b/Dockerfile index 60e76c7f2ede6beaca11659020d5991a75d5b741..fbec88c7966d6ea93495519843d6cda63f622661 100644 --- a/Dockerfile +++ b/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' && \ diff --git a/doc/design/images/asgd.gif b/doc/fluid/design/algorithm/images/asgd.gif similarity index 100% rename from doc/design/images/asgd.gif rename to doc/fluid/design/algorithm/images/asgd.gif diff --git a/doc/design/images/theta_star.gif b/doc/fluid/design/algorithm/images/theta_star.gif similarity index 100% rename from doc/design/images/theta_star.gif rename to doc/fluid/design/algorithm/images/theta_star.gif diff --git a/doc/design/parameter_average.md b/doc/fluid/design/algorithm/parameter_average.md similarity index 100% rename from doc/design/parameter_average.md rename to doc/fluid/design/algorithm/parameter_average.md diff --git a/doc/design/build_system/README.md b/doc/fluid/design/concepts/README.md similarity index 100% rename from doc/design/build_system/README.md rename to doc/fluid/design/concepts/README.md diff --git a/doc/design/block.md b/doc/fluid/design/concepts/block.md similarity index 100% rename from doc/design/block.md rename to doc/fluid/design/concepts/block.md diff --git a/doc/design/cpp_data_feeding.md b/doc/fluid/design/concepts/cpp_data_feeding.md similarity index 100% rename from doc/design/cpp_data_feeding.md rename to doc/fluid/design/concepts/cpp_data_feeding.md diff --git a/doc/design/executor.md b/doc/fluid/design/concepts/executor.md similarity index 100% rename from doc/design/executor.md rename to doc/fluid/design/concepts/executor.md diff --git a/doc/design/functions_operators_layers.md b/doc/fluid/design/concepts/functions_operators_layers.md similarity index 100% rename from doc/design/functions_operators_layers.md rename to doc/fluid/design/concepts/functions_operators_layers.md diff --git a/paddle/fluid/framework/lod_tensor.md b/doc/fluid/design/concepts/lod_tensor.md similarity index 100% rename from paddle/fluid/framework/lod_tensor.md rename to doc/fluid/design/concepts/lod_tensor.md diff --git a/doc/design/program.md b/doc/fluid/design/concepts/program.md similarity index 100% rename from doc/design/program.md rename to doc/fluid/design/concepts/program.md diff --git a/doc/design/scope.md b/doc/fluid/design/concepts/scope.md similarity index 100% rename from doc/design/scope.md rename to doc/fluid/design/concepts/scope.md diff --git a/paddle/fluid/framework/tensor.md b/doc/fluid/design/concepts/tensor.md similarity index 100% rename from paddle/fluid/framework/tensor.md rename to doc/fluid/design/concepts/tensor.md diff --git a/doc/design/tensor_array.md b/doc/fluid/design/concepts/tensor_array.md similarity index 100% rename from doc/design/tensor_array.md rename to doc/fluid/design/concepts/tensor_array.md diff --git a/doc/design/var_desc.md b/doc/fluid/design/concepts/var_desc.md similarity index 100% rename from doc/design/var_desc.md rename to doc/fluid/design/concepts/var_desc.md diff --git a/paddle/fluid/framework/variable.md b/doc/fluid/design/concepts/variable.md similarity index 100% rename from paddle/fluid/framework/variable.md rename to doc/fluid/design/concepts/variable.md diff --git a/doc/design/concurrent_programming.md b/doc/fluid/design/concurrent/concurrent_programming.md similarity index 100% rename from doc/design/concurrent_programming.md rename to doc/fluid/design/concurrent/concurrent_programming.md diff --git a/doc/design/csp.md b/doc/fluid/design/concurrent/csp.md similarity index 100% rename from doc/design/csp.md rename to doc/fluid/design/concurrent/csp.md diff --git a/doc/design/parallel_do.md b/doc/fluid/design/concurrent/parallel_do.md similarity index 100% rename from doc/design/parallel_do.md rename to doc/fluid/design/concurrent/parallel_do.md diff --git a/doc/design/float16.md b/doc/fluid/design/data_type/float16.md similarity index 100% rename from doc/design/float16.md rename to doc/fluid/design/data_type/float16.md diff --git a/doc/design/ops/images/2_level_rnn.dot b/doc/fluid/design/dynamic_rnn/2_level_rnn.dot similarity index 100% rename from doc/design/ops/images/2_level_rnn.dot rename to doc/fluid/design/dynamic_rnn/2_level_rnn.dot diff --git a/doc/design/ops/images/2_level_rnn.png b/doc/fluid/design/dynamic_rnn/2_level_rnn.png similarity index 100% rename from doc/design/ops/images/2_level_rnn.png rename to doc/fluid/design/dynamic_rnn/2_level_rnn.png diff --git a/doc/design/ops/images/rnn.dot b/doc/fluid/design/dynamic_rnn/rnn.dot similarity index 100% rename from doc/design/ops/images/rnn.dot rename to doc/fluid/design/dynamic_rnn/rnn.dot diff --git a/doc/design/ops/images/rnn.jpg b/doc/fluid/design/dynamic_rnn/rnn.jpg similarity index 100% rename from doc/design/ops/images/rnn.jpg rename to doc/fluid/design/dynamic_rnn/rnn.jpg diff --git a/doc/design/ops/rnn.md b/doc/fluid/design/dynamic_rnn/rnn.md similarity index 100% rename from doc/design/ops/rnn.md rename to doc/fluid/design/dynamic_rnn/rnn.md diff --git a/doc/design/ops/images/rnn.png b/doc/fluid/design/dynamic_rnn/rnn.png similarity index 100% rename from doc/design/ops/images/rnn.png rename to doc/fluid/design/dynamic_rnn/rnn.png diff --git a/doc/design/ops/images/rnn_2level_data.dot b/doc/fluid/design/dynamic_rnn/rnn_2level_data.dot similarity index 100% rename from doc/design/ops/images/rnn_2level_data.dot rename to doc/fluid/design/dynamic_rnn/rnn_2level_data.dot diff --git a/doc/design/ops/images/rnn_2level_data.png b/doc/fluid/design/dynamic_rnn/rnn_2level_data.png similarity index 100% rename from doc/design/ops/images/rnn_2level_data.png rename to doc/fluid/design/dynamic_rnn/rnn_2level_data.png diff --git a/paddle/fluid/operators/op_documentation/rnn_design.md b/doc/fluid/design/dynamic_rnn/rnn_design.md similarity index 100% rename from paddle/fluid/operators/op_documentation/rnn_design.md rename to doc/fluid/design/dynamic_rnn/rnn_design.md diff --git a/doc/design/if_else_op.md b/doc/fluid/design/execution/if_else_op.md similarity index 100% rename from doc/design/if_else_op.md rename to doc/fluid/design/execution/if_else_op.md diff --git a/doc/design/switch.md b/doc/fluid/design/execution/switch.md similarity index 100% rename from doc/design/switch.md rename to doc/fluid/design/execution/switch.md diff --git a/doc/design/multi_language_interface/00.why_plain_c.md b/doc/fluid/design/interface/00.why_plain_c.md similarity index 100% rename from doc/design/multi_language_interface/00.why_plain_c.md rename to doc/fluid/design/interface/00.why_plain_c.md diff --git a/doc/design/multi_language_interface/01.inference_implementation.md b/doc/fluid/design/interface/01.inference_implementation.md similarity index 100% rename from doc/design/multi_language_interface/01.inference_implementation.md rename to doc/fluid/design/interface/01.inference_implementation.md diff --git a/paddle/fluid/memory/README.md b/doc/fluid/design/memory/README.md similarity index 100% rename from paddle/fluid/memory/README.md rename to doc/fluid/design/memory/README.md diff --git a/doc/design/images/control_flow_graph.png b/doc/fluid/design/memory/images/control_flow_graph.png similarity index 100% rename from doc/design/images/control_flow_graph.png rename to doc/fluid/design/memory/images/control_flow_graph.png diff --git a/doc/design/images/dataflow_equations.png b/doc/fluid/design/memory/images/dataflow_equations.png similarity index 100% rename from doc/design/images/dataflow_equations.png rename to doc/fluid/design/memory/images/dataflow_equations.png diff --git a/doc/design/images/deep_learning.png b/doc/fluid/design/memory/images/deep_learning.png similarity index 100% rename from doc/design/images/deep_learning.png rename to doc/fluid/design/memory/images/deep_learning.png diff --git a/doc/design/memory_optimization.md b/doc/fluid/design/memory/memory_optimization.md similarity index 100% rename from doc/design/memory_optimization.md rename to doc/fluid/design/memory/memory_optimization.md diff --git a/doc/design/backward.md b/doc/fluid/design/modules/backward.md similarity index 100% rename from doc/design/backward.md rename to doc/fluid/design/modules/backward.md diff --git a/paddle/fluid/operators/op_documentation/batch_norm_op.md b/doc/fluid/design/modules/batch_norm_op.md similarity index 100% rename from paddle/fluid/operators/op_documentation/batch_norm_op.md rename to doc/fluid/design/modules/batch_norm_op.md diff --git a/doc/design/evaluator.md b/doc/fluid/design/modules/evaluator.md similarity index 100% rename from doc/design/evaluator.md rename to doc/fluid/design/modules/evaluator.md diff --git a/paddle/fluid/operators/images/batch_norm_fork.dot b/doc/fluid/design/modules/images/batch_norm_fork.dot similarity index 100% rename from paddle/fluid/operators/images/batch_norm_fork.dot rename to doc/fluid/design/modules/images/batch_norm_fork.dot diff --git a/paddle/fluid/operators/images/batch_norm_fork.png b/doc/fluid/design/modules/images/batch_norm_fork.png similarity index 100% rename from paddle/fluid/operators/images/batch_norm_fork.png rename to doc/fluid/design/modules/images/batch_norm_fork.png diff --git a/paddle/fluid/operators/images/batch_norm_op_kernel.png b/doc/fluid/design/modules/images/batch_norm_op_kernel.png similarity index 100% rename from paddle/fluid/operators/images/batch_norm_op_kernel.png rename to doc/fluid/design/modules/images/batch_norm_op_kernel.png diff --git a/doc/design/images/feed_forward.png b/doc/fluid/design/modules/images/feed_forward.png similarity index 100% rename from doc/design/images/feed_forward.png rename to doc/fluid/design/modules/images/feed_forward.png diff --git a/doc/design/images/feed_forward_regularized.png b/doc/fluid/design/modules/images/feed_forward_regularized.png similarity index 100% rename from doc/design/images/feed_forward_regularized.png rename to doc/fluid/design/modules/images/feed_forward_regularized.png diff --git a/doc/design/images/l1_regularization.png b/doc/fluid/design/modules/images/l1_regularization.png similarity index 100% rename from doc/design/images/l1_regularization.png rename to doc/fluid/design/modules/images/l1_regularization.png diff --git a/doc/design/images/l2_regularization.png b/doc/fluid/design/modules/images/l2_regularization.png similarity index 100% rename from doc/design/images/l2_regularization.png rename to doc/fluid/design/modules/images/l2_regularization.png diff --git a/doc/design/images/loss_equation.png b/doc/fluid/design/modules/images/loss_equation.png similarity index 100% rename from doc/design/images/loss_equation.png rename to doc/fluid/design/modules/images/loss_equation.png diff --git a/doc/design/infer_var_type.md b/doc/fluid/design/modules/infer_var_type.md similarity index 100% rename from doc/design/infer_var_type.md rename to doc/fluid/design/modules/infer_var_type.md diff --git a/paddle/fluid/operators/op_documentation/net_op_design.md b/doc/fluid/design/modules/net_op_design.md similarity index 100% rename from paddle/fluid/operators/op_documentation/net_op_design.md rename to doc/fluid/design/modules/net_op_design.md diff --git a/doc/design/optimizer.md b/doc/fluid/design/modules/optimizer.md similarity index 100% rename from doc/design/optimizer.md rename to doc/fluid/design/modules/optimizer.md diff --git a/doc/design/prune.md b/doc/fluid/design/modules/prune.md similarity index 100% rename from doc/design/prune.md rename to doc/fluid/design/modules/prune.md diff --git a/doc/design/python_api.md b/doc/fluid/design/modules/python_api.md similarity index 100% rename from doc/design/python_api.md rename to doc/fluid/design/modules/python_api.md diff --git a/doc/design/register_grad_op.md b/doc/fluid/design/modules/register_grad_op.md similarity index 100% rename from doc/design/register_grad_op.md rename to doc/fluid/design/modules/register_grad_op.md diff --git a/doc/design/regularization.md b/doc/fluid/design/modules/regularization.md similarity index 100% rename from doc/design/regularization.md rename to doc/fluid/design/modules/regularization.md diff --git a/doc/design/selected_rows.md b/doc/fluid/design/modules/selected_rows.md similarity index 100% rename from doc/design/selected_rows.md rename to doc/fluid/design/modules/selected_rows.md diff --git a/doc/design/api.md b/doc/fluid/design/motivation/api.md similarity index 100% rename from doc/design/api.md rename to doc/fluid/design/motivation/api.md diff --git a/doc/design/fluid-compiler.graffle b/doc/fluid/design/motivation/fluid-compiler.graffle similarity index 100% rename from doc/design/fluid-compiler.graffle rename to doc/fluid/design/motivation/fluid-compiler.graffle diff --git a/doc/design/fluid-compiler.png b/doc/fluid/design/motivation/fluid-compiler.png similarity index 100% rename from doc/design/fluid-compiler.png rename to doc/fluid/design/motivation/fluid-compiler.png diff --git a/doc/design/fluid.md b/doc/fluid/design/motivation/fluid.md similarity index 100% rename from doc/design/fluid.md rename to doc/fluid/design/motivation/fluid.md diff --git a/doc/design/fluid_compiler.md b/doc/fluid/design/motivation/fluid_compiler.md similarity index 100% rename from doc/design/fluid_compiler.md rename to doc/fluid/design/motivation/fluid_compiler.md diff --git a/doc/design/refactorization.md b/doc/fluid/design/motivation/refactorization.md similarity index 100% rename from doc/design/refactorization.md rename to doc/fluid/design/motivation/refactorization.md diff --git a/doc/design/kernel_hint_design.md b/doc/fluid/design/muti_devices/kernel_hint_design.md similarity index 100% rename from doc/design/kernel_hint_design.md rename to doc/fluid/design/muti_devices/kernel_hint_design.md diff --git a/doc/design/kernel_selection.md b/doc/fluid/design/muti_devices/kernel_selection.md similarity index 100% rename from doc/design/kernel_selection.md rename to doc/fluid/design/muti_devices/kernel_selection.md diff --git a/doc/design/operator_kernel_type.md b/doc/fluid/design/muti_devices/operator_kernel_type.md similarity index 100% rename from doc/design/operator_kernel_type.md rename to doc/fluid/design/muti_devices/operator_kernel_type.md diff --git a/doc/design/speech/deep_speech_2.md b/doc/fluid/design/network/deep_speech_2.md similarity index 98% rename from doc/design/speech/deep_speech_2.md rename to doc/fluid/design/network/deep_speech_2.md index cfdc4d6df04344c70d3334626bd38eca997c31ff..af0c6ef36feba9e0239e7a5f81a8dc9108b2471a 100644 --- a/doc/design/speech/deep_speech_2.md +++ b/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
-
+
Figure 1. Archetecture of Deep Speech 2 Network.
@@ -141,7 +141,7 @@ TODO by Assignees ### Beam Search with CTC and LM
-
+
Figure 2. Algorithm for CTC Beam Search Decoder.
diff --git a/doc/design/ops/images/LOD-and-shape-changes-during-decoding.jpg b/doc/fluid/design/network/images/LOD-and-shape-changes-during-decoding.jpg similarity index 100% rename from doc/design/ops/images/LOD-and-shape-changes-during-decoding.jpg rename to doc/fluid/design/network/images/LOD-and-shape-changes-during-decoding.jpg diff --git a/doc/design/speech/image/beam_search.png b/doc/fluid/design/network/images/beam_search.png similarity index 100% rename from doc/design/speech/image/beam_search.png rename to doc/fluid/design/network/images/beam_search.png diff --git a/doc/design/speech/image/ds2_network.png b/doc/fluid/design/network/images/ds2_network.png similarity index 100% rename from doc/design/speech/image/ds2_network.png rename to doc/fluid/design/network/images/ds2_network.png diff --git a/doc/design/ops/sequence_decoder.md b/doc/fluid/design/network/sequence_decoder.md similarity index 100% rename from doc/design/ops/sequence_decoder.md rename to doc/fluid/design/network/sequence_decoder.md diff --git a/doc/design/auto_gradient_check.md b/doc/fluid/design/others/auto_gradient_check.md similarity index 100% rename from doc/design/auto_gradient_check.md rename to doc/fluid/design/others/auto_gradient_check.md diff --git a/doc/design/dcgan.png b/doc/fluid/design/others/dcgan.png similarity index 100% rename from doc/design/dcgan.png rename to doc/fluid/design/others/dcgan.png diff --git a/doc/design/gan_api.md b/doc/fluid/design/others/gan_api.md similarity index 100% rename from doc/design/gan_api.md rename to doc/fluid/design/others/gan_api.md diff --git a/doc/design/graph.md b/doc/fluid/design/others/graph.md similarity index 100% rename from doc/design/graph.md rename to doc/fluid/design/others/graph.md diff --git a/doc/design/graph_survey.md b/doc/fluid/design/others/graph_survey.md similarity index 100% rename from doc/design/graph_survey.md rename to doc/fluid/design/others/graph_survey.md diff --git a/doc/design/images/graph_construction_example.bash b/doc/fluid/design/others/images/graph_construction_example.bash similarity index 100% rename from doc/design/images/graph_construction_example.bash rename to doc/fluid/design/others/images/graph_construction_example.bash diff --git a/doc/design/images/graph_construction_example.dot b/doc/fluid/design/others/images/graph_construction_example.dot similarity index 100% rename from doc/design/images/graph_construction_example.dot rename to doc/fluid/design/others/images/graph_construction_example.dot diff --git a/doc/design/images/graph_construction_example_all.png b/doc/fluid/design/others/images/graph_construction_example_all.png similarity index 100% rename from doc/design/images/graph_construction_example_all.png rename to doc/fluid/design/others/images/graph_construction_example_all.png diff --git a/doc/design/images/graph_construction_example_forward_backward.png b/doc/fluid/design/others/images/graph_construction_example_forward_backward.png similarity index 100% rename from doc/design/images/graph_construction_example_forward_backward.png rename to doc/fluid/design/others/images/graph_construction_example_forward_backward.png diff --git a/doc/design/images/graph_construction_example_forward_only.png b/doc/fluid/design/others/images/graph_construction_example_forward_only.png similarity index 100% rename from doc/design/images/graph_construction_example_forward_only.png rename to doc/fluid/design/others/images/graph_construction_example_forward_only.png diff --git a/doc/design/parameters_in_cpp.md b/doc/fluid/design/others/parameters_in_cpp.md similarity index 100% rename from doc/design/parameters_in_cpp.md rename to doc/fluid/design/others/parameters_in_cpp.md diff --git a/doc/design/simple_op_design.md b/doc/fluid/design/others/simple_op_design.md similarity index 100% rename from doc/design/simple_op_design.md rename to doc/fluid/design/others/simple_op_design.md diff --git a/doc/design/test.dot b/doc/fluid/design/others/test.dot similarity index 100% rename from doc/design/test.dot rename to doc/fluid/design/others/test.dot diff --git a/doc/design/test.dot.png b/doc/fluid/design/others/test.dot.png similarity index 100% rename from doc/design/test.dot.png rename to doc/fluid/design/others/test.dot.png diff --git a/doc/design/ci_build_whl.png b/doc/fluid/dev/ci_build_whl.png similarity index 100% rename from doc/design/ci_build_whl.png rename to doc/fluid/dev/ci_build_whl.png diff --git a/paddle/fluid/operators/op_documentation/name_convention.md b/doc/fluid/dev/name_convention.md similarity index 100% rename from paddle/fluid/operators/op_documentation/name_convention.md rename to doc/fluid/dev/name_convention.md diff --git a/paddle/fluid/operators/op_documentation/op_markdown_format.md b/doc/fluid/dev/op_markdown_format.md similarity index 100% rename from paddle/fluid/operators/op_documentation/op_markdown_format.md rename to doc/fluid/dev/op_markdown_format.md diff --git a/doc/design/releasing_process.md b/doc/fluid/dev/releasing_process.md similarity index 100% rename from doc/design/releasing_process.md rename to doc/fluid/dev/releasing_process.md diff --git a/doc/design/support_new_device.md b/doc/fluid/dev/support_new_device.md similarity index 100% rename from doc/design/support_new_device.md rename to doc/fluid/dev/support_new_device.md diff --git a/doc/design/reader/README.md b/doc/fluid/getstarted/concepts/reader/README.md similarity index 100% rename from doc/design/reader/README.md rename to doc/fluid/getstarted/concepts/reader/README.md diff --git a/doc/design/model_format.md b/doc/fluid/getstarted/concepts/save_model/model_format.md similarity index 100% rename from doc/design/model_format.md rename to doc/fluid/getstarted/concepts/save_model/model_format.md diff --git a/doc/design/error_clip.md b/doc/fluid/howto/performance/error_clip.md similarity index 100% rename from doc/design/error_clip.md rename to doc/fluid/howto/performance/error_clip.md diff --git a/doc/design/images/profiler.png b/doc/fluid/howto/performance/images/profiler.png similarity index 100% rename from doc/design/images/profiler.png rename to doc/fluid/howto/performance/images/profiler.png diff --git a/doc/design/profiler.md b/doc/fluid/howto/performance/profiler.md similarity index 100% rename from doc/design/profiler.md rename to doc/fluid/howto/performance/profiler.md diff --git a/doc/design/images/multigpu_allreduce.graffle b/doc/fluid/howto/third_party/images/multigpu_allreduce.graffle similarity index 100% rename from doc/design/images/multigpu_allreduce.graffle rename to doc/fluid/howto/third_party/images/multigpu_allreduce.graffle diff --git a/doc/design/images/multigpu_allreduce.png b/doc/fluid/howto/third_party/images/multigpu_allreduce.png similarity index 100% rename from doc/design/images/multigpu_allreduce.png rename to doc/fluid/howto/third_party/images/multigpu_allreduce.png diff --git a/doc/design/images/multigpu_before_convert.graffle b/doc/fluid/howto/third_party/images/multigpu_before_convert.graffle similarity index 100% rename from doc/design/images/multigpu_before_convert.graffle rename to doc/fluid/howto/third_party/images/multigpu_before_convert.graffle diff --git a/doc/design/images/multigpu_before_convert.png b/doc/fluid/howto/third_party/images/multigpu_before_convert.png similarity index 100% rename from doc/design/images/multigpu_before_convert.png rename to doc/fluid/howto/third_party/images/multigpu_before_convert.png diff --git a/doc/design/mkl/mkldnn_fluid.md b/doc/fluid/howto/third_party/mkldnn_fluid.md similarity index 100% rename from doc/design/mkl/mkldnn_fluid.md rename to doc/fluid/howto/third_party/mkldnn_fluid.md diff --git a/doc/design/paddle_nccl.md b/doc/fluid/howto/third_party/paddle_nccl.md similarity index 100% rename from doc/design/paddle_nccl.md rename to doc/fluid/howto/third_party/paddle_nccl.md diff --git a/doc/design/cluster_train/README.md b/doc/v2/design/cluster_train/README.md similarity index 100% rename from doc/design/cluster_train/README.md rename to doc/v2/design/cluster_train/README.md diff --git a/doc/design/cluster_train/checkpointing.md b/doc/v2/design/cluster_train/checkpointing.md similarity index 100% rename from doc/design/cluster_train/checkpointing.md rename to doc/v2/design/cluster_train/checkpointing.md diff --git a/doc/design/cluster_train/data_dispatch.md b/doc/v2/design/cluster_train/data_dispatch.md similarity index 100% rename from doc/design/cluster_train/data_dispatch.md rename to doc/v2/design/cluster_train/data_dispatch.md diff --git a/doc/design/cluster_train/large_model_dist_train.md b/doc/v2/design/cluster_train/large_model_dist_train.md similarity index 100% rename from doc/design/cluster_train/large_model_dist_train.md rename to doc/v2/design/cluster_train/large_model_dist_train.md diff --git a/doc/design/cluster_train/master_server.md b/doc/v2/design/cluster_train/master_server.md similarity index 100% rename from doc/design/cluster_train/master_server.md rename to doc/v2/design/cluster_train/master_server.md diff --git a/doc/design/cluster_train/pserver_client.md b/doc/v2/design/cluster_train/pserver_client.md similarity index 100% rename from doc/design/cluster_train/pserver_client.md rename to doc/v2/design/cluster_train/pserver_client.md diff --git a/doc/design/cluster_train/remote_parameter_updater.md b/doc/v2/design/cluster_train/remote_parameter_updater.md similarity index 100% rename from doc/design/cluster_train/remote_parameter_updater.md rename to doc/v2/design/cluster_train/remote_parameter_updater.md diff --git a/doc/design/cluster_train/save_model.md b/doc/v2/design/cluster_train/save_model.md similarity index 100% rename from doc/design/cluster_train/save_model.md rename to doc/v2/design/cluster_train/save_model.md diff --git a/doc/design/cluster_train/src/checkpointing.png b/doc/v2/design/cluster_train/src/checkpointing.png similarity index 100% rename from doc/design/cluster_train/src/checkpointing.png rename to doc/v2/design/cluster_train/src/checkpointing.png diff --git a/doc/design/cluster_train/src/data_dispatch.png b/doc/v2/design/cluster_train/src/data_dispatch.png similarity index 100% rename from doc/design/cluster_train/src/data_dispatch.png rename to doc/v2/design/cluster_train/src/data_dispatch.png diff --git a/doc/design/cluster_train/src/dataset.graffle b/doc/v2/design/cluster_train/src/dataset.graffle similarity index 100% rename from doc/design/cluster_train/src/dataset.graffle rename to doc/v2/design/cluster_train/src/dataset.graffle diff --git a/doc/design/cluster_train/src/dataset.png b/doc/v2/design/cluster_train/src/dataset.png similarity index 100% rename from doc/design/cluster_train/src/dataset.png rename to doc/v2/design/cluster_train/src/dataset.png diff --git a/doc/design/cluster_train/src/file_storage.graffle b/doc/v2/design/cluster_train/src/file_storage.graffle similarity index 100% rename from doc/design/cluster_train/src/file_storage.graffle rename to doc/v2/design/cluster_train/src/file_storage.graffle diff --git a/doc/design/cluster_train/src/file_storage.png b/doc/v2/design/cluster_train/src/file_storage.png similarity index 100% rename from doc/design/cluster_train/src/file_storage.png rename to doc/v2/design/cluster_train/src/file_storage.png diff --git a/doc/design/cluster_train/src/init_lock.graffle b/doc/v2/design/cluster_train/src/init_lock.graffle similarity index 100% rename from doc/design/cluster_train/src/init_lock.graffle rename to doc/v2/design/cluster_train/src/init_lock.graffle diff --git a/doc/design/cluster_train/src/init_lock.png b/doc/v2/design/cluster_train/src/init_lock.png similarity index 100% rename from doc/design/cluster_train/src/init_lock.png rename to doc/v2/design/cluster_train/src/init_lock.png diff --git a/doc/design/cluster_train/src/paddle-cloud-in-data-center.png b/doc/v2/design/cluster_train/src/paddle-cloud-in-data-center.png similarity index 100% rename from doc/design/cluster_train/src/paddle-cloud-in-data-center.png rename to doc/v2/design/cluster_train/src/paddle-cloud-in-data-center.png diff --git a/doc/design/cluster_train/src/paddle-etcd.graffle b/doc/v2/design/cluster_train/src/paddle-etcd.graffle similarity index 100% rename from doc/design/cluster_train/src/paddle-etcd.graffle rename to doc/v2/design/cluster_train/src/paddle-etcd.graffle diff --git a/doc/design/cluster_train/src/paddle-etcd.png b/doc/v2/design/cluster_train/src/paddle-etcd.png similarity index 100% rename from doc/design/cluster_train/src/paddle-etcd.png rename to doc/v2/design/cluster_train/src/paddle-etcd.png diff --git a/doc/design/cluster_train/src/paddle-model-sharding.graffle b/doc/v2/design/cluster_train/src/paddle-model-sharding.graffle similarity index 100% rename from doc/design/cluster_train/src/paddle-model-sharding.graffle rename to doc/v2/design/cluster_train/src/paddle-model-sharding.graffle diff --git a/doc/design/cluster_train/src/paddle-model-sharding.png b/doc/v2/design/cluster_train/src/paddle-model-sharding.png similarity index 100% rename from doc/design/cluster_train/src/paddle-model-sharding.png rename to doc/v2/design/cluster_train/src/paddle-model-sharding.png diff --git a/doc/design/cluster_train/src/paddle-ps-0.png b/doc/v2/design/cluster_train/src/paddle-ps-0.png similarity index 100% rename from doc/design/cluster_train/src/paddle-ps-0.png rename to doc/v2/design/cluster_train/src/paddle-ps-0.png diff --git a/doc/design/cluster_train/src/paddle-ps-1.png b/doc/v2/design/cluster_train/src/paddle-ps-1.png similarity index 100% rename from doc/design/cluster_train/src/paddle-ps-1.png rename to doc/v2/design/cluster_train/src/paddle-ps-1.png diff --git a/doc/design/cluster_train/src/paddle-ps.graffle b/doc/v2/design/cluster_train/src/paddle-ps.graffle similarity index 100% rename from doc/design/cluster_train/src/paddle-ps.graffle rename to doc/v2/design/cluster_train/src/paddle-ps.graffle diff --git a/doc/design/cluster_train/src/paddle-task-queues.graffle b/doc/v2/design/cluster_train/src/paddle-task-queues.graffle similarity index 100% rename from doc/design/cluster_train/src/paddle-task-queues.graffle rename to doc/v2/design/cluster_train/src/paddle-task-queues.graffle diff --git a/doc/design/cluster_train/src/paddle-task-queues.png b/doc/v2/design/cluster_train/src/paddle-task-queues.png similarity index 100% rename from doc/design/cluster_train/src/paddle-task-queues.png rename to doc/v2/design/cluster_train/src/paddle-task-queues.png diff --git a/doc/design/cluster_train/src/paddle-task-states.graffle b/doc/v2/design/cluster_train/src/paddle-task-states.graffle similarity index 100% rename from doc/design/cluster_train/src/paddle-task-states.graffle rename to doc/v2/design/cluster_train/src/paddle-task-states.graffle diff --git a/doc/design/cluster_train/src/paddle-task-states.png b/doc/v2/design/cluster_train/src/paddle-task-states.png similarity index 100% rename from doc/design/cluster_train/src/paddle-task-states.png rename to doc/v2/design/cluster_train/src/paddle-task-states.png diff --git a/doc/design/cluster_train/src/pserver_init.graffle b/doc/v2/design/cluster_train/src/pserver_init.graffle similarity index 100% rename from doc/design/cluster_train/src/pserver_init.graffle rename to doc/v2/design/cluster_train/src/pserver_init.graffle diff --git a/doc/design/cluster_train/src/pserver_init.png b/doc/v2/design/cluster_train/src/pserver_init.png similarity index 100% rename from doc/design/cluster_train/src/pserver_init.png rename to doc/v2/design/cluster_train/src/pserver_init.png diff --git a/doc/design/cluster_train/src/submit-job.graffle b/doc/v2/design/cluster_train/src/submit-job.graffle similarity index 100% rename from doc/design/cluster_train/src/submit-job.graffle rename to doc/v2/design/cluster_train/src/submit-job.graffle diff --git a/doc/design/cluster_train/src/submit-job.png b/doc/v2/design/cluster_train/src/submit-job.png similarity index 100% rename from doc/design/cluster_train/src/submit-job.png rename to doc/v2/design/cluster_train/src/submit-job.png diff --git a/doc/design/cluster_train/src/trainer.graffle b/doc/v2/design/cluster_train/src/trainer.graffle similarity index 100% rename from doc/design/cluster_train/src/trainer.graffle rename to doc/v2/design/cluster_train/src/trainer.graffle diff --git a/doc/design/cluster_train/src/trainer.png b/doc/v2/design/cluster_train/src/trainer.png similarity index 100% rename from doc/design/cluster_train/src/trainer.png rename to doc/v2/design/cluster_train/src/trainer.png diff --git a/doc/design/cluster_train/submit-job.md b/doc/v2/design/cluster_train/submit-job.md similarity index 100% rename from doc/design/cluster_train/submit-job.md rename to doc/v2/design/cluster_train/submit-job.md diff --git a/doc/design/mkl/image/engine.png b/doc/v2/design/mkl/image/engine.png similarity index 100% rename from doc/design/mkl/image/engine.png rename to doc/v2/design/mkl/image/engine.png diff --git a/doc/design/mkl/image/gradients.png b/doc/v2/design/mkl/image/gradients.png similarity index 100% rename from doc/design/mkl/image/gradients.png rename to doc/v2/design/mkl/image/gradients.png diff --git a/doc/design/mkl/image/layers.png b/doc/v2/design/mkl/image/layers.png similarity index 100% rename from doc/design/mkl/image/layers.png rename to doc/v2/design/mkl/image/layers.png diff --git a/doc/design/mkl/image/matrix.png b/doc/v2/design/mkl/image/matrix.png similarity index 100% rename from doc/design/mkl/image/matrix.png rename to doc/v2/design/mkl/image/matrix.png diff --git a/doc/design/mkl/image/overview.png b/doc/v2/design/mkl/image/overview.png similarity index 100% rename from doc/design/mkl/image/overview.png rename to doc/v2/design/mkl/image/overview.png diff --git a/doc/design/mkl/mkl_packed.md b/doc/v2/design/mkl/mkl_packed.md similarity index 100% rename from doc/design/mkl/mkl_packed.md rename to doc/v2/design/mkl/mkl_packed.md diff --git a/doc/design/mkl/mkldnn.md b/doc/v2/design/mkl/mkldnn.md similarity index 100% rename from doc/design/mkl/mkldnn.md rename to doc/v2/design/mkl/mkldnn.md diff --git a/doc/v2/dev/new_layer_cn.rst b/doc/v2/dev/new_layer_cn.rst index 0ded1c262adad44f4df000ef2933c7b68050f2fc..3115654b2bd87995fa63bb7828fd1b3039aea8cc 100644 --- a/doc/v2/dev/new_layer_cn.rst +++ b/doc/v2/dev/new_layer_cn.rst @@ -16,7 +16,7 @@ 下图是一个全连接层的示意图。在全连接层中，每个输出节点都连接到所有的输入节点上。 -.. image:: FullyConnected.jpg +.. image:: src/FullyConnected.jpg :align: center :scale: 60 % diff --git a/doc/v2/dev/new_layer_en.rst b/doc/v2/dev/new_layer_en.rst index 110a9fb38f890a766bb4480e91feb22d3b0838a5..b05bb45f11eb253dfb87d6283c29ec6689394d22 100644 --- a/doc/v2/dev/new_layer_en.rst +++ b/doc/v2/dev/new_layer_en.rst @@ -16,7 +16,7 @@ First we need to derive equations of the *forward* and *backward* part of the la The illustration of a fully connected layer is shown in the following figure. In a fully connected layer, all output nodes are connected to all the input nodes. -.. image:: FullyConnected.jpg +.. image:: src/FullyConnected.jpg :align: center :scale: 60 % diff --git a/doc/v2/dev/FullyConnected.jpg b/doc/v2/dev/src/FullyConnected.jpg similarity index 100% rename from doc/v2/dev/FullyConnected.jpg rename to doc/v2/dev/src/FullyConnected.jpg diff --git a/doc/v2/dev/src/doc_en.png b/doc/v2/dev/src/doc_en.png new file mode 100644 index 0000000000000000000000000000000000000000..ed6b9178fba91a3bdf45ae797a9924f84146fbc8 Binary files /dev/null and b/doc/v2/dev/src/doc_en.png differ diff --git a/doc/v2/dev/write_docs_cn.rst b/doc/v2/dev/write_docs_cn.rst index f79769b810b91c6984016d95f40b89186bfb61b0..a055bb04c0c093c9159290067e5ccbd2525cd519 100644 --- a/doc/v2/dev/write_docs_cn.rst +++ b/doc/v2/dev/write_docs_cn.rst @@ -2,20 +2,19 @@ 如何贡献文档 ############# -PaddlePaddle的文档包括英文文档 ``doc`` 和中文文档 ``doc_cn`` 两个部分。文档都是通过 `cmake`_ 驱动 `sphinx`_ 编译生成，生成后的文档分别存储在编译目录的 ``doc`` 和 ``doc_cn`` 两个子目录下。 -也可以利用PaddlePaddle 工具来编译文档，这个情况下所有的文件会存在整理过的的文件目录 .ppo_workspace/content 下 +PaddlePaddle的文档包括中英文两个部分。文档都是通过 ``cmake`` 驱动 ``sphinx`` 编译生成，也可以利用paddlepaddle.org工具来编译和预览文档。 如何构建文档 ============ -PaddlePaddle的文档构建有三种方式。 +PaddlePaddle的文档构建有两种方式，分别为使用paddlepaddle.org工具和不使用paddlepaddle.org工具，两种方式都有各自的优点，前者方便预览，后者方便开发者进行调试。这两种方式中又分别有使用docker和不使用docker的两种构建方法。 使用PaddlePaddle.org工具 --------------- -这个是目前推荐的使用方法。除了可以自动编译文档，也可以直接在网页预览文档。 +------------------------ +这个是目前推荐的使用方法。除了可以自动编译文档，还可以直接在网页中预览文档，需要注意的是，采用后续说明的其它方式虽然也可以预览文档，但是文档的样式与官网文档是不一致的，使用PaddlePaddle.org工具进行编译才能产生与官网文档样式一致的预览效果。 -文件工具是使用Docker，需要在系统里先安装好Docker工具包。Docker安装请参考Docker的官网。安装好Docker之后及可用以下命令启动工具 +PaddlePaddle.org工具可以配合Docker使用，需要在系统里先安装好Docker工具包。Docker安装请参考 `Docker的官网 `_ 。安装好Docker之后即可用以下命令启动工具 .. code-block:: bash @@ -35,7 +34,7 @@ PaddlePaddle的文档构建有三种方式。 之后再用网页连到http://localhost:8000就可以在网页上生成需要的文档 编译后的文件将被存储在工作目录 /.ppo_workspace/content。 -如果不想使用 Docker，你还可以通过运行Django框架直接激活工具的服务器。使用下面的命令来运行它。 +如果不想使用Docker，你还可以通过运行Django框架直接激活工具的服务器。使用下面的命令来运行它。 .. code-block:: bash @@ -62,37 +61,46 @@ PaddlePaddle的文档构建有三种方式。 想了解更多PaddlePaddle.org工具的详细信息，可以 `点击这里 `_ 。 -使用Docker构建 --------------- +不使用PaddlePaddle.org工具 +-------------------------- 使用Docker构建PaddlePaddle的文档，需要在系统里先安装好Docker工具包。Docker安装请参考 `Docker的官网 `_ 。安装好Docker之后可以使用源码目录下的脚本构建文档，即 -.. code-block:: bash +[TBD] - cd TO_YOUR_PADDLE_CLONE_PATH - cd paddle/scripts/tools/build_docs - sh build_docs.sh +如果不想使用Docker，也可以使用以下命令直接构建PaddlePaddle文档，即 -编译完成之后，会在当前目录生成两个子目录\: doc(英文文档目录)和 doc_cn(中文文档目录)。 -打开浏览器访问对应目录下的index.html即可访问本地文档。 +.. code-block:: bash -直接构建 --------- + mkdir paddle + cd paddle + git clone https://github.com/PaddlePaddle/Paddle.git + mkdir -p build + cd build + cmake .. -DCMAKE_BUILD_TYPE=Release -DWITH_GPU=OFF -DWITH_MKL=OFF -DWITH_DOC=ON -如果提示正确，可以执行以下命令编译生成文档，即 + # 如果只需要构建使用文档，则执行以下命令 + make -j $processors gen_proto_py + make -j $processors paddle_docs paddle_docs_cn -.. code-block:: bash + # 如果只需要构建API，则执行以下命令 + make -j $processors gen_proto_py framework_py_proto + make -j $processors copy_paddle_pybind + make -j $processors paddle_api_docs + +其中$processors代表启动和CPU核一样多的进程来并行编译，可以根据本机的CPU核数设置相应的值。 + +编译完成后，进入 ``doc/v2`` 目录，如果选择构建文档则会在该目录下生成 ``cn/html/`` 、 ``en/html`` 两个子目录，选择构建API则会生成 ``api/en/html`` 目录，分别进入这些目录下，执行以下命令： + +.. code-block:: bash - cd TO_YOUR_PADDLE_CLONE_PATH - mkdir -p build - cd build - cmake .. -DCMAKE_BUILD_TYPE=Debug -DWITH_GPU=OFF -DWITH_MKL=OFF -DWITH_DOC=ON - make gen_proto_py - make paddle_docs paddle_docs_cn + python -m SimpleHTTPServer 8088 -编译完成之后，会在当前目录生成两个子目录\: doc(英文文档目录)和 doc_cn(中文文档目录)。 -打开浏览器访问对应目录下的index.html即可访问本地文档。 +在浏览器中输入http://localhost:8088就可以看到编译生成的中/英文的文档页面和英文的API页面,下图为生成的英文文档首页示例。注意，示例中由于使用了sphinx的原始主题，所以页面的风格与官网并不一致，但这并不影响开发者进行调试。 +.. image:: src/doc_en.png + :align: center + :scale: 60 % 如何书写文档 ============ @@ -102,7 +110,7 @@ PaddlePaddle文档使用 `sphinx`_ 自动生成，用户可以参考sphinx教程 如何更新www.paddlepaddle.org ============================ -更新的文档以PR的形式提交到github中，提交方式参见 `贡献文档 `_ 。 +更新的文档以PR的形式提交到github中，提交方式参见 `如何贡献文档 `_ 。 目前PaddlePaddle的develop分支的文档是自动触发更新的，用户可以分别查看最新的 `中文文档 `_ 和 `英文文档 `_ 。 diff --git a/doc/v2/howto/cluster/index_en.rst b/doc/v2/howto/cluster/index_en.rst index 2640a09dcc904619bc97c9bd3f3d81a9dc307663..c965d30d54e71339cf10d4b05f25e740c81adbf9 100644 --- a/doc/v2/howto/cluster/index_en.rst +++ b/doc/v2/howto/cluster/index_en.rst @@ -1,8 +1,7 @@ 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 diff --git a/paddle/CMakeLists.txt b/paddle/CMakeLists.txt index a7b249d43bf3ad9924749d5e66618750f19d8bf7..d2a4b1335464f553a361728e64ed5ca177ca53da 100644 --- a/paddle/CMakeLists.txt +++ b/paddle/CMakeLists.txt @@ -1,27 +1,29 @@ -add_subdirectory(cuda) -add_subdirectory(function) -add_subdirectory(utils) -add_subdirectory(math) -add_subdirectory(gserver) -add_subdirectory(parameter) -add_subdirectory(testing) - -if(MOBILE_INFERENCE) - add_subdirectory(capi) -else() - add_subdirectory(pserver) - add_subdirectory(trainer) - add_subdirectory(scripts) +if(NOT WITH_FLUID) + add_subdirectory(cuda) + add_subdirectory(function) + add_subdirectory(utils) + add_subdirectory(math) + add_subdirectory(gserver) + add_subdirectory(parameter) - if(WITH_C_API) + if(MOBILE_INFERENCE) add_subdirectory(capi) - endif() + else() + add_subdirectory(pserver) + add_subdirectory(trainer) + add_subdirectory(scripts) - if(NOT ANDROID AND NOT IOS) - add_subdirectory(fluid) - endif() + if(WITH_C_API) + add_subdirectory(capi) + endif() - if(WITH_SWIG_PY) - add_subdirectory(api) + if(WITH_SWIG_PY) + add_subdirectory(api) + endif() endif() endif() + +add_subdirectory(testing) +if(NOT MOBILE_INFERENCE AND NOT ANDROID AND NOT IOS) + add_subdirectory(fluid) +endif() diff --git a/paddle/fluid/framework/CMakeLists.txt b/paddle/fluid/framework/CMakeLists.txt index 15e5574ecfd406b87db8370948352b7e736937ea..a4ea74a6d2fbc29dc33a6b57ee453f49ed36c7fa 100644 --- a/paddle/fluid/framework/CMakeLists.txt +++ b/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) diff --git a/paddle/fluid/framework/channel.h b/paddle/fluid/framework/channel.h index 9f8fb12098d622058a86f83c1c42a1feb1cfb2e2..adfaba26ace78f547161ad4029a741f3ca8a6764 100644 --- a/paddle/fluid/framework/channel.h +++ b/paddle/fluid/framework/channel.h @@ -15,23 +15,43 @@ limitations under the License. */ #pragma once #include // for size_t +#include #include #include "paddle/fluid/platform/enforce.h" namespace paddle { namespace framework { +enum class ChannelAction { + SEND = 0, + RECEIVE = 1, + CLOSE = 2, +}; + // Channel is the abstract class of buffered and un-buffered channels. template class Channel { public: + virtual bool CanSend() = 0; + virtual bool CanReceive() = 0; virtual bool Send(T*) = 0; virtual bool Receive(T*) = 0; virtual size_t Cap() = 0; virtual void Lock() = 0; + virtual void Unlock() = 0; + virtual bool IsClosed() = 0; virtual void Close() = 0; virtual ~Channel() {} + + virtual void AddToSendQ(const void* referrer, T* data, + std::shared_ptr cond, + std::function cb) = 0; + virtual void AddToReceiveQ(const void* referrer, T* data, + std::shared_ptr cond, + std::function cb) = 0; + virtual void RemoveFromSendQ(const void* referrer) = 0; + virtual void RemoveFromReceiveQ(const void* referrer) = 0; }; // Forward declaration of channel implementations. @@ -80,6 +100,27 @@ class ChannelHolder { return channel != nullptr ? channel->Receive(data) : false; } + bool IsClosed() { + if (IsInitialized()) { + return holder_->IsClosed(); + } + return false; + } + + bool CanSend() { + if (IsInitialized()) { + return holder_->CanSend(); + } + return false; + } + + bool CanReceive() { + if (IsInitialized()) { + return holder_->CanReceive(); + } + return false; + } + void close() { if (IsInitialized()) holder_->Close(); } @@ -97,6 +138,38 @@ class ChannelHolder { if (IsInitialized()) holder_->Unlock(); } + template + void AddToSendQ(const void* referrer, T* data, + std::shared_ptr cond, + std::function cb) { + if (IsInitialized()) { + Channel* channel = static_cast*>(holder_->Ptr()); + if (channel != nullptr) { + channel->AddToSendQ(referrer, data, cond, cb); + } + } + } + + template + void AddToReceiveQ(const void* referrer, T* data, + std::shared_ptr cond, + std::function cb) { + if (IsInitialized()) { + Channel* channel = static_cast*>(holder_->Ptr()); + if (channel != nullptr) { + channel->AddToReceiveQ(referrer, data, cond, cb); + } + } + } + + void RemoveFromSendQ(const void* referrer) { + if (IsInitialized()) holder_->RemoveFromSendQ(referrer); + } + + void RemoveFromReceiveQ(const void* referrer) { + if (IsInitialized()) holder_->RemoveFromReceiveQ(referrer); + } + inline bool IsInitialized() const { return holder_ != nullptr; } inline const std::type_index Type() { @@ -113,6 +186,11 @@ class ChannelHolder { virtual ~Placeholder() {} virtual const std::type_index Type() const = 0; virtual void* Ptr() const = 0; + 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; @@ -129,6 +207,39 @@ class ChannelHolder { virtual void* Ptr() const { return static_cast(channel_.get()); } + virtual bool IsClosed() { + if (channel_) { + return channel_->IsClosed(); + } + return false; + } + + virtual bool CanSend() { + if (channel_) { + return channel_->CanSend(); + } + return false; + } + + virtual bool CanReceive() { + if (channel_) { + return channel_->CanReceive(); + } + 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(); } diff --git a/paddle/fluid/framework/channel_impl.h b/paddle/fluid/framework/channel_impl.h index a4561031fd8c49613269e7008ce558f25f9765e4..457abbf373d4549229e8fd8bd6b2087cc6b8f5c8 100644 --- a/paddle/fluid/framework/channel_impl.h +++ b/paddle/fluid/framework/channel_impl.h @@ -29,32 +29,50 @@ class ChannelImpl : public paddle::framework::Channel { friend void paddle::framework::CloseChannel(Channel *); public: + virtual bool CanSend(); + virtual bool CanReceive(); virtual bool Send(T *); virtual bool Receive(T *); virtual size_t Cap() { return cap_; } virtual void Lock(); virtual void Unlock(); + virtual bool IsClosed(); virtual void Close(); - ChannelImpl(size_t); virtual ~ChannelImpl(); + virtual void AddToSendQ(const void *referrer, T *data, + std::shared_ptr cond, + std::function cb); + virtual void AddToReceiveQ(const void *referrer, T *data, + std::shared_ptr cond, + std::function cb); + + virtual void RemoveFromSendQ(const void *referrer); + virtual void RemoveFromReceiveQ(const void *referrer); + private: struct QueueMessage { T *data; - std::condition_variable_any cond; + std::shared_ptr cond; bool chan_closed = false; bool completed = false; + const void *referrer; // TODO(thuan): figure out better way to do this + std::function callback; - QueueMessage(T *item) : data(item) {} + QueueMessage(T *item) + : data(item), cond(std::make_shared()) {} + + QueueMessage(T *item, std::shared_ptr cond) + : data(item), cond(cond) {} void Wait(std::unique_lock &lock) { - cond.wait(lock, [this]() { return completed; }); + cond->wait(lock, [this]() { return completed; }); } void Notify() { completed = true; - cond.notify_all(); + cond->notify_all(); } }; @@ -87,6 +105,18 @@ ChannelImpl::ChannelImpl(size_t capacity) PADDLE_ENFORCE_GE(capacity, 0); } +template +bool ChannelImpl::CanSend() { + std::lock_guard lock{mu_}; + return !closed_ && (!recvq.empty() || buf_.size() < cap_); +} + +template +bool ChannelImpl::CanReceive() { + std::lock_guard lock{mu_}; + return !(closed_ && buf_.empty()) && (!sendq.empty() || buf_.size() > 0); +} + template bool ChannelImpl::Send(T *item) { send_ctr++; @@ -105,7 +135,24 @@ bool ChannelImpl::Send(T *item) { std::shared_ptr m = recvq.front(); recvq.pop_front(); // Do the data transfer - *(m->data) = std::move(*item); + // We will do this data transfer if either of the following + // cases are true + // 1. callback == nullptr // This means it was a regular channel send + // 2. callback returns true + bool do_send = true; + if (m->callback != nullptr) do_send = m->callback(ChannelAction::SEND); + if (do_send) + *(m->data) = std::move(*item); + else + // We cannot do the data transfer because + // this QueueMessage was added by Select + // and some other case was executed. + // So call the Send function again. + // We do not care about notifying other + // because they would have been notified + // by the executed select case. + return send_return(Send(item)); + // Wake up the blocked process and unlock m->Notify(); lock.unlock(); @@ -150,7 +197,25 @@ bool ChannelImpl::Receive(T *item) { std::shared_ptr m = sendq.front(); sendq.pop_front(); // Do the data transfer - *item = std::move(*(m->data)); + // We will do this data transfer if either of the following + // cases are true + // 1. callback == nullptr // This means it was a regular channel send + // 2. callback returns true + bool do_receive = true; + if (m->callback != nullptr) + do_receive = m->callback(ChannelAction::RECEIVE); + if (do_receive) + *item = std::move(*(m->data)); + else + // We cannot do the data transfer because + // this QueueMessage was added by Select + // and some other case was executed. + // So call the Receive function again. + // We do not care about notifying other + // because they would have been notified + // by the executed select case. + return recv_return(Receive(item)); + // Wake up the blocked process and unlock m->Notify(); lock.unlock(); @@ -186,6 +251,12 @@ void ChannelImpl::Unlock() { mu_.unlock(); } +template +bool ChannelImpl::IsClosed() { + std::lock_guard lock{mu_}; + return closed_; +} + template void ChannelImpl::Close() { std::unique_lock lock{mu_}; @@ -203,6 +274,12 @@ void ChannelImpl::Close() { std::shared_ptr m = recvq.front(); recvq.pop_front(); m->chan_closed = true; + + // Execute callback function (if any) + if (m->callback != nullptr) { + m->callback(ChannelAction::CLOSE); + } + m->Notify(); } @@ -211,10 +288,70 @@ void ChannelImpl::Close() { std::shared_ptr m = sendq.front(); sendq.pop_front(); m->chan_closed = true; + + // Execute callback function (if any) + if (m->callback != nullptr) { + m->callback(ChannelAction::CLOSE); + } + m->Notify(); } } +template +void ChannelImpl::AddToSendQ( + const void *referrer, T *data, + std::shared_ptr cond, + std::function cb) { + std::lock_guard lock{mu_}; + auto m = std::make_shared(data, cond); + m->referrer = referrer; + m->callback = cb; + sendq.push_back(m); +} + +template +void ChannelImpl::AddToReceiveQ( + const void *referrer, T *data, + std::shared_ptr cond, + std::function cb) { + std::lock_guard lock{mu_}; + auto m = std::make_shared(data, cond); + m->referrer = referrer; + m->callback = cb; + recvq.push_back(m); +} + +template +void ChannelImpl::RemoveFromSendQ(const void *referrer) { + std::lock_guard lock{mu_}; + + for (auto it = sendq.begin(); it != sendq.end();) { + std::shared_ptr sendMsg = (std::shared_ptr)*it; + + if (sendMsg->referrer == referrer) { + it = sendq.erase(it); + } else { + ++it; + } + } +} + +template +void ChannelImpl::RemoveFromReceiveQ(const void *referrer) { + std::lock_guard lock{mu_}; + + for (auto it = recvq.begin(); it != recvq.end();) { + std::shared_ptr recvMsg = (std::shared_ptr)*it; + + if (recvMsg->referrer == referrer) { + it = recvq.erase(it); + } else { + ++it; + } + } +} + template ChannelImpl::~ChannelImpl() { Close(); diff --git a/paddle/fluid/framework/concurrency_test.cc b/paddle/fluid/framework/concurrency_test.cc index 5770b0a5a18659e615e80a7c48113d8b543b69ec..25152054eb8452a9667bd65b4441665476c1d46d 100644 --- a/paddle/fluid/framework/concurrency_test.cc +++ b/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 -void CreateIntVariable(Scope &scope, p::CPUPlace &place, std::string name, - T value) { - // Create LoDTensor of dim [1,1] +LoDTensor *CreateVariable(Scope &scope, p::CPUPlace &place, std::string name, + T value) { + // Create LoDTensor of dim [1] auto var = scope.Var(name); auto tensor = var->GetMutable(); - tensor->Resize({1, 1}); + tensor->Resize({1}); T *expect = tensor->mutable_data(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 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 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 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 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{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(); 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(); + auto *finalData = tensor.data(); + EXPECT_EQ(finalData[0], 34); +} + } // namespace framework } // namespace paddle diff --git a/paddle/fluid/framework/executor.cc b/paddle/fluid/framework/executor.cc index 5cae38b2a857b2037f0e5ae4da50d1591da0c11a..fb7c13d36db7ebe0bdcb81a15a986a7c0ca4170f 100644 --- a/paddle/fluid/framework/executor.cc +++ b/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,6 +34,11 @@ 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) @@ -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& 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); diff --git a/paddle/fluid/framework/operator.cc b/paddle/fluid/framework/operator.cc index 1dddfcdd2e10ad91dc80aefaaf54bab04f15240b..b39a1164dbd9877d9f45cc6415d74f930921a42f 100644 --- a/paddle/fluid/framework/operator.cc +++ b/paddle/fluid/framework/operator.cc @@ -74,9 +74,6 @@ void OperatorBase::Run(const Scope& scope, const platform::Place& place) { platform::SetDeviceId(dev_id); #endif } - // profile - auto* dev_ctx = platform::DeviceContextPool::Instance().Get(place); - platform::RecordEvent record_event(Type(), dev_ctx); RunImpl(scope, place); } @@ -485,6 +482,10 @@ void OperatorWithKernel::RunImpl(const Scope& scope, this->InferShape(&infer_shape_ctx); platform::DeviceContextPool& pool = platform::DeviceContextPool::Instance(); auto* dev_ctx = pool.Get(place); + + // For profiling, don't move out of this function because that will result + // in the failure of multi-GPU profiling. + platform::RecordEvent record_event(Type(), dev_ctx); // check if op[type] has kernel registered. auto& all_op_kernels = AllOpKernels(); auto kernels_iter = all_op_kernels.find(type_); diff --git a/paddle/fluid/operators/CMakeLists.txt b/paddle/fluid/operators/CMakeLists.txt index 625e0f7561899d30b40f9daa56f743a37bdaa27f..d30124d4a3b89b802a4abaae07a33b76526f163d 100644 --- a/paddle/fluid/operators/CMakeLists.txt +++ b/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}) diff --git a/paddle/fluid/operators/assign_op.cc b/paddle/fluid/operators/assign_op.cc index 39ae3c0040d04a6d901f1d6c992d547a6778c28e..d372213e1b6008b0c4227103dd40730f86a84301 100644 --- a/paddle/fluid/operators/assign_op.cc +++ b/paddle/fluid/operators/assign_op.cc @@ -56,6 +56,7 @@ class AssignFunctor { private: void copy_tensor(const framework::LoDTensor &lod_tensor, framework::LoDTensor *out) const { + if (lod_tensor.numel() == 0) return; auto &out_tensor = *out; TensorCopy(lod_tensor, lod_tensor.place(), dev_ctx_, &out_tensor); out_tensor.set_lod(lod_tensor.lod()); diff --git a/paddle/fluid/operators/channel_recv_op.cc b/paddle/fluid/operators/channel_recv_op.cc index c12b88e7a91c4ea7044223464a2f902db494d1a8..844b3ae3b7bf87c9b253128165b3c938801d5d60 100644 --- a/paddle/fluid/operators/channel_recv_op.cc +++ b/paddle/fluid/operators/channel_recv_op.cc @@ -18,6 +18,7 @@ limitations under the License. */ #include #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()); - else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE) - return ch->Receive(var->GetMutable()); - else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY) - return ch->Receive(var->GetMutable()); - else if (type == framework::proto::VarType_Type_SELECTED_ROWS) - return ch->Receive(var->GetMutable()); - else if (type == framework::proto::VarType_Type_READER) - return ch->Receive(var->GetMutable()); - else if (type == framework::proto::VarType_Type_CHANNEL) - return ch->Receive(var->GetMutable()); - 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(); 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); diff --git a/paddle/fluid/operators/channel_send_op.cc b/paddle/fluid/operators/channel_send_op.cc index 6d7715ad229e821f02437246e3326063cb1ee757..47cf7d7efc9996e8a8db11b79c0310f77c2435a4 100644 --- a/paddle/fluid/operators/channel_send_op.cc +++ b/paddle/fluid/operators/channel_send_op.cc @@ -18,6 +18,7 @@ limitations under the License. */ #include #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()); - else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE) - return ch->Send(var->GetMutable()); - else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY) - return ch->Send(var->GetMutable()); - else if (type == framework::proto::VarType_Type_SELECTED_ROWS) - return ch->Send(var->GetMutable()); - else if (type == framework::proto::VarType_Type_READER) - return ch->Send(var->GetMutable()); - else if (type == framework::proto::VarType_Type_CHANNEL) - return ch->Send(var->GetMutable()); - 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); diff --git a/paddle/fluid/operators/concurrency/CMakeLists.txt b/paddle/fluid/operators/concurrency/CMakeLists.txt new file mode 100644 index 0000000000000000000000000000000000000000..e4617440d152b4c15d09e81cd19c76739b95b979 --- /dev/null +++ b/paddle/fluid/operators/concurrency/CMakeLists.txt @@ -0,0 +1 @@ +cc_library(concurrency SRCS channel_util.cc DEPS device_context framework_proto boost eigen3) diff --git a/paddle/fluid/operators/concurrency/channel_util.cc b/paddle/fluid/operators/concurrency/channel_util.cc new file mode 100644 index 0000000000000000000000000000000000000000..a483af7affd824da7d18676d934dc959167ef71f --- /dev/null +++ b/paddle/fluid/operators/concurrency/channel_util.cc @@ -0,0 +1,111 @@ +/* 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()); + else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE) + return ch->Send(var->GetMutable()); + else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY) + return ch->Send(var->GetMutable()); + else if (type == framework::proto::VarType_Type_SELECTED_ROWS) + return ch->Send(var->GetMutable()); + else if (type == framework::proto::VarType_Type_READER) + return ch->Send(var->GetMutable()); + else if (type == framework::proto::VarType_Type_CHANNEL) + return ch->Send(var->GetMutable()); + 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()); + else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE) + return ch->Receive(var->GetMutable()); + else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY) + return ch->Receive(var->GetMutable()); + else if (type == framework::proto::VarType_Type_SELECTED_ROWS) + return ch->Receive(var->GetMutable()); + else if (type == framework::proto::VarType_Type_READER) + return ch->Receive(var->GetMutable()); + else if (type == framework::proto::VarType_Type_CHANNEL) + return ch->Receive(var->GetMutable()); + else + PADDLE_THROW("ChannelReceive:Unsupported type"); +} + +void poc::ChannelAddToSendQ(framework::ChannelHolder *ch, const void *referrer, + framework::Variable *var, + std::shared_ptr cond, + std::function cb) { + auto type = framework::ToVarType(var->Type()); + if (type == framework::proto::VarType_Type_LOD_TENSOR) { + ch->AddToSendQ(referrer, var->GetMutable(), cond, cb); + } else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE) { + ch->AddToSendQ(referrer, var->GetMutable(), cond, + cb); + } else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY) { + ch->AddToSendQ(referrer, var->GetMutable(), cond, + cb); + } else if (type == framework::proto::VarType_Type_SELECTED_ROWS) { + ch->AddToSendQ(referrer, var->GetMutable(), cond, + cb); + } else if (type == framework::proto::VarType_Type_READER) { + ch->AddToSendQ(referrer, var->GetMutable(), cond, + cb); + } else if (type == framework::proto::VarType_Type_CHANNEL) { + ch->AddToSendQ(referrer, var->GetMutable(), cond, + cb); + } else { + PADDLE_THROW("ChannelAddToSendQ:Unsupported type"); + } +} + +void poc::ChannelAddToReceiveQ( + framework::ChannelHolder *ch, const void *referrer, + framework::Variable *var, std::shared_ptr cond, + std::function cb) { + auto type = framework::ToVarType(var->Type()); + if (type == framework::proto::VarType_Type_LOD_TENSOR) { + ch->AddToReceiveQ(referrer, var->GetMutable(), cond, + cb); + } else if (type == framework::proto::VarType_Type_LOD_RANK_TABLE) { + ch->AddToReceiveQ(referrer, var->GetMutable(), + cond, cb); + } else if (type == framework::proto::VarType_Type_LOD_TENSOR_ARRAY) { + ch->AddToReceiveQ(referrer, var->GetMutable(), + cond, cb); + } else if (type == framework::proto::VarType_Type_SELECTED_ROWS) { + ch->AddToReceiveQ(referrer, var->GetMutable(), + cond, cb); + } else if (type == framework::proto::VarType_Type_READER) { + ch->AddToReceiveQ(referrer, var->GetMutable(), + cond, cb); + } else if (type == framework::proto::VarType_Type_CHANNEL) { + ch->AddToReceiveQ(referrer, var->GetMutable(), + cond, cb); + } else { + PADDLE_THROW("ChannelAddToReceiveQ:Unsupported type"); + } +} diff --git a/paddle/fluid/operators/concurrency/channel_util.h b/paddle/fluid/operators/concurrency/channel_util.h new file mode 100644 index 0000000000000000000000000000000000000000..c3674bd9815df451751707bfa84d18dbb5fa0f6b --- /dev/null +++ b/paddle/fluid/operators/concurrency/channel_util.h @@ -0,0 +1,38 @@ +/* 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 "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 cond, + std::function cb); +void ChannelAddToReceiveQ(framework::ChannelHolder *ch, const void *referrer, + framework::Variable *var, + std::shared_ptr cond, + std::function cb); + +} // namespace concurrency +} // namespace operators +} // namespace paddle diff --git a/paddle/fluid/operators/detection_output_op.cc b/paddle/fluid/operators/detection_output_op.cc deleted file mode 100644 index f7520475917ff23535f11ccfde0ee915112bba30..0000000000000000000000000000000000000000 --- a/paddle/fluid/operators/detection_output_op.cc +++ /dev/null @@ -1,89 +0,0 @@ -/* 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("background_label_id", "(int), The background class index."); - AddAttr("num_classes", "(int), The number of the classification."); - AddAttr("nms_threshold", - "(float), The Non-maximum suppression threshold."); - AddAttr("confidence_threshold", - "(float), The classification confidence threshold."); - AddAttr("top_k", "(int), The bbox number kept of the layer’s output."); - AddAttr("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 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, - ops::DetectionOutputKernel); diff --git a/paddle/fluid/operators/detection_output_op.cu.cc b/paddle/fluid/operators/detection_output_op.cu.cc deleted file mode 100644 index 0f48765c9c67c1d3fa32b19d5e87b2acaa3c486a..0000000000000000000000000000000000000000 --- a/paddle/fluid/operators/detection_output_op.cu.cc +++ /dev/null @@ -1,21 +0,0 @@ -/* 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 ops = paddle::operators; -REGISTER_OP_CUDA_KERNEL( - detection_output, - ops::DetectionOutputKernel, - ops::DetectionOutputKernel); diff --git a/paddle/fluid/operators/detection_output_op.h b/paddle/fluid/operators/detection_output_op.h deleted file mode 100644 index af9081c93436776b6ca6ee7139e340054111e440..0000000000000000000000000000000000000000 --- a/paddle/fluid/operators/detection_output_op.h +++ /dev/null @@ -1,167 +0,0 @@ -/* 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 -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 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(shape, context.GetPlace()); - std::vector shape_axis({0, 1, 3, 4, 2}); - math::Transpose trans5; - trans5(context.template device_context(), 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(context.device_context(), in_p_tensor_transpose.data(), - src_stride, in_p_tensor_transpose.dims(), dst_stride, - dst->data() + offset); - offset += in_p_tensor_transpose.dims()[4] * src_stride[4]; - } -} -template -class DetectionOutputKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& context) const override { - const framework::Tensor* in_loc = context.Input("Loc"); - const framework::Tensor* in_conf = context.Input("Conf"); - const framework::Tensor* in_priorbox = - context.Input("PriorBox"); - auto* out = context.Output("Out"); - int num_classes = context.template Attr("num_classes"); - int top_k = context.template Attr("top_k"); - int nms_top_k = context.template Attr("nms_top_k"); - int background_label_id = context.template Attr("background_label_id"); - float nms_threshold = context.template Attr("nms_threshold"); - float confidence_threshold = - context.template Attr("confidence_threshold"); - size_t batch_size = in_conf->dims()[1]; - int conf_sum_size = in_conf->numel(); - // for softmax - std::vector 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 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 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(loc_shape, context.GetPlace()); - conf_tensor.mutable_data(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(); - transpose_fun(context, *in_loc, &loc_tensor); - transpose_fun(context, *in_conf, &conf_tensor); - conf_tensor.Resize(conf_shape_softmax); - math::SoftmaxFunctor()( - context.template device_context(), &conf_tensor, - &conf_tensor); - T* loc_data = loc_tensor.data(); - T* conf_data = conf_tensor.data(); - if (platform::is_gpu_place(context.GetPlace())) { - loc_cpu.mutable_data(loc_tensor.dims(), platform::CPUPlace()); - framework::TensorCopy(loc_tensor, platform::CPUPlace(), - context.device_context(), &loc_cpu); - loc_data = loc_cpu.data(); - conf_cpu.mutable_data(conf_tensor.dims(), platform::CPUPlace()); - framework::TensorCopy(conf_tensor, platform::CPUPlace(), - context.device_context(), &conf_cpu); - conf_data = conf_cpu.data(); - priorbox_cpu.mutable_data(in_priorbox->dims(), platform::CPUPlace()); - framework::TensorCopy(*in_priorbox, platform::CPUPlace(), - context.device_context(), &priorbox_cpu); - priorbox_data = priorbox_cpu.data(); - } - // get decode bboxes - size_t num_priors = in_priorbox->numel() / 8; - std::vector>> all_decoded_bboxes; - for (size_t n = 0; n < batch_size; ++n) { - std::vector> 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> prior_bbox_vec; - math::GetBBoxFromPriorData(priorbox_data + prior_offset, 1, - prior_bbox_vec); - std::vector> prior_bbox_var; - math::GetBBoxVarFromPriorData(priorbox_data + prior_offset, 1, - prior_bbox_var); - std::vector loc_pred_data; - for (size_t j = 0; j < 4; ++j) - loc_pred_data.push_back(*(loc_data + loc_pred_offset + j)); - math::BBox bbox = math::DecodeBBoxWithVar( - 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>> all_indices; - int num_kept = math::GetDetectionIndices( - 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 out_shape_vec({0, 0}); - framework::DDim out_shape(framework::make_ddim(out_shape_vec)); - out->Resize(out_shape); - return; - } - std::vector out_shape_vec({num_kept, 7}); - framework::DDim out_shape(framework::make_ddim(out_shape_vec)); - out->mutable_data(out_shape, context.GetPlace()); - framework::Tensor out_cpu; - T* out_data = out->data(); - if (platform::is_gpu_place(context.GetPlace())) { - out_cpu.mutable_data(out->dims(), platform::CPUPlace()); - out_data = out_cpu.data(); - } - math::GetDetectionOutput(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 diff --git a/paddle/fluid/operators/lrn_op.cc b/paddle/fluid/operators/lrn_op.cc index b0c213d637c244e4cbacbe75218537973efed047..692e85dcffa583abcb22a1629953badc67489efa 100644 --- a/paddle/fluid/operators/lrn_op.cc +++ b/paddle/fluid/operators/lrn_op.cc @@ -36,7 +36,7 @@ struct LRNFunctor { auto e_x = framework::EigenTensor::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({{m, i, 0, 0}}), @@ -92,7 +92,7 @@ struct LRNGradFunctor { Eigen::array({{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; diff --git a/paddle/fluid/operators/math/CMakeLists.txt b/paddle/fluid/operators/math/CMakeLists.txt index a181d802262d15b188060dae4330cec0e24714ab..fba1612d10f0494f4ab06fabdd0e799a74dafd53 100644 --- a/paddle/fluid/operators/math/CMakeLists.txt +++ b/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) diff --git a/paddle/fluid/operators/math/detection_util.h b/paddle/fluid/operators/math/detection_util.h deleted file mode 100644 index c31764cfaf5bbdfea2f3ed06f31f97965a8858ed..0000000000000000000000000000000000000000 --- a/paddle/fluid/operators/math/detection_util.h +++ /dev/null @@ -1,300 +0,0 @@ -/* 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 -#include "paddle/fluid/framework/selected_rows.h" -#include "paddle/fluid/platform/device_context.h" - -namespace paddle { -namespace operators { -namespace math { -template -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 -template -void GetBBoxFromPriorData(const T* prior_data, const size_t num_bboxes, - std::vector>& bbox_vec); -template -void GetBBoxVarFromPriorData(const T* prior_data, const size_t num, - std::vector>& var_vec); -template -BBox DecodeBBoxWithVar(BBox& prior_bbox, - const std::vector& prior_bbox_var, - const std::vector& loc_pred_data); -template -bool SortScorePairDescend(const std::pair& pair1, - const std::pair& pair2); -template -bool SortScorePairDescend(const std::pair>& pair1, - const std::pair>& pair2); -template -T jaccard_overlap(const BBox& bbox1, const BBox& bbox2); - -template -void ApplyNmsFast(const std::vector>& 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* indices); -template -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>>& all_decoded_bboxes, - std::vector>>* all_detection_indices); -template -BBox ClipBBox(const BBox& bbox); -template -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>>& all_indices, - const std::vector>>& all_decoded_bboxes, T* out_data); -template -void GetBBoxFromPriorData(const T* prior_data, const size_t num_bboxes, - std::vector>& 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 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 -void GetBBoxVarFromPriorData(const T* prior_data, const size_t num, - std::vector>& 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 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 -BBox DecodeBBoxWithVar(BBox& prior_bbox, - const std::vector& prior_bbox_var, - const std::vector& 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 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 -bool SortScorePairDescend(const std::pair& pair1, - const std::pair& pair2) { - return pair1.first > pair2.first; -} -template -T jaccard_overlap(const BBox& bbox1, const BBox& 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 -void ApplyNmsFast(const std::vector>& 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* indices) { - std::vector> 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); - 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(bboxes[idx], bboxes[saved_idx]); - keep = overlap <= nms_threshold; - } else { - break; - } - } - if (keep) indices->push_back(idx); - scores.erase(scores.begin()); - } -} -template -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>>& all_decoded_bboxes, - std::vector>>* all_detection_indices) { - int total_keep_num = 0; - for (size_t n = 0; n < batch_size; ++n) { - const std::vector>& decoded_bboxes = all_decoded_bboxes[n]; - size_t num_detected = 0; - std::map> 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(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> score_index_pairs; - std::vector>> score_index_pairs; - for (size_t c = 0; c < num_classes; ++c) { - const std::vector& 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>); - score_index_pairs.resize(top_k); - std::map> 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 -BBox ClipBBox(const BBox& bbox) { - T one = static_cast(1.0); - T zero = static_cast(0.0); - BBox 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 -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>>& all_indices, - const std::vector>>& all_decoded_bboxes, T* out_data) { - size_t count = 0; - for (size_t n = 0; n < batch_size; ++n) { - for (std::map>::const_iterator it = - all_indices[n].begin(); - it != all_indices[n].end(); ++it) { - size_t label = it->first; - const std::vector& indices = it->second; - const std::vector>& 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 clipped_bbox = ClipBBox(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 diff --git a/paddle/fluid/operators/math/softmax.cu b/paddle/fluid/operators/math/softmax.cu index 38e93fdf15d99eb447948378a599891074c10fc5..34ea6a91ce7743462d378cf471a5ec3a12ca51d1 100644 --- a/paddle/fluid/operators/math/softmax.cu +++ b/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 +using CudnnDataType = platform::CudnnDataType; + +template +void SoftmaxCUDNNFunctor::operator()( + const platform::CUDADeviceContext& context, const framework::Tensor* X, + framework::Tensor* Y) { + // ------------------- cudnn descriptors --------------------- + ScopedTensorDescriptor xDesc; + ScopedTensorDescriptor yDesc; + std::vector 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(layout, cudnn_tensor_dims); + cudnnTensorDescriptor_t cudnn_y_desc = + xDesc.descriptor(layout, cudnn_tensor_dims); + PADDLE_ENFORCE(platform::dynload::cudnnSoftmaxForward( + context.cudnn_handle(), CUDNN_SOFTMAX_ACCURATE, + CUDNN_SOFTMAX_MODE_INSTANCE, CudnnDataType::kOne(), cudnn_x_desc, + X->data(), CudnnDataType::kZero(), cudnn_y_desc, + Y->mutable_data(context.GetPlace()))); +} + +template +void SoftmaxGradCUDNNFunctor::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 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(layout, cudnn_tensor_dims); + cudnnTensorDescriptor_t cudnn_xgrad_desc = + dxDesc.descriptor(layout, cudnn_tensor_dims); + cudnnTensorDescriptor_t cudnn_ygrad_desc = + dyDesc.descriptor(layout, cudnn_tensor_dims); + PADDLE_ENFORCE(platform::dynload::cudnnSoftmaxBackward( + context.cudnn_handle(), CUDNN_SOFTMAX_ACCURATE, + CUDNN_SOFTMAX_MODE_INSTANCE, CudnnDataType::kOne(), cudnn_y_desc, + Y->data(), cudnn_ygrad_desc, YGrad->data(), + CudnnDataType::kZero(), cudnn_xgrad_desc, + XGrad->mutable_data(context.GetPlace()))); +} + +template class SoftmaxCUDNNFunctor; +template class SoftmaxCUDNNFunctor; +template class SoftmaxGradCUDNNFunctor; +template class SoftmaxGradCUDNNFunctor; + template class SoftmaxFunctor; template class SoftmaxFunctor; template class SoftmaxGradFunctor; diff --git a/paddle/fluid/operators/math/softmax.h b/paddle/fluid/operators/math/softmax.h index 14b2690c2a4e764058270953214a07aee8053444..da1f0b672d3a5fb5da8f4d72892be21964bdbc0d 100644 --- a/paddle/fluid/operators/math/softmax.h +++ b/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 +class SoftmaxCUDNNFunctor { + public: + void operator()(const platform::CUDADeviceContext& context, + const framework::Tensor* X, framework::Tensor* Y); +}; + +template +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 diff --git a/paddle/fluid/operators/mul_op.cc b/paddle/fluid/operators/mul_op.cc index e7bed2c39735b66c19e738c91f4977e46571143b..90af1e2d602ac039b4d98a69a889ff8b1b85ffc6 100644 --- a/paddle/fluid/operators/mul_op.cc +++ b/paddle/fluid/operators/mul_op.cc @@ -17,11 +17,14 @@ limitations under the License. */ namespace paddle { namespace operators { +using framework::OpKernelType; using framework::Tensor; -class MulOpShapeInference : public framework::InferShapeBase { +class MulOp : public framework::OperatorWithKernel { public: - void operator()(framework::InferShapeContext* ctx) const override { + using framework::OperatorWithKernel::OperatorWithKernel; + + void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of MulOp should not be null."); PADDLE_ENFORCE(ctx->HasInput("Y"), "Input(Y) of MulOp should not be null."); PADDLE_ENFORCE(ctx->HasOutput("Out"), @@ -122,7 +125,7 @@ or not. But the output only shares the LoD information with input $X$. } }; -class MulOpGrad : public framework::OperatorWithKernel { +class MulGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; @@ -156,10 +159,7 @@ class MulOpGrad : public framework::OperatorWithKernel { } // namespace paddle namespace ops = paddle::operators; -REGISTER_OPERATOR(mul, paddle::framework::OperatorWithKernel, ops::MulOpMaker, - ops::MulOpShapeInference, - paddle::framework::DefaultGradOpDescMaker); -REGISTER_OPERATOR(mul_grad, ops::MulOpGrad); +REGISTER_OP(mul, ops::MulOp, ops::MulOpMaker, mul_grad, ops::MulGradOp); REGISTER_OP_CPU_KERNEL( mul, ops::MulKernel); REGISTER_OP_CPU_KERNEL( diff --git a/paddle/fluid/operators/mul_op.cu.cc b/paddle/fluid/operators/mul_op.cu.cc index 0667530e943856576ae8c9fe4856cb6aa1448e4e..757f9c3ee2665c7ac654659416fe8dd727dca16d 100644 --- a/paddle/fluid/operators/mul_op.cu.cc +++ b/paddle/fluid/operators/mul_op.cu.cc @@ -13,9 +13,11 @@ See the License for the specific language governing permissions and limitations under the License. */ #include "paddle/fluid/operators/mul_op.h" +#include "paddle/fluid/platform/float16.h" namespace ops = paddle::operators; -REGISTER_OP_CUDA_KERNEL( - mul, ops::MulKernel); -REGISTER_OP_CUDA_KERNEL( - mul_grad, ops::MulGradKernel); +namespace plat = paddle::platform; +REGISTER_OP_CUDA_KERNEL(mul, ops::MulKernel, + ops::MulKernel); +REGISTER_OP_CUDA_KERNEL(mul_grad, + ops::MulGradKernel); diff --git a/paddle/fluid/operators/mul_op.h b/paddle/fluid/operators/mul_op.h index 38311cf87265ad0f1f815734cbf69bd682d62e62..b1260d36ebe11f65529ac274c959479dcb38ee5f 100644 --- a/paddle/fluid/operators/mul_op.h +++ b/paddle/fluid/operators/mul_op.h @@ -48,7 +48,7 @@ class MulKernel : public framework::OpKernel { } math::matmul( context.template device_context(), x_matrix, false, - y_matrix, false, 1, z, 0); + y_matrix, false, static_cast(1), z, static_cast(0)); if (z_dim.size() != 2) { z->Resize(z_dim); } diff --git a/paddle/fluid/operators/nccl_op.cu.cc b/paddle/fluid/operators/nccl_op.cu.cc index 4d83a70e7334a84bb98bd52f0172f6b7ecedb58d..ad623e1fe0f8941615b671a0c20bd3637ae6d407 100644 --- a/paddle/fluid/operators/nccl_op.cu.cc +++ b/paddle/fluid/operators/nccl_op.cu.cc @@ -106,6 +106,8 @@ class NCCLReduceKernel : public framework::OpKernel { T* recvbuffer = nullptr; if (root == gpu_id) { recvbuffer = out->mutable_data(ctx.GetPlace()); + } else { + out->Resize(framework::make_ddim({0})); } VLOG(3) << "gpu : " << gpu_id << " invoke reduce. send " << x->numel() << " recv " << out->numel(); diff --git a/paddle/fluid/operators/parallel_do_op.cc b/paddle/fluid/operators/parallel_do_op.cc index bf4d0476df32d7454d4064cb6ee454e6ad5d6fc5..4001b9a130348b4e3ea99f3017eae6d85e41fc6e 100644 --- a/paddle/fluid/operators/parallel_do_op.cc +++ b/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,11 +159,14 @@ 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] { - framework::Executor executor(place); - executor.Run(*program, cur_scope, block->ID(), - false /*create_local_scope*/); - })); + 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(place_idx) + 1); + framework::Executor executor(place); + executor.Run(*program, cur_scope, block->ID(), + false /*create_local_scope*/); + })); } for (auto &worker : workers) { worker.wait(); @@ -234,11 +238,14 @@ class ParallelDoGradOp : public framework::OperatorBase { auto *cur_scope = sub_scopes[i]; // execute - workers.emplace_back(framework::Async([program, cur_scope, place, block] { - framework::Executor executor(place); - executor.Run(*program, cur_scope, block->ID(), - false /*create_local_scope*/); - })); + 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(i) + 1); + framework::Executor executor(place); + executor.Run(*program, cur_scope, block->ID(), + false /*create_local_scope*/); + })); } for (auto &worker : workers) { worker.wait(); diff --git a/paddle/fluid/operators/scatter_op.cc b/paddle/fluid/operators/scatter_op.cc index 3fb8b56d2676f90ff7e1cefa46c459ee37f63ca8..d6fd6214711f4ee66b1daffa4db2e84aa7201e79 100644 --- a/paddle/fluid/operators/scatter_op.cc +++ b/paddle/fluid/operators/scatter_op.cc @@ -23,24 +23,24 @@ class ScatterOp : public framework::OperatorWithKernel { using framework::OperatorWithKernel::OperatorWithKernel; void InferShape(framework::InferShapeContext* ctx) const override { - PADDLE_ENFORCE(ctx->HasInput("Ref"), - "Input(Ref) of ScatterOp should not be null."); - PADDLE_ENFORCE(ctx->HasInput("Index"), - "Input(Index) of ScatterOp should not be null."); + PADDLE_ENFORCE(ctx->HasInput("X"), + "Input(X) of ScatterOp should not be null."); + PADDLE_ENFORCE(ctx->HasInput("Ids"), + "Input(Ids) of ScatterOp should not be null."); PADDLE_ENFORCE(ctx->HasInput("Updates"), "Input(Updates) of ScatterOp should not be null."); PADDLE_ENFORCE(ctx->HasOutput("Out"), "Output(Out) of ScatterOp should not be null."); auto updates_dims = ctx->GetInputDim("Updates"); - auto ref_dims = ctx->GetInputDim("Ref"); - PADDLE_ENFORCE_EQ(ctx->GetInputDim("Index").size(), 1, - "Update Index should be 1-D."); + auto ref_dims = ctx->GetInputDim("X"); + PADDLE_ENFORCE_EQ(ctx->GetInputDim("Ids").size(), 1, + "Update Ids should be 1-D."); PADDLE_ENFORCE_EQ(ref_dims.size(), updates_dims.size(), - "Reference and Updates should have the same shape size"); + "Xerence and Updates should have the same shape size"); PADDLE_ENFORCE_EQ(ctx->GetInputDim("Updates")[0], - ctx->GetInputDim("Index")[0], - "Updates and Index should have same batch-size."); + ctx->GetInputDim("Ids")[0], + "Updates and Ids should have same batch-size."); framework::DDim data_dim(updates_dims); for (int i = 1; i < data_dim.size(); ++i) { PADDLE_ENFORCE_EQ(data_dim[i], updates_dims[i]); @@ -52,7 +52,7 @@ class ScatterOp : public framework::OperatorWithKernel { framework::OpKernelType GetExpectedKernelType( const framework::ExecutionContext& ctx) const override { return framework::OpKernelType( - framework::ToDataType(ctx.Input("Ref")->type()), + framework::ToDataType(ctx.Input("X")->type()), ctx.device_context()); } }; @@ -64,14 +64,14 @@ class ScatterGradOp : public framework::OperatorWithKernel { void InferShape(framework::InferShapeContext* ctx) const override { ctx->SetOutputDim(framework::GradVarName("Updates"), ctx->GetInputDim("Updates")); - ctx->SetOutputDim(framework::GradVarName("Ref"), ctx->GetInputDim("Ref")); + ctx->SetOutputDim(framework::GradVarName("X"), ctx->GetInputDim("X")); } protected: framework::OpKernelType GetExpectedKernelType( const framework::ExecutionContext& ctx) const override { return framework::OpKernelType( - framework::ToDataType(ctx.Input("Ref")->type()), + framework::ToDataType(ctx.Input("X")->type()), ctx.device_context()); } }; @@ -80,9 +80,8 @@ class ScatterOpMaker : public framework::OpProtoAndCheckerMaker { public: ScatterOpMaker(OpProto* proto, OpAttrChecker* op_checker) : OpProtoAndCheckerMaker(proto, op_checker) { - AddInput("Ref", "The source input of scatter op"); - AddInput("Index", - "The index input of scatter op where Ref will be updated"); + AddInput("X", "The source input of scatter op"); + AddInput("Ids", "The index input of scatter op where X will be updated"); AddInput("Updates", "The updated value of updates op"); AddOutput("Out", "The output of add op"); AddComment(R"DOC( @@ -91,8 +90,8 @@ Scatter Operator. This operator obtains output by updating the input on selected indices on the first axis: $$ -Out = Ref \\ -Out[Index] = Ref[Index] + Updates +Out = X \\ +Out[Ids] = X[Ids] + Updates $$ )DOC"); diff --git a/paddle/fluid/operators/scatter_op.cu b/paddle/fluid/operators/scatter_op.cu index bdabb29fa680f8f87873b4381acf0dbd2b6195d0..ef7d700659d8d713715a10910baf739954ba0786 100644 --- a/paddle/fluid/operators/scatter_op.cu +++ b/paddle/fluid/operators/scatter_op.cu @@ -25,14 +25,14 @@ class ScatterOpCUDAKernel : public framework::OpKernel { void Compute(const framework::ExecutionContext &ctx) const override { PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()), "This kernel only runs on GPU device."); - auto *Ref = ctx.Input("Ref"); - auto *Index = ctx.Input("Index"); + auto *X = ctx.Input("X"); + auto *Ids = ctx.Input("Ids"); auto *Updates = ctx.Input("Updates"); auto *Out = ctx.Output("Out"); - Out->ShareDataWith(*Ref); + Out->ShareDataWith(*X); - GPUScatterAssign(ctx.device_context(), *Updates, *Index, Out); + GPUScatterAssign(ctx.device_context(), *Updates, *Ids, Out); } }; @@ -42,16 +42,16 @@ class ScatterGradOpCUDAKernel : public framework::OpKernel { void Compute(const framework::ExecutionContext &ctx) const override { PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()), "This kernel only runs on GPU device."); - auto *dRef = ctx.Output(framework::GradVarName("Ref")); + auto *dX = ctx.Output(framework::GradVarName("X")); auto *dUpdates = ctx.Output(framework::GradVarName("Updates")); - auto *Index = ctx.Input("Index"); + auto *Ids = ctx.Input("Ids"); auto *dOut = ctx.Input(framework::GradVarName("Out")); - // In place gradient: dRef = dO - dRef->ShareDataWith(*dOut); + // In place gradient: dX = dO + dX->ShareDataWith(*dOut); dUpdates->mutable_data(ctx.GetPlace()); - // Gradient by Gather: dUpdates = dO[Index] - GPUGather(ctx.device_context(), *dOut, *Index, dUpdates); + // Gradient by Gather: dUpdates = dO[Ids] + GPUGather(ctx.device_context(), *dOut, *Ids, dUpdates); } }; diff --git a/paddle/fluid/operators/scatter_op.h b/paddle/fluid/operators/scatter_op.h index 3c6e7ece320229e1a311ef6d7a27387d40be3c2a..2151d8a9240fc88966533f4a07d5cf56b6c1c3bc 100644 --- a/paddle/fluid/operators/scatter_op.h +++ b/paddle/fluid/operators/scatter_op.h @@ -29,15 +29,15 @@ class ScatterOpKernel : public framework::OpKernel { void Compute(const framework::ExecutionContext &ctx) const override { PADDLE_ENFORCE(platform::is_cpu_place(ctx.GetPlace()), "This kernel only runs on CPU."); - auto *Ref = ctx.Input("Ref"); - auto *Index = ctx.Input("Index"); + auto *X = ctx.Input("X"); + auto *Ids = ctx.Input("Ids"); auto *Updates = ctx.Input("Updates"); auto *Out = ctx.Output("Out"); - // In place output: Out = Ref, Out[Index] += Updates - Out->ShareDataWith(*Ref); + // In place output: Out = X, Out[Ids] += Updates + Out->ShareDataWith(*X); // Apply ScatterUpdate: Out[index] += Updates[:] - ScatterAssign(ctx.device_context(), *Updates, *Index, Out); + ScatterAssign(ctx.device_context(), *Updates, *Ids, Out); } }; @@ -47,16 +47,16 @@ class ScatterGradientOpKernel : public framework::OpKernel { void Compute(const framework::ExecutionContext &ctx) const override { PADDLE_ENFORCE(platform::is_cpu_place(ctx.GetPlace()), "This kernel only runs on CPU."); - auto *dRef = ctx.Output(framework::GradVarName("Ref")); + auto *dX = ctx.Output(framework::GradVarName("X")); auto *dUpdates = ctx.Output(framework::GradVarName("Updates")); - auto *Index = ctx.Input("Index"); + auto *Ids = ctx.Input("Ids"); auto *dOut = ctx.Input(framework::GradVarName("Out")); - // In place gradient: dRef = dO - dRef->ShareDataWith(*dOut); + // In place gradient: dX = dO + dX->ShareDataWith(*dOut); dUpdates->mutable_data(ctx.GetPlace()); - // Gradient by Gather: dUpdates += dO[Index] - CPUGather(ctx.device_context(), *dOut, *Index, dUpdates); + // Gradient by Gather: dUpdates += dO[Ids] + CPUGather(ctx.device_context(), *dOut, *Ids, dUpdates); } }; diff --git a/paddle/fluid/operators/select_op.cc b/paddle/fluid/operators/select_op.cc new file mode 100644 index 0000000000000000000000000000000000000000..8344a239df7b3fcbe91f91a17a3c5958013b55a6 --- /dev/null +++ b/paddle/fluid/operators/select_op.cc @@ -0,0 +1,414 @@ +/* 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 +#include +#include +#include +#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 casesConfigs = + Attr>(kCases); + + framework::BlockDesc *casesBlock = + Attr(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> cases = + ParseAndShuffleCases(&casesConfigs); + + // Get all unique channels involved in select + std::set channelsSet; + for (auto c : cases) { + if (!c->channelName.empty()) { + auto channelVar = scope.FindVar(c->channelName); + framework::ChannelHolder *ch = + channelVar->GetMutable(); + + if (channelsSet.find(ch) == channelsSet.end()) { + channelsSet.insert(ch); + } + } + } + + // Order all channels by their pointer address + std::vector 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(); + caseToExecuteTensor->data()[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> ParseAndShuffleCases( + std::vector *casesConfigs) const { + std::vector> cases; + std::shared_ptr defaultCase; + + if (casesConfigs != nullptr) { + boost::char_delimiters_separator sep(false, ",", ""); + for (std::vector::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(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> *cases, + std::vector channels) const { + // Lock all involved channels + lockChannels(channels); + + std::atomic caseToExecute(-1); + + std::vector>::iterator it = cases->begin(); + while (it != cases->end()) { + std::shared_ptr c = *it; + + auto chVar = scope->FindVar(c->channelName); + framework::ChannelHolder *ch = + chVar->GetMutable(); + + 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 completed(false); + std::recursive_mutex mutex; + std::unique_lock lock{mutex}; + // std::condition_variable_any selectCond; + auto selectCond = std::make_shared(); + + 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 chs) const { + std::vector::iterator it = chs.begin(); + while (it != chs.end()) { + framework::ChannelHolder *ch = *it; + ch->Lock(); + ++it; + } + } + + void unlockChannels(std::vector chs) const { + std::vector::reverse_iterator it = chs.rbegin(); + while (it != chs.rend()) { + framework::ChannelHolder *ch = *it; + ch->Unlock(); + ++it; + } + } + + void pushThreadOnChannelQueues( + const framework::Scope *scope, + std::vector> *cases, + std::shared_ptr rCond, + std::atomic &caseToExecute, std::atomic &completed, + std::recursive_mutex &callbackMutex) const { + std::vector>::iterator it = cases->begin(); + while (it != cases->end()) { + std::shared_ptr c = *it; + + auto chVar = scope->FindVar(c->channelName); + framework::ChannelHolder *ch = + chVar->GetMutable(); + + std::function cb = + [&caseToExecute, &completed, &callbackMutex, + c](framework::ChannelAction channelAction) { + std::lock_guard 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> *cases) const { + std::vector>::iterator it = cases->begin(); + while (it != cases->end()) { + std::shared_ptr c = *it; + + auto chVar = scope->FindVar(c->channelName); + framework::ChannelHolder *ch = + chVar->GetMutable(); + 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>(kCases, + "(String vector) Serialized list of" + "all cases in the select op. Each" + "case is serialized as: " + "',,,'" + "where type is 0 for default, 1 for" + "send, and 2 for receive" + "No channel and values are needed for" + "default cases."); + AddAttr(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); diff --git a/paddle/fluid/operators/sequence_softmax_cudnn_op.cu.cc b/paddle/fluid/operators/sequence_softmax_cudnn_op.cu.cc new file mode 100644 index 0000000000000000000000000000000000000000..5661f4b42f37fed7f589c515e25fd66cfcede2c7 --- /dev/null +++ b/paddle/fluid/operators/sequence_softmax_cudnn_op.cu.cc @@ -0,0 +1,105 @@ +/* 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 +class SequenceSoftmaxCUDNNKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext& ctx) const override { + auto* x = ctx.Input("X"); + auto* out = ctx.Output("Out"); + + auto lod = x->lod(); + auto dims = x->dims(); + + const size_t level = lod.size() - 1; + PADDLE_ENFORCE_EQ(dims[0], static_cast(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(ctx.GetPlace()); + for (int i = 0; i < static_cast(lod[level].size()) - 1; ++i) { + int start_pos = static_cast(lod[level][i]); + int end_pos = static_cast(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()( + ctx.template device_context(), &x_i, + &out_i); + } + } +}; + +template +class SequenceSoftmaxGradCUDNNKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext& ctx) const override { + auto* out = ctx.Input("Out"); + auto* out_grad = ctx.Input(framework::GradVarName("Out")); + auto* x = ctx.Input("X"); + auto* x_grad = ctx.Output(framework::GradVarName("X")); + + auto lod = x->lod(); + const size_t level = lod.size() - 1; + + x_grad->mutable_data(ctx.GetPlace()); + for (int i = 0; i < static_cast(lod[level].size()) - 1; ++i) { + int start_pos = static_cast(lod[level][i]); + int end_pos = static_cast(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()( + ctx.template device_context(), &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, + ops::SequenceSoftmaxCUDNNKernel) +REGISTER_OP_KERNEL(sequence_softmax_grad, CUDNN, ::paddle::platform::CUDAPlace, + ops::SequenceSoftmaxGradCUDNNKernel, + ops::SequenceSoftmaxGradCUDNNKernel) diff --git a/paddle/fluid/operators/sequence_softmax_op.cc b/paddle/fluid/operators/sequence_softmax_op.cc index 7e685eb3dc7b12ef38f06b37d99a1212cfbc992c..e8b4df04286d327f568f4c43886f9fcf89cc4a88 100644 --- a/paddle/fluid/operators/sequence_softmax_op.cc +++ b/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("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(); + 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("data_format"); + return framework::OpKernelType( + framework::ToDataType(ctx.Input("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( + "use_cudnn", + "(bool, default false) Only used in cudnn kernel, need install cudnn") + .SetDefault(false); + AddAttr( + "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("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(); + 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("data_format"); + return framework::OpKernelType( + framework::ToDataType(ctx.Input("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); + ops::SequenceSoftmaxKernel, + ops::SequenceSoftmaxKernel); REGISTER_OP_CPU_KERNEL( sequence_softmax_grad, - ops::SequenceSoftmaxGradKernel); + ops::SequenceSoftmaxGradKernel, + ops::SequenceSoftmaxGradKernel); diff --git a/paddle/fluid/operators/sequence_softmax_op.cu.cc b/paddle/fluid/operators/sequence_softmax_op.cu.cc index 295c68c5b936d6522666a4cc4e621db6f5f5f3ed..57adea3a1b9dbcbb5787d005e4d3ec595f61d4b2 100644 --- a/paddle/fluid/operators/sequence_softmax_op.cu.cc +++ b/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) + ops::SequenceSoftmaxKernel, + ops::SequenceSoftmaxKernel) REGISTER_OP_CUDA_KERNEL( sequence_softmax_grad, - ops::SequenceSoftmaxGradKernel); + ops::SequenceSoftmaxGradKernel, + ops::SequenceSoftmaxGradKernel); diff --git a/paddle/fluid/operators/softmax_cudnn_op.cu.cc b/paddle/fluid/operators/softmax_cudnn_op.cu.cc new file mode 100644 index 0000000000000000000000000000000000000000..47cb336d87f8627d86ac33d6ac32c04d5d93f753 --- /dev/null +++ b/paddle/fluid/operators/softmax_cudnn_op.cu.cc @@ -0,0 +1,62 @@ +/* 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 +class SoftmaxCUDNNKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext& context) const override { + auto* X = context.Input("X"); + auto* Out = context.Output("Out"); + + // allocate memory on device. + Out->mutable_data(context.GetPlace()); + + math::SoftmaxCUDNNFunctor()( + context.template device_context(), X, Out); + } +}; + +template +class SoftmaxGradCUDNNKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext& context) const override { + auto* Out = context.Input("Out"); + auto* dOut = context.Input(framework::GradVarName("Out")); + auto* dX = context.Output(framework::GradVarName("X")); + + // allocate memory on device. + dX->mutable_data(context.GetPlace()); + + math::SoftmaxGradCUDNNFunctor()( + context.template device_context(), Out, + dOut, dX); + } +}; + +} // namespace operators +} // namespace paddle + +namespace ops = paddle::operators; +REGISTER_OP_KERNEL(softmax, CUDNN, ::paddle::platform::CUDAPlace, + ops::SoftmaxCUDNNKernel); +REGISTER_OP_KERNEL(softmax_grad, CUDNN, ::paddle::platform::CUDAPlace, + ops::SoftmaxGradCUDNNKernel); diff --git a/paddle/fluid/operators/softmax_op.cc b/paddle/fluid/operators/softmax_op.cc index 09275ef290e8c78dc0902033e904cc4e7ccd7adb..1b63f8a499e5d20d2f10c3cd1024d1bcf78764d4 100644 --- a/paddle/fluid/operators/softmax_op.cc +++ b/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("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(); + 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("data_format"); + return framework::OpKernelType( + framework::ToDataType(ctx.Input("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( + "use_cudnn", + "(bool, default false) Only used in cudnn kernel, need install cudnn") + .SetDefault(false); + AddAttr( + "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("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(); + 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("data_format"); + return framework::OpKernelType( + framework::ToDataType(ctx.Input("X")->type()), ctx.GetPlace(), + framework::StringToDataLayout(data_format), library_); + } }; } // namespace operators diff --git a/paddle/fluid/operators/softmax_with_cross_entropy_op.cu b/paddle/fluid/operators/softmax_with_cross_entropy_op.cu index 39b246a5bedb2819fc9b7fc407cfe03e59af0b68..8f7840cee1dd95a828fd4ac8815e335a5db47e3d 100644 --- a/paddle/fluid/operators/softmax_with_cross_entropy_op.cu +++ b/paddle/fluid/operators/softmax_with_cross_entropy_op.cu @@ -23,21 +23,21 @@ using Tensor = framework::Tensor; namespace { template -__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(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(1.); } +} - __syncthreads(); - - if (tid < batch_size * class_num) { - logit_grad[tid] *= loss_grad[sample_idx]; +template +__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 { 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("soft_label")) { + int grid = (batch_size * class_num + block - 1) / block; const T* label_data = labels->data(); - SoftCrossEntropyGradientKernel< - T><<() - .stream()>>>(logit_grad_data, loss_grad_data, label_data, - batch_size, class_num); + SoftCrossEntropyGradientKernel<<>>( + 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(); - CrossEntropyGrad< - T><<() - .stream()>>>(logit_grad_data, loss_grad_data, label_data, - batch_size, class_num); + CrossEntropyGrad<<>>( + logit_grad_data, label_data, batch_size, class_num); + int num = batch_size * class_num; + grid = (num + block - 1) / block; + Scale<<>>(logit_grad_data, loss_grad_data, num, + class_num); } } }; diff --git a/paddle/fluid/platform/cudnn_helper.h b/paddle/fluid/platform/cudnn_helper.h index 1842ecd745e3f5cb75600ce00d89018f81682632..9a2ac3ff33df3f8b9e24203f9dba2130e1d16510 100644 --- a/paddle/fluid/platform/cudnn_helper.h +++ b/paddle/fluid/platform/cudnn_helper.h @@ -289,7 +289,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(); + auto& dev_ctx = ctx.device_context(); use_cudnn &= dev_ctx.cudnn_handle() != nullptr; } #endif diff --git a/paddle/fluid/platform/device_tracer.cc b/paddle/fluid/platform/device_tracer.cc index 78e00d5420bbea40c9bea4be919ec4ce5ececdcb..3b4437f576e1c2e931a86ec6d5e823ec1f344c52 100644 --- a/paddle/fluid/platform/device_tracer.cc +++ b/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 block_id_stack; +// Tracking the nested event stacks. +thread_local std::deque 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 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 diff --git a/paddle/fluid/platform/device_tracer.h b/paddle/fluid/platform/device_tracer.h index 23f7cdbdffc9f48ac5555455bf745233c81dd0cb..deb3d23f786353b8e7a2f28d094e364158885a34 100644 --- a/paddle/fluid/platform/device_tracer.h +++ b/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 diff --git a/paddle/fluid/platform/profiler.cc b/paddle/fluid/platform/profiler.cc index 28ef3e04b1c50e0d42eeb27608259c6449429da5..b25206ff35cc87dcdd363bc0de54530f629d73ed 100644 --- a/paddle/fluid/platform/profiler.cc +++ b/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 ", diff --git a/paddle/fluid/platform/profiler.h b/paddle/fluid/platform/profiler.h index 3542ce6cda87e3b013d60393e4ba93da61921940..030458f70a6630de56ef896a5730112a018ac0c3 100644 --- a/paddle/fluid/platform/profiler.h +++ b/paddle/fluid/platform/profiler.h @@ -118,6 +118,24 @@ struct RecordEvent { std::string full_name_; }; +struct RecordBlock { + explicit RecordBlock(int block_id); + ~RecordBlock(); + + private: + std::string name_; + uint64_t start_ns_; + int block_id_; +}; + +struct RecordThread { + explicit RecordThread(int thread_id); + ~RecordThread(); + + private: + uint64_t start_ns_; +}; + // 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> GetAllEvents(); diff --git a/paddle/fluid/platform/profiler.proto b/paddle/fluid/platform/profiler.proto index 71b5a9b12ef4a045ebfd3ee3d06ee25032083ff5..7b42aa785ec6ad5731e3adee1e9f189127a826a1 100644 --- a/paddle/fluid/platform/profiler.proto +++ b/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; } diff --git a/paddle/fluid/pybind/pybind.cc b/paddle/fluid/pybind/pybind.cc index d2e883caccdd34a9d662f06b83cf9a71d3d4a51e..6c05442466f5f3d8e04a8f0a2206443b1007a107 100644 --- a/paddle/fluid/pybind/pybind.cc +++ b/paddle/fluid/pybind/pybind.cc @@ -31,6 +31,7 @@ limitations under the License. */ #include "paddle/fluid/operators/cond_op.h" #include "paddle/fluid/operators/net_op.h" #include "paddle/fluid/platform/enforce.h" +#include "paddle/fluid/platform/gpu_info.h" #include "paddle/fluid/platform/place.h" #include "paddle/fluid/platform/profiler.h" #include "paddle/fluid/pybind/const_value.h" @@ -103,12 +104,14 @@ PYBIND11_PLUGIN(core) { .def("set", PyCPUTensorSetFromArray) .def("set", PyCPUTensorSetFromArray) .def("set", PyCPUTensorSetFromArray) + .def("set", PyCPUTensorSetFromArray) #ifdef PADDLE_WITH_CUDA .def("set", PyCUDATensorSetFromArray) .def("set", PyCUDATensorSetFromArray) .def("set", PyCUDATensorSetFromArray) .def("set", PyCUDATensorSetFromArray) .def("set", PyCUDATensorSetFromArray) + .def("set", PyCUDATensorSetFromArray) #endif .def("shape", [](Tensor &self) { return vectorize(self.dims()); }) .def("set_float_element", TensorSetElement) @@ -315,7 +318,6 @@ All parameter, weight, gradient are variables in Paddle. #endif }); // clang-format on - #ifdef PADDLE_WITH_CUDA py::class_(m, "Communicator").def(py::init<>()); #endif @@ -423,6 +425,12 @@ All parameter, weight, gradient are variables in Paddle. m.def("init_devices", &framework::InitDevices); m.def("is_compiled_with_cuda", IsCompiledWithCUDA); +#ifdef PADDLE_WITH_CUDA + m.def("is_float16_supported", [](const platform::CUDAPlace &place) -> bool { + // Only GPUs with Compute Capability >= 53 support float16 + return platform::GetCUDAComputeCapability(place.device) >= 53; + }); +#endif m.def("set_feed_variable", framework::SetFeedVariable); m.def("get_fetch_variable", framework::GetFetchVariable); diff --git a/paddle/fluid/pybind/tensor_py.h b/paddle/fluid/pybind/tensor_py.h index 1b0916ea0370d95a0c7dd149ee3f7b294c5e2351..6f8c597f8e610594851c318c122563523e4e7ea6 100644 --- a/paddle/fluid/pybind/tensor_py.h +++ b/paddle/fluid/pybind/tensor_py.h @@ -17,6 +17,7 @@ limitations under the License. */ #include "paddle/fluid/framework/lod_tensor.h" #include "paddle/fluid/memory/memcpy.h" #include "paddle/fluid/platform/device_context.h" +#include "paddle/fluid/platform/float16.h" #include "pybind11/numpy.h" #include "pybind11/pybind11.h" @@ -71,27 +72,39 @@ struct CastToPyBufferImpl { 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 } else if (paddle::platform::is_cpu_place(tensor.place())) { dst_tensor = tensor; } - return py::buffer_info(dst_tensor.data(), sizeof(CUR_TYPE), - py::format_descriptor::format(), - (size_t)framework::arity(dst_tensor.dims()), - dims_outside, strides); + + if (std::type_index(typeid(CUR_TYPE)) == + std::type_index(typeid(platform::float16))) { + return py::buffer_info(dst_tensor.data(), sizeof(CUR_TYPE), + "e", /* np.dtype('e') == np.float16 */ + (size_t)framework::arity(dst_tensor.dims()), + dims_outside, strides); + } else { + return py::buffer_info(dst_tensor.data(), sizeof(CUR_TYPE), + py::format_descriptor::format(), + (size_t)framework::arity(dst_tensor.dims()), + dims_outside, strides); + } } else { constexpr bool less = I + 1 < std::tuple_size>::value; return CastToPyBufferImpl()(tensor); } } }; + } // namespace details + inline py::buffer_info CastToPyBuffer(framework::Tensor &tensor) { auto buffer_info = - details::CastToPyBufferImpl()( - tensor); + details::CastToPyBufferImpl()(tensor); return buffer_info; } @@ -136,6 +149,22 @@ void PyCPUTensorSetFromArray( std::memcpy(dst, array.data(), sizeof(T) * array.size()); } +template <> +void PyCPUTensorSetFromArray( + framework::Tensor &self, + py::array_t array, + paddle::platform::CPUPlace &place) { + std::vector dims; + dims.reserve(array.ndim()); + for (size_t i = 0; i < array.ndim(); ++i) { + dims.push_back((int)array.shape()[i]); + } + + self.Resize(framework::make_ddim(dims)); + auto *dst = self.mutable_data(place); + std::memcpy(dst, array.data(), sizeof(uint16_t) * array.size()); +} + #ifdef PADDLE_WITH_CUDA template void PyCUDATensorSetFromArray( @@ -157,6 +186,28 @@ void PyCUDATensorSetFromArray( paddle::platform::GpuMemcpyAsync(dst, array.data(), sizeof(T) * array.size(), cudaMemcpyHostToDevice, dev_ctx->stream()); } + +template <> +void PyCUDATensorSetFromArray( + framework::Tensor &self, + py::array_t array, + paddle::platform::CUDAPlace &place) { + std::vector dims; + dims.reserve(array.ndim()); + for (size_t i = 0; i < array.ndim(); ++i) { + dims.push_back((int)array.shape()[i]); + } + + self.Resize(framework::make_ddim(dims)); + auto *dst = self.mutable_data(place); + + platform::DeviceContextPool &pool = platform::DeviceContextPool::Instance(); + auto dev_ctx = + static_cast(pool.Get(place)); + paddle::platform::GpuMemcpyAsync(dst, array.data(), + sizeof(uint16_t) * array.size(), + cudaMemcpyHostToDevice, dev_ctx->stream()); +} #endif } // namespace pybind diff --git a/python/CMakeLists.txt b/python/CMakeLists.txt index 6e24cbdd3f6a4f05c1691dc643d880f6f454429d..90c2dfbba78418fb7b731f5363017d70577b1ae5 100644 --- a/python/CMakeLists.txt +++ b/python/CMakeLists.txt @@ -1,27 +1,29 @@ - -file(GLOB TRAINER_PY_FILES . ./paddle/trainer/*.py) -file(GLOB HELPERS_PY_FILES . ./paddle/trainer_config_helpers/*.py) file(GLOB UTILS_PY_FILES . ./paddle/utils/*.py) -file(GLOB_RECURSE V2_PY_FILES ./paddle/v2/ *.py) file(GLOB_RECURSE FLUID_PY_FILES ./paddle/fluid/ *.py) - set(PY_FILES paddle/__init__.py - ${TRAINER_PY_FILES} - ${HELPERS_PY_FILES} ${UTILS_PY_FILES} - ${V2_PY_FILES} ${FLUID_PY_FILES}) -add_custom_target(copy_paddle_master) +if(NOT WITH_FLUID) + file(GLOB TRAINER_PY_FILES . ./paddle/trainer/*.py) + file(GLOB HELPERS_PY_FILES . ./paddle/trainer_config_helpers/*.py) + file(GLOB_RECURSE V2_PY_FILES ./paddle/v2/ *.py) + set(PY_FILES ${PY_FILES} + ${TRAINER_PY_FILES} + ${HELPERS_PY_FILES} + ${V2_PY_FILES}) -SET(COPY_PADDLE_MASTER "") -if(WITH_GOLANG) - SET(COPY_PADDLE_MASTER "copy_paddle_master") - add_custom_command(TARGET ${COPY_PADDLE_MASTER} - COMMAND cp ${paddle_master_LIB_PATH} ${PADDLE_SOURCE_DIR}/python/paddle/v2/master/ - ) - add_dependencies(copy_paddle_master paddle_master) -endif(WITH_GOLANG) + add_custom_target(copy_paddle_master) + + SET(COPY_PADDLE_MASTER "") + if(WITH_GOLANG) + SET(COPY_PADDLE_MASTER "copy_paddle_master") + add_custom_command(TARGET ${COPY_PADDLE_MASTER} + COMMAND cp ${paddle_master_LIB_PATH} ${PADDLE_SOURCE_DIR}/python/paddle/v2/master/ + ) + add_dependencies(copy_paddle_master paddle_master) + endif(WITH_GOLANG) +endif() set(MKL_SHARED_LIBS "") set(MKL_DEPENDS "") @@ -59,23 +61,28 @@ add_custom_command(OUTPUT ${PADDLE_PYTHON_BUILD_DIR}/.timestamp COMMAND ${CMAKE_COMMAND} -E copy_directory ${PADDLE_PYTHON_BUILD_DIR}/lib* ${PADDLE_PYTHON_BUILD_DIR}/lib-python DEPENDS gen_proto_py copy_paddle_pybind framework_py_proto profiler_py_proto ${PY_FILES} ${external_project_dependencies} ${COPY_PADDLE_MASTER}) -set(paddle_python_deps ${PADDLE_PYTHON_BUILD_DIR}/.timestamp paddle_pserver_main paddle_trainer paddle_merge_model ${MKL_DEPENDS}) -if(WITH_SWIG_PY) - list(APPEND paddle_python_deps python_api_wheel) +set(paddle_python_deps ${PADDLE_PYTHON_BUILD_DIR}/.timestamp ${MKL_DEPENDS}) +if(NOT WITH_FLUID) + set(paddle_python_deps ${paddle_python_deps} paddle_pserver_main paddle_trainer paddle_merge_model) + if(WITH_SWIG_PY) + list(APPEND paddle_python_deps python_api_wheel) + endif() endif() add_custom_target(paddle_python ALL DEPENDS ${paddle_python_deps}) set(PADDLE_PYTHON_PACKAGE_DIR ${CMAKE_CURRENT_BINARY_DIR}/dist/) if (WITH_TESTING) - add_subdirectory(paddle/trainer_config_helpers/tests) - if (WITH_SWIG_PY) - # enable v2 API unittest only when paddle swig api is compiled - add_subdirectory(paddle/v2/tests) - add_subdirectory(paddle/v2/reader/tests) - add_subdirectory(paddle/v2/plot/tests) - add_subdirectory(paddle/fluid/tests) + if(NOT WITH_FLUID) + add_subdirectory(paddle/trainer_config_helpers/tests) + if (WITH_SWIG_PY) + # enable v2 API unittest only when paddle swig api is compiled + add_subdirectory(paddle/v2/tests) + add_subdirectory(paddle/v2/reader/tests) + add_subdirectory(paddle/v2/plot/tests) + endif() endif() + add_subdirectory(paddle/fluid/tests) endif() install(DIRECTORY ${PADDLE_PYTHON_PACKAGE_DIR} DESTINATION opt/paddle/share/wheels diff --git a/python/paddle/fluid/__init__.py b/python/paddle/fluid/__init__.py index dcde08632a6bb4c5936c32048c2cc1dca7608b06..fcea28220485039c9daf3c5fa2688c31f9f34c42 100644 --- a/python/paddle/fluid/__init__.py +++ b/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 diff --git a/python/paddle/fluid/backward.py b/python/paddle/fluid/backward.py index b6f20daee3a585777a23255355f0a0e31328d23f..7af6ed1463ab737e871da487f2a687301652ef2d 100644 --- a/python/paddle/fluid/backward.py +++ b/python/paddle/fluid/backward.py @@ -248,12 +248,15 @@ def _callback_lookup_(op): if o_argu in self.param_grad_names: allreduce_out_name = o_argu + "__nccl_all_reduce__" op_desc = _create_op_desc_( - "ncclAllReduce", { + "ncclReduce", + { "X": [o_argu], "Communicator": ['nccl_com__do_not_change_'] - }, {"Out": [allreduce_out_name]}, - {"reduction": "ncclSum"}) + }, + {"Out": [allreduce_out_name]}, + {"reduction": "ncclSum", + "root": 0}, ) block.desc.append_op().copy_from(op_desc) op_desc = _create_op_desc_( diff --git a/python/paddle/fluid/concurrency.py b/python/paddle/fluid/concurrency.py index dec224fc886cd0739add0ebb6488625ef5063b8d..535e881c42f675198a2679cb7974af64b65cc194 100644 --- a/python/paddle/fluid/concurrency.py +++ b/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. """ diff --git a/python/paddle/fluid/framework.py b/python/paddle/fluid/framework.py index d14d6349b1bcf598e25bbeb9913d2d0da71a5054..70ecffd910a46570b5a8e576d88039fa5e22e726 100644 --- a/python/paddle/fluid/framework.py +++ b/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) diff --git a/python/paddle/fluid/layers/control_flow.py b/python/paddle/fluid/layers/control_flow.py index 1bb1aa30ee1019c6f80eb64b6dc20459e7a3073b..02cd0a05a11d8d1d52d42c2b62799f1093d0abc2 100644 --- a/python/paddle/fluid/layers/control_flow.py +++ b/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): diff --git a/python/paddle/fluid/layers/detection.py b/python/paddle/fluid/layers/detection.py index ea189749bc6cc1e37c1efc6fea424143b887cecd..a889ab6bdc6ac9494ef992a97292b7a2536c41c4 100644 --- a/python/paddle/fluid/layers/detection.py +++ b/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]) diff --git a/python/paddle/fluid/layers/nn.py b/python/paddle/fluid/layers/nn.py index f107261f3df78dfc7197d9719d0258b6ab09d487..bf161d6618b10da66f25d3f11300a4a2b10b875a 100644 --- a/python/paddle/fluid/layers/nn.py +++ b/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, diff --git a/python/paddle/fluid/layers/ops.py b/python/paddle/fluid/layers/ops.py index 0b88b639629ac73b16ec36aa5930c3d6a9665943..d7bad221c5fa7b18137bf317125195267437a644 100644 --- a/python/paddle/fluid/layers/ops.py +++ b/python/paddle/fluid/layers/ops.py @@ -59,8 +59,6 @@ __all__ = [ 'elementwise_pow', 'clip', 'clip_by_norm', - 'softmax', - 'sequence_softmax', 'logical_and', 'logical_or', 'logical_xor', @@ -70,6 +68,7 @@ __all__ = [ 'gaussian_random', 'gaussian_random_batch_size_like', 'cumsum', + 'scatter', ] + __activations__ for _OP in set(__all__): diff --git a/python/paddle/fluid/optimizer.py b/python/paddle/fluid/optimizer.py index 421963a2f9120dae3a72142681f0a30232c11166..8b8621469d856e63dfd9685d7cd3d4c1d2ada1ce 100644 --- a/python/paddle/fluid/optimizer.py +++ b/python/paddle/fluid/optimizer.py @@ -223,6 +223,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 diff --git a/python/paddle/fluid/tests/test_concurrency.py b/python/paddle/fluid/tests/test_concurrency.py index 9f7bf63c5e017251e87af94690ff32c47c538c6b..924895a9afac610059bac5f617c49712441339cc 100644 --- a/python/paddle/fluid/tests/test_concurrency.py +++ b/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() diff --git a/python/paddle/fluid/tests/unittests/CMakeLists.txt b/python/paddle/fluid/tests/unittests/CMakeLists.txt index f96c2ca4f0593b6c2624d449304f23425c69ab93..0ad273c7161977e18f91f952fd3a9dc144bf73f0 100644 --- a/python/paddle/fluid/tests/unittests/CMakeLists.txt +++ b/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 diff --git a/python/paddle/fluid/tests/unittests/test_detection_output_op.py b/python/paddle/fluid/tests/unittests/test_detection_output_op.py deleted file mode 100644 index 94681319144ee3e0d51b57944f5692183578c01b..0000000000000000000000000000000000000000 --- a/python/paddle/fluid/tests/unittests/test_detection_output_op.py +++ /dev/null @@ -1,71 +0,0 @@ -# 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() diff --git a/python/paddle/fluid/tests/unittests/test_layers.py b/python/paddle/fluid/tests/unittests/test_layers.py index 6944cca394fbc1ddde09dfeb0bc82e357a3cd225..90d70aa39fdc4d4d3f9062eb6a3eb0cdd014acfc 100644 --- a/python/paddle/fluid/tests/unittests/test_layers.py +++ b/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): diff --git a/python/paddle/fluid/tests/unittests/test_lrn_op.py b/python/paddle/fluid/tests/unittests/test_lrn_op.py index 7f2352c5882ce36d8d681a737806f3ee0e3ace98..eaff45cbb2a58798e9d55149510bec72eea370cd 100644 --- a/python/paddle/fluid/tests/unittests/test_lrn_op.py +++ b/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 diff --git a/python/paddle/fluid/tests/unittests/test_mul_op.py b/python/paddle/fluid/tests/unittests/test_mul_op.py index 9d1da420c7f70bd2a89d183a5f0a2b145f0ff475..40440bea1267112b84b66002a0bf921be3029265 100644 --- a/python/paddle/fluid/tests/unittests/test_mul_op.py +++ b/python/paddle/fluid/tests/unittests/test_mul_op.py @@ -14,6 +14,7 @@ import unittest import numpy as np +import paddle.fluid.core as core from op_test import OpTest @@ -69,5 +70,42 @@ class TestMulOp2(OpTest): ['X'], 'Out', max_relative_error=0.5, no_grad_set=set('Y')) +class TestFP16MulOp1(OpTest): + def setUp(self): + self.op_type = "mul" + x = np.random.random((32, 84)).astype("float16") + y = np.random.random((84, 100)).astype("float16") + self.inputs = {'X': x.view(np.uint16), 'Y': y.view(np.uint16)} + self.outputs = {'Out': np.dot(x, y)} + + 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-1) + + +class TestFP16MulOp2(OpTest): + def setUp(self): + self.op_type = "mul" + x = np.random.random((15, 4, 12, 10)).astype("float16") + y = np.random.random((4, 30, 8, 2, 9)).astype("float16") + self.inputs = {'X': x.view(np.uint16), 'Y': y.view(np.uint16)} + self.attrs = { + 'x_num_col_dims': 2, + 'y_num_col_dims': 2, + } + result = np.dot( + x.reshape(15 * 4, 12 * 10), y.reshape(4 * 30, 8 * 2 * 9)) + result = result.reshape(15, 4, 8, 2, 9) + self.outputs = {'Out': result} + + 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-1) + + if __name__ == "__main__": unittest.main() diff --git a/python/paddle/fluid/tests/unittests/test_parallel_op.py b/python/paddle/fluid/tests/unittests/test_parallel_op.py index 1a7551c57b26f576ab286e7b18177b9120261623..79bea148f9398152a02d70946cdc5fff1f47ba6b 100644 --- a/python/paddle/fluid/tests/unittests/test_parallel_op.py +++ b/python/paddle/fluid/tests/unittests/test_parallel_op.py @@ -15,6 +15,7 @@ import unittest import paddle.fluid as fluid +import paddle.fluid.profiler as profiler import numpy @@ -60,20 +61,23 @@ class BaseParallelForTest(unittest.TestCase): feed=feed, fetch=fetch, place=gpu, - use_parallel=False) + use_parallel=False, + use_gpu=True) result_gpu_parallel = self._run_test_impl_( callback=callback, feed=feed, fetch=fetch, place=gpu, - use_parallel=True) + use_parallel=True, + use_gpu=True) result_gpu_nccl = self._run_test_impl_( callback=callback, feed=feed, fetch=fetch, place=gpu, use_parallel=True, - use_nccl=True) + use_nccl=True, + use_gpu=True) self._assert_same_(fetch, result_cpu, result_cpu_parallel, result_gpu, result_gpu_parallel, result_gpu_nccl) else: @@ -85,7 +89,8 @@ class BaseParallelForTest(unittest.TestCase): fetch, place, use_parallel=False, - use_nccl=False): + use_nccl=False, + use_gpu=False): """ Run a single test, returns the fetch values Args: @@ -132,7 +137,12 @@ class BaseParallelForTest(unittest.TestCase): exe = fluid.Executor(place) exe.run(startup) - return exe.run(main, feed=feed, fetch_list=fetch) + if use_gpu: + profile_type = 'GPU' + else: + profile_type = 'CPU' + with profiler.profiler(profile_type, 'total', '/tmp/profiler'): + return exe.run(main, feed=feed, fetch_list=fetch) def _assert_same_(self, fetch, *args): """ diff --git a/python/paddle/fluid/tests/unittests/test_profiler.py b/python/paddle/fluid/tests/unittests/test_profiler.py index 1da6b94eea30e65913ce713e0e5e355507534161..cf6fe14a86aa1ab6ea3f60ad15f33d708e9b803a 100644 --- a/python/paddle/fluid/tests/unittests/test_profiler.py +++ b/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) diff --git a/python/paddle/fluid/tests/unittests/test_scatter_op.py b/python/paddle/fluid/tests/unittests/test_scatter_op.py index bb02a40d449860cf6c0576662e79a5e36e6e0635..fb1728743630b3ea908ae835444eff7fd71b72c8 100644 --- a/python/paddle/fluid/tests/unittests/test_scatter_op.py +++ b/python/paddle/fluid/tests/unittests/test_scatter_op.py @@ -25,7 +25,7 @@ class TestScatterOp(OpTest): updates_np = np.random.random((2, 3)).astype("float32") output_np = np.copy(ref_np) output_np[index_np] = updates_np - self.inputs = {'Ref': ref_np, 'Index': index_np, 'Updates': updates_np} + self.inputs = {'X': ref_np, 'Ids': index_np, 'Updates': updates_np} self.outputs = {'Out': output_np} def test_check_output(self): diff --git a/python/paddle/fluid/tests/unittests/test_sequence_softmax_op.py b/python/paddle/fluid/tests/unittests/test_sequence_softmax_op.py index 9e5c1e7a3d0bdf514de11e797d7139f577002c52..d6dc99bb3106feee33daa52bffb386f07cc16de5 100644 --- a/python/paddle/fluid/tests/unittests/test_sequence_softmax_op.py +++ b/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,12 +38,31 @@ 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): - self.check_output() + 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", max_relative_error=0.01) + 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__": diff --git a/python/paddle/fluid/tests/unittests/test_softmax_op.py b/python/paddle/fluid/tests/unittests/test_softmax_op.py index 8f8312edca7e2d98eb4e881f671c6afdda01c57a..4f20da2b926823db9e7ec92c95178b6d3d1feec9 100644 --- a/python/paddle/fluid/tests/unittests/test_softmax_op.py +++ b/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): - self.check_output() + 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__": diff --git a/python/paddle/fluid/tests/unittests/test_softmax_with_cross_entropy_op.py b/python/paddle/fluid/tests/unittests/test_softmax_with_cross_entropy_op.py index 889fea2ce66e64d439b51498722e571f48cd1f96..c0d9fc8f22a7c4f791d80a9cad87d003b5d54299 100644 --- a/python/paddle/fluid/tests/unittests/test_softmax_with_cross_entropy_op.py +++ b/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, diff --git a/python/setup.py.in b/python/setup.py.in index f830039a3af581d593d510326f15139377cb25f1..4cb5409524457b7bc5a99c88a0dbbfc8834923fa 100644 --- a/python/setup.py.in +++ b/python/setup.py.in @@ -62,20 +62,22 @@ write_version_py(filename='@PADDLE_SOURCE_DIR@/python/paddle/version.py') packages=['paddle', - 'paddle.proto', - 'paddle.trainer', - 'paddle.trainer_config_helpers', 'paddle.utils', - 'paddle.v2', - 'paddle.v2.dataset', - 'paddle.v2.reader', - 'paddle.v2.master', - 'paddle.v2.plot', 'paddle.fluid', 'paddle.fluid.proto', 'paddle.fluid.proto.profiler', - 'paddle.fluid.layers', - 'py_paddle'] + 'paddle.fluid.layers'] + +if '${WITH_FLUID}'== 'OFF': + packages+=['paddle.proto', + 'paddle.trainer', + 'paddle.trainer_config_helpers', + 'paddle.v2', + 'paddle.v2.dataset', + 'paddle.v2.reader', + 'paddle.v2.master', + 'paddle.v2.plot', + 'py_paddle'] with open('@PADDLE_SOURCE_DIR@/python/requirements.txt') as f: setup_requires = f.read().splitlines() @@ -84,11 +86,29 @@ if '${CMAKE_SYSTEM_PROCESSOR}' not in ['arm', 'armv7-a', 'aarch64']: setup_requires+=['opencv-python'] # the prefix is sys.prefix which should always be usr -paddle_bin_dir = 'opt/paddle/bin' -paddle_bins = ['${PADDLE_BINARY_DIR}/paddle/trainer/paddle_trainer', - '${PADDLE_BINARY_DIR}/paddle/trainer/paddle_merge_model', - '${PADDLE_BINARY_DIR}/paddle/pserver/paddle_pserver_main', - '${PADDLE_BINARY_DIR}/paddle/scripts/paddle'] +paddle_bins = '' +if '${WITH_FLUID}'== 'OFF': + paddle_bin_dir = 'opt/paddle/bin' + paddle_bins = ['${PADDLE_BINARY_DIR}/paddle/trainer/paddle_trainer', + '${PADDLE_BINARY_DIR}/paddle/trainer/paddle_merge_model', + '${PADDLE_BINARY_DIR}/paddle/pserver/paddle_pserver_main', + '${PADDLE_BINARY_DIR}/paddle/scripts/paddle'] + +package_data={'paddle.fluid': ['core.so']} +if '${WITH_FLUID}'== 'OFF': + package_data['paddle.v2.master']=['libpaddle_master.so'] + package_data['py_paddle']=['*.py','_swig_paddle.so'] + +package_dir={ + '': '${CMAKE_CURRENT_SOURCE_DIR}', + # The paddle.fluid.proto will be generated while compiling. + # So that package points to other directory. + 'paddle.fluid.proto.profiler': '${PADDLE_BINARY_DIR}/paddle/fluid/platform', + 'paddle.fluid.proto': '${PADDLE_BINARY_DIR}/paddle/fluid/framework', +} +if '${WITH_FLUID}'== 'OFF': + package_dir['py_paddle']='${PADDLE_SOURCE_DIR}/paddle/py_paddle' + paddle_rt_lib_dir = 'lib' paddle_rt_libs = ['${WARPCTC_LIBRARIES}'] @@ -101,19 +121,8 @@ setup(name='${PACKAGE_NAME}', install_requires=setup_requires, packages=packages, ext_modules=[Extension('_foo', ['stub.cc'])], - package_data={ - 'paddle.v2.master': ['libpaddle_master.so'], - 'paddle.fluid': ['core.so'], - 'py_paddle':['*.py','_swig_paddle.so'] - }, - package_dir={ - '': '${CMAKE_CURRENT_SOURCE_DIR}', - # The paddle.fluid.proto will be generated while compiling. - # So that package points to other directory. - 'paddle.fluid.proto.profiler': '${PADDLE_BINARY_DIR}/paddle/fluid/platform', - 'paddle.fluid.proto': '${PADDLE_BINARY_DIR}/paddle/fluid/framework', - 'py_paddle': '${PADDLE_SOURCE_DIR}/paddle/py_paddle' - }, + package_data=package_data, + package_dir=package_dir, scripts=paddle_bins, data_files=[(paddle_rt_lib_dir, paddle_rt_libs)] ) diff --git a/tools/timeline.py b/tools/timeline.py index ee83a1baecdd4243bb6c546486a837393980fb65..f4083c824e7333a74661d096d4954609f767c83e 100644 --- a/tools/timeline.py +++ b/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: - 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) + 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, "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()