Star583
297
代码
代码
Issues123
Pull Requests188
Actions
项目讨论
Wiki
分析
Star583
297
已合并
Revert "feat: support A5 UB-to-L1 ND-to-NZ with A2A3 path guards" #1741
zhywang创建于 9月24日
Revert "feat: support A5 UB-to-L1 ND-to-NZ with A2A3 path guards" #1741
已合并
zhywang创建于 9月24日
从zhhywang/pto-isamaster合入到cann/pto-isamaster
共 27 个文件变更+50-856
Mdocs/isa/TEXTRACT.md+0-15
@@ -138,10 +138,6 @@ canonical overload.
138 138 
139### A2A3 implementation checks139### A2A3 implementation checks
140 140 
141-For ordinary Tile inputs outside the Vec-to-Vec path, the two-Tile overload supports
142-`Mat -> Left/Right` and `Acc -> Mat`. Unmatched combinations, including `Mat -> Mat`
143-and `Vec -> Mat`, are rejected at compile time.
144- 
145For Mat-to-Left/Right layout extraction:141For Mat-to-Left/Right layout extraction:
146 142 
147- Supported element types: `int8_t`, `half`, `bfloat16_t`, `float`.143- Supported element types: `int8_t`, `half`, `bfloat16_t`, `float`.
@@ -169,15 +165,6 @@ For Mat-to-Left/Right layout extraction:
169 `mode = AccToVecMode::{SingleModeVec0, SingleModeVec1, DualModeSplitM, DualModeSplitN}`.165 `mode = AccToVecMode::{SingleModeVec0, SingleModeVec1, DualModeSplitM, DualModeSplitN}`.
170- For `TileType::Acc -> TileType::Vec` with a 32-bit destination type (`float`/`int32_t`), when using `DualModeSplitN` the `ValidCol` (before the split) must be a multiple of `32`.166- For `TileType::Acc -> TileType::Vec` with a 32-bit destination type (`float`/`int32_t`), when using `DualModeSplitN` the `ValidCol` (before the split) must be a multiple of `32`.
171 167 
172-### A5 Vec ND-to-Mat NZ extraction
173- 
174-An ND Vec source can be extracted directly into a non-compact NZ512 Mat destination.
175-The destination valid shape defines the source window at `(indexRow, indexCol)`.
176-The full window must fit the source valid shape; column width, source row stride,
177-and column offset must be 32-byte aligned. Row offsets and row tails need no fractal alignment.
178-Only the valid window is written; destination padding is preserved. See
179-[UB ND → L1 NZ](TMOV.md#ub-nd--l1-nz-a5) for the shared storage and MTE3 synchronization requirements.
180- 
181### A5 Acc-to-Mat NZ layout conversion168### A5 Acc-to-Mat NZ layout conversion
182 169 
183The following paths use an NZ destination (`BLayout::ColMajor`, `SLayout::RowMajor`):170The following paths use an NZ destination (`BLayout::ColMajor`, `SLayout::RowMajor`):
@@ -380,8 +367,6 @@ NPU ST verifies the complete instruction sequence numerically.
380 367 
381In addition to the `Mat/Acc -> ...` paths above, `TEXTRACT` supports a `TileType::Vec -> TileType::Vec` extraction path (ND and NZ layouts). A2A3 uses `CheckTExtractVecToVecCommon`; A5 checks this separately in `TEXTRACT_IMPL`:368In addition to the `Mat/Acc -> ...` paths above, `TEXTRACT` supports a `TileType::Vec -> TileType::Vec` extraction path (ND and NZ layouts). A2A3 uses `CheckTExtractVecToVecCommon`; A5 checks this separately in `TEXTRACT_IMPL`:
382 369 
383-- On A2A3, the ordinary two-Tile overload requires both Vec tiles to be ND (`RowMajor`, `NoneBox`)
384- or both NZ (`ColMajor`, `RowMajor`); other layout combinations are rejected at compile time.
385- `DstTileData::DType` must equal `SrcTileData::DType`.370- `DstTileData::DType` must equal `SrcTileData::DType`.
386- A2A3 element types: `int8_t`, `uint8_t`, `int16_t`, `uint16_t`, `int32_t`, `uint32_t`, `half`, `bfloat16_t`, `float`.371- A2A3 element types: `int8_t`, `uint8_t`, `int16_t`, `uint16_t`, `int32_t`, `uint32_t`, `half`, `bfloat16_t`, `float`.
387- A5 element types: `int8_t`, `int32_t`, `half`, `bfloat16_t`, `float`, `hifloat8_t`, `float8_e4m3_t`, `float8_e5m2_t`, `float8_e8m0_t`, `float4_e2m1x2_t`, `float4_e1m2x2_t`. A5 does not support `uint8_t`, `int16_t`, `uint16_t`, `uint32_t` or 64-bit integers on this path.372- A5 element types: `int8_t`, `int32_t`, `half`, `bfloat16_t`, `float`, `hifloat8_t`, `float8_e4m3_t`, `float8_e5m2_t`, `float8_e8m0_t`, `float4_e2m1x2_t`, `float4_e1m2x2_t`. A5 does not support `uint8_t`, `int16_t`, `uint16_t`, `uint32_t` or 64-bit integers on this path.
Mdocs/isa/TEXTRACT_zh.md+0-13
@@ -119,9 +119,6 @@ Ascend 950PR/Ascend 950DT 和 CPU 模拟器），适用范围与配对规则见
119 119 
120### Atlas A2/A3 训练系列产品/Atlas A2/A3 推理系列产品实现检查120### Atlas A2/A3 训练系列产品/Atlas A2/A3 推理系列产品实现检查
121 121 
122-对于普通 Tile 输入，双 Tile 重载在 Vec→Vec 路径外支持 `Mat -> Left/Right`
123-和 `Acc -> Mat`。未匹配的组合（包括 `Mat -> Mat` 和 `Vec -> Mat`）在编译期报错。
124- 
125对于 Mat→Left/Right 布局提取：122对于 Mat→Left/Right 布局提取：
126 123 
127- 支持的元素类型：`int8_t`、`half`、`bfloat16_t`、`float`。124- 支持的元素类型：`int8_t`、`half`、`bfloat16_t`、`float`。
@@ -148,14 +145,6 @@ Ascend 950PR/Ascend 950DT 和 CPU 模拟器），适用范围与配对规则见
148 `mode = AccToVecMode::{SingleModeVec0, SingleModeVec1, DualModeSplitM, DualModeSplitN}`。145 `mode = AccToVecMode::{SingleModeVec0, SingleModeVec1, DualModeSplitM, DualModeSplitN}`。
149- 对于 `TileType::Acc -> TileType::Vec`，当目标为32位类型（`float`/`int32_t`）且使用 `DualModeSplitN` 时，切分前的 `ValidCol` 必须是 `32` 的整数倍。146- 对于 `TileType::Acc -> TileType::Vec`，当目标为32位类型（`float`/`int32_t`）且使用 `DualModeSplitN` 时，切分前的 `ValidCol` 必须是 `32` 的整数倍。
150 147 
151-### A5 Vec ND→Mat NZ 提取
152- 
153-ND Vec 源可直接提取到非 compact 的 NZ512 Mat 目标。目标有效形状定义从
154-`(indexRow, indexCol)` 开始的源窗口，完整窗口须位于源有效形状内。
155-列宽、源物理行跨度及列偏移须 32 字节对齐，行偏移和有效行数无需分形对齐。
156-仅写入有效窗口，目标 padding 保持不变。共享的存储与 MTE3 同步要求见
157-[UB ND → L1 NZ](TMOV_zh.md#ub-nd--l1-nza5)。
158- 
159### A5 Acc→Mat NZ 布局转换148### A5 Acc→Mat NZ 布局转换
160 149 
161以下路径的目标为 NZ 布局（`BLayout::ColMajor`、`SLayout::RowMajor`）：150以下路径的目标为 NZ 布局（`BLayout::ColMajor`、`SLayout::RowMajor`）：
@@ -328,8 +317,6 @@ A5 主机测试模拟底层搬运指令并执行真实 A5 后端分发入口及
328 317 
329除上述 `Mat/Acc -> ...` 路径外，`TEXTRACT` 还支持 `TileType::Vec -> TileType::Vec` 抽取路径（ND 与 NZ 布局）。A2A3 使用 `CheckTExtractVecToVecCommon`，A5 在 `TEXTRACT_IMPL` 中单独检查：318除上述 `Mat/Acc -> ...` 路径外，`TEXTRACT` 还支持 `TileType::Vec -> TileType::Vec` 抽取路径（ND 与 NZ 布局）。A2A3 使用 `CheckTExtractVecToVecCommon`，A5 在 `TEXTRACT_IMPL` 中单独检查：
330 319 
331-- A2A3 普通双 Tile 重载要求两个 Vec Tile 同为 ND（`RowMajor`、`NoneBox`）
332- 或同为 NZ（`ColMajor`、`RowMajor`）；其他布局组合在编译期报错。
333- `DstTileData::DType` 必须等于 `SrcTileData::DType`。320- `DstTileData::DType` 必须等于 `SrcTileData::DType`。
334- A2A3 元素类型：`int8_t`、`uint8_t`、`int16_t`、`uint16_t`、`int32_t`、`uint32_t`、`half`、`bfloat16_t`、`float`。321- A2A3 元素类型：`int8_t`、`uint8_t`、`int16_t`、`uint16_t`、`int32_t`、`uint32_t`、`half`、`bfloat16_t`、`float`。
335- A5 元素类型：`int8_t`、`int32_t`、`half`、`bfloat16_t`、`float`、`hifloat8_t`、`float8_e4m3_t`、`float8_e5m2_t`、`float8_e8m0_t`、`float4_e2m1x2_t`、`float4_e1m2x2_t`。A5 不支持此路径的 `uint8_t`、`int16_t`、`uint16_t`、`uint32_t` 或 64 位整数。322- A5 元素类型：`int8_t`、`int32_t`、`half`、`bfloat16_t`、`float`、`hifloat8_t`、`float8_e4m3_t`、`float8_e5m2_t`、`float8_e8m0_t`、`float4_e2m1x2_t`、`float4_e1m2x2_t`。A5 不支持此路径的 `uint8_t`、`int16_t`、`uint16_t`、`uint32_t` 或 64 位整数。
Mdocs/isa/TINSERT.md+0-5
@@ -249,11 +249,6 @@ canonical overload.
249- **Vec → Mat** (`TileType::Vec → TileType::Mat`, UB → L1):249- **Vec → Mat** (`TileType::Vec → TileType::Mat`, UB → L1):
250 - `DstTileData::DType` must equal `SrcTileData::DType`.250 - `DstTileData::DType` must equal `SrcTileData::DType`.
251 - Supported element types: `half`, `bfloat16_t`, `float`, `int32_t`, `int8_t`, `hifloat8_t`, `float8_e4m3_t`, `float8_e5m2_t`, `float8_e8m0_t`, `float4_e2m1x2_t`, `float4_e1m2x2_t`.251 - Supported element types: `half`, `bfloat16_t`, `float`, `int32_t`, `int8_t`, `hifloat8_t`, `float8_e4m3_t`, `float8_e5m2_t`, `float8_e8m0_t`, `float4_e2m1x2_t`, `float4_e1m2x2_t`.
252- - ND-to-NZ path (A5): an ND (`RowMajor`, `NoneBox`) Vec source can be inserted into a non-compact NZ512 Mat destination.
253- Source row stride, valid column width, and `indexCol` must be 32-byte aligned; `indexRow` need not be fractal aligned.
254- Only the source valid region is written, preserving surrounding data. The window must fit the destination physical shape.
255- Packed FP4 columns count logical elements. See [UB ND → L1 NZ](TMOV.md#ub-nd--l1-nz-a5) for MTE3
256- synchronization and other shared constraints.
257 - ND path: source must be `isRowMajor`; uses `copy_ubuf_to_cbuf`. Data bytes per row must be aligned to `BLOCK_BYTE_SIZE` (32 bytes) for row-wise burst.252 - ND path: source must be `isRowMajor`; uses `copy_ubuf_to_cbuf`. Data bytes per row must be aligned to `BLOCK_BYTE_SIZE` (32 bytes) for row-wise burst.
258 - NZ path: source must be `(!isRowMajor, SFractal: RowMajor)`; uses `ComputeNZBlockParams` for fractal-block `copy_ubuf_to_cbuf`. For fp4 types (`float4_e2m1x2_t`, `float4_e1m2x2_t`), validCol and indexCol are halved for byte addressing.253 - NZ path: source must be `(!isRowMajor, SFractal: RowMajor)`; uses `ComputeNZBlockParams` for fractal-block `copy_ubuf_to_cbuf`. For fp4 types (`float4_e2m1x2_t`, `float4_e1m2x2_t`), validCol and indexCol are halved for byte addressing.
259 - ZN path: source and destination must both be `(isRowMajor, SFractal: ColMajor)`; uses `ComputeZNBlockParams` for fractal-block `copy_ubuf_to_cbuf`. ZN is the transpose-dual of NZ: rows are the fractal dimension and columns are the free dimension. `indexRow` and `validRow` must be aligned to the fractal row size (`BLOCK_BYTE_SIZE / sizeof(T)`); `indexCol` and `validCol` are unconstrained (free dimension), matching the NZ row dimension. For fp4 types (`float4_e2m1x2_t`, `float4_e1m2x2_t`), validRow and indexRow are halved for byte addressing.254 - ZN path: source and destination must both be `(isRowMajor, SFractal: ColMajor)`; uses `ComputeZNBlockParams` for fractal-block `copy_ubuf_to_cbuf`. ZN is the transpose-dual of NZ: rows are the fractal dimension and columns are the free dimension. `indexRow` and `validRow` must be aligned to the fractal row size (`BLOCK_BYTE_SIZE / sizeof(T)`); `indexCol` and `validCol` are unconstrained (free dimension), matching the NZ row dimension. For fp4 types (`float4_e2m1x2_t`, `float4_e1m2x2_t`), validRow and indexRow are halved for byte addressing.
Mdocs/isa/TINSERT_zh.md+0-4
@@ -236,10 +236,6 @@ Ascend 950PR/Ascend 950DT 和 CPU 模拟器），适用范围与配对规则见
236- **Vec → Mat**（`TileType::Vec → TileType::Mat`，UB → L1）：236- **Vec → Mat**（`TileType::Vec → TileType::Mat`，UB → L1）：
237 - `DstTileData::DType` 必须等于 `SrcTileData::DType`。237 - `DstTileData::DType` 必须等于 `SrcTileData::DType`。
238 - 支持的元素类型：`half`、`bfloat16_t`、`float`、`int32_t`、`int8_t`、`hifloat8_t`、`float8_e4m3_t`、`float8_e5m2_t`、`float8_e8m0_t`、`float4_e2m1x2_t`、`float4_e1m2x2_t`。238 - 支持的元素类型：`half`、`bfloat16_t`、`float`、`int32_t`、`int8_t`、`hifloat8_t`、`float8_e4m3_t`、`float8_e5m2_t`、`float8_e8m0_t`、`float4_e2m1x2_t`、`float4_e1m2x2_t`。
239- - ND→NZ 路径（A5）：ND（`RowMajor`、`NoneBox`）Vec 源可插入非 compact 的 NZ512 Mat 目标。
240- 源物理行跨度、有效列宽和 `indexCol` 须 32 字节对齐，`indexRow` 无需分形对齐。
241- 仅写入源有效区域，周围数据保持不变，窗口须位于目标物理形状内。
242- packed FP4 列数按逻辑元素计数。MTE3 同步及其他共享约束见 [UB ND → L1 NZ](TMOV_zh.md#ub-nd--l1-nza5)。
243 - ND路径：源必须为 `isRowMajor`；使用 `copy_ubuf_to_cbuf`。每行数据字节数必须与 `BLOCK_BYTE_SIZE`（32字节）对齐。239 - ND路径：源必须为 `isRowMajor`；使用 `copy_ubuf_to_cbuf`。每行数据字节数必须与 `BLOCK_BYTE_SIZE`（32字节）对齐。
244 - NZ路径：源必须为 `(!isRowMajor, SFractal: RowMajor)`；使用 `ComputeNZBlockParams` 进行分形块 `copy_ubuf_to_cbuf`。240 - NZ路径：源必须为 `(!isRowMajor, SFractal: RowMajor)`；使用 `ComputeNZBlockParams` 进行分形块 `copy_ubuf_to_cbuf`。
245 241 
Mdocs/isa/TMOV.md+6-40
@@ -26,17 +26,17 @@ $$ \mathrm{dst}_{i,j} = \mathrm{src}_{i,j} $$
26### ND → NZ (data repack for the Cube Unit)26### ND → NZ (data repack for the Cube Unit)
27 27 
28The Cube Unit consumes operands in **NZ** (Normal-ZigZag) fractal format: the tile is28The Cube Unit consumes operands in **NZ** (Normal-ZigZag) fractal format: the tile is
29-tiled into `16 × C0` fractals, each fractal stored with `BLayout = ColMajor` (the "N" —29+tiled into `C0 × C0` fractals, each fractal stored with `BLayout = ColMajor` (the "N" —
30column-major outer blocks) and `SLayout = RowMajor` (the "Z" — row-major within each30column-major outer blocks) and `SLayout = RowMajor` (the "Z" — row-major within each
31-fractal). `TMOV(dstNZ, src)` repacks a RowMajor `Vec` tile (`NoneBox`) into an NZ `Vec`31+fractal). `TMOV(dstNZ, src)` repacks a RowMajor `Vec`/`Mat` tile (`NoneBox`) into this
32-or, on A5, an NZ `Mat` tile. No `tmp` is required.32+NZ layout. No `tmp` is required.
33 33 
34| Operand (GM/L1 side) | `BLayout` | `SLayout` | Meaning |34| Operand (GM/L1 side) | `BLayout` | `SLayout` | Meaning |
35|----------------------|-----------|-----------|---------|35|----------------------|-----------|-----------|---------|
36| Left (A, NT) | `ColMajor` | `RowMajor` | normal NZ |36| Left (A, NT) | `ColMajor` | `RowMajor` | normal NZ |
37| Right (B, NT) | `RowMajor` | `ColMajor` | transposed NZ |37| Right (B, NT) | `RowMajor` | `ColMajor` | transposed NZ |
38 38 
39-For UB-to-UB repacking, a `CompactMode::RowPlusOne` destination (`Rows = Vec_S0 + 1`) is the canonical idiom to39+A `CompactMode::RowPlusOne` destination (`Rows = Vec_S0 + 1`) is the canonical idiom to
40avoid UB bank conflicts on the `vsstb` scatter.40avoid UB bank conflicts on the `vsstb` scatter.
41 41 
42### ND → ZN (within-fractal transpose for the transposed Cube operand)42### ND → ZN (within-fractal transpose for the transposed Cube operand)
@@ -50,7 +50,7 @@ $N \bmod 16 = 0$. Supported element types: `half`, `bfloat16_t`, `float`, `int32
50`int8_t`, `uint8_t`, `float8_e4m3_t`, `float8_e5m2_t`, `hifloat8_t`,50`int8_t`, `uint8_t`, `float8_e4m3_t`, `float8_e5m2_t`, `hifloat8_t`,
51`float4_e2m1x2_t`, `float4_e1m2x2_t` (all 1/2/4-byte storage types).51`float4_e2m1x2_t`, `float4_e1m2x2_t` (all 1/2/4-byte storage types).
52 52 
53-Unlike UB-to-UB ND→NZ (which only rearranges 32 B blocks on the block grid via `vsstb` and leaves53+Unlike ND→NZ (which only rearranges 32 B blocks on the block grid via `vsstb` and leaves
54within-block data untouched), ND→ZN must also transpose the elements **inside** each54within-block data untouched), ND→ZN must also transpose the elements **inside** each
55$K_0 \times 16$ fractal, so `vsstb` alone is insufficient. Each output fractal is the55$K_0 \times 16$ fractal, so `vsstb` alone is insufficient. Each output fractal is the
56transpose of a $K_0 \times 16$ source slice:56transpose of a $K_0 \times 16$ source slice:
@@ -109,36 +109,6 @@ Only the **X→ZZ** transforms take a `tmp` operand (the 3-arg overload). ND→Z
109the `vgather2` index buffer; DN→ZZ accepts it for interface parity but does not access it109the `vgather2` index buffer; DN→ZZ accepts it for interface parity but does not access it
110(the `vsstb` scatter needs no scratch). ND→NZ has no `tmp`.110(the `vsstb` scatter needs no scratch). ND→NZ has no `tmp`.
111 111 
112-### UB ND → L1 NZ (A5)
113- 
114-`TMOV(dst, src)` converts an ND `TileType::Vec` source (`RowMajor`, `NoneBox`) directly
115-into an NZ512 `TileType::Mat` destination (`ColMajor`, `RowMajor`, `CompactMode::Null`).
116-The data type and valid shape must match. This path preserves the bits of the supported
117-b8/b16/b32 types and packed FP4 types; it performs no numerical conversion.
118- 
119-The implementation uses multiple strided 32-byte MTE3 transfers, selecting a row or
120-column-block traversal. It requires neither a temporary UB tile nor GM staging.
121-Source row stride and valid column width must be 32-byte aligned; FP4 widths count
122-logical 4-bit elements (64 elements per 32-byte block). Addresses must be 32-byte aligned.
123-Physical source row stride and destination row count determine storage offsets, independently
124-of the valid shape. The source row stride in 32-byte blocks and the destination row count
125-must each be at most 65536 so that the DMA gaps fit in 16 bits. Both windows must fit their
126-allocated storage. Row tails need no 16-row alignment; padding outside the valid region is
127-preserved. An empty region performs no transfer. Non-aligned column tails and compact/NZ1024 destinations are not supported.
128-Static layout/stride constraints are checked at compile time; dynamic bounds and alignment
129-checks use `PTO_ASSERT` and are enabled by `_DEBUG`.
130- 
131-Issue the operation on AIV. Synchronize the producer with `PIPE_MTE3`, then notify AIC
132-after MTE3 completes and wait on `PIPE_MTE1` before consuming L1. Multiple AIVs must write
133-disjoint destination regions, and AIC must wait for every participating AIV. Calls that
134-write overlapping L1 regions require MTE3 synchronization between them.
135- 
136-`TEXTRACT(dst, src, row, col)` uses the same conversion for an ND source window; its
137-column offset must be 32-byte aligned and the window must fit the source valid shape.
138-`TINSERT(dst, src, row, col)` converts into a destination window with a 32-byte-aligned
139-column offset. The destination window must fit its physical shape. Row offsets need
140-no fractal alignment. These ordinary forms do not require a `tmp` operand.
141- 
142## Assembly Syntax112## Assembly Syntax
143 113 
144The PTO AS design recommends splitting `TMOV` into a family of ops:114The PTO AS design recommends splitting `TMOV` into a family of ops:
@@ -364,9 +334,6 @@ A5 Vec-to-Vec moves currently do not support `int64_t` / `uint64_t`; 64-bit Scal
364 334 
365### ND → NZ (data) — (128, 256) BF16335### ND → NZ (data) — (128, 256) BF16
366 336 
367-This Vec-to-Mat example targets A5. Issue it on AIV and follow the
368-[MTE3 synchronization requirements](#ub-nd--l1-nz-a5) before AIC consumes the result.
369- 
370```cpp337```cpp
371// Source: 128 rows × 256 cols BF16 RowMajor Vec tile (ND).338// Source: 128 rows × 256 cols BF16 RowMajor Vec tile (ND).
372// Destination: NZ fractal Mat tile for the Cube Unit (Left operand).339// Destination: NZ fractal Mat tile for the Cube Unit (Left operand).
@@ -377,8 +344,7 @@ SrcT src; DstT dst;
377TMOV(dst, src); // ND -> NZ, no tmp344TMOV(dst, src); // ND -> NZ, no tmp
378```345```
379 346 
380-**`tmp` for ND→NZ:** none. This Vec-to-Mat example uses strided MTE3 transfers;347+**`tmp` for ND→NZ:** none — the 2-arg overload repacks in-place via `vsstb`.
381-Vec-to-Vec repacking uses `vsstb` in UB.
382 348 
383### ND → ZZ (exponents) — `tmp` size derivation349### ND → ZZ (exponents) — `tmp` size derivation
384 350 
Mdocs/isa/TMOV_zh.md+3-32
@@ -24,14 +24,14 @@ $$ \mathrm{dst}_{i,j} = \mathrm{src}_{i,j} $$
24 24 
25### ND → NZ（为Cube Unit重新打包数据）25### ND → NZ（为Cube Unit重新打包数据）
26 26 
27-Cube Unit 以 **NZ**（Normal-ZigZag）fractal 格式消费操作数：Tile 被划分为 `16 × C0` fractal，每个 fractal 以 `BLayout = ColMajor`（"N"——列主序外层块）和 `SLayout = RowMajor`（"Z"——fractal 内行主序）存储。`TMOV(dstNZ, src)` 将 RowMajor `Vec` Tile（`NoneBox`）重新打包为 NZ `Vec` Tile，或在 A5 上转换到 NZ `Mat` Tile。无需 `tmp`。27+Cube Unit以 **NZ**（Normal-ZigZag）fractal格式消费操作数：Tile被划分为 `C0 × C0` fractal，每个fractal以 `BLayout = ColMajor`（"N"——列主序外层块）和 `SLayout = RowMajor`（"Z"——fractal内行主序）存储。`TMOV(dstNZ, src)` 将RowMajor `Vec`/`Mat` Tile（`NoneBox`）重新打包为NZ布局。无需 `tmp`。
28 28 
29| 操作数（GM/L1侧） | `BLayout` | `SLayout` | 含义 |29| 操作数（GM/L1侧） | `BLayout` | `SLayout` | 含义 |
30|-------------------|-----------|-----------|------|30|-------------------|-----------|-----------|------|
31| Left（A，NT） | `ColMajor` | `RowMajor` | 标准NZ |31| Left（A，NT） | `ColMajor` | `RowMajor` | 标准NZ |
32| Right（B，NT） | `RowMajor` | `ColMajor` | 转置NZ |32| Right（B，NT） | `RowMajor` | `ColMajor` | 转置NZ |
33 33 
34-UB 内重排时，`CompactMode::RowPlusOne` 目标（`Rows = Vec_S0 + 1`）是避免 `vsstb` scatter 上 UB bank 冲突的标准用法。34+`CompactMode::RowPlusOne` 目标（`Rows = Vec_S0 + 1`）是避免 `vsstb` scatter上UB bank冲突的标准用法。
35 35 
36### X → ZZ（microscaling指数重新打包）36### X → ZZ（microscaling指数重新打包）
37 37 
@@ -55,32 +55,6 @@ $$E_{ZZ}[c_b, p, q, \delta] = E_{DN}^{T}[16c_b + q][2p + \delta] = E_{DN}[2p + \
55 55 
56仅 **X→ZZ** 转换接受 `tmp` 操作数（3参数重载）。ND→ZZ用作 `vgather2` 索引缓冲；DN→ZZ为接口一致接受但**不访问**（`vsstb` scatter无需scratch）。ND→NZ无 `tmp`。56仅 **X→ZZ** 转换接受 `tmp` 操作数（3参数重载）。ND→ZZ用作 `vgather2` 索引缓冲；DN→ZZ为接口一致接受但**不访问**（`vsstb` scatter无需scratch）。ND→NZ无 `tmp`。
57 57 
58-### UB ND → L1 NZ（A5）
59- 
60-`TMOV(dst, src)` 将 ND `TileType::Vec` 源（`RowMajor`、`NoneBox`）直接转换到
61-NZ512 `TileType::Mat` 目标（`ColMajor`、`RowMajor`、`CompactMode::Null`）。
62-源和目标的数据类型、有效形状必须相同。此路径按位保留支持的 b8/b16/b32 类型及
63-packed FP4 类型的数据，不进行数值类型转换。
64- 
65-实现使用多条跨步 32 字节 MTE3 搬运指令，按形状选择逐行或逐列块遍历，
66-无需 UB 临时 Tile 或 GM 中转。源物理行跨度和有效列宽须 32 字节对齐；
67-FP4 列数按逻辑 4-bit 元素计数（每 32 字节块含 64 个元素），源和目标地址须 32 字节对齐。
68-地址计算使用源物理行跨度和目标物理行数，不以有效形状替代。
69-源物理行跨度（以 32 字节块为单位）和目标物理行数均不得超过 65536，以满足 DMA 间隔的 16 位限制。
70-搬运窗口必须位于已分配存储内。
71-有效行数无需 16 行对齐，有效区域外的 padding 保持不变；空区域不执行搬运。
72-不支持非对齐列尾、compact 目标或 NZ1024 目标。
73-静态布局和跨度约束在编译期检查，动态边界和对齐检查使用 `PTO_ASSERT`，仅在 `_DEBUG` 下生效。
74- 
75-此操作由 AIV 发起。数据生产者须同步到 `PIPE_MTE3`，MTE3 搬运完成后通知 AIC，
76-AIC 在 `PIPE_MTE1` 上等待后再消费 L1。多个 AIV 须写入互不重叠的目标区域，
77-AIC 须等待所有参与搬运的 AIV。连续调用写入重叠的 L1 区域时，须在调用间同步 MTE3。
78- 
79-`TEXTRACT(dst, src, row, col)` 对 ND 源窗口执行同样的转换；列偏移须 32 字节对齐，
80-窗口须位于源有效形状内。`TINSERT(dst, src, row, col)` 转换到目标窗口，
81-列偏移须 32 字节对齐，窗口须位于目标物理形状内。行偏移无需分形对齐。
82-这些普通形式均不需要 `tmp` 操作数。
83- 
84## 汇编语法58## 汇编语法
85 59 
86PTO AS设计建议将 `TMOV` 拆分为一组操作：60PTO AS设计建议将 `TMOV` 拆分为一组操作：
@@ -285,9 +259,6 @@ A5 的 Vec→Vec 移动当前不支持 `int64_t` / `uint64_t`；其他传输路
285 259 
286### ND → NZ（数据）— (128, 256) BF16260### ND → NZ（数据）— (128, 256) BF16
287 261 
288-此 Vec→Mat 示例适用于 A5，在 AIV 上执行；AIC 使用结果前须遵循
289-[MTE3 同步要求](#ub-nd--l1-nza5)。
290- 
291```cpp262```cpp
292// 源：128 行 × 256 列 BF16 RowMajor Vec Tile（ND）。263// 源：128 行 × 256 列 BF16 RowMajor Vec Tile（ND）。
293// 目标：Cube Unit 的 NZ fractal Mat Tile（Left 操作数）。264// 目标：Cube Unit 的 NZ fractal Mat Tile（Left 操作数）。
@@ -298,7 +269,7 @@ SrcT src; DstT dst;
298TMOV(dst, src); // ND -> NZ，无 tmp269TMOV(dst, src); // ND -> NZ，无 tmp
299```270```
300 271 
301-**ND→NZ 的 `tmp`：** 无。此 Vec→Mat 示例通过 MTE3 跨步搬运转换布局；Vec→Vec 路径在 UB 内使用 `vsstb` 重排。272+**ND→NZ的 `tmp`：** 无——2参数重载通过 `vsstb` 原地重新打包。
302 273 
303### ND → ZZ（指数）— `tmp` 尺寸推导274### ND → ZZ（指数）— `tmp` 尺寸推导
304 275 
Minclude/pto/npu/a2a3/TExtract.hpp+7-6
@@ -182,8 +182,6 @@ PTO_INTERNAL void TEXTRACT_TILE_IMPL(DstTileData& dst, SrcTileData& src, uint16_
182 GetCastPreQuantMode<typename SrcTileData::DType, typename DstTileData::DType>();182 GetCastPreQuantMode<typename SrcTileData::DType, typename DstTileData::DType>();
183 TExtractAccToMat<DstTileData, SrcTileData, quantPre, ReluPreMode::NoRelu>(183 TExtractAccToMat<DstTileData, SrcTileData, quantPre, ReluPreMode::NoRelu>(
184 dst.data(), src.data(), dst.GetValidRow(), dst.GetValidCol(), indexRow, indexCol);184 dst.data(), src.data(), dst.GetValidRow(), dst.GetValidCol(), indexRow, indexCol);
185- } else {
186- static_assert(sizeof(DstTileData) == 0, "TExtract: Unsupported source and destination TileType combination.");
187 }185 }
188}186}
189 187 
@@ -582,16 +580,19 @@ PTO_INTERNAL void TEXTRACT_IMPL(DstTileData& dst, SrcTileData& src, uint16_t ind
582{580{
583 if constexpr (DstTileData::Loc == TileType::Vec && SrcTileData::Loc == TileType::Vec) {581 if constexpr (DstTileData::Loc == TileType::Vec && SrcTileData::Loc == TileType::Vec) {
584 CheckTExtractVecToVecCommon<DstTileData, SrcTileData>();582 CheckTExtractVecToVecCommon<DstTileData, SrcTileData>();
585- if constexpr (583+ if constexpr (DstTileData::isRowMajor && SrcTileData::isRowMajor) {
586- DstTileData::isRowMajor && SrcTileData::isRowMajor && DstTileData::SFractal == SLayout::NoneBox &&
587- SrcTileData::SFractal == SLayout::NoneBox) {
588 TExtractVecToVecNDDispatch<DstTileData, SrcTileData>(dst, src, indexRow, indexCol);584 TExtractVecToVecNDDispatch<DstTileData, SrcTileData>(dst, src, indexRow, indexCol);
589 } else if constexpr (585 } else if constexpr (
590 !DstTileData::isRowMajor && !SrcTileData::isRowMajor && DstTileData::SFractal == SLayout::RowMajor &&586 !DstTileData::isRowMajor && !SrcTileData::isRowMajor && DstTileData::SFractal == SLayout::RowMajor &&
591 SrcTileData::SFractal == SLayout::RowMajor) {587 SrcTileData::SFractal == SLayout::RowMajor) {
592 TExtractVecToVecNZDispatch<DstTileData, SrcTileData>(dst, src, indexRow, indexCol);588 TExtractVecToVecNZDispatch<DstTileData, SrcTileData>(dst, src, indexRow, indexCol);
593 } else {589 } else {
594- static_assert(sizeof(DstTileData) == 0, "TEXTRACT Vec->Vec : Only ND-to-ND and NZ-to-NZ are supported.");590+ static_assert(
591+ DstTileData::isRowMajor == SrcTileData::isRowMajor,
592+ "TEXTRACT Vec->Vec : Source and destination layout must match (both ND or both NZ).");
593+ static_assert(
594+ DstTileData::SFractal == SrcTileData::SFractal,
595+ "TEXTRACT Vec->Vec : Source and destination SFractal must match.");
595 }596 }
596 } else if constexpr (is_conv_tile_v<SrcTileData>) {597 } else if constexpr (is_conv_tile_v<SrcTileData>) {
597 TEXTRACT_CONVTILE_IMPL(dst, src, indexRow, indexCol);598 TEXTRACT_CONVTILE_IMPL(dst, src, indexRow, indexCol);
Minclude/pto/npu/a5/TExtract.hpp+3-13
@@ -177,19 +177,9 @@ PTO_INTERNAL void TEXTRACT_TILE_IMPL(DstTileData& dst, SrcTileData& src, uint16_
177 } else if constexpr (DstTileData::Loc == TileType::Right) {177 } else if constexpr (DstTileData::Loc == TileType::Right) {
178 TExtractToRight(dst, src, indexRow, indexCol);178 TExtractToRight(dst, src, indexRow, indexCol);
179 } else if constexpr (SrcTileData::Loc == TileType::Vec && DstTileData::Loc == TileType::Mat) {179 } else if constexpr (SrcTileData::Loc == TileType::Vec && DstTileData::Loc == TileType::Mat) {
180- if constexpr (180+ TExtractVecToMat<DstTileData, SrcTileData>(
181- SrcTileData::isRowMajor && SrcTileData::SFractal == SLayout::NoneBox && !DstTileData::isRowMajor &&181+ dst.data(), src.data(), indexRow, indexCol, src.GetValidRow(), src.GetValidCol(), dst.GetValidRow(),
182- DstTileData::SFractal == SLayout::RowMajor) {182+ dst.GetValidCol());
183- PTO_ASSERT(
184- indexRow + dst.GetValidRow() <= src.GetValidRow() && indexCol + dst.GetValidCol() <= src.GetValidCol(),
185- "TEXTRACT ND-to-NZ source window exceeds valid shape.");
186- TCopyNdToNzUbToMat<DstTileData, SrcTileData>(
187- dst.data(), src.data(), dst.GetValidRow(), dst.GetValidCol(), indexRow, indexCol);
188- } else {
189- TExtractVecToMat<DstTileData, SrcTileData>(
190- dst.data(), src.data(), indexRow, indexCol, src.GetValidRow(), src.GetValidCol(), dst.GetValidRow(),
191- dst.GetValidCol());
192- }
193 } else if constexpr (DstTileData::Loc == TileType::ScaleLeft) {183 } else if constexpr (DstTileData::Loc == TileType::ScaleLeft) {
194 TExtractToAmx<DstTileData, SrcTileData>(184 TExtractToAmx<DstTileData, SrcTileData>(
195 dst.data(), src.data(), indexRow, indexCol, dst.GetValidRow(), dst.GetValidCol());185 dst.data(), src.data(), indexRow, indexCol, dst.GetValidRow(), dst.GetValidCol());
Minclude/pto/npu/a5/TExtractCommon.hpp+0-68
@@ -398,74 +398,6 @@ __tf__ PTO_INTERNAL void TExtractToBTransCompact(
398 }398 }
399}399}
400 400 
401-#if defined(PTO_NPU_ARCH_A5)
402-template <typename DstTileData, typename SrcTileData>
403-__tf__ PTO_INTERNAL void TCopyNdToNzUbToMat(
404- typename DstTileData::TileDType __out__ dst, typename SrcTileData::TileDType __in__ src, uint32_t validRows,
405- uint32_t validCols, uint32_t srcRowOffset = 0, uint32_t srcColOffset = 0, uint32_t dstRowOffset = 0,
406- uint32_t dstColOffset = 0)
407-{
408- using T = typename SrcTileData::DType;
409- static_assert(std::is_same_v<T, typename DstTileData::DType>, "UB-to-L1 ND-to-NZ requires matching data types.");
410- static_assert(sizeof(T) <= 4, "UB-to-L1 ND-to-NZ supports b4/b8/b16/b32 data.");
411- static_assert(
412- SrcTileData::Loc == TileType::Vec && SrcTileData::isRowMajor && SrcTileData::SFractal == SLayout::NoneBox,
413- "UB-to-L1 ND-to-NZ requires an ND Vec source.");
414- static_assert(
415- DstTileData::Loc == TileType::Mat && !DstTileData::isRowMajor && DstTileData::SFractal == SLayout::RowMajor &&
416- DstTileData::SFractalSize == CUBE_BLOCK_SIZE && DstTileData::Compact == CompactMode::Null,
417- "UB-to-L1 ND-to-NZ requires a non-compact NZ512 Mat destination.");
418- constexpr bool IS_FP4 = std::is_same_v<T, float4_e2m1x2_t> || std::is_same_v<T, float4_e1m2x2_t>;
419- constexpr uint32_t C0_SIZE = BLOCK_BYTE_SIZE / sizeof(T) * (IS_FP4 ? 2 : 1);
420- constexpr uint32_t SRC_ROW_STRIDE_BLOCKS = SrcTileData::RowStride / C0_SIZE;
421- constexpr uint32_t DST_COL_STRIDE_BLOCKS = DstTileData::Rows;
422- static_assert(SrcTileData::RowStride % C0_SIZE == 0, "ND source row stride must be 32-byte aligned.");
423- static_assert(
424- SRC_ROW_STRIDE_BLOCKS > 0 && SRC_ROW_STRIDE_BLOCKS <= 65536 && DST_COL_STRIDE_BLOCKS <= 65536,
425- "UB-to-L1 ND-to-NZ DMA gaps must fit in 16 bits.");
426- PTO_ASSERT(
427- validCols % C0_SIZE == 0 && srcColOffset % C0_SIZE == 0 && dstColOffset % C0_SIZE == 0,
428- "UB-to-L1 ND-to-NZ columns and column offsets must be 32-byte aligned.");
429- PTO_ASSERT(
430- srcRowOffset + validRows <= SrcTileData::Rows && srcColOffset + validCols <= SrcTileData::Cols,
431- "UB-to-L1 ND-to-NZ source window exceeds storage.");
432- PTO_ASSERT(
433- dstRowOffset + validRows <= DstTileData::Rows && dstColOffset + validCols <= DstTileData::Cols,
434- "UB-to-L1 ND-to-NZ destination window exceeds storage.");
435- if (validRows == 0 || validCols == 0) {
436- return;
437- }
438- __ubuf__ uint8_t* srcPtr = (__ubuf__ uint8_t*)__cce_get_tile_ptr(src) +
439- (srcRowOffset * SRC_ROW_STRIDE_BLOCKS + srcColOffset / C0_SIZE) * BLOCK_BYTE_SIZE;
440- __cbuf__ uint8_t* dstPtr = (__cbuf__ uint8_t*)__cce_get_tile_ptr(dst) +
441- (dstColOffset / C0_SIZE * DST_COL_STRIDE_BLOCKS + dstRowOffset) * BLOCK_BYTE_SIZE;
442- uint32_t colBlockCount = validCols / C0_SIZE;
443- constexpr uint32_t MAX_BURST_COUNT = 4095;
444- // Each burst moves one row fragment. Choose the shorter outer loop without UB scratch.
445- if (colBlockCount > validRows) {
446- for (uint32_t row = 0; row < validRows; ++row) {
447- for (uint32_t colBlock = 0; colBlock < colBlockCount; colBlock += MAX_BURST_COUNT) {
448- uint16_t burstCount = static_cast<uint16_t>(min(colBlockCount - colBlock, MAX_BURST_COUNT));
449- copy_ubuf_to_cbuf(
450- dstPtr + (colBlock * DST_COL_STRIDE_BLOCKS + row) * BLOCK_BYTE_SIZE,
451- srcPtr + (row * SRC_ROW_STRIDE_BLOCKS + colBlock) * BLOCK_BYTE_SIZE, 0, burstCount, 1, 0,
452- DST_COL_STRIDE_BLOCKS - 1);
453- }
454- }
455- } else {
456- for (uint32_t colBlock = 0; colBlock < colBlockCount; ++colBlock) {
457- for (uint32_t row = 0; row < validRows; row += MAX_BURST_COUNT) {
458- uint16_t burstCount = static_cast<uint16_t>(min(validRows - row, MAX_BURST_COUNT));
459- copy_ubuf_to_cbuf(
460- dstPtr + (colBlock * DST_COL_STRIDE_BLOCKS + row) * BLOCK_BYTE_SIZE,
461- srcPtr + (row * SRC_ROW_STRIDE_BLOCKS + colBlock) * BLOCK_BYTE_SIZE, 0, burstCount, 1,
462- SRC_ROW_STRIDE_BLOCKS - 1, 0);
463- }
464- }
465- }
466-}
467-#endif
468- 
469template <typename DstTileData, typename SrcTileData>401template <typename DstTileData, typename SrcTileData>
470__tf__ PTO_INTERNAL void TExtractVecToMat(402__tf__ PTO_INTERNAL void TExtractVecToMat(
471 typename DstTileData::TileDType __out__ dst, typename SrcTileData::TileDType __in__ src, uint16_t indexRow,403 typename DstTileData::TileDType __out__ dst, typename SrcTileData::TileDType __in__ src, uint16_t indexRow,
Minclude/pto/npu/a5/TInsert.hpp+0-10
@@ -12,9 +12,6 @@ See LICENSE in the root of the software repository for the full text of the Lice
12#define TINSERT_HPP12#define TINSERT_HPP
13#include "common.hpp"13#include "common.hpp"
14#include "utils.hpp"14#include "utils.hpp"
15-#if defined(PTO_NPU_ARCH_A5)
16-#include "TExtractCommon.hpp"
17-#endif
18#include "pto/common/arch/register/tinsert_common.hpp"15#include "pto/common/arch/register/tinsert_common.hpp"
19 16 
20namespace pto {17namespace pto {
@@ -420,13 +417,6 @@ PTO_INTERNAL void TInsertVecToMatImpl(DstTileData& dst, SrcTileData& src, uint16
420 PTO_ASSERT(validRow % irSize == 0, "TINSERT ZN : validRow must be aligned to the fractal row size!");417 PTO_ASSERT(validRow % irSize == 0, "TINSERT ZN : validRow must be aligned to the fractal row size!");
421 TInsertZNImpl<T, DstTileData, SrcTileData>(418 TInsertZNImpl<T, DstTileData, SrcTileData>(
422 dst.data(), src.data(), validRow, validCol, dstCol, indexRow, indexCol);419 dst.data(), src.data(), validRow, validCol, dstCol, indexRow, indexCol);
423-#if defined(PTO_NPU_ARCH_A5)
424- } else if constexpr (
425- SrcTileData::isRowMajor && SrcTileData::SFractal == SLayout::NoneBox && !DstTileData::isRowMajor &&
426- DstTileData::SFractal == SLayout::RowMajor) {
427- TCopyNdToNzUbToMat<DstTileData, SrcTileData>(
428- dst.data(), src.data(), src.GetValidRow(), src.GetValidCol(), 0, 0, indexRow, indexCol);
429-#endif
430 } else if constexpr (SrcTileData::isRowMajor) {420 } else if constexpr (SrcTileData::isRowMajor) {
431 uint16_t dstCols = static_cast<uint16_t>(DstTileData::Cols);421 uint16_t dstCols = static_cast<uint16_t>(DstTileData::Cols);
432 TInsertNDImpl<T, DstTileData, SrcTileData>(422 TInsertNDImpl<T, DstTileData, SrcTileData>(
Minclude/pto/npu/a5/TMov.hpp+3-13
@@ -811,19 +811,9 @@ PTO_INTERNAL void TMOV_TILE_IMPL(DstTileData& dst, SrcTileData& src)
811 }811 }
812 } else if constexpr (DstTileData::Loc == TileType::Mat) {812 } else if constexpr (DstTileData::Loc == TileType::Mat) {
813 CommonCheck<DstTileData, SrcTileData>();813 CommonCheck<DstTileData, SrcTileData>();
814- if constexpr (814+ TExtractVecToMat<DstTileData, SrcTileData>(
815- SrcTileData::isRowMajor && SrcTileData::SFractal == SLayout::NoneBox && !DstTileData::isRowMajor &&815+ dst.data(), src.data(), 0, 0, src.GetValidRow(), src.GetValidCol(), dst.GetValidRow(),
816- DstTileData::SFractal == SLayout::RowMajor) {816+ dst.GetValidCol());
817- PTO_ASSERT(
818- src.GetValidRow() == dst.GetValidRow() && src.GetValidCol() == dst.GetValidCol(),
819- "TMOV ND-to-NZ requires matching valid shapes.");
820- TCopyNdToNzUbToMat<DstTileData, SrcTileData>(
821- dst.data(), src.data(), src.GetValidRow(), src.GetValidCol());
822- } else {
823- TExtractVecToMat<DstTileData, SrcTileData>(
824- dst.data(), src.data(), 0, 0, src.GetValidRow(), src.GetValidCol(), dst.GetValidRow(),
825- dst.GetValidCol());
826- }
827 }817 }
828 }818 }
829}819}
Mtests/costmodel/st/testcase/textract/main.cpp+18-44
@@ -8,71 +8,45 @@ INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A
8See LICENSE in the root of the software repository for the full text of the License.8See LICENSE in the root of the software repository for the full text of the License.
9*/9*/
10 10 
11-#include <cstddef>
12-#include <cstdint>
13- 
14-#include <gtest/gtest.h>
15#include <pto/pto-inst.hpp>11#include <pto/pto-inst.hpp>
16#include <pto/common/constants.hpp>12#include <pto/common/constants.hpp>
13+#include <gtest/gtest.h>
14+ 
15+#include "cost_check.hpp"
17 16 
18using namespace pto;17using namespace pto;
19 18 
20namespace {19namespace {
21 20 
22template <21template <
23- typename T, size_t Rows, size_t Cols, size_t ValidRows, size_t ValidCols, uint16_t IndexRow, uint16_t IndexCol,22+ typename ST, typename DT, size_t rows, size_t cols, size_t validRows, size_t validCols, uint16_t idxRow,
24- BLayout Layout = BLayout::RowMajor, SLayout Fractal = SLayout::NoneBox>23+ uint16_t idxCol, float profiling, float accuracy>
25-void runTExtract(int expectedDmaCopies = 1, int expectedVectorCopies = 0)24+void runTExtract()
26{25{
27- constexpr int DST_VALID_ROWS = ValidRows - IndexRow;26+ constexpr int validRowsDst = validRows - idxRow;
28- constexpr int DST_VALID_COLS = ValidCols - IndexCol;27+ constexpr int validColsDst = validCols - idxCol;
29 28 
30- Tile<TileType::Vec, T, Rows, Cols, Layout, ValidRows, ValidCols, Fractal> srcTile;29+ Tile<TileType::Mat, ST, rows, cols, BLayout::RowMajor, validRows, validCols, SLayout::NoneBox, 512> srcTile;
31- Tile<TileType::Vec, T, Rows, Cols, Layout, DST_VALID_ROWS, DST_VALID_COLS, Fractal> dstTile;30+ Tile<TileType::Mat, DT, rows, cols, BLayout::RowMajor, validRowsDst, validColsDst, SLayout::NoneBox, 512> dstTile;
32 31 
33 TASSIGN(srcTile, 0x0);32 TASSIGN(srcTile, 0x0);
34 TASSIGN(dstTile, 0x10000);33 TASSIGN(dstTile, 0x10000);
35 34 
36- mocker::ResetTrace();35+ TEXTRACT(dstTile, srcTile, idxRow, idxCol);
37- TEXTRACT(dstTile, srcTile, IndexRow, IndexCol);
38 36 
39- const auto& instructions = mocker::GetTrace().executed_pto;37+ EXPECT_CYCLE_NEAR(profiling, accuracy);
40- ASSERT_EQ(instructions.size(), 1);
41- EXPECT_EQ(instructions[0].name, "TEXTRACT");
42- EXPECT_GT(instructions[0].total_cycles, 0);
43- int dmaCopies = 0;
44- int vectorCopies = 0;
45- for (const auto& call : instructions[0].cce_calls) {
46- if (call.name == "copy_ubuf_to_ubuf") {
47- ++dmaCopies;
48- } else if (call.name == "vcopy") {
49- ++vectorCopies;
50- }
51- }
52- EXPECT_EQ(dmaCopies, expectedDmaCopies);
53- EXPECT_EQ(vectorCopies, expectedVectorCopies);
54}38}
55 39 
56} // namespace40} // namespace
57 41 
58-TEST(TExtract, half_32x32_idx_0_0) { runTExtract<half, 32, 32, 32, 32, 0, 0>(); }42+TEST(TExtract, half_32x32_idx_0_0) { runTExtract<half, half, 32, 32, 32, 32, 0, 0, 0.0f, 0.0f>(); }
59 43 
60-TEST(TExtract, float_128x96_idx_0_0) { runTExtract<float, 128, 96, 128, 96, 0, 0>(); }44+TEST(TExtract, float_128x96_idx_0_0) { runTExtract<float, float, 128, 96, 128, 96, 0, 0, 0.0f, 0.0f>(); }
61 45 
62-TEST(TExtract, half_32x32_idx_8_16) { runTExtract<half, 32, 32, 32, 32, 8, 16>(); }46+TEST(TExtract, half_32x32_idx_8_16) { runTExtract<half, half, 32, 32, 32, 32, 8, 16, 0.0f, 0.0f>(); }
63 47 
64-TEST(TExtract, float_128x96_idx_8_16) { runTExtract<float, 128, 96, 128, 96, 8, 16>(); }48+TEST(TExtract, float_128x96_idx_8_16) { runTExtract<float, float, 128, 96, 128, 96, 8, 16, 0.0f, 0.0f>(); }
65 49 
66-TEST(TExtract, half_32x32_valid31_idx_8_16) { runTExtract<half, 32, 32, 31, 31, 8, 16>(0, 1); }50+TEST(TExtract, half_32x32_valid31_idx_8_16) { runTExtract<half, half, 32, 32, 31, 31, 8, 16, 0.0f, 0.0f>(); }
67 51 
68-TEST(TExtract, float_128x96_valid125_idx_8_16) { runTExtract<float, 128, 96, 125, 93, 8, 16>(1, 1); }52+TEST(TExtract, float_128x96_valid125_idx_8_16) { runTExtract<float, float, 128, 96, 125, 93, 8, 16, 0.0f, 0.0f>(); }
69- 
70-TEST(TExtract, half_nz_32x32_idx_0_0)
71-{
72- runTExtract<half, 32, 32, 32, 32, 0, 0, BLayout::ColMajor, SLayout::RowMajor>();
73-}
74- 
75-TEST(TExtract, float_nz_128x96_idx_16_16)
76-{
77- runTExtract<float, 128, 96, 128, 96, 16, 16, BLayout::ColMajor, SLayout::RowMajor>();
78-}
Mtests/cpu/st/testcase/tmov/main.cpp+0-55
@@ -207,58 +207,3 @@ TMOV_ND_TO_ZN_TEST(float, 32, 32)
207TMOV_ND_TO_ZN_TEST(float, 32, 64)207TMOV_ND_TO_ZN_TEST(float, 32, 64)
208TMOV_ND_TO_ZN_TEST(float, 64, 64)208TMOV_ND_TO_ZN_TEST(float, 64, 64)
209TMOV_ND_TO_ZN_TEST(float, 128, 128)209TMOV_ND_TO_ZN_TEST(float, 128, 128)
210- 
211-enum class Nd2NzOperation { MOV, EXTRACT, INSERT };
212- 
213-template <Nd2NzOperation Operation>
214-void testUbToL1Nd2Nz()
215-{
216- constexpr int SRC_ROWS = 48;
217- constexpr int SRC_COLS = 64;
218- constexpr int DST_ROWS = 32;
219- constexpr int DST_COLS = 48;
220- constexpr int VALID_ROWS = 17;
221- constexpr int VALID_COLS = 24;
222- constexpr int INDEX_ROW = 3;
223- constexpr int INDEX_COL = 8;
224- constexpr int C0_SIZE = 8;
225- using SrcTile = Tile<TileType::Vec, float, SRC_ROWS, SRC_COLS, BLayout::RowMajor, DYNAMIC, DYNAMIC>;
226- using DstTile =
227- Tile<TileType::Mat, float, DST_ROWS, DST_COLS, BLayout::ColMajor, DYNAMIC, DYNAMIC, SLayout::RowMajor>;
228- SrcTile src(
229- Operation == Nd2NzOperation::EXTRACT ? SRC_ROWS : VALID_ROWS,
230- Operation == Nd2NzOperation::EXTRACT ? SRC_COLS : VALID_COLS);
231- DstTile dst(
232- Operation == Nd2NzOperation::INSERT ? DST_ROWS : VALID_ROWS,
233- Operation == Nd2NzOperation::INSERT ? DST_COLS : VALID_COLS);
234- TASSIGN(src, 0);
235- TASSIGN(dst, 0);
236- for (int i = 0; i < SRC_ROWS * SRC_COLS; ++i) {
237- src.data()[i] = static_cast<float>(i);
238- }
239- std::fill(dst.data(), dst.data() + DST_ROWS * DST_COLS, -1.0f);
240- std::vector<float> expected(DST_ROWS * DST_COLS, -1.0f);
241- for (int row = 0; row < VALID_ROWS; ++row) {
242- for (int col = 0; col < VALID_COLS; ++col) {
243- const int srcRow = row + (Operation == Nd2NzOperation::EXTRACT ? INDEX_ROW : 0);
244- const int srcCol = col + (Operation == Nd2NzOperation::EXTRACT ? INDEX_COL : 0);
245- const int dstRow = row + (Operation == Nd2NzOperation::INSERT ? INDEX_ROW : 0);
246- const int dstCol = col + (Operation == Nd2NzOperation::INSERT ? INDEX_COL : 0);
247- const int offset =
248- ((dstCol / C0_SIZE * (DST_ROWS / 16) + dstRow / 16) * 16 + dstRow % 16) * C0_SIZE + dstCol % C0_SIZE;
249- expected[offset] = static_cast<float>(srcRow * SRC_COLS + srcCol);
250- }
251- }
252- if constexpr (Operation == Nd2NzOperation::EXTRACT) {
253- TEXTRACT(dst, src, INDEX_ROW, INDEX_COL);
254- } else if constexpr (Operation == Nd2NzOperation::INSERT) {
255- TINSERT(dst, src, INDEX_ROW, INDEX_COL);
256- } else {
257- TMOV(dst, src);
258- }
259- EXPECT_TRUE(std::equal(expected.begin(), expected.end(), dst.data()));
260-}
261- 
262-TEST_F(TMOVTest, nd2nz_ub_to_l1_strides) { testUbToL1Nd2Nz<Nd2NzOperation::MOV>(); }
263-TEST_F(TMOVTest, nd2nz_ub_to_l1_extract) { testUbToL1Nd2Nz<Nd2NzOperation::EXTRACT>(); }
264-TEST_F(TMOVTest, nd2nz_ub_to_l1_insert) { testUbToL1Nd2Nz<Nd2NzOperation::INSERT>(); }
Mtests/npu/a5/src/st/testcase/textract/CMakeLists.txt+0-10
@@ -19,13 +19,3 @@ target_include_directories(textract_acc2mat_kernel PRIVATE
19 ${ASCEND_HOME_PATH}/pkg_inc/runtime/runtime)19 ${ASCEND_HOME_PATH}/pkg_inc/runtime/runtime)
20target_link_options(textract_acc2mat_kernel PRIVATE --cce-fatobj-link)20target_link_options(textract_acc2mat_kernel PRIVATE --cce-fatobj-link)
21target_link_libraries(textract PRIVATE textract_acc2mat_kernel)21target_link_libraries(textract PRIVATE textract_acc2mat_kernel)
22- 
23-add_library(textract_ub2l1_kernel SHARED textract_ub2l1_kernel.cpp)
24-target_compile_options(textract_ub2l1_kernel PRIVATE
25- ${CMAKE_CCE_COMPILE_OPTIONS} --cce-aicore-arch=dav-c310 -DREGISTER_BASE -std=c++17)
26-target_include_directories(textract_ub2l1_kernel PRIVATE
27- ${ASCEND_HOME_PATH}/pkg_inc/
28- ${ASCEND_HOME_PATH}/pkg_inc/profiling/
29- ${ASCEND_HOME_PATH}/pkg_inc/runtime/runtime)
30-target_link_options(textract_ub2l1_kernel PRIVATE --cce-fatobj-link)
31-target_link_libraries(textract PRIVATE textract_ub2l1_kernel)
Mtests/npu/a5/src/st/testcase/textract/main.cpp+2-20
@@ -8,12 +8,9 @@ INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A
8See LICENSE in the root of the software repository for the full text of the License.8See LICENSE in the root of the software repository for the full text of the License.
9*/9*/
10 10 
11-#include <gtest/gtest.h>
12-#include "acl/acl.h"
13- 
14-#include "../tmov_ub2l1/ub2l1_nd2nz_test.h"
15-#include "nd2nz_cases.h"
16#include "test_common.h"11#include "test_common.h"
12+#include "acl/acl.h"
13+#include <gtest/gtest.h>
17 14 
18using namespace std;15using namespace std;
19using namespace PtoTestCommon;16using namespace PtoTestCommon;
@@ -428,18 +425,3 @@ TEST_F(TMOVTest, case15) { tmov_mx_test<15, float, int8_t, int8_t>(128, 64, 128)
428TEST_F(TMOVTest, case16) { tmov_mx_test<16, float, int8_t, int8_t>(128, 64, 128); }425TEST_F(TMOVTest, case16) { tmov_mx_test<16, float, int8_t, int8_t>(128, 64, 128); }
429 426 
430TEST_F(TMOVTest, case17) { tmov_mx_test<17, float, int8_t, int8_t>(128, 64, 128); }427TEST_F(TMOVTest, case17) { tmov_mx_test<17, float, int8_t, int8_t>(128, 64, 128); }
431- 
432-template <int32_t TestKey>
433-void launchTExtractNd2Nz(uint64_t* out, uint64_t* src, void* stream);
434- 
435-#define TEST_ND2NZ( \
436- TestKey, Name, T, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, Operation, IndexRow, \
437- IndexCol, Dynamic, SrcValidRows, SrcValidCols) \
438- TEST_F(TEXTRACTTest, Name) \
439- { \
440- PtoTestCommon::testUbToL1Nd2Nz( \
441- launchTExtractNd2Nz<TestKey>, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, \
442- Operation, IndexRow, IndexCol); \
443- }
444-TEXTRACT_UB2L1_ND2NZ_CASES(TEST_ND2NZ)
445-#undef TEST_ND2NZ
Dtests/npu/a5/src/st/testcase/textract/nd2nz_cases.h+0-39
@@ -1,39 +0,0 @@
1-/**
2-Copyright (c) 2026 Huawei Technologies Co., Ltd.
3-This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-CANN Open Software License Agreement Version 2.0 (the "License").
5-Please refer to the License for details. You may not use this file except in compliance with the License.
6-THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-See LICENSE in the root of the software repository for the full text of the License.
9-*/
10- 
11-#ifndef TEXTRACT_UB2L1_ND2NZ_CASES_H_
12-#define TEXTRACT_UB2L1_ND2NZ_CASES_H_
13- 
14-#include "../tmov_ub2l1/ub2l1_nd2nz_types.h"
15- 
16-#define TEXTRACT_UB2L1_ND2NZ_CASES(X) \
17- X(21, nd2nz_extract_offset, half, 16, 48, 96, 32, 64, 17, 48, Nd2NzOperation::EXTRACT, 3, 16, true, 48, 96) \
18- X(30, nd2nz_extract_bfloat16, bfloat16_t, 16, 48, 96, 32, 64, 17, 48, Nd2NzOperation::EXTRACT, 3, 16, true, 24, \
19- 80) \
20- X(31, nd2nz_extract_float, float, 32, 48, 48, 32, 32, 17, 24, Nd2NzOperation::EXTRACT, 3, 8, true, 24, 40) \
21- X(32, nd2nz_extract_int8, int8_t, 8, 48, 192, 32, 128, 17, 96, Nd2NzOperation::EXTRACT, 3, 32, true, 24, 160) \
22- X(33, nd2nz_extract_hifloat8, hifloat8_t, 8, 48, 192, 32, 128, 17, 96, Nd2NzOperation::EXTRACT, 3, 32, true, 24, \
23- 160) \
24- X(34, nd2nz_extract_fp8_e4m3, float8_e4m3_t, 8, 48, 192, 32, 128, 17, 96, Nd2NzOperation::EXTRACT, 3, 32, true, \
25- 24, 160) \
26- X(35, nd2nz_extract_fp8_e5m2, float8_e5m2_t, 8, 48, 192, 32, 128, 17, 96, Nd2NzOperation::EXTRACT, 3, 32, true, \
27- 24, 160) \
28- X(36, nd2nz_extract_fp8_e8m0, float8_e8m0_t, 8, 48, 192, 32, 128, 17, 96, Nd2NzOperation::EXTRACT, 3, 32, true, \
29- 24, 160) \
30- X(37, nd2nz_extract_fp4_e2m1, float4_e2m1x2_t, 4, 48, 384, 32, 256, 17, 192, Nd2NzOperation::EXTRACT, 3, 64, true, \
31- 24, 320) \
32- X(38, nd2nz_extract_fp4_e1m2, float4_e1m2x2_t, 4, 48, 384, 32, 256, 17, 192, Nd2NzOperation::EXTRACT, 3, 64, true, \
33- 24, 320) \
34- X(40, nd2nz_extract_static, half, 16, 48, 96, 32, 64, 17, 48, Nd2NzOperation::EXTRACT, 3, 16, false, 20, 64) \
35- X(41, nd2nz_extract_empty_rows, half, 16, 16, 64, 16, 32, 0, 16, Nd2NzOperation::EXTRACT, 3, 16, true, 8, 48) \
36- X(42, nd2nz_extract_empty_cols, half, 16, 16, 64, 16, 32, 7, 0, Nd2NzOperation::EXTRACT, 3, 16, true, 10, 48) \
37- X(43, nd2nz_extract_valid_edge, float, 32, 48, 64, 32, 32, 17, 24, Nd2NzOperation::EXTRACT, 3, 8, true, 20, 32)
38- 
39-#endif
Dtests/npu/a5/src/st/testcase/textract/textract_ub2l1_kernel.cpp+0-32
@@ -1,32 +0,0 @@
1-/**
2-Copyright (c) 2026 Huawei Technologies Co., Ltd.
3-This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-CANN Open Software License Agreement Version 2.0 (the "License").
5-Please refer to the License for details. You may not use this file except in compliance with the License.
6-THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-See LICENSE in the root of the software repository for the full text of the License.
9-*/
10- 
11-#include "../tmov_ub2l1/ub2l1_nd2nz_kernel.h"
12-#include "nd2nz_cases.h"
13- 
14-using namespace pto;
15- 
16-template <int32_t TestKey>
17-void launchTExtractNd2Nz(uint64_t* out, uint64_t* src, void* stream);
18- 
19-#define LAUNCH_ND2NZ( \
20- TestKey, Name, T, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, Operation, IndexRow, \
21- IndexCol, Dynamic, SrcValidRows, SrcValidCols) \
22- template <> \
23- void launchTExtractNd2Nz<TestKey>(uint64_t * out, uint64_t * src, void* stream) \
24- { \
25- PtoTestCommon::runUbToL1Nd2Nz< \
26- T, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, Operation, IndexRow, IndexCol, \
27- Dynamic, SrcValidRows, SrcValidCols><<<1, nullptr, stream>>>( \
28- reinterpret_cast<uint8_t*>(out), reinterpret_cast<uint8_t*>(src), ValidRows, ValidCols, SrcValidRows, \
29- SrcValidCols); \
30- }
31-TEXTRACT_UB2L1_ND2NZ_CASES(LAUNCH_ND2NZ)
32-#undef LAUNCH_ND2NZ
Mtests/npu/a5/src/st/testcase/tinsert/main.cpp+2-20
@@ -8,12 +8,9 @@ INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A
8See LICENSE in the root of the software repository for the full text of the License.8See LICENSE in the root of the software repository for the full text of the License.
9*/9*/
10 10 
11-#include <gtest/gtest.h>
12-#include "acl/acl.h"
13- 
14-#include "../tmov_ub2l1/ub2l1_nd2nz_test.h"
15-#include "nd2nz_cases.h"
16#include "test_common.h"11#include "test_common.h"
12+#include "acl/acl.h"
13+#include <gtest/gtest.h>
17 14 
18using namespace std;15using namespace std;
19using namespace PtoTestCommon;16using namespace PtoTestCommon;
@@ -541,18 +538,3 @@ TEST_F(TInsertTest, case_compact_rowplusone_tmov_bf16_idx0_16)
541{538{
542 testTInsertCompactTMov<4, uint16_t, 64, 48, 128, 128>(launchTInsertCompactRowPlusOneTMov<4>);539 testTInsertCompactTMov<4, uint16_t, 64, 48, 128, 128>(launchTInsertCompactRowPlusOneTMov<4>);
543}540}
544- 
545-template <int32_t TestKey>
546-void launchTInsertNd2Nz(uint64_t* out, uint64_t* src, void* stream);
547- 
548-#define TEST_ND2NZ( \
549- TestKey, Name, T, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, Operation, IndexRow, \
550- IndexCol, Dynamic, SrcValidRows, SrcValidCols) \
551- TEST_F(TInsertTest, Name) \
552- { \
553- PtoTestCommon::testUbToL1Nd2Nz( \
554- launchTInsertNd2Nz<TestKey>, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, \
555- Operation, IndexRow, IndexCol); \
556- }
557-TINSERT_UB2L1_ND2NZ_CASES(TEST_ND2NZ)
558-#undef TEST_ND2NZ
Dtests/npu/a5/src/st/testcase/tinsert/nd2nz_cases.h+0-44
@@ -1,44 +0,0 @@
1-/**
2-Copyright (c) 2026 Huawei Technologies Co., Ltd.
3-This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-CANN Open Software License Agreement Version 2.0 (the "License").
5-Please refer to the License for details. You may not use this file except in compliance with the License.
6-THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-See LICENSE in the root of the software repository for the full text of the License.
9-*/
10- 
11-#ifndef TINSERT_UB2L1_ND2NZ_CASES_H_
12-#define TINSERT_UB2L1_ND2NZ_CASES_H_
13- 
14-#include "../tmov_ub2l1/ub2l1_nd2nz_types.h"
15- 
16-#define TINSERT_UB2L1_ND2NZ_CASES(X) \
17- X(22, nd2nz_insert_offset, half, 16, 16, 64, 48, 96, 7, 48, Nd2NzOperation::INSERT, 5, 32, true, 7, 48) \
18- X(23, nd2nz_insert_wide, float, 32, 16, 128, 32, 160, 3, 128, Nd2NzOperation::INSERT, 17, 8, true, 3, 128) \
19- X(26, nd2nz_dual_insert, half, 16, 16, 64, 32, 64, 16, 64, Nd2NzOperation::DUAL_INSERT, 0, 0, true, 16, 64) \
20- X(30, nd2nz_insert_bfloat16, bfloat16_t, 16, 16, 64, 48, 96, 7, 48, Nd2NzOperation::INSERT, 5, 32, true, 7, 48) \
21- X(31, nd2nz_insert_int8, int8_t, 8, 16, 128, 48, 192, 7, 96, Nd2NzOperation::INSERT, 5, 64, true, 7, 96) \
22- X(32, nd2nz_insert_hifloat8, hifloat8_t, 8, 16, 128, 48, 192, 7, 96, Nd2NzOperation::INSERT, 5, 64, true, 7, 96) \
23- X(33, nd2nz_insert_fp8_e4m3, float8_e4m3_t, 8, 16, 128, 48, 192, 7, 96, Nd2NzOperation::INSERT, 5, 64, true, 7, \
24- 96) \
25- X(34, nd2nz_insert_fp8_e5m2, float8_e5m2_t, 8, 16, 128, 48, 192, 7, 96, Nd2NzOperation::INSERT, 5, 64, true, 7, \
26- 96) \
27- X(35, nd2nz_insert_fp8_e8m0, float8_e8m0_t, 8, 16, 128, 48, 192, 7, 96, Nd2NzOperation::INSERT, 5, 64, true, 7, \
28- 96) \
29- X(36, nd2nz_insert_fp4_e2m1, float4_e2m1x2_t, 4, 16, 256, 48, 384, 7, 192, Nd2NzOperation::INSERT, 5, 128, true, \
30- 7, 192) \
31- X(37, nd2nz_insert_fp4_e1m2, float4_e1m2x2_t, 4, 16, 256, 48, 384, 7, 192, Nd2NzOperation::INSERT, 5, 128, true, \
32- 7, 192) \
33- X(38, nd2nz_insert_int32, int32_t, 32, 16, 32, 48, 48, 7, 24, Nd2NzOperation::INSERT, 5, 16, true, 7, 24) \
34- X(40, nd2nz_insert_static, half, 16, 16, 64, 48, 96, 7, 48, Nd2NzOperation::INSERT, 5, 32, false, 7, 48) \
35- X(41, nd2nz_insert_empty_rows, half, 16, 16, 64, 32, 64, 0, 32, Nd2NzOperation::INSERT, 5, 16, true, 0, 32) \
36- X(42, nd2nz_insert_empty_cols, float, 32, 16, 32, 32, 32, 7, 0, Nd2NzOperation::INSERT, 3, 8, true, 7, 0) \
37- X(43, nd2nz_dual_insert_fp4, float4_e2m1x2_t, 4, 16, 128, 32, 128, 16, 128, Nd2NzOperation::DUAL_INSERT, 0, 0, \
38- true, 16, 128) \
39- X(44, nd2nz_insert_fp4_65536_static, float4_e2m1x2_t, 4, 1, 65536, 16, 65536, 1, 65536, Nd2NzOperation::INSERT, 5, \
40- 0, false, 1, 65536) \
41- X(45, nd2nz_insert_fp4_65536_dynamic, float4_e1m2x2_t, 4, 1, 65600, 16, 65536, 1, 65536, Nd2NzOperation::INSERT, \
42- 5, 0, true, 1, 65536)
43- 
44-#endif
Mtests/npu/a5/src/st/testcase/tinsert/tinsert_kernel.cpp+0-21
@@ -13,9 +13,6 @@ See LICENSE in the root of the software repository for the full text of the Lice
13#include <pto/common/constants.hpp>13#include <pto/common/constants.hpp>
14#include <pto/npu/a5/TInsert.hpp>14#include <pto/npu/a5/TInsert.hpp>
15 15 
16-#include "../tmov_ub2l1/ub2l1_nd2nz_kernel.h"
17-#include "nd2nz_cases.h"
18- 
19using namespace pto;16using namespace pto;
20 17 
21template <int M, int N, typename OutType>18template <int M, int N, typename OutType>
@@ -2376,21 +2373,3 @@ template void launchTInsertNDVec<10>(uint8_t* out, uint8_t* srcIn, uint8_t* dstI
2376template void launchTInsertNDVec<11>(uint8_t* out, uint8_t* srcIn, uint8_t* dstIn, void* stream);2373template void launchTInsertNDVec<11>(uint8_t* out, uint8_t* srcIn, uint8_t* dstIn, void* stream);
2377template void launchTInsertND<1>(uint64_t* out, uint64_t* src, void* stream);2374template void launchTInsertND<1>(uint64_t* out, uint64_t* src, void* stream);
2378template void launchTInsertND<2>(uint64_t* out, uint64_t* src, void* stream);2375template void launchTInsertND<2>(uint64_t* out, uint64_t* src, void* stream);
2379- 
2380-template <int32_t TestKey>
2381-void launchTInsertNd2Nz(uint64_t* out, uint64_t* src, void* stream);
2382- 
2383-#define LAUNCH_ND2NZ( \
2384- TestKey, Name, T, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, Operation, IndexRow, \
2385- IndexCol, Dynamic, SrcValidRows, SrcValidCols) \
2386- template <> \
2387- void launchTInsertNd2Nz<TestKey>(uint64_t * out, uint64_t * src, void* stream) \
2388- { \
2389- PtoTestCommon::runUbToL1Nd2Nz< \
2390- T, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, Operation, IndexRow, IndexCol, \
2391- Dynamic, SrcValidRows, SrcValidCols><<<1, nullptr, stream>>>( \
2392- reinterpret_cast<uint8_t*>(out), reinterpret_cast<uint8_t*>(src), ValidRows, ValidCols, SrcValidRows, \
2393- SrcValidCols); \
2394- }
2395-TINSERT_UB2L1_ND2NZ_CASES(LAUNCH_ND2NZ)
2396-#undef LAUNCH_ND2NZ
Mtests/npu/a5/src/st/testcase/tmov_ub2l1/CMakeLists.txt+1-7
@@ -8,10 +8,4 @@
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# --------------------------------------------------------------------------------9# --------------------------------------------------------------------------------
10 10 
11-pto_mix_st(tmov_ub2l1)11+pto_mix_st(tmov_ub2l1)
12- 
13-if(AUTO_MODE)
14- # The ND-to-NZ tests alias raw and typed Tiles using manual TASSIGN addresses.
15- target_compile_definitions(tmov_ub2l1 PRIVATE PTO_SKIP_UB2L1_ND2NZ_ST)
16- target_compile_definitions(tmov_ub2l1_kernel_mix PRIVATE PTO_SKIP_UB2L1_ND2NZ_ST)
17-endif()
Mtests/npu/a5/src/st/testcase/tmov_ub2l1/main.cpp+4-24
@@ -8,15 +8,9 @@ INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A
8See LICENSE in the root of the software repository for the full text of the License.8See LICENSE in the root of the software repository for the full text of the License.
9*/9*/
10 10 
11-#include <cstdint>
12-#include <cstdlib>
13- 
14-#include <gtest/gtest.h>
15-#include "acl/acl.h"
16- 
17-#include "nd2nz_cases.h"
18#include "test_common.h"11#include "test_common.h"
19-#include "ub2l1_nd2nz_test.h"12+#include "acl/acl.h"
13+#include <gtest/gtest.h>
20 14 
21using namespace std;15using namespace std;
22using namespace PtoTestCommon;16using namespace PtoTestCommon;
@@ -43,9 +37,7 @@ template <int32_t testKey, typename dType>
43void testTMovUb2L1(int32_t srcRows, int32_t srcCols, int32_t dstRows, int32_t dstCols)37void testTMovUb2L1(int32_t srcRows, int32_t srcCols, int32_t dstRows, int32_t dstCols)
44{38{
45 aclInit(nullptr);39 aclInit(nullptr);
46- const char* deviceEnv = std::getenv("PTO_DEVICE_ID");40+ aclrtSetDevice(0);
47- const int deviceId = deviceEnv ? std::atoi(deviceEnv) : 0;
48- aclrtSetDevice(deviceId);
49 aclrtStream stream;41 aclrtStream stream;
50 aclrtCreateStream(&stream);42 aclrtCreateStream(&stream);
51 43 
@@ -73,7 +65,7 @@ void testTMovUb2L1(int32_t srcRows, int32_t srcCols, int32_t dstRows, int32_t ds
73 aclrtFreeHost(dstHost);65 aclrtFreeHost(dstHost);
74 aclrtFreeHost(srcHost);66 aclrtFreeHost(srcHost);
75 aclrtDestroyStream(stream);67 aclrtDestroyStream(stream);
76- aclrtResetDevice(deviceId);68+ aclrtResetDevice(0);
77 aclFinalize();69 aclFinalize();
78 70 
79 std::vector<dType> golden(dstByteSize / sizeof(dType));71 std::vector<dType> golden(dstByteSize / sizeof(dType));
@@ -101,15 +93,3 @@ TEST_F(TMovUb2l1Test, case7) { testTMovUb2L1<7, uint16_t>(64, 64, 48, 48); }
101TEST_F(TMovUb2l1Test, case8) { testTMovUb2L1<8, float>(128, 128, 64, 64); }93TEST_F(TMovUb2l1Test, case8) { testTMovUb2L1<8, float>(128, 128, 64, 64); }
102 94 
103TEST_F(TMovUb2l1Test, case9) { testTMovUb2L1<9, int8_t>(256, 256, 32, 32); }95TEST_F(TMovUb2l1Test, case9) { testTMovUb2L1<9, int8_t>(256, 256, 32, 32); }
104- 
105-#define TEST_ND2NZ( \
106- TestKey, Name, T, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, Operation, IndexRow, \
107- IndexCol, Dynamic, SrcValidRows, SrcValidCols) \
108- TEST_F(TMovUb2l1Test, Name) \
109- { \
110- PtoTestCommon::testUbToL1Nd2Nz( \
111- launchTmovUb2l1<TestKey>, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, \
112- Operation, IndexRow, IndexCol); \
113- }
114-TMOV_UB2L1_ND2NZ_CASES(TEST_ND2NZ)
115-#undef TEST_ND2NZ
Dtests/npu/a5/src/st/testcase/tmov_ub2l1/nd2nz_cases.h+0-36
@@ -1,36 +0,0 @@
1-/**
2-Copyright (c) 2026 Huawei Technologies Co., Ltd.
3-This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-CANN Open Software License Agreement Version 2.0 (the "License").
5-Please refer to the License for details. You may not use this file except in compliance with the License.
6-THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-See LICENSE in the root of the software repository for the full text of the License.
9-*/
10- 
11-#ifndef TMOV_UB2L1_ND2NZ_CASES_H_
12-#define TMOV_UB2L1_ND2NZ_CASES_H_
13- 
14-#include "ub2l1_nd2nz_types.h"
15- 
16-#ifdef PTO_SKIP_UB2L1_ND2NZ_ST
17-#define TMOV_UB2L1_ND2NZ_CASES(X)
18-#else
19-#define TMOV_UB2L1_ND2NZ_CASES(X) \
20- X(10, nd2nz_half_static, half, 16, 16, 32, 16, 32, 16, 32, Nd2NzOperation::MOV, 0, 0, false, 16, 32) \
21- X(11, nd2nz_half_wide, half, 16, 16, 512, 16, 512, 3, 512, Nd2NzOperation::MOV, 0, 0, true, 3, 512) \
22- X(12, nd2nz_float_strides, float, 32, 32, 64, 32, 32, 19, 24, Nd2NzOperation::MOV, 0, 0, true, 19, 24) \
23- X(13, nd2nz_bfloat16_strides, bfloat16_t, 16, 32, 96, 32, 64, 17, 48, Nd2NzOperation::MOV, 0, 0, true, 17, 48) \
24- X(14, nd2nz_int8_strides, int8_t, 8, 32, 160, 32, 128, 31, 96, Nd2NzOperation::MOV, 0, 0, true, 31, 96) \
25- X(15, nd2nz_fp8_e4m3, float8_e4m3_t, 8, 16, 64, 16, 64, 16, 64, Nd2NzOperation::MOV, 0, 0, false, 16, 64) \
26- X(16, nd2nz_fp8_e5m2, float8_e5m2_t, 8, 16, 64, 16, 64, 16, 64, Nd2NzOperation::MOV, 0, 0, false, 16, 64) \
27- X(17, nd2nz_hifloat8, hifloat8_t, 8, 16, 64, 16, 64, 16, 64, Nd2NzOperation::MOV, 0, 0, false, 16, 64) \
28- X(18, nd2nz_fp8_e8m0, float8_e8m0_t, 8, 16, 64, 16, 64, 16, 64, Nd2NzOperation::MOV, 0, 0, false, 16, 64) \
29- X(19, nd2nz_fp4_e2m1, float4_e2m1x2_t, 4, 32, 256, 32, 192, 17, 128, Nd2NzOperation::MOV, 0, 0, true, 17, 128) \
30- X(20, nd2nz_fp4_e1m2, float4_e1m2x2_t, 4, 16, 128, 16, 128, 16, 128, Nd2NzOperation::MOV, 0, 0, false, 16, 128) \
31- X(24, nd2nz_burst_split, half, 16, 4112, 16, 4112, 16, 4101, 16, Nd2NzOperation::MOV, 0, 0, true, 4101, 16) \
32- X(25, nd2nz_empty, half, 16, 16, 32, 16, 32, 0, 32, Nd2NzOperation::MOV, 0, 0, true, 0, 32) \
33- X(27, nd2nz_empty_cols, float, 32, 16, 32, 16, 32, 7, 0, Nd2NzOperation::MOV, 0, 0, true, 7, 0)
34-#endif
35- 
36-#endif
Mtests/npu/a5/src/st/testcase/tmov_ub2l1/tmov_ub2l1_kernel.cpp+1-20
@@ -9,14 +9,10 @@ See LICENSE in the root of the software repository for the full text of the Lice
9*/9*/
10 10 
11#include <iostream>11#include <iostream>
12- 
13#include <pto/pto-inst.hpp>12#include <pto/pto-inst.hpp>
14#include <pto/common/pto_tile.hpp>13#include <pto/common/pto_tile.hpp>
15#include <pto/common/constants.hpp>14#include <pto/common/constants.hpp>
16 15 
17-#include "nd2nz_cases.h"
18-#include "ub2l1_nd2nz_kernel.h"
19- 
20using namespace std;16using namespace std;
21using namespace pto;17using namespace pto;
22 18 
@@ -25,7 +21,7 @@ __tf__ PTO_INTERNAL void tf_copy_cbuf_to_ubuf(
25 typename DstTileData::TileDType __out__ dst, typename SrcTileData::TileDType __in__ src, uint16_t vector,21 typename DstTileData::TileDType __out__ dst, typename SrcTileData::TileDType __in__ src, uint16_t vector,
26 uint16_t blockLen)22 uint16_t blockLen)
27{23{
28- __cbuf__ T* srcMatAddr = (__cbuf__ T*)__cce_get_tile_ptr(src);24+ __cbuf__ T* srcMatAddr = __cce_get_tile_ptr(src);
29 __ubuf__ T* dstUbAddr = __cce_get_tile_ptr(dst);25 __ubuf__ T* dstUbAddr = __cce_get_tile_ptr(dst);
30 copy_cbuf_to_ubuf(26 copy_cbuf_to_ubuf(
31 (__ubuf__ void*)dstUbAddr, (__cbuf__ void*)srcMatAddr, vector, 1, blockLen, 0,27 (__ubuf__ void*)dstUbAddr, (__cbuf__ void*)srcMatAddr, vector, 1, blockLen, 0,
@@ -162,18 +158,3 @@ template void launchTmovUb2l1<6>(uint64_t* out, uint64_t* src, void* stream);
162template void launchTmovUb2l1<7>(uint64_t* out, uint64_t* src, void* stream);158template void launchTmovUb2l1<7>(uint64_t* out, uint64_t* src, void* stream);
163template void launchTmovUb2l1<8>(uint64_t* out, uint64_t* src, void* stream);159template void launchTmovUb2l1<8>(uint64_t* out, uint64_t* src, void* stream);
164template void launchTmovUb2l1<9>(uint64_t* out, uint64_t* src, void* stream);160template void launchTmovUb2l1<9>(uint64_t* out, uint64_t* src, void* stream);
165- 
166-#define LAUNCH_ND2NZ( \
167- TestKey, Name, T, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, Operation, IndexRow, \
168- IndexCol, Dynamic, SrcValidRows, SrcValidCols) \
169- template <> \
170- void launchTmovUb2l1<TestKey>(uint64_t * out, uint64_t * src, void* stream) \
171- { \
172- PtoTestCommon::runUbToL1Nd2Nz< \
173- T, ElementBits, SrcRows, SrcCols, DstRows, DstCols, ValidRows, ValidCols, Operation, IndexRow, IndexCol, \
174- Dynamic, SrcValidRows, SrcValidCols><<<1, nullptr, stream>>>( \
175- reinterpret_cast<uint8_t*>(out), reinterpret_cast<uint8_t*>(src), ValidRows, ValidCols, SrcValidRows, \
176- SrcValidCols); \
177- }
178-TMOV_UB2L1_ND2NZ_CASES(LAUNCH_ND2NZ)
179-#undef LAUNCH_ND2NZ
Dtests/npu/a5/src/st/testcase/tmov_ub2l1/ub2l1_nd2nz_kernel.h+0-166
@@ -1,166 +0,0 @@
1-/**
2-Copyright (c) 2026 Huawei Technologies Co., Ltd.
3-This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-CANN Open Software License Agreement Version 2.0 (the "License").
5-Please refer to the License for details. You may not use this file except in compliance with the License.
6-THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-See LICENSE in the root of the software repository for the full text of the License.
9-*/
10- 
11-#ifndef UB2L1_ND2NZ_KERNEL_H_
12-#define UB2L1_ND2NZ_KERNEL_H_
13- 
14-#include <pto/pto-inst.hpp>
15-#include <pto/common/constants.hpp>
16-#include <pto/common/pto_tile.hpp>
17- 
18-#include "ub2l1_nd2nz_types.h"
19- 
20-namespace PtoTestCommon {
21- 
22-using namespace pto;
23- 
24-template <typename T, typename DstTileData, typename SrcTileData>
25-__tf__ PTO_INTERNAL void readbackNzMat(
26- typename DstTileData::TileDType __out__ dst, typename SrcTileData::TileDType __in__ src, uint16_t vectorId,
27- uint16_t blockCount, uint32_t srcByteOffset = 0)
28-{
29- __cbuf__ T* srcMatAddr = (__cbuf__ T*)((__cbuf__ uint8_t*)__cce_get_tile_ptr(src) + srcByteOffset);
30- __ubuf__ T* dstUbAddr = __cce_get_tile_ptr(dst);
31- copy_cbuf_to_ubuf((__ubuf__ void*)dstUbAddr, (__cbuf__ void*)srcMatAddr, vectorId, 1, blockCount, 0, 0);
32-}
33- 
34-template <typename MatTile, typename UbTile>
35-__tf__ PTO_INTERNAL void initNzMat(
36- typename MatTile::TileDType __out__ dst, typename UbTile::TileDType __in__ src, uint32_t byteCount,
37- uint32_t dstByteOffset)
38-{
39- copy_ubuf_to_cbuf(
40- (__cbuf__ uint8_t*)__cce_get_tile_ptr(dst) + dstByteOffset, (__ubuf__ void*)__cce_get_tile_ptr(src), 0, 1,
41- byteCount / BLOCK_BYTE_SIZE, 0, 0);
42-}
43- 
44-template <
45- typename T, int ElementBits, int SrcRows, int SrcCols, int DstRows, int DstCols, int ValidRows, int ValidCols,
46- Nd2NzOperation Operation, int IndexRow, int IndexCol, bool Dynamic, int SrcValidRows, int SrcValidCols>
47-__global__ AICORE void runUbToL1Nd2Nz(
48- __gm__ uint8_t* out, __gm__ uint8_t* input, uint32_t validRows, uint32_t validCols, uint32_t srcValidRows,
49- uint32_t srcValidCols)
50-{
51- constexpr uint32_t L1_READY = 0;
52- constexpr uint32_t UB_READY = 1;
53- constexpr uint32_t L1_INITIALIZED = 2;
54- constexpr uint32_t INIT_ACK = 3;
55- constexpr uint32_t AIV1_EVENT_OFFSET = 16;
56- constexpr uint32_t DST_BYTES = DstRows * DstCols * ElementBits / 8;
57- static_assert(SrcRows * SrcCols * ElementBits / 8 <= 256 * 1024, "ST source exceeds A5 UB capacity.");
58- static_assert(DST_BYTES <= 512 * 1024, "ST destination exceeds A5 L1 capacity.");
59- constexpr uint32_t MAX_READBACK_BYTES = 128 * 1024;
60- constexpr uint32_t READBACK_BYTES = DST_BYTES > MAX_READBACK_BYTES ? MAX_READBACK_BYTES : DST_BYTES;
61- constexpr uint32_t READBACK_COUNT = (DST_BYTES + READBACK_BYTES - 1) / READBACK_BYTES;
62- constexpr bool IS_INSERT = Operation == Nd2NzOperation::INSERT || Operation == Nd2NzOperation::DUAL_INSERT;
63- using SrcTile = Tile<
64- TileType::Vec, T, SrcRows, SrcCols, BLayout::RowMajor, Dynamic ? DYNAMIC : SrcValidRows,
65- Dynamic ? DYNAMIC : SrcValidCols>;
66- using DstTile = Tile<
67- TileType::Mat, T, DstRows, DstCols, BLayout::ColMajor, Dynamic ? DYNAMIC : (IS_INSERT ? DstRows : ValidRows),
68- Dynamic ? DYNAMIC : (IS_INSERT ? DstCols : ValidCols), SLayout::RowMajor>;
69- using RawSrc = Tile<TileType::Vec, uint8_t, SrcRows, SrcCols * ElementBits / 8>;
70- using RawDst = Tile<TileType::Vec, uint8_t, 1, READBACK_BYTES, BLayout::RowMajor, DYNAMIC, DYNAMIC>;
71- using SrcGlobal = GlobalTensor<
72- uint8_t, Shape<1, 1, 1, SrcRows, SrcCols * ElementBits / 8>,
73- pto::Stride<1, 1, 1, SrcCols * ElementBits / 8, 1>>;
74- using DstGlobal = GlobalTensor<uint8_t, Shape<1, 1, 1, 1, DYNAMIC>, pto::Stride<1, 1, 1, READBACK_BYTES, 1>>;
75- SrcTile src;
76- DstTile dst;
77- if constexpr (Dynamic) {
78- src.SetValidShape(srcValidRows, srcValidCols);
79- dst.SetValidShape(IS_INSERT ? DstRows : validRows, IS_INSERT ? DstCols : validCols);
80- }
81- RawSrc rawSrc;
82- RawDst rawDst(1, READBACK_BYTES);
83- TASSIGN(src, 0);
84- TASSIGN(dst, 0);
85- TASSIGN(rawSrc, 0);
86- TASSIGN(rawDst, 0);
87- SrcGlobal srcGlobal(
88- input +
89- (Operation == Nd2NzOperation::DUAL_INSERT ? get_subblockid() * SrcRows * SrcCols * ElementBits / 8 : 0));
90-#if defined(__DAV_VEC__)
91- // Initialize L1 once; the dual-AIV case writes disjoint row windows.
92- if (get_subblockid() == 0) {
93- for (uint32_t chunk = 0; chunk < READBACK_COUNT; ++chunk) {
94- const uint32_t offset = chunk * READBACK_BYTES;
95- const uint32_t bytes = min(READBACK_BYTES, DST_BYTES - offset);
96- rawDst.SetValidShape(1, bytes);
97- DstGlobal dstGlobal(out + offset, Shape<1, 1, 1, 1, DYNAMIC>(bytes));
98- TLOAD(rawDst, dstGlobal);
99- set_flag(PIPE_MTE2, PIPE_MTE3, EVENT_ID0);
100- wait_flag(PIPE_MTE2, PIPE_MTE3, EVENT_ID0);
101- initNzMat<DstTile, RawDst>(dst.data(), rawDst.data(), bytes, offset);
102- set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
103- wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
104- }
105- if constexpr (Operation == Nd2NzOperation::DUAL_INSERT) {
106- set_intra_block(PIPE_MTE3, L1_INITIALIZED);
107- }
108- }
109- if constexpr (Operation == Nd2NzOperation::DUAL_INSERT) {
110- wait_intra_block(PIPE_MTE2, INIT_ACK);
111- }
112- if (Operation == Nd2NzOperation::DUAL_INSERT || get_subblockid() == 0) {
113- TLOAD(rawSrc, srcGlobal);
114- set_flag(PIPE_MTE2, PIPE_MTE3, EVENT_ID0);
115- wait_flag(PIPE_MTE2, PIPE_MTE3, EVENT_ID0);
116- if constexpr (Operation == Nd2NzOperation::EXTRACT) {
117- TEXTRACT(dst, src, IndexRow, IndexCol);
118- } else if constexpr (IS_INSERT) {
119- TINSERT(
120- dst, src, IndexRow + (Operation == Nd2NzOperation::DUAL_INSERT ? get_subblockid() * validRows : 0),
121- IndexCol);
122- } else {
123- TMOV(dst, src);
124- }
125- set_intra_block(PIPE_MTE3, L1_READY);
126- if (get_subblockid() == 0) {
127- for (uint32_t chunk = 0; chunk < READBACK_COUNT; ++chunk) {
128- const uint32_t offset = chunk * READBACK_BYTES;
129- const uint32_t bytes = min(READBACK_BYTES, DST_BYTES - offset);
130- rawDst.SetValidShape(1, bytes);
131- DstGlobal dstGlobal(out + offset, Shape<1, 1, 1, 1, DYNAMIC>(bytes));
132- wait_intra_block(PIPE_MTE3, UB_READY);
133- TSTORE(dstGlobal, rawDst);
134- if (chunk + 1 < READBACK_COUNT) {
135- set_intra_block(PIPE_MTE3, L1_READY);
136- }
137- }
138- }
139- }
140-#endif
141-#if defined(__DAV_CUBE__)
142- if constexpr (Operation == Nd2NzOperation::DUAL_INSERT) {
143- wait_intra_block(PIPE_MTE1, L1_INITIALIZED);
144- set_intra_block(PIPE_MTE1, INIT_ACK);
145- set_intra_block(PIPE_MTE1, INIT_ACK + AIV1_EVENT_OFFSET);
146- }
147- wait_intra_block(PIPE_MTE1, L1_READY);
148- if constexpr (Operation == Nd2NzOperation::DUAL_INSERT) {
149- wait_intra_block(PIPE_MTE1, L1_READY + AIV1_EVENT_OFFSET);
150- }
151- // Wait for GM stores before reusing the bounded UB readback buffer.
152- for (uint32_t chunk = 0; chunk < READBACK_COUNT; ++chunk) {
153- if (chunk > 0) {
154- wait_intra_block(PIPE_MTE1, L1_READY);
155- }
156- const uint32_t offset = chunk * READBACK_BYTES;
157- const uint32_t bytes = min(READBACK_BYTES, DST_BYTES - offset);
158- readbackNzMat<uint8_t, RawDst, DstTile>(rawDst.data(), dst.data(), 0, bytes / BLOCK_BYTE_SIZE, offset);
159- set_intra_block(PIPE_MTE1, UB_READY);
160- }
161-#endif
162-}
163- 
164-} // namespace PtoTestCommon
165- 
166-#endif
Dtests/npu/a5/src/st/testcase/tmov_ub2l1/ub2l1_nd2nz_test.h+0-83
@@ -1,83 +0,0 @@
1-/**
2-Copyright (c) 2026 Huawei Technologies Co., Ltd.
3-This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-CANN Open Software License Agreement Version 2.0 (the "License").
5-Please refer to the License for details. You may not use this file except in compliance with the License.
6-THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-See LICENSE in the root of the software repository for the full text of the License.
9-*/
10- 
11-#ifndef UB2L1_ND2NZ_TEST_H_
12-#define UB2L1_ND2NZ_TEST_H_
13- 
14-#include <cstddef>
15-#include <cstdint>
16-#include <cstdlib>
17-#include <vector>
18- 
19-#include <gtest/gtest.h>
20-#include "acl/acl.h"
21- 
22-#include "ub2l1_nd2nz_types.h"
23- 
24-namespace PtoTestCommon {
25- 
26-inline void testUbToL1Nd2Nz(
27- void (*launch)(uint64_t*, uint64_t*, void*), int elementBits, int srcRows, int srcCols, int dstRows, int dstCols,
28- int validRows, int validCols, Nd2NzOperation operation, int indexRow, int indexCol)
29-{
30- const bool isInsert = operation == Nd2NzOperation::INSERT || operation == Nd2NzOperation::DUAL_INSERT;
31- // Compare bytes, including padding sentinels, so packed and floating formats retain every bit.
32- const size_t srcTileBytes = srcRows * srcCols * elementBits / 8;
33- const int partCount = operation == Nd2NzOperation::DUAL_INSERT ? 2 : 1;
34- const size_t srcBytes = partCount * srcTileBytes;
35- const size_t dstBytes = dstRows * dstCols * elementBits / 8;
36- std::vector<uint8_t> input(srcBytes);
37- std::vector<uint8_t> output(dstBytes, 0xa5);
38- std::vector<uint8_t> golden = output;
39- for (size_t i = 0; i < input.size(); ++i) {
40- input[i] = static_cast<uint8_t>((i * 17 + i / 7) % 251);
41- }
42- for (int part = 0; part < partCount; ++part) {
43- for (int row = 0; row < validRows; ++row) {
44- for (int colByte = 0; colByte < validCols * elementBits / 8; ++colByte) {
45- const int srcRow = row + (operation == Nd2NzOperation::EXTRACT ? indexRow : 0);
46- const int srcColByte =
47- colByte + (operation == Nd2NzOperation::EXTRACT ? indexCol * elementBits / 8 : 0);
48- const int dstRow = row + (isInsert ? indexRow : 0) + part * validRows;
49- const int dstColByte = colByte + (isInsert ? indexCol * elementBits / 8 : 0);
50- const int blockRow = dstRow / 16;
51- const int blockCol = dstColByte / 32;
52- const size_t offset =
53- ((blockCol * (dstRows / 16) + blockRow) * 16 + dstRow % 16) * 32 + dstColByte % 32;
54- golden[offset] = input[part * srcTileBytes + srcRow * srcCols * elementBits / 8 + srcColByte];
55- }
56- }
57- }
58- ASSERT_TRUE(aclInit(nullptr) == ACL_SUCCESS);
59- const char* deviceEnv = std::getenv("PTO_DEVICE_ID");
60- const int deviceId = deviceEnv ? std::atoi(deviceEnv) : 0;
61- ASSERT_TRUE(aclrtSetDevice(deviceId) == ACL_SUCCESS);
62- aclrtStream stream;
63- ASSERT_TRUE(aclrtCreateStream(&stream) == ACL_SUCCESS);
64- void* srcDevice = nullptr;
65- void* dstDevice = nullptr;
66- ASSERT_TRUE(aclrtMalloc(&srcDevice, srcBytes, ACL_MEM_MALLOC_HUGE_FIRST) == ACL_SUCCESS);
67- ASSERT_TRUE(aclrtMalloc(&dstDevice, dstBytes, ACL_MEM_MALLOC_HUGE_FIRST) == ACL_SUCCESS);
68- ASSERT_TRUE(aclrtMemcpy(srcDevice, srcBytes, input.data(), srcBytes, ACL_MEMCPY_HOST_TO_DEVICE) == ACL_SUCCESS);
69- ASSERT_TRUE(aclrtMemcpy(dstDevice, dstBytes, output.data(), dstBytes, ACL_MEMCPY_HOST_TO_DEVICE) == ACL_SUCCESS);
70- launch(static_cast<uint64_t*>(dstDevice), static_cast<uint64_t*>(srcDevice), stream);
71- EXPECT_TRUE(aclrtSynchronizeStream(stream) == ACL_SUCCESS);
72- EXPECT_TRUE(aclrtMemcpy(output.data(), dstBytes, dstDevice, dstBytes, ACL_MEMCPY_DEVICE_TO_HOST) == ACL_SUCCESS);
73- EXPECT_TRUE(aclrtFree(srcDevice) == ACL_SUCCESS);
74- EXPECT_TRUE(aclrtFree(dstDevice) == ACL_SUCCESS);
75- EXPECT_TRUE(aclrtDestroyStream(stream) == ACL_SUCCESS);
76- EXPECT_TRUE(aclrtResetDevice(deviceId) == ACL_SUCCESS);
77- EXPECT_TRUE(aclFinalize() == ACL_SUCCESS);
78- EXPECT_TRUE(output == golden);
79-}
80- 
81-} // namespace PtoTestCommon
82- 
83-#endif
Dtests/npu/a5/src/st/testcase/tmov_ub2l1/ub2l1_nd2nz_types.h+0-16
@@ -1,16 +0,0 @@
1-/**
2-Copyright (c) 2026 Huawei Technologies Co., Ltd.
3-This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-CANN Open Software License Agreement Version 2.0 (the "License").
5-Please refer to the License for details. You may not use this file except in compliance with the License.
6-THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-See LICENSE in the root of the software repository for the full text of the License.
9-*/
10- 
11-#ifndef UB2L1_ND2NZ_TYPES_H_
12-#define UB2L1_ND2NZ_TYPES_H_
13- 
14-enum class Nd2NzOperation { MOV, EXTRACT, INSERT, DUAL_INSERT };
15- 
16-#endif