[wangchen965]

Summary

LegalizePTOVMIMemoryAccess is not enabled by default. Due to the shortcomings of the analysis ability of community tvm, LegalizePTOvmiMemoryAccess is added at present, and the writing method that may lead to physical cross-border reading is rejected, but this will affect the usability of VMI.
Take x_ub = T.alloc_shared((heads_per_task, hidden_size), T.bfloat16) as an example.
Original writing：
x_f32 = V.vcvt(V.vload(x_ub[head_id, offset], size=vec_size), 'float32')
LegalizePTOVMIMemoryAccess will reject this writing, and you need to write:

x_flat = T.view(x_ub, (heads_per_task * hidden_size,))
x_f32 = V.vcvt(V.vload(x_flat[head_id * hidden_size + offset], size=vec_size), 'float32')

Now, LegalizePTOvmiMemoryAccess is set to be disabled by default, allowing the use of the original writing method.

Usage

@tilelang.jit(target='pto',
              pass_configs={"tl.enable_pto_vmi_coverage_check": True})

or

TILELANG_ENABLE_PTO_VMI_COVERAGE_CHECK=1 python ...

[KurrinQu]

新增 Python PassConfigKey 不能代替 C++ PassContext 配置注册。当前代码没有注册 tl.enable_pto_vmi_coverage_check，所以文档中的 pass_configs={...: True} 在建立 PassContext 时就会报 Invalid config option，还未进入 pipeline；设置 False 同样失败。请参照 tl.enable_auto_schedule 注册 Bool 配置，增加对应的TVM_REGISTER_PASS_CONFIG_OPTION(..., Bool)，沿用现有配置声明方式。验证显式配置能够覆盖环境变量默认值。

[wangchen965]

Fixed in src/ascend/op/builtin.cc
TVM_REGISTER_PASS_CONFIG_OPTION(kEnablePtoVmiCoverageCheck, Bool);