[ONNX] Add CumSum operator to ONNX frontend

[echuraev]

Thanks for contributing to TVM! Please refer to guideline https://tvm.apache.org/docs/contribute/ for useful information and tips. After the pull request is submitted, please request code reviews from Reviewers by @ them in the pull request thread.

[mbrookhart]

Fix lint?

I'm a little confused, I'm looking at the ONNX spec, and I only see 2 inputs, not three. What is the third input you're using for datatype?
https://github.com/onnx/onnx/blob/master/docs/Operators.md#cumsum

We should probably either handle or error when the exclusive or reverse attributes are true.

[mbrookhart]

Oh, I just noticed it's a draft :) You're probably not ready for review, sorry.

[echuraev]

@mbrookhart, thank you for you questions. Now I added error generation for exclusive and reverse attributes, but I'm planning to implement them in this PR.

[masahi]

exclusive can be easily supported by cumsum(data) - data. Similarly, I expect reverse can also be supported by composing with reverse op. But I think this attribute is weird (why would anyone one want to do reverse cumsum), so it might not be worth it.

It's strange ONNX Cumsum doesn't have output dtype attribute. Without it, there could be overflow issues for int32 input. Also it wouldn't work for boolean input, which is a common input to cumsum. I guess they expect casting to be done by users beforehand.

[mbrookhart]

Looking very close.

It looks like you have a bit of a hybrid approach with the two options, one you're supporting in the relay constructor, the other you support in the TOPI CPU kernel, but not the GPU kernel?

I wonder if it makes more sense to do the reverse/exclusive changes explicitly in the importer so we don't have to pipe it through the multiple layers? I'm not sure if support the attributes in topi or doing it via composition in the importer makes more sense.

[masahi]

I think reverse should be done in the frontend, since this seems ONNX specific. For exclusive, if you want to avoid the overhead of subtraction, we can do cumsum with exclusive scan. See

tvm/python/tvm/topi/cuda/scan.py

Line 517 in 1e0d356

return inclusive_scan(data, axis, output_dtype=dtype, binop=tvm.tir.generic.add)

if exclusive=True, we can use exclusive_scan instead of inclusive_scan.

[echuraev]

@mbrookhart, @masahi, I have one question about implementation of exclusive attribute. In current implementation, I added this attribute to topi kernels, but also I could do the same thing only for ONNX frontend (by adding subtract operation after cumsum). What do you think, what is the better approach? In first case we don't have overhead of subtract operation, but in the second we should write less code.

[masahi]

I'm happy with the current implementation. It is faster, and exclusive cumsum could be useful for other purpose (just as exclusive_scan in topi cuda is used in multiple places).

[mbrookhart]

Thanks @echuraev @masahi