tvm.driver

Namespace for driver APIs

tvm.lower(sch, args, name='main', binds=None, simple_mode=False)

Lowering step before build into target.

Parameters

  • sch (tvm.te.schedule.Schedule) – The schedule to be built

  • args (list of Buffer or Tensor or Var) – The argument lists to the function.

  • name (str, optional) – The name of result function.

  • binds (dict of Tensor to Buffer, optional) – Dictionary that maps the Tensor to Buffer which specified the data layout requirement of the function. By default, a new compact buffer is created for each tensor in the argument.

  • simple_mode (bool, optional) – Whether only output simple and compact statement, this will skip LoopPartition, api wrapper generation and Unrolling.

Returns

m – The result IRModule, if simple_mode=False Then the Stmt before make api is returned.

Return type

IRModule or Stmt

tvm.build(inputs, args=None, target=None, target_host=None, name='default_function', binds=None)

Build a function with arguments as signature. Code will be generated for devices coupled with target information.

Parameters

  • inputs (tvm.te.Schedule, IRModule, or dict of target to IRModule) – The schedule to be built

  • args (list of Buffer or Tensor or Var, optional) – The argument lists to the function.

  • target (str or tvm.target.Target, optional) – The target and option of the compilation.

  • target_host (str or tvm.target.Target optional) – Host compilation target, if target is device. When TVM compiles device specific program such as CUDA, we also need host(CPU) side code to interact with the driver setup the dimensions and parameters correctly. target_host is used to specify the host side codegen target. By default, llvm is used if it is enabled, otherwise a stackvm intepreter is used.

  • name (str, optional) – The name of result function.

  • binds (dict, optional) – Dictionary that maps the binding of symbolic buffer to Tensor. By default, a new buffer is created for each tensor in the argument.

Returns

ret – A module that combines both host and device code.

Return type

tvm.module

Examples

There are two typical example uses of this function depending on the type of the argument inputs: 1. it is an IRModule.

n = 2
A = te.placeholder((n,), name='A')
B = te.placeholder((n,), name='B')
C = te.compute(A.shape, lambda *i: A(*i) + B(*i), name='C')
s = tvm.te.create_schedule(C.op)
m = tvm.lower(s, [A, B, C], name="test_add")
rt_mod = tvm.build(m, target="llvm")

  1. it is a dict of compilation target to IRModule.

n = 2
A = te.placeholder((n,), name='A')
B = te.placeholder((n,), name='B')
C = te.compute(A.shape, lambda *i: A(*i) + B(*i), name='C')
s1 = tvm.te.create_schedule(C.op)
with tvm.target.cuda() as cuda_tgt:
  s2 = topi.cuda.schedule_injective(cuda_tgt, [C])
  m1 = tvm.lower(s1, [A, B, C], name="test_add1")
  m2 = tvm.lower(s2, [A, B, C], name="test_add2")
  rt_mod = tvm.build({"llvm": m1, "cuda": m2}, target_host="llvm")

Note

See the note on tvm.target on target string format.