diff --git a/lib/PTO/Transforms/PTOViewToMemref.cpp b/lib/PTO/Transforms/PTOViewToMemref.cpp index 053638a6..94c51277 100644 --- a/lib/PTO/Transforms/PTOViewToMemref.cpp +++ b/lib/PTO/Transforms/PTOViewToMemref.cpp @@ -617,21 +617,16 @@ struct PTOViewToMemrefPass auto configAttr = tbTy.getConfigAttr(); if (!configAttr) configAttr = pto::TileBufConfigAttr::getDefault(ctx); - // 6. If alloc_tile provides an explicit address, keep the original - // pointer_cast lowering intact and additionally rebind through - // pto.bind_tile. PointerCastOp continues to carry the tile metadata - // used by existing lowering paths, while BindTileOp provides the - // unified anchor EmitC uses to recover tile_buf information. + // 6. If alloc_tile provides an explicit address, lower directly to + // pto.pointer_cast. Rebinding through pto.bind_tile here is redundant + // and can produce an extra tile rewrap in EmitC for dynamic valid + // shapes (double TASSIGN pattern). if (Value addr = op.getAddr()) { auto pc = rewriter.create( loc, targetType, ValueRange{addr}, vRow ? vRow : Value(), vCol ? vCol : Value(), configAttr); markForceDynamicValidShape(pc, tbTy.hasDynamicValid(), ctx); - auto bindOp = rewriter.create( - loc, targetType, pc.getResult(), vRow ? vRow : Value(), - vCol ? vCol : Value(), configAttr); - markForceDynamicValidShape(bindOp, tbTy.hasDynamicValid(), ctx); - rewriter.replaceOp(op, bindOp.getResult()); + rewriter.replaceOp(op, pc.getResult()); continue; } diff --git a/test/basic/alloc_tile_addr_dynamic_no_rebind.pto b/test/basic/alloc_tile_addr_dynamic_no_rebind.pto new file mode 100644 index 00000000..a7cc1893 --- /dev/null +++ b/test/basic/alloc_tile_addr_dynamic_no_rebind.pto @@ -0,0 +1,23 @@ +// RUN: ptoas --pto-level=level3 %s | FileCheck %s + +module { + func.func @print_alloc_addr_dyn(%arg0: index, %arg1: index) attributes {pto.entry} { + %c0_i64 = arith.constant 0 : i64 + + %0 = pto.alloc_tile addr = %c0_i64 valid_row = %arg0 valid_col = %arg1 + : !pto.tile_buf + + pto.tprint ins(%0 : !pto.tile_buf) + return + } +} + +// CHECK-LABEL: __global__ AICORE void print_alloc_addr_dyn( +// CHECK: Tile [[TILE:v[0-9]+]] = Tile( +// CHECK: TASSIGN([[TILE]], [[ADDR:v[0-9]+]]); +// CHECK-NOT: .data() +// CHECK-NOT: reinterpret_cast +// CHECK: TPRINT([[TILE]]);