[TVMScript] Encourage using T.Buffer directly

Previously there are two equivalent ways of declaring a buffer in
TVMScript:

```python
buffer = T.buffer_decl(...)
buffer = T.Buffer(...)
```

The two approaches are aliases to each other and are essentially the
same in implementation. Therefore, this PR encourages to use `T.Buffer`
as the recommended approach as it's a bit shorter. Meanwhile,
`T.buffer_decl` will continue to be valid in TVMScript, but a
deprecation warning will be emitted if its used.

Author: junrushao

python/tvm/script/ir_builder/tir/__init__.py

@@ -17,4 +17,4 @@
 """Package tvm.script.ir_builder.tir"""
 from .ir import *  # pylint: disable=wildcard-import,redefined-builtin
 from .ir import boolean as bool  # pylint: disable=redefined-builtin
-from .ir import buffer_decl as Buffer
+from .ir import buffer as Buffer

python/tvm/script/ir_builder/tir/ir.py

@@ -86,7 +86,7 @@
 # pylint: enable=unused-import
 
 
-def buffer_decl(
+def buffer(
     shape: Union[List[PrimExpr], Tuple[PrimExpr], PrimExpr, Integral],
     dtype: str = "float32",
     data: Var = None,
@@ -138,7 +138,7 @@ def buffer_decl(
         The declared buffer.
     """
     shape = (shape,) if isinstance(shape, (PrimExpr, Integral)) else shape
-    return _ffi_api.BufferDecl(  # type: ignore[attr-defined] # pylint: disable=no-member
+    return _ffi_api.Buffer(  # type: ignore[attr-defined] # pylint: disable=no-member
         shape,
         dtype,
         "",
@@ -153,6 +153,11 @@ def buffer_decl(
     )
 
 
+@deprecated("T.buffer_decl(...)", "T.Buffer(...)")
+def buffer_decl(*args, **kwargs):
+    return buffer(*args, **kwargs)
+
+
 def prim_func() -> frame.PrimFuncFrame:
     """The primitive function statement.
 
@@ -1431,7 +1436,7 @@ def ptr(dtype: str, storage_scope: str = "global") -> Var:
     return _ffi_api.Ptr(dtype, storage_scope)  # type: ignore[attr-defined] # pylint: disable=no-member
 
 
-@deprecated("T.buffer_var", "T.Ptr")
+@deprecated("T.buffer_var", "T.handle")
 def buffer_var(dtype: str, storage_scope: str = "global") -> Var:
     """The pointer declaration function.
 
@@ -1814,6 +1819,7 @@ def wrapped(*args, **kwargs):
     "float16x64",
     "float32x64",
     "float64x64",
+    "buffer",
     "buffer_decl",
     "prim_func",
     "arg",

python/tvm/script/parser/tir/entry.py

@@ -21,7 +21,7 @@
 from tvm.ir.base import deprecated
 from tvm.tir import Buffer, PrimFunc
 
-from ...ir_builder.tir import buffer_decl, ptr
+from ...ir_builder.tir import buffer, ptr
 from .._core import parse, utils
 
 
@@ -66,7 +66,7 @@ def __call__(
         buffer_type="",
         axis_separators=None,
     ) -> Buffer:
-        return buffer_decl(
+        return buffer(
             shape,
             dtype=dtype,
             data=data,

src/script/ir_builder/tir/ir.cc

@@ -593,8 +593,7 @@ TVM_STATIC_IR_FUNCTOR(Namer, vtable)
       Namer::Name(var->var, name);
     });
 
-TVM_REGISTER_GLOBAL("script.ir_builder.tir.BufferDecl").set_body_typed(BufferDecl);
-
+TVM_REGISTER_GLOBAL("script.ir_builder.tir.Buffer").set_body_typed(BufferDecl);
 TVM_REGISTER_GLOBAL("script.ir_builder.tir.PrimFunc").set_body_typed(PrimFunc);
 TVM_REGISTER_GLOBAL("script.ir_builder.tir.Arg")
     .set_body_typed([](String name, ObjectRef obj) -> ObjectRef {

tests/python/unittest/test_auto_scheduler_feature.py

@@ -21,9 +21,8 @@
 import tempfile
 
 import tvm
-from tvm import te, auto_scheduler, relay
+from tvm import auto_scheduler, relay, te
 from tvm.script import tir as T
-
 from tvm.testing.auto_scheduler import matmul_auto_scheduler_test
 
 
@@ -78,8 +77,8 @@ def test_cpu_matmul():
     """
 
     # check touched memory in bytes, touched unique memory in bytes, reuse distance, etc.
-    assert fequal(fea_dict[c_name + ".bytes"], math.log2(512**3 * 4 + 1))
-    assert fequal(fea_dict[b_name + ".unique_bytes"], math.log2(512**2 * 4 + 1))
+    assert fequal(fea_dict[c_name + ".bytes"], math.log2(512 ** 3 * 4 + 1))
+    assert fequal(fea_dict[b_name + ".unique_bytes"], math.log2(512 ** 2 * 4 + 1))
     assert fequal(fea_dict[c_name + ".reuse_dis_iter"], math.log2(8 * 16 + 1))
     assert fequal(fea_dict[c_name + ".reuse_dis_bytes"], math.log2((8 * 16 + 8 + 16) * 4 + 1))
     assert fequal(fea_dict[c_name + ".reuse_ct"], math.log2(512 + 1))
@@ -209,9 +208,9 @@ def tir_matmul(
 ) -> None:
     # function attr dict
     T.func_attr({"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True})
-    A_flat = T.buffer_decl([16384], dtype="float32", data=A.data)
-    B_flat = T.buffer_decl([16384], dtype="float32", data=B.data)
-    C_flat = T.buffer_decl([16384], dtype="float32", data=C.data)
+    A_flat = T.Buffer([16384], dtype="float32", data=A.data)
+    B_flat = T.Buffer([16384], dtype="float32", data=B.data)
+    C_flat = T.Buffer([16384], dtype="float32", data=C.data)
     # body
     for x, y in T.grid(128, 128):
         C_flat[x * 128 + y] = T.float32(0)