We are currently solving the following issue within our repository. Here is the issue text:
--- BEGIN ISSUE ---
[TOPI] Typo in operator key
https://github.com/apache/incubator-tvm/blob/bdfefbb03f5aab96ee677ee28a166dd6ab5dbf3f/python/tvm/topi/bifrost/dense.py#L26

"biforst" should be "bifrost". This bug makes the op totally unavailable in Relay.

I can fix this bug if expected, but I don't know how to add a proper test.

--- END ISSUE ---

Below are some code segments, each from a relevant file. One or more of these files may contain bugs.
--- BEGIN FILES ---
Path: `python/tvm/topi/bifrost/dense.py`
Content:
```
1 # Licensed to the Apache Software Foundation (ASF) under one
2 # or more contributor license agreements.  See the NOTICE file
3 # distributed with this work for additional information
4 # regarding copyright ownership.  The ASF licenses this file
5 # to you under the Apache License, Version 2.0 (the
6 # "License"); you may not use this file except in compliance
7 # with the License.  You may obtain a copy of the License at
8 #
9 #   http://www.apache.org/licenses/LICENSE-2.0
10 #
11 # Unless required by applicable law or agreed to in writing,
12 # software distributed under the License is distributed on an
13 # "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
14 # KIND, either express or implied.  See the License for the
15 # specific language governing permissions and limitations
16 # under the License.
17 # pylint: disable=invalid-name,unused-variable
18 """dense schedule on ARM Mali Biforst GPU"""
19 from tvm import te
20 from tvm import autotvm
21 
22 from .. import nn
23 from ..util import traverse_inline
24 
25 
26 @autotvm.register_topi_compute("dense.biforst")
27 def dense(_, data, weight, bias=None, out_dtype=None):
28     """Dense operator on Biforst"""
29     return nn.dense(data, weight, bias, out_dtype)
30 
31 
32 @autotvm.register_topi_schedule("dense.bifrost")
33 def schedule_dense(cfg, outs):
34     """Schedule for dense operator.
35 
36     Parameters
37     ----------
38     cfg: ConfigEntity
39         The config entity for this template
40     outs: Array of Tensor
41         The computation graph description of dense
42         in the format of an array of tensors.
43 
44     Returns
45     -------
46     s: Schedule
47         The computation schedule for dense.
48     """
49     outs = [outs] if isinstance(outs, te.tensor.Tensor) else outs
50     s = te.create_schedule([x.op for x in outs])
51 
52     def _callback(op):
53         if op.tag == "dense":
54             vec_size = [1, 2, 4, 8, 16]
55             max_unroll = 32
56 
57             dense_out = op.output(0)
58             output = outs[0]
59 
60             y, x = s[output].op.axis
61             c = s[dense_out].op.reduce_axis[0]
62 
63             ##### space definition begin #####
64             cfg.define_split("tile_y", y, num_outputs=3)
65             cfg.define_split("tile_x", x, num_outputs=3)
66             cfg.define_split("c_unroll", c, num_outputs=2, max_factor=64)
67 
68             # fallback support
69             if cfg.is_fallback:
70                 ref_log = autotvm.tophub.load_reference_log("mali", "rk3399", "dense.bifrost")
71                 cfg.fallback_with_reference_log(ref_log)
72             ##### space definition end #####
73 
74             if dense_out.op in s.outputs:
75                 dense_out = s.cache_write(output, "local")
76 
77             by, ty, yi = cfg["tile_y"].apply(s, output, y)
78             bx, tx, xi = cfg["tile_x"].apply(s, output, x)
79 
80             s[output].bind(by, te.thread_axis("blockIdx.y"))
81             s[output].bind(bx, te.thread_axis("blockIdx.x"))
82             s[output].bind(ty, te.thread_axis("threadIdx.y"))
83             s[output].bind(tx, te.thread_axis("threadIdx.x"))
84 
85             if cfg["tile_y"].size[-1] < max_unroll:
86                 s[output].unroll(yi)
87             if cfg["tile_x"].size[-1] in vec_size:
88                 s[output].vectorize(xi)
89             s[dense_out].compute_at(s[output], tx)
90 
91             k = s[dense_out].op.reduce_axis[0]
92             y, x = s[dense_out].op.axis
93             k, k_unroll = cfg["c_unroll"].apply(s, dense_out, k)
94             s[dense_out].reorder(k, k_unroll, y, x)
95             s[dense_out].unroll(k_unroll)
96             if cfg["tile_y"].size[-1] < max_unroll:
97                 s[dense_out].unroll(y)
98             if cfg["tile_x"].size[-1] in vec_size:
99                 s[dense_out].vectorize(x)
100 
101     traverse_inline(s, outs[0].op, _callback)
102     return s
103 
104 
105 def fuse_and_bind(s, tensor, axis=None, num_thread=None):
106     """ fuse all the axis and bind to GPU threads """
107     axis = axis or s[tensor].op.axis
108     fused = s[tensor].fuse(*axis)
109     bx, tx = s[tensor].split(fused, num_thread)
110     s[tensor].bind(bx, te.thread_axis("blockIdx.x"))
111     s[tensor].bind(tx, te.thread_axis("threadIdx.x"))
112     return bx, tx
113
```
--- END FILES ---

Please first localize the bug based on the issue statement, and then generate a patch according to the `git diff` format fenced by three backticks.

Here is an example:

```diff
diff --git a/examples/server_async.py b/examples/server_async.py
--- a/examples/server_async.py
+++ b/examples/server_async.py
@@ -313,4 +313,4 @@
 
 
 if __name__ == "__main__":
-    asyncio.run(run_async_server("."), debug=True)
+    asyncio.run(run_async_server(), debug=True)
diff --git a/examples/server_sync.py b/examples/server_sync.py
--- a/examples/server_sync.py
+++ b/examples/server_sync.py
@@ -313,5 +313,5 @@
 
 
 if __name__ == "__main__":
-    server = run_sync_server(".")
+    server = run_sync_server()
     server.shutdown()

```

diff --git a/python/tvm/topi/bifrost/dense.py b/python/tvm/topi/bifrost/dense.py
--- a/python/tvm/topi/bifrost/dense.py
+++ b/python/tvm/topi/bifrost/dense.py
@@ -23,7 +23,7 @@
 from ..util import traverse_inline
 
 
-@autotvm.register_topi_compute("dense.biforst")
+@autotvm.register_topi_compute("dense.bifrost")
 def dense(_, data, weight, bias=None, out_dtype=None):
     """Dense operator on Biforst"""
     return nn.dense(data, weight, bias, out_dtype)

{"golden_diff": "diff --git a/python/tvm/topi/bifrost/dense.py b/python/tvm/topi/bifrost/dense.py\n--- a/python/tvm/topi/bifrost/dense.py\n+++ b/python/tvm/topi/bifrost/dense.py\n@@ -23,7 +23,7 @@\n from ..util import traverse_inline\n \n \n-@autotvm.register_topi_compute(\"dense.biforst\")\n+@autotvm.register_topi_compute(\"dense.bifrost\")\n def dense(_, data, weight, bias=None, out_dtype=None):\n     \"\"\"Dense operator on Biforst\"\"\"\n     return nn.dense(data, weight, bias, out_dtype)\n", "issue": "[TOPI] Typo in operator key\nhttps://github.com/apache/incubator-tvm/blob/bdfefbb03f5aab96ee677ee28a166dd6ab5dbf3f/python/tvm/topi/bifrost/dense.py#L26\r\n\r\n\"biforst\" should be \"bifrost\". This bug makes the op totally unavailable in Relay.\r\n\r\nI can fix this bug if expected, but I don't know how to add a proper test.\n", "before_files": [{"content": "# Licensed to the Apache Software Foundation (ASF) under one\n# or more contributor license agreements.  See the NOTICE file\n# distributed with this work for additional information\n# regarding copyright ownership.  The ASF licenses this file\n# to you under the Apache License, Version 2.0 (the\n# \"License\"); you may not use this file except in compliance\n# with the License.  You may obtain a copy of the License at\n#\n#   http://www.apache.org/licenses/LICENSE-2.0\n#\n# Unless required by applicable law or agreed to in writing,\n# software distributed under the License is distributed on an\n# \"AS IS\" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY\n# KIND, either express or implied.  See the License for the\n# specific language governing permissions and limitations\n# under the License.\n# pylint: disable=invalid-name,unused-variable\n\"\"\"dense schedule on ARM Mali Biforst GPU\"\"\"\nfrom tvm import te\nfrom tvm import autotvm\n\nfrom .. import nn\nfrom ..util import traverse_inline\n\n\n@autotvm.register_topi_compute(\"dense.biforst\")\ndef dense(_, data, weight, bias=None, out_dtype=None):\n    \"\"\"Dense operator on Biforst\"\"\"\n    return nn.dense(data, weight, bias, out_dtype)\n\n\n@autotvm.register_topi_schedule(\"dense.bifrost\")\ndef schedule_dense(cfg, outs):\n    \"\"\"Schedule for dense operator.\n\n    Parameters\n    ----------\n    cfg: ConfigEntity\n        The config entity for this template\n    outs: Array of Tensor\n        The computation graph description of dense\n        in the format of an array of tensors.\n\n    Returns\n    -------\n    s: Schedule\n        The computation schedule for dense.\n    \"\"\"\n    outs = [outs] if isinstance(outs, te.tensor.Tensor) else outs\n    s = te.create_schedule([x.op for x in outs])\n\n    def _callback(op):\n        if op.tag == \"dense\":\n            vec_size = [1, 2, 4, 8, 16]\n            max_unroll = 32\n\n            dense_out = op.output(0)\n            output = outs[0]\n\n            y, x = s[output].op.axis\n            c = s[dense_out].op.reduce_axis[0]\n\n            ##### space definition begin #####\n            cfg.define_split(\"tile_y\", y, num_outputs=3)\n            cfg.define_split(\"tile_x\", x, num_outputs=3)\n            cfg.define_split(\"c_unroll\", c, num_outputs=2, max_factor=64)\n\n            # fallback support\n            if cfg.is_fallback:\n                ref_log = autotvm.tophub.load_reference_log(\"mali\", \"rk3399\", \"dense.bifrost\")\n                cfg.fallback_with_reference_log(ref_log)\n            ##### space definition end #####\n\n            if dense_out.op in s.outputs:\n                dense_out = s.cache_write(output, \"local\")\n\n            by, ty, yi = cfg[\"tile_y\"].apply(s, output, y)\n            bx, tx, xi = cfg[\"tile_x\"].apply(s, output, x)\n\n            s[output].bind(by, te.thread_axis(\"blockIdx.y\"))\n            s[output].bind(bx, te.thread_axis(\"blockIdx.x\"))\n            s[output].bind(ty, te.thread_axis(\"threadIdx.y\"))\n            s[output].bind(tx, te.thread_axis(\"threadIdx.x\"))\n\n            if cfg[\"tile_y\"].size[-1] < max_unroll:\n                s[output].unroll(yi)\n            if cfg[\"tile_x\"].size[-1] in vec_size:\n                s[output].vectorize(xi)\n            s[dense_out].compute_at(s[output], tx)\n\n            k = s[dense_out].op.reduce_axis[0]\n            y, x = s[dense_out].op.axis\n            k, k_unroll = cfg[\"c_unroll\"].apply(s, dense_out, k)\n            s[dense_out].reorder(k, k_unroll, y, x)\n            s[dense_out].unroll(k_unroll)\n            if cfg[\"tile_y\"].size[-1] < max_unroll:\n                s[dense_out].unroll(y)\n            if cfg[\"tile_x\"].size[-1] in vec_size:\n                s[dense_out].vectorize(x)\n\n    traverse_inline(s, outs[0].op, _callback)\n    return s\n\n\ndef fuse_and_bind(s, tensor, axis=None, num_thread=None):\n    \"\"\" fuse all the axis and bind to GPU threads \"\"\"\n    axis = axis or s[tensor].op.axis\n    fused = s[tensor].fuse(*axis)\n    bx, tx = s[tensor].split(fused, num_thread)\n    s[tensor].bind(bx, te.thread_axis(\"blockIdx.x\"))\n    s[tensor].bind(tx, te.thread_axis(\"threadIdx.x\"))\n    return bx, tx\n", "path": "python/tvm/topi/bifrost/dense.py"}], "after_files": [{"content": "# Licensed to the Apache Software Foundation (ASF) under one\n# or more contributor license agreements.  See the NOTICE file\n# distributed with this work for additional information\n# regarding copyright ownership.  The ASF licenses this file\n# to you under the Apache License, Version 2.0 (the\n# \"License\"); you may not use this file except in compliance\n# with the License.  You may obtain a copy of the License at\n#\n#   http://www.apache.org/licenses/LICENSE-2.0\n#\n# Unless required by applicable law or agreed to in writing,\n# software distributed under the License is distributed on an\n# \"AS IS\" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY\n# KIND, either express or implied.  See the License for the\n# specific language governing permissions and limitations\n# under the License.\n# pylint: disable=invalid-name,unused-variable\n\"\"\"dense schedule on ARM Mali Biforst GPU\"\"\"\nfrom tvm import te\nfrom tvm import autotvm\n\nfrom .. import nn\nfrom ..util import traverse_inline\n\n\n@autotvm.register_topi_compute(\"dense.bifrost\")\ndef dense(_, data, weight, bias=None, out_dtype=None):\n    \"\"\"Dense operator on Biforst\"\"\"\n    return nn.dense(data, weight, bias, out_dtype)\n\n\n@autotvm.register_topi_schedule(\"dense.bifrost\")\ndef schedule_dense(cfg, outs):\n    \"\"\"Schedule for dense operator.\n\n    Parameters\n    ----------\n    cfg: ConfigEntity\n        The config entity for this template\n    outs: Array of Tensor\n        The computation graph description of dense\n        in the format of an array of tensors.\n\n    Returns\n    -------\n    s: Schedule\n        The computation schedule for dense.\n    \"\"\"\n    outs = [outs] if isinstance(outs, te.tensor.Tensor) else outs\n    s = te.create_schedule([x.op for x in outs])\n\n    def _callback(op):\n        if op.tag == \"dense\":\n            vec_size = [1, 2, 4, 8, 16]\n            max_unroll = 32\n\n            dense_out = op.output(0)\n            output = outs[0]\n\n            y, x = s[output].op.axis\n            c = s[dense_out].op.reduce_axis[0]\n\n            ##### space definition begin #####\n            cfg.define_split(\"tile_y\", y, num_outputs=3)\n            cfg.define_split(\"tile_x\", x, num_outputs=3)\n            cfg.define_split(\"c_unroll\", c, num_outputs=2, max_factor=64)\n\n            # fallback support\n            if cfg.is_fallback:\n                ref_log = autotvm.tophub.load_reference_log(\"mali\", \"rk3399\", \"dense.bifrost\")\n                cfg.fallback_with_reference_log(ref_log)\n            ##### space definition end #####\n\n            if dense_out.op in s.outputs:\n                dense_out = s.cache_write(output, \"local\")\n\n            by, ty, yi = cfg[\"tile_y\"].apply(s, output, y)\n            bx, tx, xi = cfg[\"tile_x\"].apply(s, output, x)\n\n            s[output].bind(by, te.thread_axis(\"blockIdx.y\"))\n            s[output].bind(bx, te.thread_axis(\"blockIdx.x\"))\n            s[output].bind(ty, te.thread_axis(\"threadIdx.y\"))\n            s[output].bind(tx, te.thread_axis(\"threadIdx.x\"))\n\n            if cfg[\"tile_y\"].size[-1] < max_unroll:\n                s[output].unroll(yi)\n            if cfg[\"tile_x\"].size[-1] in vec_size:\n                s[output].vectorize(xi)\n            s[dense_out].compute_at(s[output], tx)\n\n            k = s[dense_out].op.reduce_axis[0]\n            y, x = s[dense_out].op.axis\n            k, k_unroll = cfg[\"c_unroll\"].apply(s, dense_out, k)\n            s[dense_out].reorder(k, k_unroll, y, x)\n            s[dense_out].unroll(k_unroll)\n            if cfg[\"tile_y\"].size[-1] < max_unroll:\n                s[dense_out].unroll(y)\n            if cfg[\"tile_x\"].size[-1] in vec_size:\n                s[dense_out].vectorize(x)\n\n    traverse_inline(s, outs[0].op, _callback)\n    return s\n\n\ndef fuse_and_bind(s, tensor, axis=None, num_thread=None):\n    \"\"\" fuse all the axis and bind to GPU threads \"\"\"\n    axis = axis or s[tensor].op.axis\n    fused = s[tensor].fuse(*axis)\n    bx, tx = s[tensor].split(fused, num_thread)\n    s[tensor].bind(bx, te.thread_axis(\"blockIdx.x\"))\n    s[tensor].bind(tx, te.thread_axis(\"threadIdx.x\"))\n    return bx, tx\n", "path": "python/tvm/topi/bifrost/dense.py"}]}