prompt: 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()
  
  ```
  
golden_diff: 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)
  
verification_info: {"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"}]}