code: def CPUAccessRewrite():
      """Detect CPU access to VTA buffer and get address correctly.
  
      VTA's buffer is an opaque handle that do not
      correspond to address in CPU.
      This pass detect CPU access and rewrite to use pointer
      returned VTABufferCPUPtr for CPU access.
  
      Returns
      -------
      fpass : tvm.transform.Pass
          The pass
      """
  
      def _ftransform(f, mod, ctx):
          env = get_env()
  
          var_remap = {}
          buf_remap = {}
  
          def find_var_remap(old_var):
              if old_var in var_remap:
                  return var_remap[old_var]
  
              new_var = tvm.tir.Var(
                  old_var.name + "_ptr", dtype=old_var.type_annotation
              )
              var_remap[old_var] = new_var
              return new_var
  
          def find_buf_remap(old_buf):
              if old_buf in buf_remap:
                  return buf_remap[old_buf]
  
              new_var = find_var_remap(old_buf.data)
              new_buf = tvm.tir.decl_buffer(
                  shape=old_buf.shape,
                  dtype=old_buf.dtype,
                  data=new_var,
                  strides=old_buf.strides,
                  elem_offset=old_buf.elem_offset,
                  scope=old_buf.scope,
                  data_alignment=old_buf.data_alignment,
                  offset_factor=old_buf.offset_factor,
                  buffer_type="auto_broadcast"
                  if (old_buf.buffer_type == 2)
                  else "",
                  axis_separators=old_buf.axis_separators,
              )
              buf_remap[old_buf] = new_buf
              return new_buf
  
          def _post_order(op):
              if isinstance(op, tvm.tir.Allocate):
                  buffer_var = op.buffer_var
                  if buffer_var not in var_remap:
                      return None
                  new_var = var_remap[buffer_var]
                  let_stmt = tvm.tir.LetStmt(
                      new_var,
                      tvm.tir.call_extern(
                          "handle",
                          "VTABufferCPUPtr",
                          env.dev.command_handle,
                          buffer_var,
                      ),
                      op.body,
                  )
                  alloc = tvm.tir.Allocate(
                      buffer_var,
                      op.dtype,
                      op.extents,
                      op.condition,
                      let_stmt,
                  )
                  del var_remap[buffer_var]
                  bufs_to_delete = [
                      old_buf
                      for old_buf in buf_remap
                      if old_buf.data.same_as(buffer_var)
                  ]
                  for buf in bufs_to_delete:
                      del buf_remap[buf]
                  return alloc
  
              if isinstance(op, tvm.tir.BufferLoad):
                  return tvm.tir.BufferLoad(
                      find_buf_remap(op.buffer), op.indices
                  )
  
              if isinstance(op, tvm.tir.BufferStore):
                  return tvm.tir.BufferStore(
                      find_buf_remap(op.buffer), op.value, op.indices
                  )
  
              raise RuntimeError("not reached")
  
          stmt_in = f.body
          stmt = tvm.tir.stmt_functor.ir_transform(
              stmt_in,
              None,
              _post_order,
              ["tir.Allocate", "tir.BufferLoad", "tir.BufferStore"],
          )
  
          for old_var, new_var in var_remap.items():
              stmt = tvm.tir.LetStmt(
                  new_var,
                  tvm.tir.call_extern(
                      "handle",
                      "VTABufferCPUPtr",
                      env.dev.command_handle,
                      old_var,
                  ),
                  stmt,
              )
          return f.with_body(stmt)
  
      return tvm.tir.transform.prim_func_pass(
          _ftransform, opt_level=0, name="tir.vta.CPUAccessRewrite"
      )
  
file_path: tvm/vta/python/vta/transform.py
imports: import tvm
  from tvm import te
  from tvm.topi import utils
  from .environment import get_env
label_desc: Use image node_id 3 for calling a global function with example usage: CPUAccessRewrite() and returns: tvm, f, new_var, new_buf, var_remap, buf_remap, alloc, tvm, tvm, None