Skip to content
This repository has been archived by the owner on Nov 17, 2023. It is now read-only.

[MXNET-807] Support integer label type in ctc_loss operator #12468

Merged
merged 17 commits into from
Sep 12, 2018
Merged

[MXNET-807] Support integer label type in ctc_loss operator #12468

Show file tree
Hide file tree
Changes from 11 commits
Commits
Show all changes
17 commits
Select commit Hold shift + click to select a range
f0a757b
Support integer type in ctc_loss
apeforest Sep 5, 2018
7af7274
Support any data type in ctc_loss operator
apeforest Sep 5, 2018
5e99e7e
Enable integer type in labels and fix lint errors
apeforest Sep 5, 2018
eb30964
Fix compilation error in GPU
apeforest Sep 5, 2018
774c61b
Add unit tests
apeforest Sep 6, 2018
d9dc6e6
Merge remote-tracking branch 'upstream/master' into bugfix/ctc_loss_i…
apeforest Sep 6, 2018
59f48f2
Undo indentation
apeforest Sep 6, 2018
1b3d141
Undo blank line
apeforest Sep 6, 2018
299b1e7
Undo blank line
apeforest Sep 6, 2018
ec5cc3c
Add unit test for large number of classes
apeforest Sep 6, 2018
59e5d7c
Merge branch 'bugfix/ctc_loss_integer' of https://github.com/apefores…
apeforest Sep 6, 2018
c8b7cd4
move unit tests to test_operator.py per reviewer advice
apeforest Sep 6, 2018
973daca
update unit test
apeforest Sep 6, 2018
217069e
update unit test
apeforest Sep 6, 2018
4574c7c
update unit test using random seed
apeforest Sep 6, 2018
fa61a0a
Update unit test
apeforest Sep 8, 2018
3fbb3f5
Fix unit test difference Python2 and Python3
apeforest Sep 11, 2018
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
152 changes: 84 additions & 68 deletions src/operator/contrib/ctc_loss-inl.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -256,66 +256,69 @@ class CTCLossOp : public Operator {
exceed_cudnn_limit = false;
Stream<xpu> *s = ctx.get_stream<xpu>();

Tensor<xpu, 3, real_t> data =
MSHADOW_TYPE_SWITCH(in_data[ctc_loss::kLabel].type_flag_, DType, {
Tensor<xpu, 3, real_t> data =
in_data[ctc_loss::kData].get<xpu, 3, real_t>(s);
Tensor<xpu, 2, real_t> labels =
in_data[ctc_loss::kLabel].get<xpu, 2, real_t>(s);
Tensor<xpu, 2, DType> labels =
in_data[ctc_loss::kLabel].get<xpu, 2, DType>(s);

Tensor<xpu, 1, real_t> costs =
Tensor<xpu, 1, real_t> costs =
out_data[ctc_loss::kOut].get<xpu, 1, real_t>(s);
Tensor<xpu, 3, real_t> grad =
Tensor<xpu, 3, real_t> grad =
out_data[ctc_loss::kGrad].get<xpu, 3, real_t>(s);

int max_seq_len = data.size(0);
int batch_size = data.size(1);
int alphabet_size = data.size(2);

// data_lengths
std::vector<int> data_lengths(batch_size, max_seq_len);
if (param_.use_data_lengths) {
int kInputLength = 2;
IndexTensorToVector(in_data[kInputLength].get<xpu, 1, real_t>(s), &data_lengths);
}

// label_lengths
std::vector<int> packed_labels;
std::vector<int> label_lengths(batch_size);

if (param_.use_label_lengths) {
int kLabelLength = 2+param_.use_data_lengths;
exceed_cudnn_limit = PackLabelByLength(labels, in_data[kLabelLength].get<xpu, 1, real_t>(s),
&packed_labels, &label_lengths);
} else {
exceed_cudnn_limit = LabelTensorToPackedVector(labels, param_.blank_label == 0?0:-1,
&packed_labels, &label_lengths);
}

// CUDNN is disabled due to lack of support for input lengths
/* #if defined(__CUDACC__) && MXNET_USE_CUDNN == 1 && CUDNN_MAJOR >= 7 */
/* if (!exceed_cudnn_limit) { */
/* cudnn_forward(ctx, s, data, costs, grad, */
/* &data_lengths, &label_lengths, &packed_labels, */
/* max_seq_len, batch_size, alphabet_size, */
/* req[ctc_loss::kGrad] != mxnet::kNullOp); */
/* } else { */
/* baidu_forward(ctx, s, data, costs, grad, */
/* &data_lengths, &label_lengths, &packed_labels, */
/* batch_size, alphabet_size, req[ctc_loss::kGrad] != mxnet::kNullOp); */
/* } */
/* #else */

baidu_forward(ctx, s, data, costs, grad,
&data_lengths, &label_lengths, &packed_labels,
batch_size, alphabet_size, req[ctc_loss::kGrad] != mxnet::kNullOp);

if (param_.use_data_lengths) {
// baidu warp CTC implementation sometimes includes undefined gradients
// for data outside of length mask. Setting to 0 to make it consistent
// with CPU implementation.
int kInputLength = 2;
mxnet_op::SequenceMask(grad, in_data[kInputLength].get<xpu, 1, real_t>(s),
static_cast<real_t>(0));
}
int max_seq_len = data.size(0);
int batch_size = data.size(1);
int alphabet_size = data.size(2);

// data_lengths
std::vector<int> data_lengths(batch_size, max_seq_len);
if (param_.use_data_lengths) {
int kInputLength = 2;
IndexTensorToVector(in_data[kInputLength].get<xpu, 1, real_t>(s), &data_lengths);
}

// label_lengths
std::vector<int> packed_labels;
std::vector<int> label_lengths(batch_size);

if (param_.use_label_lengths) {
int kLabelLength = 2 + param_.use_data_lengths;
exceed_cudnn_limit =
PackLabelByLength(labels, in_data[kLabelLength].get<xpu, 1, DType>(s),
&packed_labels, &label_lengths);
} else {
exceed_cudnn_limit = LabelTensorToPackedVector(labels, param_.blank_label == 0 ? 0 : -1,
&packed_labels, &label_lengths);
}

// CUDNN is disabled due to lack of support for input lengths
/* #if defined(__CUDACC__) && MXNET_USE_CUDNN == 1 && CUDNN_MAJOR >= 7 */
/* if (!exceed_cudnn_limit) { */
/* cudnn_forward(ctx, s, data, costs, grad, */
/* &data_lengths, &label_lengths, &packed_labels, */
/* max_seq_len, batch_size, alphabet_size, */
/* req[ctc_loss::kGrad] != mxnet::kNullOp); */
/* } else { */
/* baidu_forward(ctx, s, data, costs, grad, */
/* &data_lengths, &label_lengths, &packed_labels, */
/* batch_size, alphabet_size, req[ctc_loss::kGrad] != mxnet::kNullOp);*/
/* } */
/* #else */

baidu_forward(ctx, s, data, costs, grad,
&data_lengths, &label_lengths, &packed_labels,
batch_size, alphabet_size, req[ctc_loss::kGrad] != mxnet::kNullOp);

if (param_.use_data_lengths) {
// baidu warp CTC implementation sometimes includes undefined gradients
// for data outside of length mask. Setting to 0 to make it consistent
// with CPU implementation.
int kInputLength = 2;
mxnet_op::SequenceMask(grad, in_data[kInputLength].get<xpu, 1, real_t>(s),
static_cast<real_t>(0));
}
});
}

virtual void Backward(const OpContext &ctx,
Expand Down Expand Up @@ -434,17 +437,17 @@ class CTCLossOp : public Operator {
}
#endif // __CUDACC__ && CUDNN

inline virtual void baidu_forward(const OpContext &ctx,
mshadow::Stream<xpu>* s,
mshadow::Tensor<xpu, 3, real_t> data,
mshadow::Tensor<xpu, 1, real_t> costs,
mshadow::Tensor<xpu, 3, real_t> grad,
std::vector<int>* data_lengths,
std::vector<int>* label_lengths,
std::vector<int>* packed_labels,
int batch_size,
int alphabet_size,
bool req_grad) {
inline void baidu_forward(const OpContext &ctx,
mshadow::Stream<xpu>* s,
mshadow::Tensor<xpu, 3, real_t> data,
mshadow::Tensor<xpu, 1, real_t> costs,
mshadow::Tensor<xpu, 3, real_t> grad,
std::vector<int>* data_lengths,
std::vector<int>* label_lengths,
std::vector<int>* packed_labels,
int batch_size,
int alphabet_size,
bool req_grad) {
using namespace mshadow;
// allocate temporary workspace
size_t size_bytes;
Expand All @@ -461,7 +464,7 @@ class CTCLossOp : public Operator {
compute_ctc_cost(data, costs.dptr_, grad.dptr_, packed_labels->data(),
label_lengths->data(), data_lengths->data(),
workspace.dptr_, req_grad,
param_.blank_label == 0?0:(alphabet_size-1));
param_.blank_label == 0 ? 0 : (alphabet_size-1));
}
}; // class CTCLossOp

Expand Down Expand Up @@ -534,11 +537,24 @@ class CTCLossProp : public OperatorProperty {
TShape oshape(1);
oshape[0] = dshape[1]; // batch size
out_shape->clear();
out_shape->push_back(oshape);
out_shape->push_back(oshape); // forward output
out_shape->push_back(dshape); // grad output
return true;
}

bool InferType(std::vector<int> *in_type,
std::vector<int> *out_type,
std::vector<int> *aux_type) const override {
CHECK_LE(in_type->size(), this->ListArguments().size());
int dtype = (*in_type)[ctc_loss::kData];
CHECK_NE(dtype, -1) << "Input data must have specified type";

out_type->clear();
out_type->push_back(dtype); // forward output
out_type->push_back(dtype); // grad output
return true;
}

OperatorProperty *Copy() const override {
auto ptr = new CTCLossProp();
ptr->param_ = param_;
Expand Down
58 changes: 58 additions & 0 deletions tests/python/unittest/test_contrib_operator.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -244,6 +244,64 @@ def assert_match(inputs, x, y, threshold, is_ascend=False):
assert_match([[0.5, 0.6], [0.1, 0.2], [0.3, 0.4]], [1, -1, 0], [2, 0], 1e-12, False)
assert_match([[0.5, 0.6], [0.1, 0.2], [0.3, 0.4]], [-1, 0, 1], [1, 2], 100, True)

def test_ctc_loss_op():
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

CTC loss tests can be found at https://github.com/apache/incubator-mxnet/blob/master/tests/python/unittest/test_operator.py#L4500, and integration at https://github.com/apache/incubator-mxnet/blob/master/tests/python/unittest/test_loss.py#L186. Test cases are from hand calculated examples.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Feel free to add test cases for large labels there.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for the reference. Moving the unit test there.

batch_size = 10
seq_len = 5
label_len = 3
num_classes = 6
np.random.seed(1)
x = np.random.uniform(size=(seq_len, batch_size, num_classes))
y = np.random.randint(0, num_classes, size=(batch_size, label_len))

def test_cpu(x, y):
data = mx.nd.array(x, ctx=mx.cpu(0))
label = mx.nd.array(y, ctx=mx.cpu(0))
loss = mx.nd.contrib.ctc_loss(data=data, label=label)
loss = mx.nd.make_loss(loss)
expected_output = [9.604521, 7.096151, 4.906869, 5.5237527, 5.9895644, 5.584548,
5.528411, 5.765914, 6.740701, 5.2625823]
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This testing strategy (i.e. compare the output from random input and labels with fixed seed from recorded output) is not meaningful and does not guarantee anything. It merely increases the line coverage.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Did not notice the unit test in test_operator.py. I have removed this one.

assert np.isclose(loss.asnumpy(), expected_output).all()

def test_gpu(x, y):
data = mx.nd.array(x, ctx=mx.gpu(0))
label = mx.nd.array(y, ctx=mx.gpu(0))
loss = mx.nd.contrib.ctc_loss(data=data, label=label)
loss = mx.nd.make_loss(loss)
expected_output = [9.604521, 7.096151, 4.906869, 5.5237527, 5.9895644, 5.584548,
5.528411, 5.765914, 6.740701, 5.2625823]
assert np.isclose(loss.asnumpy(), expected_output).all()

def test_integer_label(x, y):
data = mx.nd.array(x, ctx=mx.cpu(0))
label = mx.nd.array(y, ctx=mx.cpu(0), dtype=np.int32)
loss = mx.nd.contrib.ctc_loss(data=data, label=label)
loss = mx.nd.make_loss(loss)
expected_output = [9.604521, 7.096151, 4.906869, 5.5237527, 5.9895644, 5.584548,
5.528411, 5.765914, 6.740701, 5.2625823]
assert np.isclose(loss.asnumpy(), expected_output).all()

def test_large_classes():
batch_size = 1024
seq_len = 35
label_len = 10
num_classes = 6000
x = np.random.uniform(size=(seq_len, batch_size, num_classes))
y = np.random.randint(0, num_classes, size=(batch_size, label_len))

data = mx.nd.array(x, ctx=mx.gpu(0))
label = mx.nd.array(y, ctx=mx.gpu(0))
loss = mx.nd.contrib.ctc_loss(data=data, label=label)
loss = mx.nd.make_loss(loss)
expected_output_sum = 282733.95318603516
assert np.isclose(sum(loss.asnumpy(), expected_output_sum))

test_cpu(x, y)
test_integer_label(x, y)
if default_context().device_type == 'gpu':
test_gpu(x, y)
test_large_classes()


if __name__ == '__main__':
import nose
nose.runmodule()