tilelang.ascend.language.print¶

Ascend dialect T.print.

Same debug externs as the common implementation (debug_print_var, debug_print_buffer_value, emitted by the Ascend codegen via src/tl_templates/ascend/debug.h), but with the NPU’s execution model: a kernel body runs once per AI core and UB/L1 buffers are core-local, so buffer prints need no CUDA-style single-thread gating. Keeping the policy here (rather than probing hardware inside the CUDA dialect module) keeps each dialect self-contained.

Functions¶

print([obj, msg])

Print a buffer, a primitive expression, or a plain message.

Module Contents¶

tilelang.ascend.language.print.print(obj=None, msg='')¶

Print a buffer, a primitive expression, or a plain message.

Parameters:

  • obj (Any) – A tirx.Buffer (local, shared, shared.dyn or global scope), a tirx.PrimExpr, or None for a message-only print.

  • msg (str) – Optional message prefix.

Return type:

None