tilelang.ascend.debug¶

Ascend dialect device assert.

The Ascend toolkit’s assert() macro works in both aicore and SIMT code, so device asserts lower through the same tl.device_assert builtins as CUDA. Mirroring the CUDA dialect’s policy, the assert compiles to a no-op on hosts without an NPU toolkit so host-only tracing and tests still run.

Functions¶

device_assert(condition[, msg, no_stack_info])

Device-side assert emulation for Ascend targets.

Module Contents¶

tilelang.ascend.debug.device_assert(condition, msg='', no_stack_info=False)¶

Device-side assert emulation for Ascend targets.

Parameters:

  • condition (tvm.tirx.PrimExpr)

  • msg (str)