Skip to content

Commit 10f9644

Browse files
zh794390558zh794390558
authored
Align CTC grad scale same with ESPNet (PaddlePaddle#34729)
* dygraph support more ctc grad scale * scale for 1.x * fix unitest * fix unitest * format code * fix unittest * fix log info * unittest cov * fix format;notest,test=cpu,coverage * skip ctc_loss egs;test=cpu * warpctc grad cov;test=coverage * add dygraph test;test=coverage * format;test=cpu,coverage * format;test=cpu * add api compat;test=cpu * add cpu test * rename * rename * fix * fix test * format * eigen cpu * eigen gpu grad pass * cuda gpu pass * format * fix ci
1 parent 8046e33 commit 10f9644

File tree

15 files changed

+602
-104
lines changed

15 files changed

+602
-104
lines changed

‎paddle/fluid/operators/CMakeLists.txt‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -81,10 +81,10 @@ op_library(run_program_op SRCS run_program_op.cc run_program_op.cu.cc DEPS execu
8181

8282
if (WITH_GPU OR WITH_ROCM)
8383
if(WITH_ROCM)
84-
op_library(warpctc_op DEPS dynload_warpctc sequence_padding sequence_scale SRCS warpctc_op.cc warpctc_op.cu.cc)
84+
op_library(warpctc_op DEPS dynload_warpctc sequence_padding sequence_scale SRCS warpctc_op.cc warpctc_op.cu)
8585
# warpctc_op needs cudnn 7 above
8686
elseif(${CUDNN_MAJOR_VERSION} VERSION_LESS 7)
87-
op_library(warpctc_op DEPS dynload_warpctc sequence_padding sequence_scale SRCS warpctc_op.cc warpctc_op.cu.cc)
87+
op_library(warpctc_op DEPS dynload_warpctc sequence_padding sequence_scale SRCS warpctc_op.cc warpctc_op.cu)
8888
else()
8989
op_library(warpctc_op DEPS dynload_warpctc sequence_padding sequence_scale)
9090
endif()

‎paddle/fluid/operators/math/sequence_padding.cc‎

Lines changed: 27 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,8 @@ void CopyValidData(framework::Tensor* dst_tensor,
3333
const framework::Tensor* src_tensor,
3434
const framework::Vector<size_t>& seq_offsets,
3535
int pad_seq_len, int step_width, bool norm_by_len,
36-
CopyType type, PadLayout layout) {
36+
bool norm_by_batchsize, bool norm_by_total_logits_len,
37+
int total_logits_len, CopyType type, PadLayout layout) {
3738
int seq_num = seq_offsets.size() - 1;
3839
const T* src_data = src_tensor->data<T>();
3940
T* dst_data = dst_tensor->data<T>();
@@ -54,7 +55,21 @@ void CopyValidData(framework::Tensor* dst_tensor,
5455
int pad_data_offset = layout == kBatchLengthWidth
5556
? seq_idx * pad_seq_len * step_width
5657
: seq_idx * step_width;
57-
float scale = 1.0f / static_cast<float>(valid_seq_len);
58+
59+
float scale = 1.0f;
60+
if (norm_by_total_logits_len) {
61+
scale = 1.0f / static_cast<float>(total_logits_len);
62+
VLOG(3) << "[warpctc grad][norm_by_total_logits_len]: scale " << scale
63+
<< "total_logits_len " << total_logits_len;
64+
} else if (norm_by_batchsize) {
65+
scale = 1.0f / static_cast<float>(seq_num);
66+
VLOG(3) << "[warpctc grad][norm_by_batchsize]: scale " << scale << "B "
67+
<< seq_num;
68+
} else if (norm_by_len) {
69+
scale = 1.0f / static_cast<float>(valid_seq_len);
70+
VLOG(3) << "[warpctc grad][norm_by_len]: scale " << scale << "T "
71+
<< valid_seq_len;
72+
}
5873

5974
for (int step_idx = 0; step_idx < valid_seq_len; ++step_idx) {
6075
const T* src =
@@ -97,6 +112,8 @@ class PaddingLoDTensorFunctor<platform::CPUDeviceContext, T> {
97112
framework::LoDTensor* pad_tensor,
98113
const framework::LoDTensor& pad_value, int pad_seq_len = -1,
99114
int lod_level = 0, bool norm_by_times = false,
115+
bool norm_by_batchsize = false,
116+
bool norm_by_total_logits_len = false,
100117
const PadLayout layout = kBatchLengthWidth) {
101118
auto seq_lod = seq_tensor.lod();
102119
const auto seq_offsets = framework::ToAbsOffset(seq_lod)[lod_level];
@@ -131,7 +148,8 @@ class PaddingLoDTensorFunctor<platform::CPUDeviceContext, T> {
131148
}
132149

133150
CopyValidData<T>(pad_tensor, &seq_tensor, seq_offsets, pad_seq_len,
134-
step_width, norm_by_times, kSeqToPad, layout);
151+
step_width, norm_by_times, false, false, 0, kSeqToPad,
152+
layout);
135153
}
136154
};
137155

@@ -142,20 +160,25 @@ class UnpaddingLoDTensorFunctor<platform::CPUDeviceContext, T> {
142160
const framework::LoDTensor& pad_tensor,
143161
framework::LoDTensor* seq_tensor, int pad_seq_len = -1,
144162
int lod_level = 0, bool norm_by_times = false,
163+
bool norm_by_batchsize = false,
164+
bool norm_by_total_logits_len = false,
145165
const PadLayout layout = kBatchLengthWidth) {
146166
auto seq_offsets = framework::ToAbsOffset(seq_tensor->lod())[lod_level];
147167
const auto& seq_tensor_dims = seq_tensor->dims();
148168
const auto& pad_tensor_dims = pad_tensor.dims();
149169
if (pad_seq_len == -1) {
150170
pad_seq_len = MaximumSequenceLength(seq_offsets);
151171
}
172+
int total_logits_len = TotalSequenceLength(seq_offsets);
152173
int step_width = seq_tensor->numel() / seq_tensor_dims[0];
153174

154175
CheckDims(seq_tensor_dims, pad_tensor_dims, seq_offsets, pad_seq_len,
155176
step_width, layout);
156177

157178
CopyValidData<T>(seq_tensor, &pad_tensor, seq_offsets, pad_seq_len,
158-
step_width, norm_by_times, kPadToSeq, layout);
179+
step_width, norm_by_times, norm_by_batchsize,
180+
norm_by_total_logits_len, total_logits_len, kPadToSeq,
181+
layout);
159182
}
160183
};
161184

‎paddle/fluid/operators/math/sequence_padding.cu‎

Lines changed: 20 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,9 @@ template <typename T, CopyType Type>
2323
__global__ void SequencePaddingKernel(
2424
T* dst, const T* src, const T* pad_value, bool is_constant_pad,
2525
const size_t* seq_offsets, const size_t seq_num, const size_t pad_seq_len,
26-
const size_t step_width, bool norm_by_len, const PadLayout layout) {
26+
const size_t step_width, bool norm_by_len, bool norm_by_batchsize,
27+
bool norm_by_total_logits_len, int total_logits_len,
28+
const PadLayout layout) {
2729
size_t seq_idx = blockIdx.y;
2830
size_t seq_len = seq_offsets[seq_idx + 1] - seq_offsets[seq_idx];
2931

@@ -38,7 +40,15 @@ __global__ void SequencePaddingKernel(
3840
src + (Type == kSeqToPad ? seq_data_offset : pad_data_offset);
3941

4042
if (step_idx < seq_len) {
41-
float scale = norm_by_len ? (1.0f / static_cast<float>(seq_len)) : 1.0f;
43+
float scale = 1.0f;
44+
if (norm_by_total_logits_len) {
45+
scale = 1.0f / static_cast<float>(total_logits_len);
46+
} else if (norm_by_batchsize) {
47+
scale = 1.0f / static_cast<float>(seq_num);
48+
} else if (norm_by_len) {
49+
scale = norm_by_len ? (1.0f / static_cast<float>(seq_len)) : 1.0f;
50+
}
51+
4252
for (size_t i = threadIdx.x; i < step_width; i += blockDim.x) {
4353
dst_data[i] = scale * src_data[i];
4454
}
@@ -57,6 +67,8 @@ class PaddingLoDTensorFunctor<platform::CUDADeviceContext, T> {
5767
framework::LoDTensor* pad_tensor,
5868
const framework::LoDTensor& pad_value, int pad_seq_len = -1,
5969
int lod_level = 0, bool norm_by_times = false,
70+
bool norm_by_batchsize = false,
71+
bool norm_by_total_logits_len = false,
6072
const PadLayout layout = kBatchLengthWidth) {
6173
auto seq_lod = seq_tensor.lod();
6274
const auto seq_offsets = framework::ToAbsOffset(seq_lod)[lod_level];
@@ -107,7 +119,7 @@ class PaddingLoDTensorFunctor<platform::CUDADeviceContext, T> {
107119
SequencePaddingKernel<T, kSeqToPad><<<grid, threads, 0, context.stream()>>>(
108120
pad_data, seq_data, pad_value_data, pad_value.numel() == 1,
109121
seq_offsets.CUDAData(context.GetPlace()), seq_num, pad_seq_len,
110-
step_width, norm_by_times, layout);
122+
step_width, norm_by_times, false, false, 0, layout);
111123
}
112124
};
113125

@@ -118,6 +130,8 @@ class UnpaddingLoDTensorFunctor<platform::CUDADeviceContext, T> {
118130
const framework::LoDTensor& pad_tensor,
119131
framework::LoDTensor* seq_tensor, int pad_seq_len = -1,
120132
int lod_level = 0, bool norm_by_times = false,
133+
bool norm_by_batchsize = false,
134+
bool norm_by_total_logits_len = false,
121135
const PadLayout layout = kBatchLengthWidth) {
122136
auto seq_offsets = framework::ToAbsOffset(seq_tensor->lod())[lod_level];
123137
const auto& seq_tensor_dims = seq_tensor->dims();
@@ -126,6 +140,7 @@ class UnpaddingLoDTensorFunctor<platform::CUDADeviceContext, T> {
126140
if (pad_seq_len == -1) {
127141
pad_seq_len = max_seq_len;
128142
}
143+
int total_logits_len = TotalSequenceLength(seq_offsets);
129144
int step_width = seq_tensor->numel() / seq_tensor_dims[0];
130145
int seq_num = seq_offsets.size() - 1;
131146

@@ -159,7 +174,8 @@ class UnpaddingLoDTensorFunctor<platform::CUDADeviceContext, T> {
159174
SequencePaddingKernel<T, kPadToSeq><<<grid, threads, 0, context.stream()>>>(
160175
seq_data, pad_data, nullptr, false,
161176
seq_offsets.CUDAData(context.GetPlace()), seq_num, pad_seq_len,
162-
step_width, norm_by_times, layout);
177+
step_width, norm_by_times, norm_by_batchsize, norm_by_total_logits_len,
178+
total_logits_len, layout);
163179
}
164180
};
165181

‎paddle/fluid/operators/math/sequence_padding.h‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -107,6 +107,8 @@ class PaddingLoDTensorFunctor {
107107
framework::LoDTensor* pad_tensor,
108108
const framework::LoDTensor& pad_value, int pad_seq_len = -1,
109109
int lod_level = 0, bool norm_by_times = false,
110+
bool norm_by_batchsize = false,
111+
bool norm_by_total_logits_len = false,
110112
const PadLayout layout = kBatchLengthWidth);
111113
};
112114

@@ -117,6 +119,8 @@ class UnpaddingLoDTensorFunctor {
117119
const framework::LoDTensor& pad_tensor,
118120
framework::LoDTensor* seq_tensor, int pad_seq_len = -1,
119121
int lod_level = 0, bool norm_by_times = false,
122+
bool norm_by_batchsize = false,
123+
bool norm_by_total_logits_len = false,
120124
const PadLayout layout = kBatchLengthWidth);
121125
};
122126

‎paddle/fluid/operators/math/sequence_padding_test.cc‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -66,13 +66,13 @@ void TestSequencePadding(const DeviceContext &context,
6666
}
6767

6868
paddle::operators::math::PaddingLoDTensorFunctor<DeviceContext, T>()(
69-
context, seq, &padding, pad_value, -1, 0, false,
69+
context, seq, &padding, pad_value, -1, 0, false, false, false,
7070
paddle::operators::math::kLengthBatchWidth);
7171

7272
seq_back.set_lod(lod);
7373
seq_back.mutable_data<T>(seq_dims, place);
7474
paddle::operators::math::UnpaddingLoDTensorFunctor<DeviceContext, T>()(
75-
context, padding, &seq_back, -1, 0, false,
75+
context, padding, &seq_back, -1, 0, false, false, false,
7676
paddle::operators::math::kLengthBatchWidth);
7777

7878
if (paddle::platform::is_cpu_place(place)) {

‎paddle/fluid/operators/sequence_ops/sequence_pad_op.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -46,7 +46,7 @@ class SequencePadOpKernel : public framework::OpKernel<T> {
4646

4747
math::PaddingLoDTensorFunctor<DeviceContext, T>()(
4848
ctx.template device_context<DeviceContext>(), *x, out, *pad_value,
49-
padded_length, 0, false, math::kBatchLengthWidth);
49+
padded_length, 0, false, false, false, math::kBatchLengthWidth);
5050

5151
LoDTensor seq_len;
5252
seq_len.Resize(len_t->dims());
@@ -72,7 +72,7 @@ class SequencePadGradOpKernel : public framework::OpKernel<T> {
7272

7373
math::UnpaddingLoDTensorFunctor<DeviceContext, T>()(
7474
ctx.template device_context<DeviceContext>(), *d_out, d_x,
75-
padded_length, 0, false, math::kBatchLengthWidth);
75+
padded_length, 0, false, false, false, math::kBatchLengthWidth);
7676
}
7777
}
7878
};

‎paddle/fluid/operators/sequence_ops/sequence_unpad_op.h‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -69,7 +69,8 @@ class SequenceUnpadOpKernel : public framework::OpKernel<T> {
6969

7070
int64_t padded_length = x_t->dims()[1];
7171
math::UnpaddingLoDTensorFunctor<DeviceContext, T>()(
72-
dev_ctx, *x_t, out_t, padded_length, 0, false, math::kBatchLengthWidth);
72+
dev_ctx, *x_t, out_t, padded_length, 0, false, false, false,
73+
math::kBatchLengthWidth);
7374
}
7475
};
7576

@@ -93,7 +94,7 @@ class SequenceUnpadGradOpKernel : public framework::OpKernel<T> {
9394

9495
math::PaddingLoDTensorFunctor<DeviceContext, T>()(
9596
ctx.template device_context<DeviceContext>(), *d_out, d_x, zero_pads,
96-
padded_length, 0, false, math::kBatchLengthWidth);
97+
padded_length, 0, false, false, false, math::kBatchLengthWidth);
9798
}
9899
}
99100
};

‎paddle/fluid/operators/warpctc_op.cc‎

Lines changed: 29 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -125,6 +125,17 @@ class WarpCTCOpMaker : public framework::OpProtoAndCheckerMaker {
125125
"normalize the gradients by the number of time-step, "
126126
"which is also the sequence's length.")
127127
.SetDefault(false);
128+
AddAttr<bool>(
129+
"norm_by_batchsize",
130+
"(bool, default: false), normalize the loss by the batch size."
131+
"If True, supersedes norm_by_times")
132+
.SetDefault(false);
133+
AddAttr<bool>(
134+
"norm_by_total_logits_len",
135+
"(bool, default: false), normalize the loss by the total number of "
136+
"frames"
137+
"in the batch. If True, supersedes norm_by_batchsize and norm_by_times")
138+
.SetDefault(false);
128139
AddComment(R"DOC(
129140
An operator integrating the open-source
130141
[warp-ctc](https://github.com/baidu-research/warp-ctc) library, which is used in
@@ -206,3 +217,21 @@ REGISTER_OP_CPU_KERNEL(
206217
warpctc_grad,
207218
ops::WarpCTCGradKernel<paddle::platform::CPUDeviceContext, float>,
208219
ops::WarpCTCGradKernel<paddle::platform::CPUDeviceContext, double>);
220+
221+
REGISTER_OP_VERSION(warpctc)
222+
.AddCheckpoint(
223+
R"ROC(
224+
Upgrade warpctc add a new attribute [norm_by_batchsize] and [norm_by_total_logits_len])ROC",
225+
paddle::framework::compatible::OpVersionDesc()
226+
.NewAttr(
227+
"norm_by_batchsize",
228+
"(bool, default: false), normalize the loss by the batch size."
229+
"If True, supersedes norm_by_times",
230+
false)
231+
.NewAttr("norm_by_total_logits_len",
232+
"(bool, default: false), normalize the loss by the total "
233+
"number of "
234+
"frames"
235+
"in the batch. If True, supersedes norm_by_batchsize and "
236+
"norm_by_times",
237+
false));

0 commit comments

Comments
 (0)