modified reduce_max reduce_min reduce_prod for higher_performance and fix a bug in reduce_op.cuh #32974
Conversation
|
Thanks for your contribution! |
…iling123/Paddle into reduce_max_min_prod_all_any
…iling123/Paddle into reduce_max_min_prod_all_any
|
|
AnnaTrainingG
force-pushed
the
reduce_max_min_prod_all_any
branch
from
3069b09
to
2e8ad8f
Compare
AnnaTrainingG
changed the title
AnnaTrainingG
force-pushed
the
reduce_max_min_prod_all_any
branch
from
4ce20bb
to
4bd9644
Compare
AnnaTrainingG
force-pushed
the
reduce_max_min_prod_all_any
branch
from
d2dea81
to
469e0a5
Compare
AnnaTrainingG
force-pushed
the
reduce_max_min_prod_all_any
branch
from
469e0a5
to
8700894
Compare
|
LGTM |
1 similar comment
|
LGTM |
| } | ||
| } | ||
|
|
||
| // module function designed for global function |
There was a problem hiding this comment.
感觉模板可以再简化一下，一些参数没有必要通过模板传，比如ReduceType。关于ReduceType的if判断只执行一次，并没有在循环里面，所以通过输入参数传也不会多影响性能。减少一些模板，应该能够剪短一些编译时间。
对于TransformOp，感觉LaunchReduceKernel和LaunchKernel这两个函数是不需要将TransformOp作为模板的？ReduceKernelFunction看起来是需要的。另外，LaunchReduceKernel和LaunchKernel函数命名缺乏辨识度，不能准确地表达函数的功能。
| @@ -141,21 +174,24 @@ struct ReduceConfig { | |||
| void Run() { | |||
There was a problem hiding this comment.
L170的comment：输入参数都建议改成const std::vector &类型。
| @@ -523,22 +606,22 @@ static void launchKernel(const Tx* x_data, Ty* y_data, | |||
| ReduceKernelFunction< | |||
| Ty, Ty, ReduceOp, detail::IdentityFunctor<Ty>, 128, kRank, kReduceRank, | |||
| ReduceType::kReduceHigherDim><<<grid, block, 0, stream>>>( | |||
There was a problem hiding this comment.
L597 - L599 comment：直接写成CUB_REDUCE_TYPE_CASE(ReduceType::kReduceLastDim)这样？若ReduceType不作为模板，也就不需要这个swith case了。
| // SetOutputData for ReduceHigherDim when should_reduce_again is true, | ||
| // temp_output should be stored temp_data in output_data space or stored in | ||
| // y_data; | ||
| config.SetOutputData(y_data, x.place(), tmp); | ||
| framework::Tensor tmp; | ||
| config.SetOutputData(y_data, x.place(), &tmp); | ||
|
|
||
| if (config.reduce_num == 1) { | ||
| auto out_dims = y->dims(); | ||
| framework::TensorCopy(x, y->place(), y); | ||
| y->Resize(out_dims); | ||
| return; | ||
| } |
| } | ||
| }; | ||
|
|
||
| template <typename T, template <typename, typename> class ReduceOp> |
There was a problem hiding this comment.
上面的实现都可能需要复用到别的算子里面（比如broadcast反向），但ReduceCudaKernel只用于reduce_xxx算子的实现，所以L749 - L771最好不要放到这个头文件里面。
| int, ops::ProdFunctor>, | ||
| ops::ReduceKernel<paddle::platform::CUDADeviceContext, | ||
| int64_t, ops::ProdFunctor>); | ||
| REGISTER_OP_CUDA_KERNEL( |
There was a problem hiding this comment.
从注释来看，原来之所以加这个ifdef是因为原来的reduce采用Eigen实现，而Eigen对double的支持有问题。我们已经全部改成了cuda+cub的方式，或许这个ifdef可以去掉。
PR types
Function optimization
PR changes
OPs
Describe
modified reduce_min reduce_max reduce_prod reduce_all reduce_any
ctest结果：
Test project /paddle_test/commit/Paddle/build
Start 709: test_max_op
100% tests passed, 0 tests failed out of 1
Total Test time (real) = 7.51 sec
Test project /paddle_test/commit/Paddle/build
Start 719: test_min_op
100% tests passed, 0 tests failed out of 1
Total Test time (real) = 6.81 sec
Test project /paddle_test/commit/Paddle/build
Start 826: test_prod_op
100% tests passed, 0 tests failed out of 1
Total Test time (real) = 7.78 sec
以max 为例进行性能比对：
benchmark性能数据如下：