Skip to content

Commit

Permalink
Mark VTA's CPU-side allocations with disable_lower_builtin
Browse files Browse the repository at this point in the history
Because VTA's codegen is allowed to call host-side, it uses the
`"cpu"` tag.  Therefore, the allocations that are already handled with
`VTABufferCPUPtr` should opt-out of using the device API from
`LowerTVMBuiltin`.
  • Loading branch information
Lunderberg committed Jul 5, 2023
1 parent 2aa1971 commit af2964d
Showing 1 changed file with 8 additions and 1 deletion.
9 changes: 8 additions & 1 deletion vta/python/vta/transform.py
Original file line number Diff line number Diff line change
Expand Up @@ -202,7 +202,14 @@ def _post_order(op):
),
op.body,
)
alloc = tvm.tir.Allocate(buffer_var, op.dtype, op.extents, op.condition, let_stmt)
alloc = tvm.tir.Allocate(
buffer_var,
op.dtype,
op.extents,
op.condition,
let_stmt,
annotations={"disable_lower_builtin": True},
)
del var_remap[buffer_var]
bufs_to_delete = [
old_buf for old_buf in buf_remap if old_buf.data.same_as(buffer_var)
Expand Down

0 comments on commit af2964d

Please sign in to comment.