[TE] Bind a rank-0 tensor in decl_tensor_intrin

[leeexyz]

Hi All,

This PR addressed issue 1 in #9829. This can solve the following error.

~/project/apache/tvm/tests/python/unittest$ pytest test_te_schedule.py -k test_tensor_intrin_red
[14:40:10] ../src/target/target_kind.cc:163: Warning: Unable to detect CUDA version, default to "-arch=sm_20" instead
enabled targets: llvm
pytest marker: 
======================================= test session starts ========================================
platform linux -- Python 3.7.5, pytest-6.2.3, py-1.10.0, pluggy-0.13.1
rootdir: ~/project/apache/tvm/tests/python/unittest
collected 18 items / 17 deselected / 1 selected                                                    

test_te_schedule.py F                                                                        [100%]

============================================= FAILURES =============================================
______________________________________ test_tensor_intrin_red ______________________________________

    def test_tensor_intrin_red():
        # not keepdims
        n = 16
        x = te.placeholder((n,), name="x")
        k = te.reduce_axis((0, n), "k")
        z = te.compute((), lambda : te.sum(x[k], axis=k), name="z")
    
        def intrin_func(ins, outs):
            assert isinstance(ins[0], tvm.te.schedule.Buffer)
            assert ins[0].shape[0].value == n
            return tvm.tir.call_packed("vsum", ins[0].data, outs[0].data, ins[0].shape[0])
    
        intrin = te.decl_tensor_intrin(z.op, intrin_func,
            binds={
                x : tvm.tir.decl_buffer(x.shape, x.dtype),
>               z : tvm.tir.decl_buffer(z.shape, z.dtype)
            }
        )

test_te_schedule.py:285: 
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ 
../../../python/tvm/te/tensor_intrin.py:131: in decl_tensor_intrin
    if t in binds
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ 

self = Tensor(shape=[], op.name=z), other = Tensor(shape=[], op.name=z)

    def __eq__(self, other):
        if not isinstance(other, Tensor):
            if isinstance(other, _expr.ExprOp):
                return _expr.EqualOp(self, other)
            return False
        if self.ndim == 0 and other.ndim == 0:
            raise ValueError(
>               "Equal == comparison among rank-0 tensor is ambiguous, "
                "use Tensor.equal for content expression equvalence, "
                "use Tensor.same_as for exact reference comparison"
            )
E           ValueError: Equal == comparison among rank-0 tensor is ambiguous, use Tensor.equal for content expression equvalence, use Tensor.same_as for exact reference comparison

../../../python/tvm/te/tensor.py:91: ValueError
===================================== short test summary info ======================================
FAILED test_te_schedule.py::test_tensor_intrin_red - ValueError: Equal == comparison among rank-0...
================================= 1 failed, 17 deselected in 0.09s =================================

[junrushao]

CC: @Hzfengsy

[Hzfengsy]

I'm not sure. Can you explain why "comparison among rank-0 tensor is ambiguous"?

[leeexyz]

I'm not sure. Can you explain why "comparison among rank-0 tensor is ambiguous"?

I guess here, the tensor.equal(other) and tensor.same_as(other) is not suitable for non rank-0 tensors, so it adds the check.

1. Reference check makes no sense because the address of the tensor is changed every time when invoking op.output(id).
2. Content checks have some constraints, such as Tensor dimension needs to be matched.

All these two ways cannot be put into the __eq__ method. But they may make sense for rank-0 tensors. A possible fix for this would be simply to delete it.

    def __eq__(self, other):
        if not isinstance(other, Tensor):
            if isinstance(other, _expr.ExprOp):
                return _expr.EqualOp(self, other)
            return False
        return _ffi_api.TensorEqual(self, other)

[areusch]

@Hzfengsy @leeexyz what's the next step with this one?