From 4feae253786cf826d677ec68ec0f99785661d9e3 Mon Sep 17 00:00:00 2001 From: Qiao Longfei Date: Sun, 13 Jan 2019 22:05:18 +0800 Subject: [PATCH] fix build problem test=develop --- .../operators/math/detail/gru_cpu_kernel.h | 25 ++++++++----------- .../operators/math/detail/gru_gpu_kernel.h | 5 ++-- .../fluid/operators/math/detail/gru_kernel.h | 5 ++-- paddle/fluid/operators/math/gru_compute.cc | 3 +-- python/paddle/fluid/layers/nn.py | 7 +++--- 5 files changed, 20 insertions(+), 25 deletions(-) diff --git a/paddle/fluid/operators/math/detail/gru_cpu_kernel.h b/paddle/fluid/operators/math/detail/gru_cpu_kernel.h index f3f5cad0cb4..c6dd972e12b 100644 --- a/paddle/fluid/operators/math/detail/gru_cpu_kernel.h +++ b/paddle/fluid/operators/math/detail/gru_cpu_kernel.h @@ -298,8 +298,7 @@ void hl_naive_gru_backward_reset_grad(OpResetGrad op_reset_grad, T *gate_value, T *gate_grad, T *prev_out_value, T *prev_out_grad, T *reset_output_grad, int frame_size, - ActivationType active_gate, - bool origin_mode) { + ActivationType active_gate) { T r_update_gate_value; T r_update_gate_grad; T r_reset_gate_value; @@ -329,8 +328,7 @@ void hl_naive_gru_backward_reset_grad(OpResetGrad op_reset_grad, T *gate_value, op_reset_grad(&r_update_gate_value, &r_update_gate_grad, &r_reset_gate_value, &r_reset_gate_grad, &r_prev_out_value, - &r_prev_out_grad, &r_reset_output_grad, active_gate, - origin_mode); + &r_prev_out_grad, &r_reset_output_grad, active_gate); update_gate_grad[i] = r_update_gate_grad; reset_gate_grad[i] = r_reset_gate_grad; @@ -389,8 +387,8 @@ template void hl_avx_gru_backward_reset_grad(OpResetGrad op_reset_grad, T *gate_value, T *gate_grad, T *prev_out_value, T *prev_out_grad, T *reset_output_grad, - int frame_size, ActivationType active_gate, - bool origin_mode) { + int frame_size, + ActivationType active_gate) { #ifdef __AVX__ __m256 r_update_gate_value; __m256 r_update_gate_grad; @@ -422,8 +420,7 @@ void hl_avx_gru_backward_reset_grad(OpResetGrad op_reset_grad, T *gate_value, op_reset_grad(&r_update_gate_value, &r_update_gate_grad, &r_reset_gate_value, &r_reset_gate_grad, &r_prev_out_value, - &r_prev_out_grad, &r_reset_output_grad, active_gate, - origin_mode); + &r_prev_out_grad, &r_reset_output_grad, active_gate); update_gate_grad[i] = r_update_gate_grad; reset_gate_grad[i] = r_reset_gate_grad; @@ -469,18 +466,16 @@ template inline void backward_reset_grad(OpResetGrad op_reset_grad, GRUMetaValue value, GRUMetaGrad grad, int frame_size, int batch_size, - ActivationType active_gate, bool origin_mode) { + ActivationType active_gate) { for (int b = 0; b < batch_size; b++) { if (OpResetGrad::avx && !(frame_size & (8 - 1)) && (sizeof(T) == 4)) { - hl_avx_gru_backward_reset_grad(op_reset_grad, value.gate_value, - grad.gate_grad, value.prev_out_value, - grad.prev_out_grad, grad.reset_output_grad, - frame_size, active_gate, origin_mode); + hl_avx_gru_backward_reset_grad( + op_reset_grad, value.gate_value, grad.gate_grad, value.prev_out_value, + grad.prev_out_grad, grad.reset_output_grad, frame_size, active_gate); } else { hl_naive_gru_backward_reset_grad( op_reset_grad, value.gate_value, grad.gate_grad, value.prev_out_value, - grad.prev_out_grad, grad.reset_output_grad, frame_size, active_gate, - origin_mode); + grad.prev_out_grad, grad.reset_output_grad, frame_size, active_gate); } value.gate_value += frame_size * 3; diff --git a/paddle/fluid/operators/math/detail/gru_gpu_kernel.h b/paddle/fluid/operators/math/detail/gru_gpu_kernel.h index 6b57da1046a..af501a6188c 100644 --- a/paddle/fluid/operators/math/detail/gru_gpu_kernel.h +++ b/paddle/fluid/operators/math/detail/gru_gpu_kernel.h @@ -159,7 +159,8 @@ __global__ void KeGruBackwardResetGrad(OpResetGrad op_reset_grad, T *gate_value, T *gate_grad, T *prev_out_value, T *prev_out_grad, T *reset_output_grad, int frame_size, int batch_size, - ActivationType active_gate) { + ActivationType active_gate, + bool origin_mode) { const int frame_idx = blockIdx.x * blockDim.x + threadIdx.x; if (frame_idx >= frame_size) return; int batch_idx = 0; @@ -189,7 +190,7 @@ __global__ void KeGruBackwardResetGrad(OpResetGrad op_reset_grad, T *gate_value, op_reset_grad(&r_update_gate_value, &r_update_gate_grad, &r_reset_gate_value, &r_reset_gate_grad, &r_prev_out_value, &r_prev_out_grad, - &r_reset_output_grad, active_gate); + &r_reset_output_grad, active_gate, origin_mode); gate_grad[frame_idx + frame_size * 0] = r_update_gate_grad; gate_grad[frame_idx + frame_size * 1] = r_reset_gate_grad; diff --git a/paddle/fluid/operators/math/detail/gru_kernel.h b/paddle/fluid/operators/math/detail/gru_kernel.h index c464d9cec4b..894f5f04d24 100644 --- a/paddle/fluid/operators/math/detail/gru_kernel.h +++ b/paddle/fluid/operators/math/detail/gru_kernel.h @@ -163,8 +163,7 @@ class gru_resetGrad { HOSTDEVICE void operator()(T *value_update_gate, T *grad_update_gate, T *value_reset_gate, T *grad_reset_gate, T *value_prev_out, T *grad_prev_out, - T *grad_reset_output, ActivationType act_gate, - bool origin_mode) { + T *grad_reset_output, ActivationType act_gate) { *grad_reset_gate = (*grad_reset_output * (*value_prev_out)); *grad_prev_out += (*grad_reset_output * (*value_reset_gate)); *grad_update_gate = @@ -181,7 +180,7 @@ class gru_resetGrad { __m256 *grad_update_gate, __m256 *value_reset_gate, __m256 *grad_reset_gate, __m256 *value_prev_out, __m256 *grad_prev_out, __m256 *grad_reset_output, - ActivationType act_gate, bool origin_mode) { + ActivationType act_gate) { *grad_reset_gate = _mm256_mul_ps(*grad_reset_output, *value_prev_out); *grad_prev_out = _mm256_add_ps( *grad_prev_out, _mm256_mul_ps(*grad_reset_output, *value_reset_gate)); diff --git a/paddle/fluid/operators/math/gru_compute.cc b/paddle/fluid/operators/math/gru_compute.cc index b875f7d4f4b..07c5cbf3337 100644 --- a/paddle/fluid/operators/math/gru_compute.cc +++ b/paddle/fluid/operators/math/gru_compute.cc @@ -78,8 +78,7 @@ struct GRUUnitGradFunctor { } detail::backward_reset_grad(detail::backward::gru_resetGrad(), value, - grad, frame_size, batch_size, active_gate, - origin_mode); + grad, frame_size, batch_size, active_gate); if (grad.prev_out_grad && value.prev_out_value) { blas.GEMM(false, true, batch_size, frame_size, frame_size * 2, 1, grad.gate_grad, frame_size * 3, value.gate_weight, diff --git a/python/paddle/fluid/layers/nn.py b/python/paddle/fluid/layers/nn.py index 3352ff58315..b78ae63c645 100644 --- a/python/paddle/fluid/layers/nn.py +++ b/python/paddle/fluid/layers/nn.py @@ -885,8 +885,9 @@ def dynamic_gru(input, h_t & = (1-u_t) \odot h_{t-1} + u_t \odot \\tilde{h_t} - if origin_mode is True, then the equation is from paper - `Learning Phrase Representations using RNN Encoder–Decoder for Statistical + + if origin_mode is True then the equation is from paper + Learning Phrase Representations using RNN Encoder-Decoder for Statistical Machine Translation `_ .. math:: @@ -1014,7 +1015,7 @@ def gru_unit(input, **GRU unit layer** if origin_mode is True, then the equation of a gru step is from paper - `Learning Phrase Representations using RNN Encoder–Decoder for Statistical + `Learning Phrase Representations using RNN Encoder-Decoder for Statistical Machine Translation `_ .. math:: -- GitLab