diff --git a/docs/designs/vpto-integer-address-canonicalization-design-zh.md b/docs/designs/vpto-integer-address-canonicalization-design-zh.md
new file mode 100644
index 0000000000..9713b0acf8
--- /dev/null
+++ b/docs/designs/vpto-integer-address-canonicalization-design-zh.md
@@ -0,0 +1,673 @@
+# VPTO 整数地址规范化设计
+
+## 1. 结论
+
+Issue #591 暴露的是 pointer SSA 形态问题：地址的循环变化量仍藏在
+`pto.castptr` 的整数输入中，而现有 `PTOAddressAnalysis` 只会沿
+`pto.addptr` 累积 element offset。
+
+本设计只规范化 pointer，不解释使用该 pointer 的 memory operation：
+
+```mlir
+%byte_address = ... : i64
+%ptr = pto.castptr %byte_address : i64 -> !pto.ptr<T, ub>
+```
+
+规范化的目标形态是「规范根 + addptr(element_offset)」：
+
+```mlir
+%root_addr = ...            // 规范根整数输入：常量 0，或不可归零原子（见 3.2）
+%base = pto.castptr %root_addr : i64 -> !pto.ptr<T, ub>
+%ptr = pto.addptr %base, %element_offset
+    : !pto.ptr<T, ub> -> !pto.ptr<T, ub>
+```
+
+`%element_offset` 必须是与原字节表达式「可归零部分」等价的精确商，并按
+`pto.addptr` 要求物化为 `index`。规范根取 `castptr(0)` 是退化特例：整个字节
+表达式全部归入 element offset；含运行时原子基址（如 kernel 参数）时，原子保留
+为根（见 3.2 与 C15）。变换不改 memory operation，不感知 LLVM intrinsic、offset
+字段宽度、Bisheng post-update 编码、访问 footprint 或 post-update 行为。
+
+设计由两个彼此独立的 canonical rewrite 组成：
+
+1. **整数地址规范化**：`castptr(byte_address)` 变为
+   `castptr(规范根) + addptr(element_offset)`（规范根为 `castptr(0)` 或不可归零
+   原子，见 3.2/3.5）；这是核心设计。
+2. **addptr 吸收**：对本来就有 element offset 的 operation，将
+   `op(addptr(base, A), O)` 变为 `op(base, A + O)`；这是面向现有后端形态的普通
+   canonical fold，不参与地址等价分析。
+
+第一步已经让 `PTOAddressAnalysis` 和它的 consumers 看见地址递推。第二步是否执行，
+不影响第一步的正确性。
+
+## 2. 问题边界
+
+### 2.1 当前主线已有的能力
+
+当前 `PTOAddressAnalysis` 已经统一表示 typed pointer 地址：
+
+- `PTOAddressExpr` 保存 root/base、`pto.addptr` element offset 和 operation offset；
+- `getAddresses()` 沿相同 element type 的 `pto.addptr` 链向上遍历；
+- `getDeltaBytes()`、`getDifferenceBytes()` 将 pointer offset 和 operation offset
+  统一换算为 byte difference；
+- `convertDeltaToUnit()` 对线性表达式做精确单位换算，不能整除时返回
+  `InexactUnitConversion`；
+- `PTOValueEvolutionAnalysis` 已经表示常量、加减、常量乘法和不同种类的整数 cast。
+
+缺口只有一个：`getAddresses()` 到达 `pto.castptr(integer)` 后，把它作为 opaque root，
+不会把整数输入中的循环变化量变成 element offset。
+
+因此不新增额外的地址模型、终端 footprint 或 operation adapter。需要新增的是一个使用现有 typed expression
+能力的 rewrite，以及它所需的精确商和 SSA 物化 helper。
+
+### 2.2 `pto.addptr` 是 PTO 层的语义目标
+
+`pto.addptr` 是 Pure op，offset 以 pointer element 为单位。规范化在 PTO IR 内证明：
+
+```text
+integer byte address
+  == zero-origin typed pointer + element offset
+```
+
+LLVM 最终使用 GEP、整数加法还是 target intrinsic，不属于该证明。特别是 LLVM lowering
+中出现的 `i32 trunc` 只是当前 emitter 的结果，不能成为地址分析的输入、拒绝条件或
+正确性依据。若某个 operation 的 lowering 字段过窄，那是 operation verifier、target
+legality 或 emitter 的问题。
+
+**实测（2026-08-26）**：`castptr %x : i32 -> ptr<f16, ub>` 当前经
+`ConvertPtoCastPtrOp` 直接生成 `inttoptr i32 %x to ptr addrspace(6)`（无位宽扩展），
+LLVM 合法且 bisheng 可编译。窄整数到 64 位指针的扩展在 LLVM 层**有书面定义**
+（`LangRef.rst` 的 `inttoptr` Semantics：窄→宽 zero-extend、宽→窄 truncate），
+bisheng 行为与之一致；PTO 层只是未文档化继承该语义。LLVM 真正的未定义在
+provenance——纯整数 `inttoptr` 形成的指针不基于任何指针，按 pointer aliasing
+rules 解引用是 UB，但 PTO 的 zero-origin 数值地址模型（§2.3）本就按整数位模式
+处理 UB 空间，规范化不改变这一性质。lowering 事实既不能作为规范化拒绝条件，
+也不能作为 round-trip 证明依据（§3.3 的实现近似因此保持"仅 64 位输入"）。
+
+### 2.3 地址空间前提不是用户配置
+
+该变换只适用于 PTO 语义明确规定为 **zero-origin integral address space** 的空间：
+
+- 地址由编译器或开发者以整数管理；
+- 地址 0 有效；
+- `castptr(integer)` 按整数 bit pattern 形成地址；
+- `addptr` 的 element scaling 与该整数地址使用相同的 byte 单位。
+
+A5 UB 满足这些条件。这个判定应来自 PTO dialect/target 的权威语义，不引入可由用户任意
+组合的地址配置对象。其他空间只有在其 PTO 语义同样满足四项条件时才能启用；
+LLVM object/provenance 管理的空间和语义未知的空间保持原样。
+
+这一区分也解释了为什么 LLVM 关于 `null + GEP` 的通用限制不直接否定 UB 变换：UB 的
+0 是实际地址原点，不是 C/LLVM object model 中不可解引用的 null object。
+
+## 3. 规范化语义
+
+### 3.1 输入和输出
+
+设原 op 为：
+
+```text
+P = castptr(B) : integer -> ptr<T, S>
+```
+
+其中：
+
+- `B` 是整数 byte-address expression；
+- `E = sizeof_storage(T)`，单位为 byte；
+- `S` 是已确认的 zero-origin integral address space；
+- `Waddr` 是该空间的地址位宽；
+- `Windex` 是当前 PTO target 中 `index` 的位宽。
+
+把 `B` 线性化为 `C + sum(Ki * Xi)` 后，规范根与商分别定义：
+
+- **规范根整数输入** `R`：所有系数 `Ki` 都能被 `E` 整除时，`R = 0`（退化特例，
+  根为 `castptr(0)`）；否则取「系数为 1 且不可被 `E` 整除的原子叶子」为 `R`
+  （见 3.2 的原子根规则，最常见是运行时基址参数；仅支持单个此类原子，多原子
+  场景整体拒绝）。
+- **商** `Q`：`B - R` 的可整除部分除以 `E` 的精确商，可物化为 `index`，满足：
+
+```text
+CastToAddressWidth(B - R)
+  == ScaleElementOffset(CastIndex(Q), E)       (mod 2^Waddr)
+```
+
+然后生成：
+
+```text
+P' = addptr(castptr(R), Q) : ptr<T, S>
+```
+
+只要上式成立，`P` 和 `P'` 作为 pointer value 等价，与它们是否被 load、store、pointer
+cast、循环递推或多个 operation 使用无关。这就是本设计不需要分析地址终端的原因。
+
+### 3.2 精确商与原子根
+
+对线性 typed expression，默认规范形式是对整个 `B` 求 `E` 的精确商，不先选择 base，
+也不把表达式拆成 invariant 和 variant 部分：
+
+```text
+B = C + sum(Ki * Xi)
+Q = C/E + sum((Ki/E) * Xi)
+```
+
+只有常数项 `C` 和每个系数 `Ki` 都能被 `E` 整除时，线性精确商存在，规范根
+`R = 0`。cast、block argument 和其他不能穿透的 value 作为原子 `Xi` 保留，因此不会
+为了整除而把 cast 沿数据流移动。
+
+**原子根例外**：当存在系数 `Ki` 不可被 `E` 整除的原子叶子时（最常见是系数 1 的
+运行时基址参数，如 kernel 参数的 `castptr(%param)`），该原子不能归入 offset，
+保留为规范根：
+
+```text
+B = R + (C' + sum(Kj * Xj))         // R 是不可归零原子部分
+Q = C'/E + sum((Kj/E) * Xj)
+P' = addptr(castptr(R), Q)
+```
+
+闭合边界：只支持系数为 1 的原子叶子作根（`%param` 本身）；系数 `|Ki| != 1` 且
+不可整除的原子（如 `%param * 3` 对 `E = 2`）没有唯一的「根 + 偏移」分离，整体
+拒绝，保持原样。多个系数为 1 的原子叶子同时存在时，它们组成的和表达式作为根
+（根仍唯一）；该场景当前按保守策略整体拒绝，留给后续扩展。
+
+现有 `normalizePTOLinearExpr()` 与 `dividePTOLinearExprExact()` 已经实现这项代数能力。
+新 rewrite 可以复用它们，但还必须完成 typed round-trip proof；仅有 `int64_t` 系数整除
+不等于跨位宽转换一定正确。
+
+不允许为了扩大覆盖范围在 IR 中生成动态 `div`/`rem`。规范化的目标是更直接的地址
+表达式；新增运行时除法既有成本，也把“静态可证明”变成了运行时条件。
+
+### 3.3 cast 和有限位宽
+
+精确商保留原表达式中的 cast 边界。以下输入可以把 cast 结果作为一个原子：
+
+```mlir
+%x64 = arith.extui %x : i32 to i64
+%bytes = arith.muli %x64, %c4096 : i64
+```
+
+商为 `%x64 * 2048`，没有移动 `%x64`。相反，默认不把：
+
+```text
+trunc(mul(x, 4096))
+```
+
+改写成：
+
+```text
+mul(trunc(x), 2048)
+```
+
+因为截断发生点改变了。
+
+从商表达式到 `index` 还必须验证 round trip：
+
+```text
+Q-expression -> index -> addptr element scaling -> Waddr
+```
+
+保持原地址 bit pattern。`Windex` 和 `Waddr` 来自 PTO target/data-layout 语义，不从
+LLVM lowering 产物反推。当前 `PTOValueEvolutionAnalysis` 中固定 64 位的 index 处理是
+实现现状，不应被提升为跨 target 合同。
+
+**实现（C14 的 round-trip proof）**：同宽（64 位）输入平凡无损，直接接受。更窄
+输入（如 i32）按 LLVM `inttoptr` 的 zero-extend 语义（§2.2 实测）做 round-trip
+证明：字节表达式 `B` 可证 `0 <= B < 2^inputWidth` 时，`Q = B/E < 2^inputWidth` 且
+`zext(Q) * E == zext(B)` 在 64 位 index 域精确成立。证明路径：常量表达式用
+`foldPTOConstant`；循环递推用 `PTOValueEvolutionAnalysis::getRange`（要求非负且
+signed 上界留有余量）。已知边界：分析器的 Mul range 目前仅支持 `index` 类型
+（`getSyntheticEvolutionImpl` 的 `hasIndexType` 检查），因此 i32 循环递推（如
+`muli(iv32, 4096)`）即使数学上不溢出也会返回 `PossibleWrap` 而拒绝——扩展现有
+分析器支持非 index 算术 range 是后续工作，不影响本设计边界。宽于 index 的输入
+（如 i128）会截断，拒绝。
+
+带 `nuw`/`nsw` 的原算术只要求在原程序有定义的输入上保持地址，并且新表达式不新增
+poison。reifier 不复制无法证明的 overflow flag。无法证明 refinement 时保持原样。
+
+### 3.4 非整除地址
+
+例如：
+
+```mlir
+%bytes = arith.addi %multiple_of_2, %c1 : i64
+%ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+```
+
+不存在 f16 element offset `Q` 使 `2 * Q` 表示所有输入，因此不规范化。
+
+本设计不退化成：
+
+```text
+castptr(0) -> ptr<i8>
+addptr(byte_offset)
+castptr -> ptr<T>
+```
+
+该形态虽然可能保持数值地址，却重新引入 ptr-to-ptr cast，并且当前
+`PTOAddressAnalysis` 不会透明地穿过最后一个 cast。它不能形成唯一 canonical form，
+也不能解决本设计要解决的分析入口问题。拒绝非整除输入是闭合的语义边界。
+
+这里的「非整除」指常数项或系数不可整除且不适用 3.2 的原子根例外：常数项不可整除
+（如上面的 `+1`）没有 typed canonical form，整体拒绝；单个系数为 1 的原子叶子不可
+归零则是「保留为根」而非拒绝（C15，见 3.2 闭合边界），两者拒绝原因不同（与 6.3
+的区分一致）。
+
+### 3.5 收敛性、幂等性和共享
+
+规范化必须收敛到唯一的规范形（NF），不是若干条互不相关的 peephole：
+
+```text
+NF = castptr(R)                     // Q 恒为 0 时（纯根：castptr(0) 或 castptr(%原子)）
+   | castptr(R) + addptr(Q)         // R 是常量 0 或不可归零原子输入，Q 非平凡
+```
+
+rewrite 规则 `R` 只匹配 integer-to-pointer `pto.castptr`，且要求**可归零部分的精确商
+`Q = (B - R) / E` 不恒为 0**（即 `B - R` 非平凡）：
+
+- **终止**：每次应用 `R` 消灭一个「Q 非平凡」的 `castptr`；`R` 的产物中唯一的
+  `castptr` 输入是常量 0 或不可归零原子，它们的 `B - R = 0`、`Q` 恒为 0，不满足
+  `R` 的前提，不会被 canonicalizer worklist 再次命中。应用次数 ≤ 输入中可规范化
+  `castptr` 的个数，有限。
+  （若只排除常量 0 而不排除纯原子输入，`castptr(%param)` 会被反复包成
+  `addptr(castptr(%param), 0)`，永不终止——所以「Q 不恒为 0」才是完整的排除条件，
+  它同时覆盖常量 0 与纯原子两类输入。）
+- **唯一性**：根由不可归零原子部分唯一确定（无原子时为常量 0；仅支持单个原子，
+  见 3.2），`Q` 是 `B - R` 的唯一精确商（3.2）。同一字节表达式只有一种 NF。
+- **幂等**：`R` 的输出是 NF，再次运行不匹配、不改 IR。
+
+若原 `%ptr` 有多个 users，统一替换其定义即可；不按 memory operation 分别 clone 地址
+表达式。若商表达式不能在原 `castptr` 位置合法物化，整个 rewrite 失败，不在 users 附近
+复制计算。
+
+## 4. 实现结构
+
+### 4.1 分析，不新建第二套地址体系
+
+rewrite 使用现有 `PTOValueEvolutionAnalysis`：
+
+```text
+castptr integer input
+  -> getExpr()
+  -> normalizePTOLinearExpr()
+  -> dividePTOLinearExprExact(elementBytes)
+  -> typed/index round-trip proof
+```
+
+分析结果可以是 pass 内部的小型值对象：
+
+```c++
+struct ExactElementQuotient {
+  PTOTypedExprRef quotient;
+  Type sourceIntegerType;
+  unsigned addressWidth;
+  unsigned indexWidth;
+};
+```
+
+它不是新的 public address model，也不保存 operation、builder 或待创建 SSA。element
+storage size 应复用 `PTOAddressAnalysis` 现有规则；实现时把当前 file-local
+`getElementBytes()` 提升为 Analysis 层共享 helper，避免复制类型规则。
+
+PR #1260 增加的 typed address/alignment/contiguity 查询与这里互补，但不替代精确商。
+截至本文核对的当前 checkout，它仍不是可直接依赖的主线 API；本设计不以该 PR 合入为
+前提，也不复制它的 alignment 或 stream-fusion 能力。PR 合入后，规范化输出自然成为
+这些查询的普通 `addptr` 输入。
+
+### 4.2 SSA 物化
+
+reifier 负责：
+
+1. 复用支配原 `castptr` 的 SSA leaves；
+2. 按 typed expression 重建常量、加减和常量乘法；
+3. 保留原子 cast value，不跨 cast 重排；
+4. 生成经过 round-trip proof 的 `index` value；
+5. 在原 `castptr` 位置生成规范根整数输入（常量 0 或原子叶子，见 3.2）、同 pointer
+   type 的 `castptr(R)` 和 `pto.addptr`；
+6. 原子地替换所有 users，失败时不改 IR。
+
+只重建 Pure、可推测执行且已由 proof 覆盖的算术。load、call、带副作用 op 或 region
+外不可捕获的值都只可作为支配位置合法的叶子，不能 clone。
+
+**规范根的提升约束**：规范根 `castptr(R)` 是 loop-invariant（`R` 为常量 0 或函数
+参数等原子），reifier 把它提升到最近 `scf.for` 之外，使 post-update consumer 看到
+循环外 base。提升仅当 `R` 对**所有**被跳过的循环都是 loop-invariant 时合法；若原子根
+定义在循环内（如 `index_cast(iv)` 本身不可整除），提升会违反 SSA 支配，此时 `castptr(R)`
+留在原位（IR 保持合法，只是不触发 post-update）。商 `Q` 始终在原 `castptr` 位置物化
+（依赖循环变量的部分留在循环内）。
+
+### 4.3 与 `PTOAddressAnalysis` 的接入
+
+规范化后无需给 `PTOAddressAnalysis` 增加 raw-integer root：
+
+```text
+castptr(integer)                         opaque root
+        |
+        | canonicalize
+        v
+castptr(R) -> addptr(element_offset)     existing analyzable form
+                                          (R = 常量 0 或不可归零原子)
+```
+
+`getAddresses()` 按现有逻辑沿 `addptr` 得到 element offset；
+`getDeltaBytes()` 按现有逻辑得到 loop step；SoftPostUpdate、alignment、contiguity 等
+consumer 只消费这一套结果。pass 改写 IR 后正常失效并重建 analysis cache。
+
+## 5. 独立的 addptr 吸收 fold
+
+实际 Bisheng 验证表明，下面两种 PTO 形态并不等价地触发后端 post-update：
+
+```mlir
+// 地址分析可见，但当前 Bisheng 仍生成 VLDI + SADD。
+%ptr = pto.addptr %base, %a
+%v = pto.vldsx2 %ptr[%zero], ...
+
+// 当前 Bisheng 能生成 VLDS post-update。
+%v = pto.vldsx2 %base[%a], ...
+```
+
+因此需要时可再执行一个普通 canonical fold：
+
+```text
+op(addptr(base, A), O) -> op(base, A + O)
+```
+
+其合法性条件仅来自现有 `VPTOAddressSemantics`：
+
+- operation 当前访问恰有对应 base 和 offset operand；
+- offset unit 是 Element，且 element type 与 `addptr` 相同；
+- operation 不是已经带 `updatedBase` 的 post-update form；
+- `A + O` 能在原 offset 类型中无损物化；
+- 替换不改变其他 base users；
+- **`addptr` 的 base 必须是 integer-backed `castptr`**（规范化的产物形态）。任意
+  `addptr` 链（base 是用户指针或另一层 addptr）不折叠——`VPTOSoftPostUpdate`
+  对这类链有自己的 sequential base-chain post-update 处理，提前折叠会改变或破坏
+  该结构（lit 回归：`soft_postupdate_sequential-base-chain`）。
+
+这个 fold 不读取 LLVM intrinsic 字段，不推测 post-update step，也不分析 footprint。
+它只是把同单位的两级加法折叠到 operation 已有的地址 operand。对于 offset unit 为 Byte、
+Block 或 Alignment 的 operation，除非现有 semantics 提供精确的同单位换算，否则不做。
+
+operation 语义建模复用主线的 `VPTOAddressSemanticsOpInterface`（14 个 load/store op
+已声明接口：`vlds`/`vldsx2`/`vldus`/`plds`/`pldi`/`vsts`/`vstus`/`psts`/`psti`/
+`sprsts`/`sprsti`/`vstas`/`vsldb`/`vsstb`，实现见 `lib/PTO/IR/VPTOAddressSemantics.cpp`
+的 `getDefaultVPTOAddressSemantics`；`VPTOSoftPostUpdate` 已按同一接口消费
+`postUpdate`）。C11 fold 只读 `currentAccesses` 的 base/offset/unit，不需要新增语义
+建模。已知缺口：`vstsx2` 未声明该接口（与 `vldsx2` 不对称），`PTOAddressAnalysis` 与
+C11 fold 对它均不可用；实现时补一行接口声明 + `getDefaultVPTOAddressSemantics` case
+（与 `vsts` 对称），或在 fold 中显式声明暂不支持。
+
+## 6. Case 驱动的设计收敛
+
+case 不是实现后的补充，而是设计合同的可执行边界。每个接受 case 必须同时检查输出形态
+和地址等价；每个拒绝 case 必须检查 IR 保持不变。核心矩阵如下。
+
+| Case | 输入关键形态 | 期望 | 锁定的设计结论 |
+| --- | --- | --- | --- |
+| C01 | Issue #591: `index_cast(iv) * 4096 -> ptr<f16,ub>` | `castptr(0) + addptr(iv * 2048)` | 整条 byte expression 精确除 element size |
+| C02 | `8192 + index_cast(iv) * 4096 -> ptr<f16,ub>` | offset 为 `4096 + iv * 2048` | 不需要挑选 invariant base，所有地址统一归到 offset |
+| C03 | `index_cast(iv) * 4096 -> ptr<f32,ub>` | offset 为 `iv * 1024` | element storage size 来自 pointer type |
+| C04 | 任意线性 i64 byte address -> `ptr<i8,ub>` | 同值 element offset | element size 1 是恒等换算 |
+| C05 | `(extui x) * 4096 -> ptr<f16,ub>` | 保留 `extui x`，系数变 2048 | cast 作为原子，不移动 cast |
+| C06 | `trunc(x * 4096) -> ptr<f16,ub>` | 不改 | 不跨截断做代数重排 |
+| C07 | `iv * 4096 + 1 -> ptr<f16,ub>` | 不改 | 非整除 byte address 没有 typed element canonical form |
+| C08 | 地址空间不是 zero-origin integral | 不改 | 不把 UB 数值地址规则扩散到 object pointer |
+| C09 | 已是规范形：`castptr(0) + addptr(E)`、`castptr(%原子) + addptr(E)`，或裸 `castptr(0)` / `castptr(%原子)` | 不改 | 规则要求 `Q` 非平凡；收敛到规范形，无重复包装（见 3.5） |
+| C10 | 一个 integer-backed pointer 有多个 users | 生成一个共享 addptr | rewrite pointer producer，不按终端复制 |
+| C11 | `%ptr = addptr(base,A); op %ptr[O]` | 可选 fold 为 `op base[A+O]` | 后端形态 fold 与地址规范化分离 |
+| C12 | 已有 post-update op | 不做 C11 fold | 当前访问与访问后更新不混合 |
+| C13 | `%raw = castptr %p : ptr -> i64; %bytes = raw * 2` | 保持不变 | pointer-derived integer 不进入本规范化；不引入 provenance parser |
+| C14 | 商不能无损 round-trip 到 target index | 不改 | index/address width 是 PTO target 证明条件 |
+| C15 | `castptr(%param + iv * 4096)`（`%param` 为 kernel 参数） | `castptr(%param)` 为根，offset 只含 `iv * 2048` | 运行时原子基址是规范根；分析与发射以「根循环不变」为前提 |
+
+### 6.1 C01：Issue #591
+
+```mlir
+%iv64 = arith.index_cast %iv : index to i64
+%bytes = arith.muli %iv64, %c4096 : i64
+%ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+```
+
+线性式只有一个系数 4096，`sizeof(f16)=2`，精确商是：
+
+```text
+Q = index_cast(iv) * 2048
+```
+
+A5 的 PTO target 证明该商可无损进入 index 后，输出：
+
+```mlir
+%zero = arith.constant 0 : i64
+%base = pto.castptr %zero : i64 -> !pto.ptr<f16, ub>
+%elements = ... // 与 index_cast(iv) * 2048 bit-equivalent
+%ptr = pto.addptr %base, %elements
+```
+
+memory operation 原样使用 `%ptr`。如果再需要当前 Bisheng 的 post-update 形态，C11 fold
+把这个 `addptr` 吸收到 operation offset；这是第二个独立测试。
+
+### 6.2 C02：常量和动态项一起归一
+
+```text
+B = 8192 + iv * 4096
+E = 2
+Q = 4096 + iv * 2048
+```
+
+不需要设计“新 base 选 8192 还是 0”的策略。唯一 canonical base 是 0（常量场景；
+运行时原子基址见 6.5/C15），完整商进入 `addptr`。这消除了 invariant-term
+selection、profitability search 和多个等价 normal form。
+
+### 6.3 C06/C07：两类拒绝原因不同
+
+C06 可能在数学上总是偶数，但当前表达式模型不能穿过 trunc 证明同一个有限宽函数，属于
+**proof unavailable**。C07 的常数项 1 明确不可被 2 整除，属于 **no exact typed form**。
+诊断和单元测试应区分两者，避免以后把分析能力不足误认为语义上不可能。
+
+### 6.4 C13：pointer-derived integer
+
+即使某个 PTO 输入在数值上看起来可以整除，也不把 pointer-to-integer-to-pointer
+round-trip 纳入本规范化。它需要单独的 PTO provenance/round-trip 合同，而不是依赖
+LLVM lowering 的事实。当前设计明确保持该 case 不变；这不会影响 Issue #591，因为
+Issue #591 的整数表达式没有 pointer-derived leaf。
+
+### 6.5 C15：运行时原子基址
+
+**转换输入**——`castptr` 的整数输入含不可归零的运行时原子：
+
+```mlir
+%param : i64                        // kernel 参数，UB 基址
+%bytes = arith.muli %iv, %c4096 : i64
+%addr  = arith.addi %param, %bytes : i64
+%base  = pto.castptr %addr : i64 -> !pto.ptr<f16, ub>
+```
+
+`%param` 是原子叶子（block argument），系数 1 不能被 `E = 2` 整除，按 3.2 的原子根
+例外保留为规范根：输出为 `addptr(castptr(%param), iv * 2048)`。该形态的等价性不需要
+证明 `%param` 本身可整除；分析与发射都只依赖「根循环不变」：
+
+- `PTOAddressAnalysis::getPointerDelta` 对循环外定义的 root 返回 0
+  （`PTOAddressAnalysis.cpp` 的 `isDefinedOutsideOfLoop` 分支）；
+- `VPTOSoftPostUpdate` 的 post-update 前提是 base 循环外（`VPTOSoftPostUpdate.cpp`
+  的 `isDefinedOutsideOfLoop` 检查），不要求根是常量 0。
+
+这证明规范形是「规范根 + addptr(唯一商)」而不仅是「castptr(0) + addptr」：
+`castptr(0)` 是整条可整除时的退化特例，`castptr(%param)` 是纯原子根退化特例。
+C15 同时是 §9.4 规范形表述的收口用例。
+
+**C15 实测（2026-08-25，A5 bisheng 工具链）**：以下验证针对**规范形态的发射行为**
+（`addptr(castptr(%param), iv*2048)` 经 C11 fold 前/后的两个变体；转换本身由
+C01-C14 的 lit 覆盖）。`.work/c15-investigation/` 记录了 `castptr(%param)` 根两种
+形态的 LLVM IR 与 device 二进制对比：
+
+- `c15_op_offset.ll`（根 + operation offset，SoftPostUpdate 开启）循环体为
+  `@llvm.hivm.vldsx2.post.v128f16(ptr %p, i32 4096, ...)`，返回三元组含更新后指针
+  并被 `scf.for` phi 循环携带；与 `castptr(0)` 根对照 `ctrl0_root.ll` 的 intrinsic
+  完全同构，仅 base 装载（`inttoptr %param` vs `null`）不同。
+- 两者经 bisheng（`--cce-aicore-arch=dav-c310-vec --cce-aicore-only
+  -cce-bitcode-is-aicore`）编出的 device `.text` 同为 0x120 字节，prologue 逐字节
+  相同，循环体结构一致 → `castptr(%param)` 根 **同样触发 VLDS post-update**。
+- 对照 `c15_addptr.ll`（根 + addptr，未 fold）为 `@llvm.hivm.vldsx2.v128f16`
+  （非 post）+ 循环内 `mul`/`getelementptr`，device `.text` 为 0x128 字节（多 8
+  字节 = 一条地址更新指令，对应 `VLDI + SADD` 形态）→ 不触发 post-update。
+
+结论：post-update 触发只依赖 base 循环不变，与根是否为常量 0 无关；C11 fold
+仍是触发 post-update 的必要步骤。
+
+## 7. 测试分层
+
+### 7.1 Analysis/helper 单元测试
+
+直接测试 exact quotient 和 typed round trip，不经过 memory operation：
+
+- 常数、正负系数、多个线性项全部可整除；
+- 常数项或任一系数不可整除；
+- cast atom 保留；
+- 原子根例外：系数 1 的原子叶子不可整除时保留为规范根；系数 `|K| != 1` 且不可
+  整除时整体拒绝；
+- trunc 外层保持 opaque；
+- i32/i64/index 的可证明和不可证明 round trip；
+- 带 overflow flag 的 source-defined-domain refinement；
+- element size 1、2、4、8；
+- pointer-derived leaf 必须保持 opaque 并被 canonicalizer 拒绝。
+
+这些测试锁定证明器，而不是某个后端输出。
+
+### 7.2 Canonicalization lit 测试
+
+实现 pass 时以 C01-C15 构造一个 lit 文件，检查：
+
+```text
+accepted:
+  integer castptr 消失
+  只生成一个规范根 castptr（常量 0 或不可归零原子输入）
+  quotient 只生成一次
+  addptr element type/address space 不变
+
+rejected:
+  原 castptr 和整数表达式逐字形态保持
+  不残留部分创建的 zero/base/arithmetic
+```
+
+同一文件至少包含一个无 memory user 的 pointer cast round-trip 和一个多 users case，防止
+实现重新滑向“按 memory operation 做地址 rebase”。
+
+### 7.3 地址分析测试
+
+规范化后运行 `-pto-print-address-analysis`：
+
+- C01 的 root 是 `castptr(0)`；
+- C15 的 root 是 `castptr(%param)`，element offset 是 `iv * 2048`；
+- element offset 是 `iv * 2048`；
+- f16 byte delta 是 element delta 的 2 倍；
+- operation 自身已有的 offset 仍独立存在；
+- SoftPostUpdate 不需要 raw-integer matcher。
+
+### 7.4 当前可运行的 baseline
+
+仓库中的 `test/lit/vpto/integer_address_forms_baseline.pto` 已构造并保留三种输入：
+
+1. integer-backed `castptr`；
+2. `castptr(0) + typed addptr`；
+3. `castptr(0) + operation offset`。
+
+它检查当前 LLVM handoff 的真实差异，并用 store sink 防止被测 load/address chain 被 DCE。
+该测试是实现前 baseline，不是地址规范化正确性的证明。
+
+### 7.5 Bisheng 与运行时集成测试
+
+后端观察单独测试，不写入 analysis 合同：
+
+| PTO 输入 | 当前 Bisheng 观察 |
+| --- | --- |
+| integer-backed `castptr` | `VLDI + SADD` |
+| `castptr(0) + addptr` | 仍为 `VLDI + SADD` |
+| `castptr(0) + operation offset` | `VLDS ... post-update` |
+| `castptr(%param) + operation offset`（C15） | `VLDS ... post-update`（2026-08-25 实测，见 6.5） |
+| `castptr(%param) + addptr`（C15 未 fold） | `VLDI + SADD`（2026-08-25 实测，见 6.5） |
+
+此外用边界输入对原/新 kernel 做 differential runtime test：0、1、最大合法 loop index、
+含常量 base 的 C02、含运行时基址的 C15、多个 element type。汇编选择是收益回归；
+访问结果相同才是语义回归。
+
+## 8. Pass 位置
+
+地址规范化必须在 integer-backed `castptr` 已形成之后、所有 typed-address consumers 之前：
+
+```text
+形成 VPTO pointer SSA
+  -> integer address canonicalization
+  -> optional addptr absorption canonicalization
+  -> PTOAddressAnalysis consumers
+  -> LLVM lowering
+```
+
+这是数据依赖，不把任何具体 post-update pass 写进设计合同。rewrite 修改 pointer SSA 后，
+MLIR analysis preservation 应如实声明失效；下游自然重建 `PTOAddressAnalysis`。
+
+闭环 Issue #591 需要 C11 fold 的输出进入 post-update 阶段，因此以
+`VPTOSoftPostUpdate` 默认启用为前提（主线 #1330 已将 `--enable-vpto-soft-postupdate`
+默认开启）。C11 fold 必须在 post-update pass 之前执行；若 post-update 被显式关闭，
+第一步规范化仍改善分析可见性，但不改变发射形态（§7.5 第二行证据）。
+
+## 9. 实现完成条件
+
+实现只有同时满足以下条件才算闭合：
+
+1. C01-C15 均有 executable lit/unit case，正例和反例都有；
+2. exact quotient 与 typed round-trip proof 分开测试；
+3. canonicalizer 不查询 memory operation 或 LLVM lowering；
+4. 输出只有「规范根 + addptr」这一种 pointer normal form（规范根为 `castptr(0)`
+   或不可归零原子输入，见 3.2；`Q` 恒为 0 时退化为纯根 `castptr(R)`）；
+5. 规则要求 `Q` 非平凡（同时排除常量 0 与纯原子输入），规范化终止于规范形；
+   非整除和证明不足均原子失败，不留下半成品 IR；
+6. `PTOAddressAnalysis` 无第二套 integer-backed root/matcher；
+7. addptr absorption 是独立 pattern，可单独开关和测试；
+8. Issue #591 的 PTO IR、address-analysis 输出、Bisheng 汇编和 differential result 四层
+   证据均通过；
+9. 性能验收与 Issue #591 的目标绑定：rolled tile 循环体内 0 个 `RV_SADD`/`RV_SMOVK`
+   标量地址 op（对照 fully-unrolled 形态），且 rolled 形态的 RVEC cycles 不劣于
+   unrolled 形态（目标 ≤1.1×）；WHT N=512 的 RVEC cycle 对比作为回归记录。
+
+**性能验收实测（2026-08-27，camodel/dav_3510 仿真）**：`test/vpto/cases/micro-op/
+issue591-perf-rolled` 与 `issue591-perf-unrolled` 两个 case（16 tile，f16，
+`castptr(t*4096)` 读 UB + vadd + 写回，跑 `run_host_vpto_validation.sh`）：
+
+| 指标 | rolled（规范化后） | unrolled（基线） |
+| --- | --- | --- |
+| `kernal total ticks` | 3959 | 3957 |
+| `rvec_veccore0_busy_cycle` | 70 | 71 |
+| 循环内 `RV_SADD`/`RV_SMOVK` | **0** | 2（循环外 prologue） |
+| `RV_VLDS` / `RV_VLDI` | 16 / 0 | 15 / 1 |
+
+rolled 与 unrolled 性能持平（1.0005× ≤ 1.1×），循环体内零标量地址 op——验收
+达成。注意：本地 build 的 `VPTOSoftPostUpdate` 对**动态 trip count** 的 `scf.for`
+不转换（`UnknownIterationDomain`，缺 hw-native-sys #1330 的修复），验收 case 使用
+常量 trip count；动态 trip 的 post-update 依赖 #1330 或后续改进。
+
+## 10. 调研依据与事实证据
+
+- LLVM [GetElementPtr FAQ](https://llvm.org/docs/GetElementPtr.html) 说明 GEP 与整数地址算术
+  的 object/provenance 边界；本设计据此只在 PTO 已声明的数值地址空间中使用 zero base。
+- LLVM [LangRef pointer aliasing rules](https://llvm.org/docs/LangRef.html#pointer-aliasing-rules)
+  说明通用 LLVM pointer 不能仅凭数值相等任意重建；这也是 C08 的来源。
+- MLIR [Data Layout](https://mlir.llvm.org/docs/DataLayout/) 将 index bitwidth 作为 target
+  语义；本设计从 PTO target/data layout 取得位宽，不从 lowering 结果反推。
+- LLVM ScalarEvolution 与 `SCEVExpander` 分离符号证明和 SSA 物化；本设计同样把 exact
+  quotient proof 与 reifier 分开。
+- 当前仓库事实来自 `PTOValueEvolutionAnalysis`、`PTOAddressAnalysis`、
+  `VPTOAddressSemantics` 和 `VPTOLLVMEmitterHelper` 的实际实现。
+- `.work/issue591-postupdate/` 中的 `original`、`addptr`、`normalized` LLVM/汇编产物记录了
+  三种形态的 Bisheng A5 实测结果；它们是诊断证据，不作为提交测试。
+- `.work/c15-investigation/` 记录了 C15（`castptr(%param)` 运行时根）两种形态的
+  LLVM IR 与 device 二进制对比（2026-08-25 实测，见 6.5）；同为诊断证据。
+- 设备编译参数来自 `tools/ptoas/ObjectEmission.cpp` 的
+  `compileDeviceLLVMToObject()`（`bisheng --cce-aicore-arch=dav-c310-vec
+  --cce-aicore-only -cce-bitcode-is-aicore -c -x ir`）。
+
+设计的最终边界可以概括为：
+
+```text
+地址规范化只证明 pointer value 等价；
+地址分析只消费 canonical pointer SSA；
+operation fold 只合并同单位 offset；
+lowering 和 post-update 只作为独立集成观察。
+```
diff --git a/include/PTO/Analysis/PTOAddressAnalysis.h b/include/PTO/Analysis/PTOAddressAnalysis.h
index f20b7b78c3..ce03949e53 100644
--- a/include/PTO/Analysis/PTOAddressAnalysis.h
+++ b/include/PTO/Analysis/PTOAddressAnalysis.h
@@ -32,6 +32,13 @@ class AnalysisManager;
 
 namespace pto {
 
+/// Element storage size in bytes for an integer/float element type, or for a
+/// pointer/view-typed SSA value. Shared by address analysis and the integer
+/// address canonicalization rewrite so the element-size rule lives in one
+/// place.
+std::optional<int64_t> getPTOElementBytes(Type elementType);
+std::optional<int64_t> getPTOElementBytes(Value pointer);
+
 struct PTOTypedAddressOffset {
   Value sourceValue;
   PTOTypedExprRef value;
diff --git a/include/PTO/IR/VPTOOps.td b/include/PTO/IR/VPTOOps.td
index bda2b16ab5..8e44ca68d5 100644
--- a/include/PTO/IR/VPTOOps.td
+++ b/include/PTO/IR/VPTOOps.td
@@ -3659,7 +3659,8 @@ def PTO_VmulconvOp : PTO_VectorMicroOp<"vmulconv", [Pure]> {
 }
 
 def PTO_Vstsx2Op : PTO_VectorMicroOp<"vstsx2", [
-    DeclareOpInterfaceMethods<MemoryEffectsOpInterface>
+    DeclareOpInterfaceMethods<MemoryEffectsOpInterface>,
+    VPTOAddressSemanticsOpInterface
   ]> {
   let arguments = (ins
     PTO_VectorType:$low,
diff --git a/include/PTO/Transforms/Passes.h b/include/PTO/Transforms/Passes.h
index 1a8dc14307..cbc430d6d4 100644
--- a/include/PTO/Transforms/Passes.h
+++ b/include/PTO/Transforms/Passes.h
@@ -102,6 +102,8 @@ std::unique_ptr<Pass> createPTOInferVPTOVecScopePass();
 std::unique_ptr<Pass> createVPTOExpandWrapperOpsPass();
 std::unique_ptr<Pass> createVPTOSoftPostUpdatePass();
 std::unique_ptr<Pass> createPTOPrintAddressAnalysisPass();
+std::unique_ptr<Pass> createPTOIntegerAddressCanonicalizationPass();
+std::unique_ptr<Pass> createPTOAbsorbAddPtrPass();
 std::unique_ptr<Pass> createPTOVPTOPtrBoundaryPass();
 std::unique_ptr<Pass>
 createPTOLowLevelLoopFusionPass(const PTOLowLevelLoopFusionOptions &options = {});
diff --git a/include/PTO/Transforms/Passes.td b/include/PTO/Transforms/Passes.td
index 13a084cbc4..74e91928ff 100644
--- a/include/PTO/Transforms/Passes.td
+++ b/include/PTO/Transforms/Passes.td
@@ -1141,8 +1141,8 @@ def VMILowerUnifiedToLegacy : Pass<"vmi-lower-unified-to-legacy", "ModuleOp"> {
     Ops lowered (Category A–C6):
       A:  vci, vinterpret_cast, vsel, vbrc  →  iota, bitcast, select, broadcast/group_broadcast
       B:  vadd/vsub/vmul/vdiv/vmin/vmax/vand/vor/vxor/vshl/vshr (masked binary)
-          vneg/vabs/vsqrt/vexp/vln/vrelu/vnot (masked unary)
-          →  binary ops discard mask/pmode; unary ops preserve zero mode with select;
+          vneg/vsqrt/vexp/vln/vrelu/vnot (masked unary)
+          →  mask-less legacy operations;
              vshr selects shrui for explicit unsigned elements and shrsi otherwise
       C1: vcmp/vcmps  →  legacy cmp + select
       C2: vcvt        →  legacy extf/truncf/fptosi/sitofp/extsi/extui/trunci
@@ -1150,8 +1150,8 @@ def VMILowerUnifiedToLegacy : Pass<"vmi-lower-unified-to-legacy", "ModuleOp"> {
       C4: pset/pge    →  create_mask/create_group_mask
       C6: vcadd/vcmax/vcmin  →  legacy reduce variants
 
-    Ops NOT lowered (no legacy equivalent — require direct VMIToVPTO 1:N patterns):
-      plt, vadds, vmuls, vmaxs, vmins, vshls, vshrs, vhist, vintlv, vdintlv, vselr,
+    Ops kept unified for direct VMIToVPTO 1:N patterns:
+      plt, vabs, vadds, vmuls, vmaxs, vmins, vshls, vshrs, vhist, vintlv, vdintlv, vselr,
       vgather, vgatherb, vscatter,
       vexpdif, vaxpy, vlrelu, vprelu, vmull, vmula
   }];
@@ -1300,6 +1300,44 @@ def PTOPrintAddressAnalysis
                            "mlir::scf::SCFDialect"];
 }
 
+def PTOIntegerAddressCanonicalization
+    : Pass<"pto-canonicalize-integer-address", "func::FuncOp"> {
+  let summary = "Canonicalize integer-backed pto.castptr into castptr(root) + addptr(quotient)";
+  let description = [{
+    Rewrites `castptr(B)` with an integer byte-address input `B` into the
+    canonical pointer form `addptr(castptr(R), Q)` where `R` is the canonical
+    root (constant 0, or a single non-divisible atom leaf such as a runtime
+    base-address kernel parameter) and `Q` is the exact element quotient
+    `(B - R) / sizeof(element)`. The rewrite only matches zero-origin integral
+    address spaces (A5 UB). It refuses non-linear inputs, non-exact quotients,
+    non-unit-coefficient or multiple non-divisible atoms, and inputs whose
+    width does not round-trip into index. The rule requires a non-trivial
+    quotient (constant-zero and pure-atom inputs stay untouched) so the
+    rewrite converges to the normal form in one pass.
+  }];
+  let constructor = "mlir::pto::createPTOIntegerAddressCanonicalizationPass()";
+  let dependentDialects = ["mlir::func::FuncDialect",
+                           "mlir::pto::PTODialect",
+                           "mlir::arith::ArithDialect"];
+}
+
+def PTOAbsorbAddPtr
+    : Pass<"pto-absorb-addptr", "func::FuncOp"> {
+  let summary = "Absorb addptr element offsets into VPTO memory op offsets";
+  let description = [{
+    Canonical fold `op(addptr(base, A), O) -> op(base, A + O)` for VPTO memory
+    operations whose offset unit is Element and whose element type matches the
+    addptr. Legality comes only from VPTOAddressSemanticsOpInterface:
+    current access with an Element offset, no updated-base post-update form,
+    and a no-loss index addition. This is the backend-shape fold that lets the
+    post-update consumer see the affine offset directly on the operation.
+  }];
+  let constructor = "mlir::pto::createPTOAbsorbAddPtrPass()";
+  let dependentDialects = ["mlir::func::FuncDialect",
+                           "mlir::pto::PTODialect",
+                           "mlir::arith::ArithDialect"];
+}
+
 def PTOVPTOPtrBoundary
     : Pass<"pto-vpto-ptr-boundary", "ModuleOp"> {
   let summary =
diff --git a/lib/PTO/Analysis/PTOAddressAnalysis.cpp b/lib/PTO/Analysis/PTOAddressAnalysis.cpp
index 104b9c2259..0cd8619f87 100644
--- a/lib/PTO/Analysis/PTOAddressAnalysis.cpp
+++ b/lib/PTO/Analysis/PTOAddressAnalysis.cpp
@@ -21,25 +21,6 @@ namespace {
 
 static constexpr int64_t kBlockSizeBytes = 32;
 
-static std::optional<int64_t> getElementBytes(Value pointer) {
-  Type elementType;
-  if (auto pointerType = dyn_cast<PtrType>(pointer.getType())) {
-    elementType = pointerType.getElementType();
-  } else if (auto memrefType = dyn_cast<BaseMemRefType>(pointer.getType())) {
-    elementType = memrefType.getElementType();
-  } else {
-    return std::nullopt;
-  }
-  if (!elementType || !elementType.isIntOrFloat()) {
-    return std::nullopt;
-  }
-  unsigned bitWidth = elementType.getIntOrFloatBitWidth();
-  if (bitWidth == 0 || bitWidth % mlir::pto::kValue8 != 0) {
-    return std::nullopt;
-  }
-  return static_cast<int64_t>(bitWidth / mlir::pto::kValue8);
-}
-
 static std::optional<int64_t> getUnitBytes(Operation *operation,
                                            VPTOAddressUnit unit,
                                            int64_t elementBytes) {
@@ -91,6 +72,29 @@ static bool isZero(const PTOTypedExprRef &expression) {
 
 } // namespace
 
+std::optional<int64_t> mlir::pto::getPTOElementBytes(Type elementType) {
+  if (!elementType || !elementType.isIntOrFloat()) {
+    return std::nullopt;
+  }
+  unsigned bitWidth = elementType.getIntOrFloatBitWidth();
+  if (bitWidth == 0 || bitWidth % mlir::pto::kValue8 != 0) {
+    return std::nullopt;
+  }
+  return static_cast<int64_t>(bitWidth / mlir::pto::kValue8);
+}
+
+std::optional<int64_t> mlir::pto::getPTOElementBytes(Value pointer) {
+  Type elementType;
+  if (auto pointerType = dyn_cast<PtrType>(pointer.getType())) {
+    elementType = pointerType.getElementType();
+  } else if (auto memrefType = dyn_cast<BaseMemRefType>(pointer.getType())) {
+    elementType = memrefType.getElementType();
+  } else {
+    return std::nullopt;
+  }
+  return getPTOElementBytes(elementType);
+}
+
 PTOAddressAnalysis::PTOAddressAnalysis(func::FuncOp func,
                                        AnalysisManager &analysisManager)
     : func(func),
@@ -118,7 +122,7 @@ PTOAddressAnalysis::getAddresses(Operation *operation) {
   VPTOAddressSemantics contract = semantics.getVPTOAddressSemantics();
   for (const VPTOAddressAccess &access : contract.currentAccesses) {
     Value base = access.baseOperand->get();
-    auto elementBytes = getElementBytes(base);
+    auto elementBytes = getPTOElementBytes(base);
     if (!elementBytes) {
       return PTOAnalysisResult<SmallVector<PTOAddressExpr>>::unknown(
           PTOAnalysisUnknownReason::UnknownElementSize);
@@ -132,7 +136,7 @@ PTOAddressAnalysis::getAddresses(Operation *operation) {
 
     while (auto addPointer =
                address.rootOrBase.getDefiningOp<AddPtrOp>()) {
-      auto parentElementBytes = getElementBytes(addPointer.getPtr());
+      auto parentElementBytes = getPTOElementBytes(addPointer.getPtr());
       if (!parentElementBytes || *parentElementBytes != *elementBytes) {
         break;
       }
diff --git a/lib/PTO/IR/VPTOAddressSemantics.cpp b/lib/PTO/IR/VPTOAddressSemantics.cpp
index 34b629251e..423dbfed2a 100644
--- a/lib/PTO/IR/VPTOAddressSemantics.cpp
+++ b/lib/PTO/IR/VPTOAddressSemantics.cpp
@@ -95,6 +95,15 @@ mlir::pto::getDefaultVPTOAddressSemantics(Operation *operation) {
             postUpdate(base, &offset, VPTOAddressUnit::Element,
                        op.getUpdatedBase())};
       })
+      .Case<Vstsx2Op>([](Vstsx2Op op) {
+        // vstsx2 has no result, so it cannot carry a post-update updated base;
+        // only the current access (destination + element offset) is modeled.
+        OpOperand &base = op.getDestinationMutable();
+        OpOperand &offset = op.getOffsetMutable();
+        return VPTOAddressSemantics{
+            {oneAccess(base, offset, VPTOAddressUnit::Element)},
+            std::nullopt};
+      })
       .Case<VstusOp>([](VstusOp op) {
         OpOperand &base = op.getBaseMutable();
         return VPTOAddressSemantics{
diff --git a/lib/PTO/Transforms/CMakeLists.txt b/lib/PTO/Transforms/CMakeLists.txt
index eb92bf1edb..eb6889e141 100644
--- a/lib/PTO/Transforms/CMakeLists.txt
+++ b/lib/PTO/Transforms/CMakeLists.txt
@@ -48,6 +48,8 @@ add_mlir_dialect_library(PTOTransforms
   VPTOMaskSimplify.cpp
   VPTOExpandWrapperOps.cpp
   VPTOSoftPostUpdate.cpp
+  PTOIntegerAddressCanonicalization.cpp
+  PTOAbsorbAddPtr.cpp
   PTOPrintAddressAnalysis.cpp
   VPTOScheduler/VPTORegPressureTracker.cpp
   VPTOScheduler/VPTOSchedBoundary.cpp
diff --git a/lib/PTO/Transforms/PTOAbsorbAddPtr.cpp b/lib/PTO/Transforms/PTOAbsorbAddPtr.cpp
new file mode 100644
index 0000000000..1db00d72a5
--- /dev/null
+++ b/lib/PTO/Transforms/PTOAbsorbAddPtr.cpp
@@ -0,0 +1,123 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+//===- PTOAbsorbAddPtr.cpp -----------------------------------------------===//
+//
+// Backend-shape canonical fold for VPTO memory operations:
+//
+//   op(addptr(base, A), O)  ->  op(base, A + O)
+//
+// Bisheng only emits a post-update load/store when the affine offset sits on
+// the operation itself; an offset hidden inside an `addptr` is re-materialized
+// as `VLDI + SADD` per iteration. This fold is the "addptr absorption" rule of
+// docs/designs/vpto-integer-address-canonicalization-design-zh.md and must run
+// before the post-update consumer (VPTOSoftPostUpdate).
+//
+// Legality comes only from VPTOAddressSemanticsOpInterface: a current access
+// with an Element-unit offset, a base that is an `addptr` result, no
+// updated-base post-update form, and a no-loss index addition.
+//
+//===----------------------------------------------------------------------===//
+
+#include "PTO/IR/VPTOAddressSemantics.h"
+#include "PTO/Transforms/Passes.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/IR/PatternMatch.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
+
+namespace mlir {
+namespace pto {
+#define GEN_PASS_DEF_PTOABSORBADDPTR
+#include "PTO/Transforms/Passes.h.inc"
+} // namespace pto
+} // namespace mlir
+
+using namespace mlir;
+using namespace mlir::pto;
+
+namespace {
+
+struct AbsorbAddPtrIntoOpOffset final : public RewritePattern {
+  AbsorbAddPtrIntoOpOffset(MLIRContext *context)
+      : RewritePattern(MatchAnyOpTypeTag(), /*benefit=*/1, context) {}
+
+  LogicalResult matchAndRewrite(Operation *op,
+                                PatternRewriter &rewriter) const override {
+    auto semantics = dyn_cast<pto::VPTOAddressSemanticsOpInterface>(op);
+    if (!semantics) {
+      return rewriter.notifyMatchFailure(op, "no address semantics");
+    }
+    VPTOAddressSemantics contract = semantics.getVPTOAddressSemantics();
+    if (contract.currentAccesses.empty()) {
+      return rewriter.notifyMatchFailure(op, "no current access");
+    }
+    const VPTOAddressAccess &access = contract.currentAccesses.front();
+    if (!access.offset || access.offset->unit != VPTOAddressUnit::Element) {
+      return rewriter.notifyMatchFailure(op, "offset is not element-unit");
+    }
+    // Post-update form: the offset operand denotes the after-access advance,
+    // not a current access; never fold those.
+    if (contract.postUpdate && contract.postUpdate->updatedBase) {
+      return rewriter.notifyMatchFailure(op, "already in post-update form");
+    }
+
+    Value base = access.baseOperand->get();
+    auto addptr = base.getDefiningOp<pto::AddPtrOp>();
+    if (!addptr) {
+      return rewriter.notifyMatchFailure(op, "base is not an addptr");
+    }
+    // Only fold addptr whose base is itself an integer-backed castptr — the
+    // canonical shape produced by pto-canonicalize-integer-address. Arbitrary
+    // addptr chains over user pointers are left alone: VPTOSoftPostUpdate has
+    // its own sequential base-chain handling for those, and folding them here
+    // would change (or destroy) that post-update structure.
+    if (!addptr.getPtr().getDefiningOp<pto::CastPtrOp>()) {
+      return rewriter.notifyMatchFailure(
+          op, "addptr base is not a castptr (sequential chain handled by "
+              "soft post-update)");
+    }
+    if (!addptr.getOffset().getType().isIndex() ||
+        !access.offset->operand->get().getType().isIndex()) {
+      return rewriter.notifyMatchFailure(op, "offsets are not index-typed");
+    }
+    // addptr and the op share the same pointer element type because the op's
+    // base operand *is* the addptr result (AllTypesMatch on AddPtrOp).
+
+    Value combined = rewriter.create<arith::AddIOp>(
+        op->getLoc(), addptr.getOffset(), access.offset->operand->get());
+
+    rewriter.modifyOpInPlace(op, [&]() {
+      access.baseOperand->set(addptr.getPtr());
+      access.offset->operand->set(combined);
+    });
+    if (addptr->use_empty()) {
+      rewriter.eraseOp(addptr);
+    }
+    return success();
+  }
+};
+
+struct PTOAbsorbAddPtrPass final
+    : public pto::impl::PTOAbsorbAddPtrBase<PTOAbsorbAddPtrPass> {
+  void runOnOperation() override {
+    func::FuncOp func = getOperation();
+    RewritePatternSet patterns(&getContext());
+    patterns.add<AbsorbAddPtrIntoOpOffset>(&getContext());
+    if (failed(applyPatternsAndFoldGreedily(func, std::move(patterns)))) {
+      signalPassFailure();
+    }
+  }
+};
+
+} // namespace
+
+std::unique_ptr<Pass> mlir::pto::createPTOAbsorbAddPtrPass() {
+  return std::make_unique<PTOAbsorbAddPtrPass>();
+}
diff --git a/lib/PTO/Transforms/PTOIntegerAddressCanonicalization.cpp b/lib/PTO/Transforms/PTOIntegerAddressCanonicalization.cpp
new file mode 100644
index 0000000000..5f7ae94cf3
--- /dev/null
+++ b/lib/PTO/Transforms/PTOIntegerAddressCanonicalization.cpp
@@ -0,0 +1,433 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+//===- PTOIntegerAddressCanonicalization.cpp -----------------------------===//
+//
+// Canonicalizes integer-backed `pto.castptr` into the canonical pointer form
+//
+//   castptr(byte_address)  ->  addptr(castptr(R), Q)
+//
+// where `R` is the canonical root (constant 0, or a single non-divisible atom
+// leaf such as a runtime base-address kernel parameter) and `Q` is the exact
+// element quotient `(B - R) / sizeof(element)`. The rewrite only matches
+// zero-origin integral address spaces (A5 UB) and requires a non-trivial
+// quotient so it converges to the normal form in one pass (constant-zero and
+// pure-atom inputs stay untouched).
+//
+// This is the "integer address canonicalization" rule of
+// docs/designs/vpto-integer-address-canonicalization-design-zh.md. The addptr
+// absorption fold into memory-op offsets lives in PTOAbsorbAddPtr.cpp.
+//
+//===----------------------------------------------------------------------===//
+
+#include "PTO/Analysis/PTOAddressAnalysis.h"
+#include "PTO/Analysis/PTOValueEvolutionAnalysis.h"
+#include "PTO/IR/PTO.h"
+#include "PTO/Transforms/Passes.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/SCF/IR/SCF.h"
+#include "mlir/IR/BuiltinAttributes.h"
+#include "mlir/IR/PatternMatch.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
+#include "llvm/ADT/STLExtras.h"
+
+namespace mlir {
+namespace pto {
+#define GEN_PASS_DEF_PTOINTEGERADDRESSCANONICALIZATION
+#include "PTO/Transforms/Passes.h.inc"
+} // namespace pto
+} // namespace mlir
+
+using namespace mlir;
+using namespace mlir::pto;
+
+namespace {
+
+// Current PTO targets fix index at 64 bits (PTOValueEvolutionAnalysis uses the
+// same assumption); the round-trip proof below requires the input width to
+// equal this width so CastIndex is a same-width conversion.
+static constexpr unsigned kIndexBitWidth = 64;
+
+/// Convert a signedness-carrying integer to the equivalent signless carrier
+/// required by arith operations. Reuse the source of an existing no-op bridge
+/// when possible so canonicalization does not grow redundant casts.
+static Value getSignlessIntegerCarrier(PatternRewriter &rewriter, Location loc,
+                                       Value value) {
+  auto integerType = dyn_cast<IntegerType>(value.getType());
+  if (!integerType || integerType.isSignless()) {
+    return value;
+  }
+  auto carrierType = rewriter.getIntegerType(integerType.getWidth());
+  if (auto cast = value.getDefiningOp<UnrealizedConversionCastOp>()) {
+    if (cast.getInputs().size() == 1 &&
+        cast.getInputs().front().getType() == carrierType) {
+      return cast.getInputs().front();
+    }
+  }
+  return rewriter
+      .create<UnrealizedConversionCastOp>(loc, carrierType, value)
+      .getResult(0);
+}
+
+/// Rebuild a cast with `kind` into `targetType`. Same-type casts are skipped
+/// so absorbed index extensions do not leave a redundant index_cast behind.
+/// arith's ext/si/trunc ops require fixed-width targets, so an index target
+/// is reached through index_cast/index_castui (which preserve the same
+/// zero/sign extension semantics).
+static Value createCastToTarget(PatternRewriter &rewriter, Location loc,
+                                PTOCastKind kind, Value input,
+                                Type targetType) {
+  if (input.getType() == targetType) {
+    return input;
+  }
+  if (targetType.isIndex()) {
+    switch (kind) {
+    case PTOCastKind::IndexCast:
+    case PTOCastKind::ExtSI:
+      return rewriter.create<arith::IndexCastOp>(loc, targetType, input);
+    case PTOCastKind::IndexCastUI:
+    case PTOCastKind::ExtUI:
+      return rewriter.create<arith::IndexCastUIOp>(loc, targetType, input);
+    case PTOCastKind::TruncI:
+      // Truncation into index is ill-defined; refuse.
+      return {};
+    }
+  }
+  switch (kind) {
+  case PTOCastKind::IndexCast:
+    return rewriter.create<arith::IndexCastOp>(loc, targetType, input);
+  case PTOCastKind::IndexCastUI:
+    return rewriter.create<arith::IndexCastUIOp>(loc, targetType, input);
+  case PTOCastKind::ExtSI:
+    return rewriter.create<arith::ExtSIOp>(loc, targetType, input);
+  case PTOCastKind::ExtUI:
+    return rewriter.create<arith::ExtUIOp>(loc, targetType, input);
+  case PTOCastKind::TruncI:
+    return rewriter.create<arith::TruncIOp>(loc, targetType, input);
+  }
+  return {};
+}
+
+/// Materialize a synthetic typed expression as SSA values directly in
+/// `targetType` (the index domain). Leaves (nodes carrying a sourceValue) are
+/// reused verbatim when their type already matches; an index<->int extension
+/// cast leaf is rebuilt from its cast input so the emitted offset stays a
+/// direct affine index expression (the post-update consumer needs that shape).
+/// Only synthetic constant/add/sub/mul nodes built by buildPTOTypedExpr are
+/// reconstructed.
+static Value materializeQuotient(PatternRewriter &rewriter, Location loc,
+                                 const PTOTypedExprRef &expr,
+                                 Type targetType) {
+  if (!expr) {
+    return {};
+  }
+  if (expr->sourceValue) {
+    Value src = expr->sourceValue;
+    if (src.getType() == targetType) {
+      return src;
+    }
+    if (expr->kind == PTOTypedExpr::Kind::Cast) {
+      // Rebuild the cast into the target type from its input, materializing
+      // the input in its own type first (preserves zero/sign semantics).
+      Type innerType = expr->lhs ? expr->lhs->type : Type();
+      Value inner = innerType ? materializeQuotient(rewriter, loc, expr->lhs,
+                                                    innerType)
+                              : Value();
+      if (!inner) {
+        return {};
+      }
+      return createCastToTarget(rewriter, loc, expr->castKind, inner,
+                                targetType);
+    }
+    // Plain integer leaf (e.g. an i64 atom inside the quotient): same-width
+    // conversion into index keeps the bit pattern.
+    if (src.getType().isIntOrIndex() && targetType.isIndex() &&
+        src.getType().getIntOrFloatBitWidth() == kIndexBitWidth) {
+      src = getSignlessIntegerCarrier(rewriter, loc, src);
+      return rewriter.create<arith::IndexCastOp>(loc, targetType, src);
+    }
+    return {};
+  }
+  switch (expr->kind) {
+  case PTOTypedExpr::Kind::Constant:
+    return rewriter.create<arith::ConstantOp>(
+        loc, rewriter.getIntegerAttr(targetType, expr->constant));
+  case PTOTypedExpr::Kind::Add: {
+    Value lhs = materializeQuotient(rewriter, loc, expr->lhs, targetType);
+    Value rhs = materializeQuotient(rewriter, loc, expr->rhs, targetType);
+    if (!lhs || !rhs) {
+      return {};
+    }
+    return rewriter.create<arith::AddIOp>(loc, lhs, rhs);
+  }
+  case PTOTypedExpr::Kind::Sub: {
+    Value lhs = materializeQuotient(rewriter, loc, expr->lhs, targetType);
+    Value rhs = materializeQuotient(rewriter, loc, expr->rhs, targetType);
+    if (!lhs || !rhs) {
+      return {};
+    }
+    return rewriter.create<arith::SubIOp>(loc, lhs, rhs);
+  }
+  case PTOTypedExpr::Kind::Mul: {
+    Value lhs = materializeQuotient(rewriter, loc, expr->lhs, targetType);
+    Value rhs = materializeQuotient(rewriter, loc, expr->rhs, targetType);
+    if (!lhs || !rhs) {
+      return {};
+    }
+    return rewriter.create<arith::MulIOp>(loc, lhs, rhs);
+  }
+  case PTOTypedExpr::Kind::Cast:
+  case PTOTypedExpr::Kind::Opaque:
+    // buildPTOTypedExpr never emits cast/opaque nodes; a source-backed node
+    // would have been returned above.
+    return {};
+  }
+  return {};
+}
+
+/// Strip the sourceValue from structural (add/sub/mul) nodes so
+/// normalizePTOLinearExpr expands them instead of treating the whole
+/// expression as one opaque atom. Cast/opaque/constant leaves keep their
+/// sourceValue and stay atomic, matching the design's cast-as-atom rule.
+static PTOTypedExprRef stripStructuralSourceValues(const PTOTypedExprRef &expr) {
+  if (!expr) {
+    return expr;
+  }
+  switch (expr->kind) {
+  case PTOTypedExpr::Kind::Add:
+    return makePTOAddExpr(stripStructuralSourceValues(expr->lhs),
+                          stripStructuralSourceValues(expr->rhs), expr->type);
+  case PTOTypedExpr::Kind::Sub:
+    return makePTOSubExpr(stripStructuralSourceValues(expr->lhs),
+                          stripStructuralSourceValues(expr->rhs), expr->type);
+  case PTOTypedExpr::Kind::Mul:
+    return makePTOMulExpr(stripStructuralSourceValues(expr->lhs),
+                          stripStructuralSourceValues(expr->rhs), expr->type);
+  default:
+    // Constant / Cast / Opaque stay atomic (and keep their sourceValue).
+    return expr;
+  }
+}
+
+struct CanonicalizeIntegerCastPtr final
+    : public OpRewritePattern<pto::CastPtrOp> {
+  CanonicalizeIntegerCastPtr(MLIRContext *context,
+                             PTOValueEvolutionAnalysis &valueEvolution)
+      : OpRewritePattern<pto::CastPtrOp>(context),
+        valueEvolution(valueEvolution) {}
+
+  LogicalResult matchAndRewrite(pto::CastPtrOp castOp,
+                                PatternRewriter &rewriter) const override {
+    auto ptrType = dyn_cast<pto::PtrType>(castOp.getResult().getType());
+    if (!ptrType) {
+      return rewriter.notifyMatchFailure(castOp, "result is not a ptr type");
+    }
+    // Only zero-origin integral address spaces are eligible (A5 UB prints as
+    // "ub" and maps to AddressSpace::VEC).
+    if (ptrType.getMemorySpace().getAddressSpace() != pto::AddressSpace::VEC) {
+      return rewriter.notifyMatchFailure(castOp, "not a zero-origin space");
+    }
+    Value input = castOp.getInput();
+    Type inputType = input.getType();
+    if (!inputType.isIntOrIndex()) {
+      return rewriter.notifyMatchFailure(castOp, "input is not an integer");
+    }
+    unsigned inputWidth =
+        inputType.isIndex() ? kIndexBitWidth
+                            : inputType.getIntOrFloatBitWidth();
+
+    // Design C14 requires the *quotient* to round-trip losslessly into index.
+    // Same-width (64-bit) inputs are trivially lossless. Narrower inputs
+    // (e.g. i32) are accepted only when the byte expression provably does not
+    // wrap in its own width: with the zero-extension semantics of LLVM
+    // inttoptr (design §2.2), 0 <= B < 2^inputWidth implies Q = B / E <
+    // 2^inputWidth and zext(Q) * E == zext(B) exactly in the 64-bit index
+    // domain. Wider inputs would truncate and are rejected.
+    auto roundTripsLosslessly = [&]() -> bool {
+      if (inputWidth == kIndexBitWidth) {
+        return true;
+      }
+      if (inputWidth > kIndexBitWidth) {
+        return false;
+      }
+      if (auto folded = foldPTOConstant(valueEvolution.getExpr(input))) {
+        // Constant byte expression: non-negative and fits the input width.
+        return *folded >= 0 &&
+               static_cast<uint64_t>(*folded) <
+                   (uint64_t{1} << (inputWidth - 1));
+      }
+      // Loop-carried byte expression: require a provable non-negative range
+      // whose signed upper bound leaves room in the input width.
+      scf::ForOp loop = castOp->getParentOfType<scf::ForOp>();
+      if (!loop) {
+        return false;
+      }
+      auto range = valueEvolution.getRange(input, loop);
+      if (!range) {
+        return false;
+      }
+      unsigned rangeWidth = std::max(1U, inputWidth);
+      return range.value->lowerInclusive.isNonNegative() &&
+             range.value->upperInclusive.ult(
+                 APInt(rangeWidth, 1).shl(inputWidth - 1));
+    };
+    if (!roundTripsLosslessly()) {
+      return rewriter.notifyMatchFailure(
+          castOp, "byte expression does not provably round-trip into index");
+    }
+
+    std::optional<int64_t> elementBytes =
+        getPTOElementBytes(ptrType.getElementType());
+    if (!elementBytes || *elementBytes <= 0) {
+      return rewriter.notifyMatchFailure(castOp, "unknown element size");
+    }
+
+    auto linear = normalizePTOLinearExpr(
+        stripStructuralSourceValues(valueEvolution.getExpr(input)));
+    if (!linear) {
+      return rewriter.notifyMatchFailure(castOp, "input is not linear");
+    }
+
+    // Pointer-derived integer leaves (ptr-to-int castptr / ptrtoint results)
+    // need a separate provenance contract and stay untouched (design C13).
+    for (const PTOLinearTerm &term : linear->terms) {
+      if (!term.atom || !term.atom->sourceValue) {
+        continue;
+      }
+      Operation *leafDef = term.atom->sourceValue.getDefiningOp();
+      if (isa_and_nonnull<pto::PtrToIntOp>(leafDef)) {
+        return rewriter.notifyMatchFailure(castOp,
+                                           "pointer-derived integer leaf");
+      }
+      if (auto ptrCast = dyn_cast_or_null<pto::CastPtrOp>(leafDef)) {
+        if (!isa<pto::PtrType>(ptrCast.getResult().getType())) {
+          return rewriter.notifyMatchFailure(
+              castOp, "pointer-derived integer leaf");
+        }
+      }
+    }
+
+    // Canonical-root selection: a coefficient not divisible by the element
+    // size cannot move into the offset. Only a single unit-coefficient atom
+    // leaf is supported as the root; everything else is rejected.
+    Value rootInteger = nullptr;
+    PTOLinearExpr quotient = *linear;
+    SmallVector<PTOLinearTerm> nonDivisible;
+    for (const PTOLinearTerm &term : quotient.terms) {
+      if (term.coefficient % *elementBytes != 0) {
+        nonDivisible.push_back(term);
+      }
+    }
+    if (!nonDivisible.empty()) {
+      if (nonDivisible.size() > 1) {
+        return rewriter.notifyMatchFailure(
+            castOp, "multiple non-divisible atoms have no unique root");
+      }
+      const PTOLinearTerm &rootTerm = nonDivisible.front();
+      // Design contract: only a unit-coefficient (+1) atom is the canonical
+      // root; negative or scaled atoms have no unique root/offset split.
+      if (rootTerm.coefficient != 1) {
+        return rewriter.notifyMatchFailure(
+            castOp, "non-unit-coefficient atom has no root/offset split");
+      }
+      if (!rootTerm.atom || !rootTerm.atom->sourceValue ||
+          !rootTerm.atom->sourceValue.getType().isIntOrIndex()) {
+        return rewriter.notifyMatchFailure(
+            castOp, "atom leaf is not a materializable integer");
+      }
+      rootInteger = rootTerm.atom->sourceValue;
+      Value rootLeaf = rootInteger;
+      llvm::erase_if(quotient.terms, [&](const PTOLinearTerm &term) {
+        return term.coefficient == rootTerm.coefficient &&
+               term.atom && term.atom->sourceValue == rootLeaf;
+      });
+    }
+
+    if (!dividePTOLinearExprExact(quotient, *elementBytes)) {
+      return rewriter.notifyMatchFailure(castOp, "no exact element quotient");
+    }
+    if (isZeroPTOLinearExpr(quotient)) {
+      return rewriter.notifyMatchFailure(
+          castOp, "quotient is trivial (already canonical)");
+    }
+
+    // Round-trip proof: inputWidth == index width makes CastIndex a
+    // same-width conversion, and every quotient leaf comes from the input
+    // expression, so `Q -> index -> element scaling` preserves the address
+    // bit pattern (mod 2^Waddr) because Q * E == B - R is an exact identity.
+    // The quotient is materialized directly in the index domain so the
+    // emitted offset is a direct affine index expression that the post-update
+    // consumer can analyze.
+    PTOTypedExprRef quotientExpr = buildPTOTypedExpr(quotient, inputType);
+    rewriter.setInsertionPoint(castOp);
+    Value indexValue = materializeQuotient(
+        rewriter, castOp.getLoc(), quotientExpr, rewriter.getIndexType());
+    if (!indexValue) {
+      return rewriter.notifyMatchFailure(castOp, "cannot materialize quotient");
+    }
+
+    // The canonical root (castptr(0) or castptr(%atom)) is hoisted above any
+    // enclosing scf.for so the post-update consumer sees a base defined
+    // outside the loop (VPTOSoftPostUpdate requires that). Hoisting is only
+    // legal when the root value itself is loop-invariant; an atom-root defined
+    // inside the loop must stay in place or the hoisted castptr would violate
+    // SSA dominance.
+    Operation *anchor = castOp.getOperation();
+    bool rootIsLoopInvariant = true;
+    while (scf::ForOp forOp = anchor->getParentOfType<scf::ForOp>()) {
+      if (rootInteger && !forOp.isDefinedOutsideOfLoop(rootInteger)) {
+        rootIsLoopInvariant = false;
+      }
+      anchor = forOp.getOperation();
+    }
+    rewriter.setInsertionPoint(rootIsLoopInvariant ? anchor
+                                                   : castOp.getOperation());
+    Value rootValue = rootInteger;
+    if (!rootValue) {
+      Type rootType = inputType;
+      if (auto integerType = dyn_cast<IntegerType>(inputType)) {
+        rootType = rewriter.getIntegerType(integerType.getWidth());
+      }
+      rootValue = rewriter.create<arith::ConstantOp>(
+          castOp.getLoc(), rewriter.getIntegerAttr(rootType, 0));
+    }
+    Value base =
+        rewriter.create<pto::CastPtrOp>(castOp.getLoc(), ptrType, rootValue);
+
+    rewriter.setInsertionPoint(castOp);
+    Value canonical = rewriter.create<pto::AddPtrOp>(castOp.getLoc(), ptrType,
+                                                     base, indexValue);
+    rewriter.replaceOp(castOp, canonical);
+    return success();
+  }
+
+  PTOValueEvolutionAnalysis &valueEvolution;
+};
+
+struct PTOIntegerAddressCanonicalizationPass final
+    : public pto::impl::PTOIntegerAddressCanonicalizationBase<
+          PTOIntegerAddressCanonicalizationPass> {
+  void runOnOperation() override {
+    func::FuncOp func = getOperation();
+    auto &valueEvolution = getAnalysis<pto::PTOValueEvolutionAnalysis>();
+    RewritePatternSet patterns(&getContext());
+    patterns.add<CanonicalizeIntegerCastPtr>(&getContext(), valueEvolution);
+    if (failed(applyPatternsAndFoldGreedily(func, std::move(patterns)))) {
+      signalPassFailure();
+    }
+  }
+};
+
+} // namespace
+
+std::unique_ptr<Pass>
+mlir::pto::createPTOIntegerAddressCanonicalizationPass() {
+  return std::make_unique<PTOIntegerAddressCanonicalizationPass>();
+}
diff --git a/test/dsl/issue_591_repro.pto b/test/dsl/issue_591_repro.pto
new file mode 100644
index 0000000000..62d210131e
--- /dev/null
+++ b/test/dsl/issue_591_repro.pto
@@ -0,0 +1,35 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+module attributes {
+  pto.target_arch = "a5",
+  pto.kernel_kind = #pto.kernel_kind<vector>
+} {
+  // Streams ntile UB tiles; tile t lives at UB byte offset t * stride.
+  func.func @affine_base_in_scf_for(%x: !pto.ptr<f16, gm>, %sink: !pto.ptr<f16, ub>, %ntile: i32)
+      attributes {pto.kernel} {
+    %c0 = arith.constant 0 : index
+    %c1 = arith.constant 1 : index
+    %e0 = arith.constant 0 : index
+    %stride = arith.constant 4096 : i64
+    %nt = arith.index_cast %ntile : i32 to index
+    %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+    scf.for %t = %c0 to %nt step %c1 {
+      %ti = arith.index_cast %t : index to i64
+      %off = arith.muli %ti, %stride : i64
+      %ub = pto.castptr %off : i64 -> !pto.ptr<f16, ub>
+      %lo, %hi = pto.vldsx2 %ub[%e0], "DINTLV_B16"
+          : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+      pto.vsts %lo, %sink[%e0], %mask
+          : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+      pto.vsts %hi, %sink[%e0], %mask
+          : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    }
+    return
+  }
+}
diff --git a/test/lit/vpto/absorb_addptr_offset.pto b/test/lit/vpto/absorb_addptr_offset.pto
new file mode 100644
index 0000000000..d669ddd5f6
--- /dev/null
+++ b/test/lit/vpto/absorb_addptr_offset.pto
@@ -0,0 +1,120 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// RUN: pto-test-opt %s -pto-absorb-addptr -canonicalize | FileCheck %s
+
+// C11: load through addptr(base, A) absorbs A into the op offset.
+// CHECK-LABEL: func.func @absorb_load
+// CHECK-NOT: pto.addptr
+// CHECK: %{{.*}}, %{{.*}} = pto.vldsx2 %{{.*}}[%arg0], "DINTLV_B16"
+func.func @absorb_load(%off: index, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c0_i64 = arith.constant 0 : i64
+  %base = pto.castptr %c0_i64 : i64 -> !pto.ptr<f16, ub>
+  %ptr = pto.addptr %base, %off : !pto.ptr<f16, ub> -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C11: store through addptr(base, A) absorbs A into the op offset.
+// CHECK-LABEL: func.func @absorb_store
+// CHECK-NOT: pto.addptr
+// CHECK: pto.vsts %{{.*}}, %{{.*}}[%arg0], %{{.*}}
+func.func @absorb_store(%off: index, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c0_i64 = arith.constant 0 : i64
+  %base = pto.castptr %c0_i64 : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %sink[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  %ptr = pto.addptr %base, %off : !pto.ptr<f16, ub> -> !pto.ptr<f16, ub>
+  pto.vsts %lo, %ptr[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C11: non-zero op offset adds to the absorbed addptr offset.
+// CHECK-LABEL: func.func @absorb_existing_offset
+// CHECK-NOT: pto.addptr
+// CHECK: %[[SUM:.*]] = arith.addi %arg0, %arg1
+// CHECK: pto.vldsx2 %{{.*}}[%[[SUM]]]
+func.func @absorb_existing_offset(%a: index, %o: index, %sink: !pto.ptr<f16, ub>) {
+  %c0_i64 = arith.constant 0 : i64
+  %base = pto.castptr %c0_i64 : i64 -> !pto.ptr<f16, ub>
+  %ptr = pto.addptr %base, %a : !pto.ptr<f16, ub> -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%o], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%o], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C10: addptr with multiple users folds every user; the addptr dies once
+// its last use is absorbed.
+// CHECK-LABEL: func.func @absorb_multi_user
+// CHECK-NOT: pto.addptr
+// CHECK: pto.vldsx2 %{{.*}}[%arg0]
+// CHECK: pto.vldsx2 %{{.*}}[%arg0]
+func.func @absorb_multi_user(%off: index, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c0_i64 = arith.constant 0 : i64
+  %base = pto.castptr %c0_i64 : i64 -> !pto.ptr<f16, ub>
+  %ptr = pto.addptr %base, %off : !pto.ptr<f16, ub> -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  %lo2, %hi2 = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  pto.vsts %lo2, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// vstsx2 (interface added with this change) absorbs addptr too.
+// CHECK-LABEL: func.func @absorb_vstsx2
+// CHECK-NOT: pto.addptr
+// CHECK: pto.vstsx2 %{{.*}}, %{{.*}}, %{{.*}}[%arg0]
+func.func @absorb_vstsx2(%off: index, %src: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c0_i64 = arith.constant 0 : i64
+  %base = pto.castptr %c0_i64 : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %src[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  %ptr = pto.addptr %base, %off : !pto.ptr<f16, ub> -> !pto.ptr<f16, ub>
+  pto.vstsx2 %lo, %hi, %ptr[%c0], "INTLV_B16", %mask
+      : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.ptr<f16, ub>, index, !pto.mask<b16>
+  return
+}
+
+// Chained addptr over a user pointer is NOT absorbed: sequential base
+// chains are VPTOSoftPostUpdate's job, folding them here would destroy that
+// post-update structure.
+// CHECK-LABEL: func.func @absorb_chained
+// CHECK: %[[P1:.*]] = pto.addptr %arg0, %arg1
+// CHECK: %[[P2:.*]] = pto.addptr %[[P1]], %arg2
+// CHECK: pto.vldsx2 %[[P2]][%c0]
+func.func @absorb_chained(%base: !pto.ptr<f16, ub>, %a: index, %b: index,
+                          %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %p1 = pto.addptr %base, %a : !pto.ptr<f16, ub> -> !pto.ptr<f16, ub>
+  %p2 = pto.addptr %p1, %b : !pto.ptr<f16, ub> -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %p2[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
diff --git a/test/lit/vpto/auto_vecscope_infer_shared_ptr_capture.pto b/test/lit/vpto/auto_vecscope_infer_shared_ptr_capture.pto
index 42b5852578..a3e30f2a9c 100644
--- a/test/lit/vpto/auto_vecscope_infer_shared_ptr_capture.pto
+++ b/test/lit/vpto/auto_vecscope_infer_shared_ptr_capture.pto
@@ -28,11 +28,12 @@ module attributes {pto.target_arch = "a5"} {
 }
 
 // CHECK-LABEL: func.func @auto_vecscope_infer_shared_ptr_capture
-// CHECK: %[[UBOUT:.*]] = pto.castptr %c2048_i64 : i64 -> !pto.ptr<ui16, ub>
+// CHECK: %[[UBASE:.*]] = pto.castptr %c0_i64 : i64 -> !pto.ptr<ui16, ub>
+// CHECK: %[[UBOUT:.*]] = pto.addptr %[[UBASE]], %c1024 : <ui16, ub> -> <ui16, ub>
 // CHECK: pto.vecscope {
 // CHECK: pto.vcvt
 // CHECK: pto.vpack
-// CHECK: pto.vsts %{{.*}}, %[[UBOUT]][%c0], %{{.*}}
+// CHECK: pto.vsts %{{.*}}, %[[UBASE]][%c1024], %{{.*}}
 // CHECK: }
 // CHECK: pto.barrier
 // CHECK: pto.copy_ubuf_to_gm %[[UBOUT]], %{{.*}}, %c0_i64, %c1_i64, %c256_i64, %c0_i64, %c256_i64, %c256_i64
diff --git a/test/lit/vpto/integer_address_canonicalization.pto b/test/lit/vpto/integer_address_canonicalization.pto
new file mode 100644
index 0000000000..37d0690b92
--- /dev/null
+++ b/test/lit/vpto/integer_address_canonicalization.pto
@@ -0,0 +1,459 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// RUN: pto-test-opt %s -pto-canonicalize-integer-address | FileCheck %s
+// RUN: pto-test-opt %s -pto-canonicalize-integer-address -pto-canonicalize-integer-address | FileCheck %s --check-prefix=IDEMPOTENT
+
+// C01: Issue #591 minimal pattern. Whole byte expression is exactly divisible.
+// CHECK-LABEL: func.func @c01_affine_iv
+// CHECK: %[[ZERO:.*]] = arith.constant 0 : i64
+// CHECK: %[[Q:.*]] = arith.muli %{{.*}}, %c2048 : index
+// CHECK: %[[BASE:.*]] = pto.castptr %[[ZERO]] : i64 -> !pto.ptr<f16, ub>
+// CHECK: %[[PTR:.*]] = pto.addptr %[[BASE]], %[[Q]]  : <f16, ub> -> <f16, ub>
+// CHECK: pto.vldsx2 %[[PTR]][%c0]
+func.func @c01_affine_iv(%iv: index, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i64
+  %iv64 = arith.index_cast %iv : index to i64
+  %bytes = arith.muli %iv64, %c4096 : i64
+  %ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C02: Constant base and dynamic term normalize together (root stays 0).
+// CHECK-LABEL: func.func @c02_const_plus_iv
+// CHECK: %[[ZERO:.*]] = arith.constant 0 : i64
+// CHECK: %[[Q:.*]] = arith.muli %{{.*}}, %c2048 : index
+// CHECK: pto.castptr %[[ZERO]] : i64 -> !pto.ptr<f16, ub>
+// CHECK: pto.addptr
+func.func @c02_const_plus_iv(%iv: index, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c8192 = arith.constant 8192 : i64
+  %c4096 = arith.constant 4096 : i64
+  %iv64 = arith.index_cast %iv : index to i64
+  %dyn = arith.muli %iv64, %c4096 : i64
+  %bytes = arith.addi %c8192, %dyn : i64
+  %ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C03: f32 element -> quotient scales by 4.
+// CHECK-LABEL: func.func @c03_f32
+// CHECK: arith.muli %{{.*}}, %c1024 : index
+func.func @c03_f32(%iv: index, %sink: !pto.ptr<f32, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i64
+  %iv64 = arith.index_cast %iv : index to i64
+  %bytes = arith.muli %iv64, %c4096 : i64
+  %ptr = pto.castptr %bytes : i64 -> !pto.ptr<f32, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f32, ub>, index -> !pto.vreg<64xf32>, !pto.vreg<64xf32>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<64xf32>, !pto.ptr<f32, ub>, !pto.mask<b16>
+  return
+}
+
+// C05: cast leaf is preserved as an atom (rebuilt as index_castui into the
+// index domain, preserving zero-extension); coefficient is divided.
+// CHECK-LABEL: func.func @c05_extui_atom
+// CHECK: %[[X:.*]] = arith.index_castui %arg0 : i32 to index
+// CHECK: arith.muli %[[X]], %c2048{{.*}} : index
+func.func @c05_extui_atom(%x: i32, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i64
+  %x64 = arith.extui %x : i32 to i64
+  %bytes = arith.muli %x64, %c4096 : i64
+  %ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C07: non-divisible constant -> unchanged.
+// CHECK-LABEL: func.func @c07_nondivisible_const
+// CHECK: %[[BYTES:.*]] = arith.addi %{{.*}}, %c1{{.*}} : i64
+// CHECK: %[[PTR:.*]] = pto.castptr %[[BYTES]] : i64 -> !pto.ptr<f16, ub>
+// CHECK: pto.vldsx2 %[[PTR]][%c0]
+func.func @c07_nondivisible_const(%iv: index, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i64
+  %c1 = arith.constant 1 : i64
+  %iv64 = arith.index_cast %iv : index to i64
+  %dyn = arith.muli %iv64, %c4096 : i64
+  %bytes = arith.addi %dyn, %c1 : i64
+  %ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C08: non-zero-origin space (gm) -> unchanged.
+// CHECK-LABEL: func.func @c08_gm_space
+// CHECK: %[[PTR:.*]] = pto.castptr %{{.*}} : i64 -> !pto.ptr<f16, gm>
+// CHECK: %[[BACK:.*]] = pto.castptr %[[PTR]] : !pto.ptr<f16, gm> -> i64
+// CHECK: return %[[BACK]]
+func.func @c08_gm_space(%iv: index) -> i64 {
+  %c4096 = arith.constant 4096 : i64
+  %iv64 = arith.index_cast %iv : index to i64
+  %bytes = arith.muli %iv64, %c4096 : i64
+  %ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, gm>
+  %back = pto.castptr %ptr : !pto.ptr<f16, gm> -> i64
+  return %back : i64
+}
+
+// C09a: bare castptr(0) is canonical -> unchanged.
+// CHECK-LABEL: func.func @c09a_bare_zero
+// CHECK: %[[ZERO:.*]] = arith.constant 0 : i64
+// CHECK: %[[PTR:.*]] = pto.castptr %[[ZERO]] : i64 -> !pto.ptr<f16, ub>
+// CHECK-NOT: pto.addptr
+func.func @c09a_bare_zero(%sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c0i = arith.constant 0 : i64
+  %ptr = pto.castptr %c0i : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C09b: pure-atom root is canonical -> unchanged.
+// CHECK-LABEL: func.func @c09b_bare_atom
+// CHECK: %[[PTR:.*]] = pto.castptr %arg0 : i64 -> !pto.ptr<f16, ub>
+// CHECK-NOT: pto.addptr
+func.func @c09b_bare_atom(%param: i64, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %ptr = pto.castptr %param : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C10: multiple users share one rewritten pointer.
+// CHECK-LABEL: func.func @c10_multi_user
+// CHECK: %[[ZERO:.*]] = arith.constant 0 : i64
+// CHECK: %[[BASE:.*]] = pto.castptr %[[ZERO]] : i64 -> !pto.ptr<f16, ub>
+// CHECK: %[[PTR:.*]] = pto.addptr %[[BASE]], %{{.*}}  : <f16, ub> -> <f16, ub>
+// CHECK: pto.vldsx2 %[[PTR]][%c0]
+// CHECK: pto.vldsx2 %[[PTR]][%c0]
+func.func @c10_multi_user(%iv: index, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i64
+  %iv64 = arith.index_cast %iv : index to i64
+  %bytes = arith.muli %iv64, %c4096 : i64
+  %ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  %lo2, %hi2 = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  pto.vsts %lo2, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C13: pointer-derived integer stays untouched.
+// CHECK-LABEL: func.func @c13_ptr_derived
+// CHECK: %[[RAW:.*]] = pto.castptr %arg0 : !pto.ptr<f16, ub> -> i64
+// CHECK: %[[BYTES:.*]] = arith.muli %[[RAW]], %c2{{.*}} : i64
+// CHECK: %[[PTR:.*]] = pto.castptr %[[BYTES]] : i64 -> !pto.ptr<f16, ub>
+func.func @c13_ptr_derived(%base: !pto.ptr<f16, ub>, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c2 = arith.constant 2 : i64
+  %raw = pto.castptr %base : !pto.ptr<f16, ub> -> i64
+  %bytes = arith.muli %raw, %c2 : i64
+  %ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C14: input width (i32) does not round-trip into index -> unchanged.
+// CHECK-LABEL: func.func @c14_i32_input
+// CHECK: %[[PTR:.*]] = pto.castptr %{{.*}} : i32 -> !pto.ptr<f16, ub>
+func.func @c14_i32_input(%iv: i32, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i32
+  %bytes = arith.muli %iv, %c4096 : i32
+  %ptr = pto.castptr %bytes : i32 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C15: runtime atom base -> root stays castptr(%param).
+// CHECK-LABEL: func.func @c15_runtime_atom
+// CHECK: %[[Q:.*]] = arith.muli %{{.*}}, %c2048 : index
+// CHECK: %[[BASE:.*]] = pto.castptr %arg0 : i64 -> !pto.ptr<f16, ub>
+// CHECK: %[[PTR:.*]] = pto.addptr %[[BASE]], %[[Q]]  : <f16, ub> -> <f16, ub>
+func.func @c15_runtime_atom(%param: i64, %iv: index, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i64
+  %iv64 = arith.index_cast %iv : index to i64
+  %bytes = arith.muli %iv64, %c4096 : i64
+  %addr = arith.addi %param, %bytes : i64
+  %ptr = pto.castptr %addr : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C16: a signedness-carrying address uses signless carriers for the synthesized
+// zero root and index offset.
+// CHECK-LABEL: func.func @c16_unsigned_input
+// CHECK-DAG: %[[ZERO:.*]] = arith.constant 0 : i64
+// CHECK-DAG: %[[OFFSET:.*]] = arith.constant 4096 : index
+// CHECK: %[[BASE:.*]] = pto.castptr %[[ZERO]] : i64 -> !pto.ptr<ui8, ub>
+// CHECK: %[[PTR:.*]] = pto.addptr %[[BASE]], %[[OFFSET]]
+// CHECK: return %[[PTR]]
+func.func @c16_unsigned_input() -> !pto.ptr<ui8, ub> {
+  %c4096 = arith.constant 4096 : i64
+  %bytes = builtin.unrealized_conversion_cast %c4096 : i64 to ui64
+  %ptr = pto.castptr %bytes : ui64 -> !pto.ptr<ui8, ub>
+  return %ptr : !pto.ptr<ui8, ub>
+}
+
+// Rejected: multiple non-divisible atoms have no unique root.
+// CHECK-LABEL: func.func @reject_multi_atom
+// CHECK: %[[PTR:.*]] = pto.castptr %{{.*}} : i64 -> !pto.ptr<f16, ub>
+func.func @reject_multi_atom(%p1: i64, %p2: i64, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i64
+  %iv64 = arith.constant 16 : i64
+  %bytes = arith.muli %iv64, %c4096 : i64
+  %addr = arith.addi %p1, %p2 : i64
+  %addr2 = arith.addi %addr, %bytes : i64
+  %ptr = pto.castptr %addr2 : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// Rejected: non-unit-coefficient atom has no root/offset split.
+// CHECK-LABEL: func.func @reject_nonunit_coeff
+// CHECK: %[[PTR:.*]] = pto.castptr %{{.*}} : i64 -> !pto.ptr<f16, ub>
+func.func @reject_nonunit_coeff(%p: i64, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c3 = arith.constant 3 : i64
+  %c4096 = arith.constant 4096 : i64
+  %iv64 = arith.constant 16 : i64
+  %scaled = arith.muli %p, %c3 : i64
+  %dyn = arith.muli %iv64, %c4096 : i64
+  %addr = arith.addi %scaled, %dyn : i64
+  %ptr = pto.castptr %addr : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// i32 constant byte expression round-trips losslessly into index (LLVM
+// inttoptr zero-extension semantics); narrower constant inputs are accepted.
+// CHECK-LABEL: func.func @i32_const_accept
+// CHECK: pto.castptr %{{.*}} : i32 -> !pto.ptr<f16, ub>
+// CHECK: pto.addptr %{{.*}}, %c2048
+func.func @i32_const_accept(%sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i32
+  %ptr = pto.castptr %c4096 : i32 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// C04: element size 1 (i8) is an identity quotient.
+// CHECK-LABEL: func.func @c04_i8
+// CHECK: %[[ZERO:.*]] = arith.constant 0 : i64
+// CHECK: %[[Q:.*]] = arith.muli %{{.*}}, %c4096 : index
+// CHECK: pto.castptr %[[ZERO]] : i64 -> !pto.ptr<i8, ub>
+// CHECK: pto.addptr
+func.func @c04_i8(%iv: index, %sink: !pto.ptr<i8, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i64
+  %iv64 = arith.index_cast %iv : index to i64
+  %bytes = arith.muli %iv64, %c4096 : i64
+  %ptr = pto.castptr %bytes : i64 -> !pto.ptr<i8, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<i8, ub>, index -> !pto.vreg<256xi8>, !pto.vreg<256xi8>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<256xi8>, !pto.ptr<i8, ub>, !pto.mask<b16>
+  return
+}
+
+// C06: trunc in the byte expression cannot be proven -> unchanged.
+// CHECK-LABEL: func.func @c06_trunc
+// CHECK: %[[TRUNC:.*]] = arith.trunci %{{.*}} : i64 to i32
+// CHECK: %[[PTR:.*]] = pto.castptr %[[TRUNC]] : i32 -> !pto.ptr<f16, ub>
+func.func @c06_trunc(%iv: i64, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i64
+  %wide = arith.muli %iv, %c4096 : i64
+  %bytes = arith.trunci %wide : i64 to i32
+  %ptr = pto.castptr %bytes : i32 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// Subtraction input: B = iv*4096 - 8192 normalizes like addition.
+// CHECK-LABEL: func.func @sub_input
+// CHECK: %[[ZERO:.*]] = arith.constant 0 : i64
+// CHECK: pto.castptr %[[ZERO]] : i64 -> !pto.ptr<f16, ub>
+// CHECK: pto.addptr
+func.func @sub_input(%iv: index, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c8192 = arith.constant 8192 : i64
+  %c4096 = arith.constant 4096 : i64
+  %iv64 = arith.index_cast %iv : index to i64
+  %dyn = arith.muli %iv64, %c4096 : i64
+  %bytes = arith.subi %dyn, %c8192 : i64
+  %ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// Non-linear input (srem) cannot be linearized -> unchanged.
+// CHECK-LABEL: func.func @nonlinear
+// CHECK: %[[PTR:.*]] = pto.castptr %{{.*}} : i64 -> !pto.ptr<f16, ub>
+func.func @nonlinear(%iv: i64, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i64
+  %c7 = arith.constant 7 : i64
+  %r = arith.remsi %iv, %c7 : i64
+  %bytes = arith.muli %r, %c4096 : i64
+  %ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// Atom root defined inside the loop: the base must NOT be hoisted (hoisting
+// would violate SSA dominance); the IR stays legal with the base in place.
+// CHECK-LABEL: func.func @atom_root_in_loop
+// CHECK: scf.for
+// CHECK: pto.castptr %{{.*}} : i64 -> !pto.ptr<f16, ub>
+// CHECK: pto.addptr %{{.*}}, %{{.*}}
+func.func @atom_root_in_loop(%sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c1 = arith.constant 1 : index
+  %c2 = arith.constant 2 : i64
+  %c1024 = arith.constant 1024 : index
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  scf.for %t = %c0 to %c1024 step %c1 {
+    %t64 = arith.index_cast %t : index to i64
+    %addr = arith.addi %t64, %c2 : i64
+    %ptr = pto.castptr %addr : i64 -> !pto.ptr<f16, ub>
+    %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    pto.vsts %lo, %sink[%c0], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  }
+  return
+}
+
+// Rejected: negative-coefficient atom has no root/offset split.
+// CHECK-LABEL: func.func @reject_neg_coeff
+// CHECK: %[[PTR:.*]] = pto.castptr %{{.*}} : i64 -> !pto.ptr<f16, ub>
+func.func @reject_neg_coeff(%p: i64, %sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c4096 = arith.constant 4096 : i64
+  %neg = arith.subi %c4096, %p : i64
+  %ptr = pto.castptr %neg : i64 -> !pto.ptr<f16, ub>
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+      : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+  pto.vsts %lo, %sink[%c0], %mask
+      : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+  return
+}
+
+// Nested loops: the canonical base is hoisted above the OUTERMOST loop.
+// CHECK-LABEL: func.func @nested_loop_hoist
+// CHECK: %[[BASE:.*]] = pto.castptr %{{.*}} : i64 -> !pto.ptr<f16, ub>
+// CHECK: scf.for
+// CHECK: scf.for
+// CHECK: pto.addptr %[[BASE]], %{{.*}}
+func.func @nested_loop_hoist(%sink: !pto.ptr<f16, ub>) {
+  %c0 = arith.constant 0 : index
+  %c1 = arith.constant 1 : index
+  %c64 = arith.constant 64 : index
+  %c4096 = arith.constant 4096 : i64
+  %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+  scf.for %i = %c0 to %c64 step %c1 {
+    scf.for %j = %c0 to %c64 step %c1 {
+      %j64 = arith.index_cast %j : index to i64
+      %bytes = arith.muli %j64, %c4096 : i64
+      %ptr = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+      %lo, %hi = pto.vldsx2 %ptr[%c0], "DINTLV_B16"
+          : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+      pto.vsts %lo, %sink[%c0], %mask
+          : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    }
+  }
+  return
+}
+
+// Second run leaves the canonical form untouched (no double addptr wrap).
+// IDEMPOTENT-LABEL: func.func @c01_affine_iv
+// IDEMPOTENT: pto.castptr %{{.*}} : i64 -> !pto.ptr<f16, ub>
+// IDEMPOTENT: pto.addptr %{{.*}}, %{{.*}} : <f16, ub> -> <f16, ub>
+// IDEMPOTENT-NOT: pto.addptr %{{.*}}, %{{.*}} : <f16, ub> -> <f16, ub>
+// IDEMPOTENT-LABEL: func.func @c02_const_plus_iv
+// IDEMPOTENT-LABEL: func.func @c15_runtime_atom
+// IDEMPOTENT: pto.castptr %arg0 : i64 -> !pto.ptr<f16, ub>
diff --git a/test/lit/vpto/soft_postupdate_delta-nonzero-lb.pto b/test/lit/vpto/soft_postupdate_delta-nonzero-lb.pto
index 9468257fe5..d81547321f 100644
--- a/test/lit/vpto/soft_postupdate_delta-nonzero-lb.pto
+++ b/test/lit/vpto/soft_postupdate_delta-nonzero-lb.pto
@@ -2,8 +2,8 @@
 
 // IV starts at 64 (not 0). The initial pointer must be base + 64, not base.
 // CHECK-LABEL: func.func @nonzero_lb
-// CHECK: %[[INIT:.*]] = pto.addptr
-// CHECK: scf.for {{.*}} iter_args({{.*}}, %{{.*}} = %[[INIT]]
+// CHECK: scf.for {{.*}} iter_args({{.*}} = %{{.*}}, %{{.*}} = %{{.*}}, %{{.*}} = %{{.*}})
+// CHECK: %{{.*}}, %{{.*}} = pto.vlds %{{.*}}[%c64]
 
 module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind<vector>} {
   func.func @nonzero_lb(%input: !pto.ptr<f32, gm>, %output: !pto.ptr<f32, gm>) attributes {pto.kernel} {
diff --git a/test/vpto/cases/micro-op/issue591-perf-rolled/compare.py b/test/vpto/cases/micro-op/issue591-perf-rolled/compare.py
new file mode 100644
index 0000000000..c556cd0054
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-rolled/compare.py
@@ -0,0 +1,30 @@
+#!/usr/bin/python3
+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+# CANN Open Software License Agreement Version 2.0 (the "License").
+# Please refer to the License for details. You may not use this file except in compliance with the License.
+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+# See LICENSE in the root of the software repository for the full text of the License.
+
+# Issue #591 performance acceptance: the numeric result is intentionally not
+# gated here; the acceptance evidence is the simulator instruction log
+# (zero RV_SADD/RV_SMOVK in the rolled tile loop) and the RVEC cycle count
+# from *_summary_log (rolled <= 1.1x unrolled). Check the output exists.
+
+import os
+import sys
+
+PATH = "v3.bin"
+
+
+def main():
+    if not os.path.exists(PATH) or os.path.getsize(PATH) == 0:
+        print(f"[ERROR] output missing or empty: {PATH}")
+        sys.exit(2)
+    print("[INFO] compare passed (output present; perf evidence is in "
+          "core*_summary_log and core*.veccore*.instr_log.dump)")
+
+
+if __name__ == "__main__":
+    main()
diff --git a/test/vpto/cases/micro-op/issue591-perf-rolled/golden.py b/test/vpto/cases/micro-op/issue591-perf-rolled/golden.py
new file mode 100644
index 0000000000..0642c7bc85
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-rolled/golden.py
@@ -0,0 +1,35 @@
+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+# CANN Open Software License Agreement Version 2.0 (the "License").
+# Please refer to the License for details. You may not use this file except in compliance with the License.
+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+# See LICENSE in the root of the software repository for the full text of the License.
+
+# Issue #591 performance acceptance: generate the input payload only. The
+# acceptance criterion is the CA-model instruction mix and RVEC cycle count,
+# not the numeric result (correctness is covered by lit + differential tests).
+
+import argparse
+from pathlib import Path
+
+import numpy as np
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--output-dir", type=Path, default=Path("."))
+    parser.add_argument("--seed", type=int, default=0)
+    args = parser.parse_args()
+    rng = np.random.default_rng(args.seed)
+    out = args.output_dir
+    out.mkdir(parents=True, exist_ok=True)
+    elems = 16 * 128
+    v1 = rng.integers(0, 1024, size=elems, dtype=np.uint16)
+    v1.tofile(out / "v1.bin")
+    golden = np.zeros(elems, dtype=np.uint16)
+    golden.tofile(out / "golden_v3.bin")
+
+
+if __name__ == "__main__":
+    main()
diff --git a/test/vpto/cases/micro-op/issue591-perf-rolled/kernel.pto b/test/vpto/cases/micro-op/issue591-perf-rolled/kernel.pto
new file mode 100644
index 0000000000..4d0ecceb12
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-rolled/kernel.pto
@@ -0,0 +1,61 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// Issue #591 performance acceptance case (rolled form).
+// The tile loop reads UB through castptr(t*4096) — the affine base-pointer
+// pattern. With pto-canonicalize-integer-address + pto-absorb-addptr the loop
+// body must emit a VLDS post-update load with zero scalar address ops
+// (RV_SADD/RV_SMOVK), measured on the CA model simulator.
+module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind<vector>} {
+  func.func @issue591_perf_rolled(%in: !pto.ptr<f16, gm>, %out: !pto.ptr<f16, gm>)
+      attributes {pto.kernel} {
+    %c0 = arith.constant 0 : index
+    %c1 = arith.constant 1 : index
+    %c16 = arith.constant 16 : index
+    %c128 = arith.constant 128 : index
+    %c0_i64 = arith.constant 0 : i64
+    %c4096 = arith.constant 4096 : i64
+    %c65536 = arith.constant 65536 : i64
+    %c128_i64 = arith.constant 128 : i64
+    %c32_i64 = arith.constant 32 : i64
+    %c16_i64 = arith.constant 16 : i64
+    %c2048_i64 = arith.constant 2048 : i64
+
+    %ub_in = pto.castptr %c0_i64 : i64 -> !pto.ptr<f16, ub>
+    %ub_out = pto.castptr %c65536 : i64 -> !pto.ptr<f16, ub>
+
+    pto.mte_gm_ub %in, %ub_in, %c0_i64, %c2048_i64
+      nburst(%c16_i64, %c128_i64, %c128_i64)
+      : !pto.ptr<f16, gm>, !pto.ptr<f16, ub>, i64, i64, i64, i64, i64
+    pto.set_flag["PIPE_MTE2", "PIPE_V", "EVENT_ID0"]
+    pto.wait_flag["PIPE_MTE2", "PIPE_V", "EVENT_ID0"]
+
+    pto.vecscope {
+      %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+      scf.for %t = %c0 to %c16 step %c1 {
+        %t64 = arith.index_cast %t : index to i64
+        %bytes = arith.muli %t64, %c4096 : i64
+        %p = pto.castptr %bytes : i64 -> !pto.ptr<f16, ub>
+        %lo, %hi = pto.vldsx2 %p[%c0], "DINTLV_B16"
+            : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+        %sum = pto.vadd %lo, %hi, %mask
+            : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+        %outoff = arith.muli %t, %c128 : index
+        pto.vsts %sum, %ub_out[%outoff], %mask
+            : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+      }
+    }
+    pto.set_flag["PIPE_V", "PIPE_MTE3", "EVENT_ID0"]
+    pto.wait_flag["PIPE_V", "PIPE_MTE3", "EVENT_ID0"]
+    pto.mte_ub_gm %ub_out, %out, %c2048_i64
+      nburst(%c16_i64, %c128_i64, %c128_i64)
+      : !pto.ptr<f16, ub>, !pto.ptr<f16, gm>, i64, i64, i64, i64
+    pto.barrier #pto.pipe<PIPE_ALL>
+    return
+  }
+}
diff --git a/test/vpto/cases/micro-op/issue591-perf-rolled/launch.cpp b/test/vpto/cases/micro-op/issue591-perf-rolled/launch.cpp
new file mode 100644
index 0000000000..9579fa9207
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-rolled/launch.cpp
@@ -0,0 +1,32 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+#ifndef __VEC_SCOPE__
+#define __VEC_SCOPE__
+#endif
+
+#if defined(__CCE_AICORE__) && defined(__NPU_ARCH__) && (__NPU_ARCH__ == 2201)
+typedef struct { unsigned char v; } hifloat8_t;
+typedef struct { unsigned char v; } float8_e4m3_t;
+typedef struct { unsigned char v; } float8_e5m2_t;
+typedef struct { unsigned char v; } float8_e8m0_t;
+typedef struct { unsigned char v; } float4_e1m2x2_t;
+typedef struct { unsigned char v; } float4_e2m1x2_t;
+#endif
+#include <stdint.h>
+#ifndef __CPU_SIM
+#include "acl/acl.h"
+#endif
+
+extern "C" __global__ [aicore] void issue591_perf_rolled(__gm__ half *in,
+                                                         __gm__ half *out);
+
+void LaunchIssue591PerfRolled(uint16_t *in, uint16_t *out, void *stream) {
+  issue591_perf_rolled<<<1, nullptr, stream>>>((__gm__ half *)in,
+                                               (__gm__ half *)out);
+}
diff --git a/test/vpto/cases/micro-op/issue591-perf-rolled/main.cpp b/test/vpto/cases/micro-op/issue591-perf-rolled/main.cpp
new file mode 100644
index 0000000000..d5aa7fd144
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-rolled/main.cpp
@@ -0,0 +1,91 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// Issue #591 performance acceptance host runner (rolled form).
+
+#include "test_common.h"
+#include "acl/acl.h"
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+
+using namespace PtoTestCommon;
+
+#define ACL_CHECK(expr)                                                          \
+  do {                                                                           \
+    const aclError _ret = (expr);                                                \
+    if (_ret != ACL_SUCCESS) {                                                   \
+      std::fprintf(stderr, "[ERROR] %s failed: %d (%s:%d)\n", #expr,            \
+                   (int)_ret, __FILE__, __LINE__);                               \
+      rc = 1;                                                                    \
+      goto cleanup;                                                              \
+    }                                                                            \
+  } while (0)
+
+#define FILE_CHECK(expr, path)                                                   \
+  do {                                                                           \
+    if (!(expr)) {                                                               \
+      std::fprintf(stderr, "[ERROR] file operation failed: %s (%s:%d)\n",       \
+                   path, __FILE__, __LINE__);                                    \
+      rc = 1;                                                                    \
+      goto cleanup;                                                              \
+    }                                                                            \
+  } while (0)
+
+void LaunchIssue591PerfRolled(uint16_t *in, uint16_t *out, void *stream);
+
+int main() {
+  constexpr size_t ELEMS = 16 * 128;
+  constexpr size_t SZ = ELEMS * sizeof(uint16_t);
+
+  uint16_t *h_in = nullptr, *h_out = nullptr;
+  uint16_t *d_in = nullptr, *d_out = nullptr;
+
+  int rc = 0;
+  bool aclInited = false;
+  bool deviceSet = false;
+  int deviceId = 0;
+  aclrtStream stream = nullptr;
+  size_t fsize = 0;
+
+  ACL_CHECK(aclInit(nullptr));
+  aclInited = true;
+  if (const char *envDevice = std::getenv("ACL_DEVICE_ID"))
+    deviceId = std::atoi(envDevice);
+  ACL_CHECK(aclrtSetDevice(deviceId));
+  deviceSet = true;
+  ACL_CHECK(aclrtCreateStream(&stream));
+
+  ACL_CHECK(aclrtMallocHost((void **)&h_in, SZ));
+  ACL_CHECK(aclrtMallocHost((void **)&h_out, SZ));
+  ACL_CHECK(aclrtMalloc((void **)&d_in, SZ, ACL_MEM_MALLOC_HUGE_FIRST));
+  ACL_CHECK(aclrtMalloc((void **)&d_out, SZ, ACL_MEM_MALLOC_HUGE_FIRST));
+
+  fsize = SZ;
+  FILE_CHECK(ReadFile("./v1.bin", fsize, h_in, SZ) && fsize == SZ, "./v1.bin");
+  ACL_CHECK(aclrtMemcpy(d_in, SZ, h_in, SZ, ACL_MEMCPY_HOST_TO_DEVICE));
+
+  LaunchIssue591PerfRolled(d_in, d_out, stream);
+  ACL_CHECK(aclrtSynchronizeStream(stream));
+
+  ACL_CHECK(aclrtMemcpy(h_out, SZ, d_out, SZ, ACL_MEMCPY_DEVICE_TO_HOST));
+  FILE_CHECK(WriteFile("./v3.bin", h_out, SZ), "./v3.bin");
+
+cleanup:
+  aclrtFree(d_in);
+  aclrtFree(d_out);
+  aclrtFreeHost(h_in);
+  aclrtFreeHost(h_out);
+  if (stream != nullptr)
+    aclrtDestroyStream(stream);
+  if (deviceSet)
+    aclrtResetDevice(deviceId);
+  if (aclInited)
+    aclFinalize();
+  return rc;
+}
diff --git a/test/vpto/cases/micro-op/issue591-perf-rolled/ptoas.flags b/test/vpto/cases/micro-op/issue591-perf-rolled/ptoas.flags
new file mode 100644
index 0000000000..28c4691e94
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-rolled/ptoas.flags
@@ -0,0 +1 @@
+--pto-arch a5 --pto-backend=vpto --enable-vpto-soft-postupdate
diff --git a/test/vpto/cases/micro-op/issue591-perf-unrolled/compare.py b/test/vpto/cases/micro-op/issue591-perf-unrolled/compare.py
new file mode 100644
index 0000000000..c556cd0054
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-unrolled/compare.py
@@ -0,0 +1,30 @@
+#!/usr/bin/python3
+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+# CANN Open Software License Agreement Version 2.0 (the "License").
+# Please refer to the License for details. You may not use this file except in compliance with the License.
+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+# See LICENSE in the root of the software repository for the full text of the License.
+
+# Issue #591 performance acceptance: the numeric result is intentionally not
+# gated here; the acceptance evidence is the simulator instruction log
+# (zero RV_SADD/RV_SMOVK in the rolled tile loop) and the RVEC cycle count
+# from *_summary_log (rolled <= 1.1x unrolled). Check the output exists.
+
+import os
+import sys
+
+PATH = "v3.bin"
+
+
+def main():
+    if not os.path.exists(PATH) or os.path.getsize(PATH) == 0:
+        print(f"[ERROR] output missing or empty: {PATH}")
+        sys.exit(2)
+    print("[INFO] compare passed (output present; perf evidence is in "
+          "core*_summary_log and core*.veccore*.instr_log.dump)")
+
+
+if __name__ == "__main__":
+    main()
diff --git a/test/vpto/cases/micro-op/issue591-perf-unrolled/golden.py b/test/vpto/cases/micro-op/issue591-perf-unrolled/golden.py
new file mode 100644
index 0000000000..0642c7bc85
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-unrolled/golden.py
@@ -0,0 +1,35 @@
+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+# CANN Open Software License Agreement Version 2.0 (the "License").
+# Please refer to the License for details. You may not use this file except in compliance with the License.
+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+# See LICENSE in the root of the software repository for the full text of the License.
+
+# Issue #591 performance acceptance: generate the input payload only. The
+# acceptance criterion is the CA-model instruction mix and RVEC cycle count,
+# not the numeric result (correctness is covered by lit + differential tests).
+
+import argparse
+from pathlib import Path
+
+import numpy as np
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--output-dir", type=Path, default=Path("."))
+    parser.add_argument("--seed", type=int, default=0)
+    args = parser.parse_args()
+    rng = np.random.default_rng(args.seed)
+    out = args.output_dir
+    out.mkdir(parents=True, exist_ok=True)
+    elems = 16 * 128
+    v1 = rng.integers(0, 1024, size=elems, dtype=np.uint16)
+    v1.tofile(out / "v1.bin")
+    golden = np.zeros(elems, dtype=np.uint16)
+    golden.tofile(out / "golden_v3.bin")
+
+
+if __name__ == "__main__":
+    main()
diff --git a/test/vpto/cases/micro-op/issue591-perf-unrolled/kernel.pto b/test/vpto/cases/micro-op/issue591-perf-unrolled/kernel.pto
new file mode 100644
index 0000000000..832ca3f95a
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-unrolled/kernel.pto
@@ -0,0 +1,186 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// Issue #591 performance acceptance case (fully-unrolled form).
+module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind<vector>} {
+  func.func @issue591_perf_unrolled(%in: !pto.ptr<f16, gm>, %out: !pto.ptr<f16, gm>)
+      attributes {pto.kernel} {
+    %c0 = arith.constant 0 : index
+    %c128 = arith.constant 128 : index
+    %c0_i64 = arith.constant 0 : i64
+    %c65536 = arith.constant 65536 : i64
+    %c128_i64 = arith.constant 128 : i64
+    %c32_i64 = arith.constant 32 : i64
+    %c16_i64 = arith.constant 16 : i64
+    %c2048_i64 = arith.constant 2048 : i64
+
+    %ub_in = pto.castptr %c0_i64 : i64 -> !pto.ptr<f16, ub>
+    %ub_out = pto.castptr %c65536 : i64 -> !pto.ptr<f16, ub>
+
+    pto.mte_gm_ub %in, %ub_in, %c0_i64, %c2048_i64
+      nburst(%c16_i64, %c128_i64, %c128_i64)
+      : !pto.ptr<f16, gm>, !pto.ptr<f16, ub>, i64, i64, i64, i64, i64
+    pto.set_flag["PIPE_MTE2", "PIPE_V", "EVENT_ID0"]
+    pto.wait_flag["PIPE_MTE2", "PIPE_V", "EVENT_ID0"]
+
+    pto.vecscope {
+      %mask = pto.pset_b16 "PAT_ALL" : !pto.mask<b16>
+    %ct0_i64 = arith.constant 0 : i64
+    %p0 = pto.castptr %ct0_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo0, %hi0 = pto.vldsx2 %p0[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum0 = pto.vadd %lo0, %hi0, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o0 = arith.constant 0 : index
+    pto.vsts %sum0, %ub_out[%o0], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct1_i64 = arith.constant 4096 : i64
+    %p1 = pto.castptr %ct1_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo1, %hi1 = pto.vldsx2 %p1[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum1 = pto.vadd %lo1, %hi1, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o1 = arith.constant 128 : index
+    pto.vsts %sum1, %ub_out[%o1], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct2_i64 = arith.constant 8192 : i64
+    %p2 = pto.castptr %ct2_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo2, %hi2 = pto.vldsx2 %p2[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum2 = pto.vadd %lo2, %hi2, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o2 = arith.constant 256 : index
+    pto.vsts %sum2, %ub_out[%o2], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct3_i64 = arith.constant 12288 : i64
+    %p3 = pto.castptr %ct3_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo3, %hi3 = pto.vldsx2 %p3[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum3 = pto.vadd %lo3, %hi3, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o3 = arith.constant 384 : index
+    pto.vsts %sum3, %ub_out[%o3], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct4_i64 = arith.constant 16384 : i64
+    %p4 = pto.castptr %ct4_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo4, %hi4 = pto.vldsx2 %p4[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum4 = pto.vadd %lo4, %hi4, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o4 = arith.constant 512 : index
+    pto.vsts %sum4, %ub_out[%o4], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct5_i64 = arith.constant 20480 : i64
+    %p5 = pto.castptr %ct5_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo5, %hi5 = pto.vldsx2 %p5[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum5 = pto.vadd %lo5, %hi5, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o5 = arith.constant 640 : index
+    pto.vsts %sum5, %ub_out[%o5], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct6_i64 = arith.constant 24576 : i64
+    %p6 = pto.castptr %ct6_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo6, %hi6 = pto.vldsx2 %p6[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum6 = pto.vadd %lo6, %hi6, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o6 = arith.constant 768 : index
+    pto.vsts %sum6, %ub_out[%o6], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct7_i64 = arith.constant 28672 : i64
+    %p7 = pto.castptr %ct7_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo7, %hi7 = pto.vldsx2 %p7[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum7 = pto.vadd %lo7, %hi7, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o7 = arith.constant 896 : index
+    pto.vsts %sum7, %ub_out[%o7], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct8_i64 = arith.constant 32768 : i64
+    %p8 = pto.castptr %ct8_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo8, %hi8 = pto.vldsx2 %p8[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum8 = pto.vadd %lo8, %hi8, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o8 = arith.constant 1024 : index
+    pto.vsts %sum8, %ub_out[%o8], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct9_i64 = arith.constant 36864 : i64
+    %p9 = pto.castptr %ct9_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo9, %hi9 = pto.vldsx2 %p9[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum9 = pto.vadd %lo9, %hi9, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o9 = arith.constant 1152 : index
+    pto.vsts %sum9, %ub_out[%o9], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct10_i64 = arith.constant 40960 : i64
+    %p10 = pto.castptr %ct10_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo10, %hi10 = pto.vldsx2 %p10[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum10 = pto.vadd %lo10, %hi10, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o10 = arith.constant 1280 : index
+    pto.vsts %sum10, %ub_out[%o10], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct11_i64 = arith.constant 45056 : i64
+    %p11 = pto.castptr %ct11_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo11, %hi11 = pto.vldsx2 %p11[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum11 = pto.vadd %lo11, %hi11, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o11 = arith.constant 1408 : index
+    pto.vsts %sum11, %ub_out[%o11], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct12_i64 = arith.constant 49152 : i64
+    %p12 = pto.castptr %ct12_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo12, %hi12 = pto.vldsx2 %p12[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum12 = pto.vadd %lo12, %hi12, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o12 = arith.constant 1536 : index
+    pto.vsts %sum12, %ub_out[%o12], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct13_i64 = arith.constant 53248 : i64
+    %p13 = pto.castptr %ct13_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo13, %hi13 = pto.vldsx2 %p13[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum13 = pto.vadd %lo13, %hi13, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o13 = arith.constant 1664 : index
+    pto.vsts %sum13, %ub_out[%o13], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct14_i64 = arith.constant 57344 : i64
+    %p14 = pto.castptr %ct14_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo14, %hi14 = pto.vldsx2 %p14[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum14 = pto.vadd %lo14, %hi14, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o14 = arith.constant 1792 : index
+    pto.vsts %sum14, %ub_out[%o14], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    %ct15_i64 = arith.constant 61440 : i64
+    %p15 = pto.castptr %ct15_i64 : i64 -> !pto.ptr<f16, ub>
+    %lo15, %hi15 = pto.vldsx2 %p15[%c0], "DINTLV_B16"
+        : !pto.ptr<f16, ub>, index -> !pto.vreg<128xf16>, !pto.vreg<128xf16>
+    %sum15 = pto.vadd %lo15, %hi15, %mask
+        : !pto.vreg<128xf16>, !pto.vreg<128xf16>, !pto.mask<b16> -> !pto.vreg<128xf16>
+    %o15 = arith.constant 1920 : index
+    pto.vsts %sum15, %ub_out[%o15], %mask
+        : !pto.vreg<128xf16>, !pto.ptr<f16, ub>, !pto.mask<b16>
+    }
+    pto.set_flag["PIPE_V", "PIPE_MTE3", "EVENT_ID0"]
+    pto.wait_flag["PIPE_V", "PIPE_MTE3", "EVENT_ID0"]
+    pto.mte_ub_gm %ub_out, %out, %c2048_i64
+      nburst(%c16_i64, %c128_i64, %c128_i64)
+      : !pto.ptr<f16, ub>, !pto.ptr<f16, gm>, i64, i64, i64, i64
+    pto.barrier #pto.pipe<PIPE_ALL>
+    return
+  }
+}
diff --git a/test/vpto/cases/micro-op/issue591-perf-unrolled/launch.cpp b/test/vpto/cases/micro-op/issue591-perf-unrolled/launch.cpp
new file mode 100644
index 0000000000..1d25c94b0c
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-unrolled/launch.cpp
@@ -0,0 +1,32 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+#ifndef __VEC_SCOPE__
+#define __VEC_SCOPE__
+#endif
+
+#if defined(__CCE_AICORE__) && defined(__NPU_ARCH__) && (__NPU_ARCH__ == 2201)
+typedef struct { unsigned char v; } hifloat8_t;
+typedef struct { unsigned char v; } float8_e4m3_t;
+typedef struct { unsigned char v; } float8_e5m2_t;
+typedef struct { unsigned char v; } float8_e8m0_t;
+typedef struct { unsigned char v; } float4_e1m2x2_t;
+typedef struct { unsigned char v; } float4_e2m1x2_t;
+#endif
+#include <stdint.h>
+#ifndef __CPU_SIM
+#include "acl/acl.h"
+#endif
+
+extern "C" __global__ [aicore] void issue591_perf_unrolled(__gm__ half *in,
+                                                           __gm__ half *out);
+
+void LaunchIssue591PerfUnrolled(uint16_t *in, uint16_t *out, void *stream) {
+  issue591_perf_unrolled<<<1, nullptr, stream>>>((__gm__ half *)in,
+                                                 (__gm__ half *)out);
+}
diff --git a/test/vpto/cases/micro-op/issue591-perf-unrolled/main.cpp b/test/vpto/cases/micro-op/issue591-perf-unrolled/main.cpp
new file mode 100644
index 0000000000..887fa76bba
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-unrolled/main.cpp
@@ -0,0 +1,91 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// Issue #591 performance acceptance host runner (unrolled form).
+
+#include "test_common.h"
+#include "acl/acl.h"
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+
+using namespace PtoTestCommon;
+
+#define ACL_CHECK(expr)                                                          \
+  do {                                                                           \
+    const aclError _ret = (expr);                                                \
+    if (_ret != ACL_SUCCESS) {                                                   \
+      std::fprintf(stderr, "[ERROR] %s failed: %d (%s:%d)\n", #expr,            \
+                   (int)_ret, __FILE__, __LINE__);                               \
+      rc = 1;                                                                    \
+      goto cleanup;                                                              \
+    }                                                                            \
+  } while (0)
+
+#define FILE_CHECK(expr, path)                                                   \
+  do {                                                                           \
+    if (!(expr)) {                                                               \
+      std::fprintf(stderr, "[ERROR] file operation failed: %s (%s:%d)\n",       \
+                   path, __FILE__, __LINE__);                                    \
+      rc = 1;                                                                    \
+      goto cleanup;                                                              \
+    }                                                                            \
+  } while (0)
+
+void LaunchIssue591PerfUnrolled(uint16_t *in, uint16_t *out, void *stream);
+
+int main() {
+  constexpr size_t ELEMS = 16 * 128;
+  constexpr size_t SZ = ELEMS * sizeof(uint16_t);
+
+  uint16_t *h_in = nullptr, *h_out = nullptr;
+  uint16_t *d_in = nullptr, *d_out = nullptr;
+
+  int rc = 0;
+  bool aclInited = false;
+  bool deviceSet = false;
+  int deviceId = 0;
+  aclrtStream stream = nullptr;
+  size_t fsize = 0;
+
+  ACL_CHECK(aclInit(nullptr));
+  aclInited = true;
+  if (const char *envDevice = std::getenv("ACL_DEVICE_ID"))
+    deviceId = std::atoi(envDevice);
+  ACL_CHECK(aclrtSetDevice(deviceId));
+  deviceSet = true;
+  ACL_CHECK(aclrtCreateStream(&stream));
+
+  ACL_CHECK(aclrtMallocHost((void **)&h_in, SZ));
+  ACL_CHECK(aclrtMallocHost((void **)&h_out, SZ));
+  ACL_CHECK(aclrtMalloc((void **)&d_in, SZ, ACL_MEM_MALLOC_HUGE_FIRST));
+  ACL_CHECK(aclrtMalloc((void **)&d_out, SZ, ACL_MEM_MALLOC_HUGE_FIRST));
+
+  fsize = SZ;
+  FILE_CHECK(ReadFile("./v1.bin", fsize, h_in, SZ) && fsize == SZ, "./v1.bin");
+  ACL_CHECK(aclrtMemcpy(d_in, SZ, h_in, SZ, ACL_MEMCPY_HOST_TO_DEVICE));
+
+  LaunchIssue591PerfUnrolled(d_in, d_out, stream);
+  ACL_CHECK(aclrtSynchronizeStream(stream));
+
+  ACL_CHECK(aclrtMemcpy(h_out, SZ, d_out, SZ, ACL_MEMCPY_DEVICE_TO_HOST));
+  FILE_CHECK(WriteFile("./v3.bin", h_out, SZ), "./v3.bin");
+
+cleanup:
+  aclrtFree(d_in);
+  aclrtFree(d_out);
+  aclrtFreeHost(h_in);
+  aclrtFreeHost(h_out);
+  if (stream != nullptr)
+    aclrtDestroyStream(stream);
+  if (deviceSet)
+    aclrtResetDevice(deviceId);
+  if (aclInited)
+    aclFinalize();
+  return rc;
+}
diff --git a/test/vpto/cases/micro-op/issue591-perf-unrolled/ptoas.flags b/test/vpto/cases/micro-op/issue591-perf-unrolled/ptoas.flags
new file mode 100644
index 0000000000..28c4691e94
--- /dev/null
+++ b/test/vpto/cases/micro-op/issue591-perf-unrolled/ptoas.flags
@@ -0,0 +1 @@
+--pto-arch a5 --pto-backend=vpto --enable-vpto-soft-postupdate
diff --git a/tools/ptoas/ptoas.cpp b/tools/ptoas/ptoas.cpp
index 9ca7c61b1c..8a7a7f35e1 100644
--- a/tools/ptoas/ptoas.cpp
+++ b/tools/ptoas/ptoas.cpp
@@ -3090,6 +3090,14 @@ static void prepareVPTOForEmission(PassManager &pm) {
       createVPTOExpandWrapperOpsPass());
   kernelModulePM.addNestedPass<func::FuncOp>(
       pto::createPTOInferVPTOVecScopePass());
+  // Integer address canonicalization must run after integer-backed castptr
+  // forms are produced and before the typed-address consumers. The addptr
+  // absorption fold then exposes the affine offset on the operation itself so
+  // the post-update consumer can strength-reduce it (Issue #591).
+  kernelModulePM.addNestedPass<func::FuncOp>(
+      pto::createPTOIntegerAddressCanonicalizationPass());
+  kernelModulePM.addNestedPass<func::FuncOp>(
+      pto::createPTOAbsorbAddPtrPass());
   if (enableSoftPostUpdate) {
     kernelModulePM.addPass(pto::createVPTOSoftPostUpdatePass());
   }