diff --git a/docs/isa/micro-isa/18-special-scalar.md b/docs/isa/micro-isa/18-special-scalar.md new file mode 100644 index 0000000000..a9d4370c0e --- /dev/null +++ b/docs/isa/micro-isa/18-special-scalar.md @@ -0,0 +1,393 @@ +# 18. Special Scalar Operations + +> **Category:** PTO scalar query, pointer/address, and scalar-memory operations +> **Dialect:** `pto` + +Special Scalar operations provide the PTO-specific scalar facilities used +around vector and tile code. They query the current kernel execution instance, +construct and adjust typed pointers, access one scalar element through the +scalar pipeline, and perform ordinary AICore GM accesses that bypass the local +L1 data cache. + +This group does not include shared scalar arithmetic, which remains in +[Arith](14-shared-arith.md), or SIMT workitem operations, which remain in +[SIMT Ops](17-simt.md). An operation with a scalar operand but a vector result, +such as `pto.vadds`, belongs to [Vec-Scalar Ops](08-vec-scalar-ops.md). + +--- + +## Operation Summary + +| Family | Operations | Purpose | +|--------|------------|---------| +| Kernel execution queries | `pto.get_block_idx`, `pto.get_subblock_idx`, `pto.get_block_num`, `pto.get_subblock_num` | Query the block or subblock identity and launch extent visible to the current kernel instance | +| Typed pointer/address operations | `pto.castptr`, `pto.addptr` | Construct, reinterpret, and offset `!pto.ptr` values | +| Scalar-pipeline memory | `pto.load_scalar`, `pto.store_scalar` | Read or write one element through the general scalar-memory interface | +| AICore scalar GM L1-bypass | `pto.ld_dev`, `pto.st_dev` | Read or write one integer GM element while bypassing the local L1 data cache | + +--- + +## Common Pointer and Offset Rules + +Memory operations in this chapter use the typed pointer form +`!pto.ptr`: + +- `T` is the element type stored at the pointed-to location; +- `space` identifies the memory space, such as `gm` or `ub`; +- a pointer carries an address, element type, and memory-space interpretation, + but no tensor shape or stride metadata; +- every `%offset` operand in this chapter has type `index` and is measured in + elements of `T`, not bytes. + +For a base address `base`, element type `T`, and element offset `offset`, the +effective byte address is: + +```text +effective_address = base + offset * sizeof(T) +``` + +For example, offset `3` on `!pto.ptr` selects the element beginning +12 bytes after the base address. + +--- + +## Kernel Execution Query Operations + +These nullary, side-effect-free operations expose the block-level execution +state visible to the current PTO kernel instance. They return `i64` values and +do not perform memory access, synchronization, tiling, or work partitioning by +themselves. + +They are distinct from the `pto.get_tid_*`, `pto.get_block_idx_*`, and related +SIMT workitem queries documented in [SIMT Ops](17-simt.md). + +### `pto.get_block_idx` + +- **Purpose:** Return the linear block index of the current kernel instance. +- **Syntax:** + + ```mlir + %block = pto.get_block_idx + ``` + +- **Operands and attributes:** None. +- **Result:** One `i64` block index. +- **Semantics:** The result identifies the current block in the range + `[0, block_num)`, where `block_num` is the value returned by + `pto.get_block_num` for the same launch. + +```text +block = current_block_index +0 <= block < block_num +``` + +### `pto.get_subblock_idx` + +- **Purpose:** Return the subblock index visible to the current kernel + instance. +- **Syntax:** + + ```mlir + %subblock = pto.get_subblock_idx + ``` + +- **Operands and attributes:** None. +- **Result:** One `i64` subblock index. +- **Semantics:** The result identifies the current subblock in the range + `[0, subblock_num)`, where `subblock_num` is returned by + `pto.get_subblock_num` for the same execution instance. + +```text +subblock = current_subblock_index +0 <= subblock < subblock_num +``` + +### `pto.get_block_num` + +- **Purpose:** Return the total number of blocks in the current kernel launch. +- **Syntax:** + + ```mlir + %block_num = pto.get_block_num + ``` + +- **Operands and attributes:** None. +- **Result:** One `i64` block count. +- **Semantics:** The result is the launch-wide block count used to interpret + `pto.get_block_idx`. + +### `pto.get_subblock_num` + +- **Purpose:** Return the number of subblocks visible to the current execution + instance. +- **Syntax:** + + ```mlir + %subblock_num = pto.get_subblock_num + ``` + +- **Operands and attributes:** None. +- **Result:** One `i64` subblock count. +- **Semantics:** The result is the subblock count used to interpret + `pto.get_subblock_idx`. + +### Block Partitioning Example + +The following example assigns a disjoint 2048-element GM window to each +block: + +```mlir +%block = pto.get_block_idx +%block_num = pto.get_block_num +%block_len = arith.constant 2048 : index +%block_as_index = arith.index_cast %block : i64 to index +%block_offset = arith.muli %block_as_index, %block_len : index +%block_in = pto.addptr %gm_in, %block_offset + : !pto.ptr -> !pto.ptr +%block_out = pto.addptr %gm_out, %block_offset + : !pto.ptr -> !pto.ptr +``` + +The query operations report the launch state; the surrounding arithmetic and +pointer operations define the actual partitioning policy. + +--- + +## Typed Pointer and Address Operations + +### `pto.castptr` + +- **Purpose:** Explicitly convert between an integer address, a typed PTO + pointer, or a memref base address without moving data. +- **Syntax:** + + ```mlir + %result = pto.castptr %input : input-type -> result-type + ``` + +- **Operands:** `%input` is an integer, a memref, or `!pto.ptr`. +- **Result:** An integer or `!pto.ptr` according to the selected form. +- **Attributes:** None. +- **Legal forms:** + + | Input | Result | Meaning | + |-------|--------|---------| + | integer | `!pto.ptr` | Interpret the integer as an address in `space` | + | `!pto.ptr` | integer | Expose the pointer address as an integer | + | `!pto.ptr` | `!pto.ptr` | Reinterpret the element type while preserving the address and memory space | + | `memref<..., space>` | `!pto.ptr` | Extract the aligned base address and represent it as a PTO pointer | + +- **Constraints:** Integer-to-integer and memref-to-integer forms are invalid. + Pointer-to-pointer casts must preserve the PTO memory space. A memref with an + explicit PTO memory space must be cast to the same space. The operation does + not dereference the address and does not change the referenced bytes. + +```text +result.address = input.address +result.space = requested_space +result.element_type = requested_element_type +``` + +Examples: + +```mlir +%gm_i32 = pto.castptr %addr : i64 -> !pto.ptr +%gm_i8 = pto.castptr %gm_i32 + : !pto.ptr -> !pto.ptr +%addr_again = pto.castptr %gm_i8 : !pto.ptr -> i64 +``` + +### `pto.addptr` + +- **Purpose:** Produce a pointer displaced from a typed base pointer. +- **Syntax:** + + ```mlir + %result = pto.addptr %ptr, %offset + : !pto.ptr -> !pto.ptr + ``` + +- **Operands:** `%ptr` is `!pto.ptr` and `%offset` is `index`. +- **Result:** A pointer with exactly the same element type and memory space as + `%ptr`. +- **Attributes:** None. +- **Semantics:** `%offset` is a signed element displacement. Positive values + advance the pointer and negative values move it toward lower addresses. + +```text +result.address = ptr.address + offset * sizeof(T) +result.element_type = T +result.space = space +``` + +- **Constraints:** The result type must exactly match the input pointer type. + `pto.addptr` does not access memory and does not perform bounds checking. + +Example: + +```mlir +%c16 = arith.constant 16 : index +%tail = pto.addptr %base, %c16 + : !pto.ptr -> !pto.ptr +``` + +`%tail` points 16 `f32` elements, or 64 bytes, after `%base`. + +--- + +## Scalar-Pipeline Memory Operations + +`pto.load_scalar` and `pto.store_scalar` access one element through the general +scalar-memory interface. The pointer element type and memory space determine +the accessed value type and storage domain. + +### `pto.load_scalar` + +- **Purpose:** Read one scalar element from a typed PTO pointer. +- **Syntax:** + + ```mlir + %value = pto.load_scalar %ptr[%offset] + : !pto.ptr -> T + ``` + +- **Operands:** `%ptr` is `!pto.ptr` and `%offset` is an element + offset of type `index`. +- **Result:** One scalar value of type `T`. +- **Attributes:** None. +- **Semantics:** Read the element at `ptr + offset` through the scalar pipeline. + +```text +value = memory[ptr.address + offset * sizeof(T)] as T +``` + +- **Constraints:** The result type must exactly match the pointer element type. + This op returns a scalar, not a `!pto.vreg` value, and has no vector load + distribution or mask clauses. + +### `pto.store_scalar` + +- **Purpose:** Write one scalar element through a typed PTO pointer. +- **Syntax:** + + ```mlir + pto.store_scalar %value, %ptr[%offset] + : !pto.ptr, T + ``` + +- **Operands:** `%value` has type `T`, `%ptr` is `!pto.ptr`, and + `%offset` is an element offset of type `index`. +- **Results:** None. +- **Attributes:** None. +- **Semantics:** Write `%value` to the element at `ptr + offset` through the + scalar pipeline. + +```text +memory[ptr.address + offset * sizeof(T)] = value +``` + +- **Constraints:** `%value` must exactly match the pointer element type. This + op writes one scalar element and has no vector store distribution or mask + clauses. + +Example round trip in UB: + +```mlir +%c7 = arith.constant 7 : index +%value = pto.load_scalar %ub[%c7] : !pto.ptr -> i32 +pto.store_scalar %value, %ub_out[%c7] : !pto.ptr, i32 +``` + +--- + +## AICore Scalar GM L1-Bypass Operations + +`pto.ld_dev` and `pto.st_dev` are the ordinary AICore scalar GM access pair +for accesses that must bypass the local L1 data cache. They are not SIMT +operations and must not be substituted with `pto.ldg` or `pto.stg`, whose +execution scope and cache-control contract are different. + +### Common Contract + +- the pointer must be `!pto.ptr`; +- `T` must be one of `i8`, `i16`, `i32`, or `i64`; +- `%offset` has type `index` and is measured in elements of `T`; +- load result and store value types must exactly match `T`; +- no `l1cache` or `l2cache` policy attribute is accepted; +- the op must appear in an ordinary AICore entry function, outside both a + `pto.simt_entry` function and `pto.section.simt`; +- the supported target profile is A5 with CANN output version 9.0.0 official + or newer. + +Both operations are non-atomic. They do not imply synchronization, memory +ordering, cache invalidation, cache writeback, or an L2 cache policy. Programs +that combine these accesses with another cached path must provide any required +synchronization and cache maintenance separately. Cache behavior beyond the +local L1 data cache is target-defined. + +### `pto.ld_dev` + +- **Purpose:** Read one integer scalar from GM while bypassing the local L1 + data cache. +- **Syntax:** + + ```mlir + %value = pto.ld_dev %ptr[%offset] : !pto.ptr -> T + ``` + +- **Operands:** `%ptr` is `!pto.ptr` and `%offset` is an element offset + of type `index`. +- **Result:** One value of type `T` containing exactly the bytes read from GM. +- **Attributes:** None. +- **Semantics:** Read `sizeof(T)` bytes from the selected GM element. No sign + extension, zero extension, truncation, or numeric conversion is part of the + observable operation semantics. + +```text +address = ptr.address + offset * sizeof(T) +value = GM[address : address + sizeof(T)] as T +``` + +### `pto.st_dev` + +- **Purpose:** Write one integer scalar to GM while bypassing the local L1 data + cache. +- **Syntax:** + + ```mlir + pto.st_dev %value, %ptr[%offset] : !pto.ptr, T + ``` + +- **Operands:** `%value` has type `T`, `%ptr` is `!pto.ptr`, and + `%offset` is an element offset of type `index`. +- **Results:** None. +- **Attributes:** None. +- **Semantics:** Write exactly `sizeof(T)` bytes from `%value` to the selected + GM element. + +```text +address = ptr.address + offset * sizeof(T) +GM[address : address + sizeof(T)] = value as bytes +``` + +### Nonzero-Offset Example + +```mlir +%c3 = arith.constant 3 : index +%value = pto.ld_dev %src[%c3] : !pto.ptr -> i32 +pto.st_dev %value, %dst[%c3] : !pto.ptr, i32 +``` + +Both operations access element 3, which begins 12 bytes after the corresponding +`i32` base address. The load and store bypass the local L1 data cache; they do +not establish ordering with other memory operations. + +--- + +## Choosing a Scalar Memory Operation + +| Requirement | Operation family | +|-------------|------------------| +| General typed scalar access through the scalar-memory interface | `pto.load_scalar`, `pto.store_scalar` | +| Ordinary AICore integer GM access that bypasses local L1 | `pto.ld_dev`, `pto.st_dev` | +| SIMT workitem scalar memory access | See [SIMT Ops](17-simt.md) | diff --git a/docs/vpto-spec.md b/docs/vpto-spec.md index 4c13bc72ba..6ba16b0400 100644 --- a/docs/vpto-spec.md +++ b/docs/vpto-spec.md @@ -738,45 +738,10 @@ At the PTO micro Instruction level, these runtime-query ops are pure scalar prod In this pattern, all blocks execute the same kernel body, but each block sees a different `%block` value and therefore computes a different GM window. -#### `pto.get_block_idx` - -- **syntax:** `%block = pto.get_block_idx` -- **result:** `i64` -- **semantics:** Return the current block ID in the range `[0, pto.get_block_num())`. - -```c -block = block_idx(); -``` - -#### `pto.get_subblock_idx` - -- **syntax:** `%subblock = pto.get_subblock_idx` -- **result:** `i64` -- **semantics:** Return the current subblock ID in the range `[0, pto.get_subblock_num())`. - -```c -subblock = subblock_idx(); -``` - -#### `pto.get_block_num` - -- **syntax:** `%block_num = pto.get_block_num` -- **result:** `i64` -- **semantics:** Return the total number of launched blocks visible to the current kernel instance. - -```c -block_num = block_num(); -``` - -#### `pto.get_subblock_num` - -- **syntax:** `%subblock_num = pto.get_subblock_num` -- **result:** `i64` -- **semantics:** Return the total number of visible subblocks for the current execution instance. - -```c -subblock_num = subblock_num(); -``` +The complete syntax, result types, constraints, semantics, pseudocode, and +partitioning example for `pto.get_block_idx`, `pto.get_subblock_idx`, +`pto.get_block_num`, and `pto.get_subblock_num` are documented in +[Special Scalar Operations](isa/micro-isa/18-special-scalar.md#kernel-execution-query-operations). #### `pto.store_vfsimt_info` @@ -974,64 +939,9 @@ Typical examples: ### Pointer Operations -#### `pto.castptr` - -- **syntax:** `%result = pto.castptr %addr : i64 -> !pto.ptr` -- **semantics:** Reinterpret a scalar address value as a typed PTO pointer in the target memory space. - -```c -result = (ptr)addr; -``` - -`pto.castptr` is a pointer-construction operation. It does not perform data movement and does not by itself imply any load/store side effect. - -#### `pto.addptr` - -- **syntax:** `%result = pto.addptr %ptr, %offset : !pto.ptr -> !pto.ptr` -- **semantics:** Compute a new pointer by advancing the base pointer by an element offset. - -```c -result = ptr + offset; // offset counted in elements, not bytes -``` - -`pto.addptr` preserves both the element type `T` and the memory-space tag `space`. - -#### `pto.load_scalar` - -- **syntax:** `%value = pto.load_scalar %ptr[%offset] : !pto.ptr -> T` -- **semantics:** Load one scalar element from a pointer-like operand. - -```c -value = ptr[offset]; -``` - -- **inputs:** - `%ptr` is a typed PTO pointer `!pto.ptr`, and `%offset` is an - `index` displacement counted in elements. -- **outputs:** - `%value` is the loaded scalar element. -- **constraints and limitations:** - The result type MUST match the element type of `%ptr`. This op is a scalar - memory helper; unlike `pto.vlds`, it does not produce a `vreg` result and - does not participate in vector load `dist` families. - -#### `pto.store_scalar` - -- **syntax:** `pto.store_scalar %value, %ptr[%offset] : !pto.ptr, T` -- **semantics:** Store one scalar element to a pointer-like operand. - -```c -ptr[offset] = value; -``` - -- **inputs:** - `%value` is the scalar value to store. `%ptr` is a typed PTO pointer - `!pto.ptr`, and `%offset` is an `index` displacement counted in - elements. -- **constraints and limitations:** - The stored value type MUST match the element type of `%ptr`. This op is a - scalar memory helper; unlike `pto.vsts`, it does not consume a mask and does - not target vector-store `dist` families. +The complete contracts for `pto.castptr`, `pto.addptr`, `pto.load_scalar`, and +`pto.store_scalar` are documented in +[Special Scalar Operations](isa/micro-isa/18-special-scalar.md#typed-pointer-and-address-operations). #### `pto.load` @@ -1070,6 +980,11 @@ ptr[offset] = value; The stored value type MUST match the element type of `%ptr`. This is the preferred scalar memory op for VPTO/SIMT authoring. +The complete syntax, type restrictions, execution-scope rules, cache behavior, +target availability, and examples for `pto.ld_dev` and `pto.st_dev` are +documented in +[Special Scalar Operations](isa/micro-isa/18-special-scalar.md#aicore-scalar-gm-l1-bypass-operations). + #### Pointer-Based Vector Access Example The following lowered-style fragment shows how typed PTO pointers flow through @@ -1408,6 +1323,7 @@ This section provides a categorized overview of all PTO micro Instruction operat | 15 | [SCF (Shared MLIR Dialect)](isa/micro-isa/15-shared-scf.md) | Structured loops, branches, and loop-carried state around PTO regions | 5 | `scf.for`, `scf.if`, `scf.while`, `scf.condition`, `scf.yield` | | 16 | [Cube Matrix Multiply](isa/micro-isa/16-cube-matmul.md) | GM↔L1 (`l1`/cbuf) staging, L1 (`l1`)↔UB/BT/FB side moves, L1→L0A/L0B loads, L0C (`l0c`) matmul, and FIXPIPE MTE writeback | 19 | `pto.mte_gm_l1`, `pto.mte_l1_ub`, `pto.mte_gm_l1_frac`, `pto.mte_l1_bt`, `pto.mte_l1_fb`, `pto.mte_l1_l0a`, `pto.mte_l1_l0b`, `pto.mte_l1_l0a_mx`, `pto.mte_l1_l0b_mx`, `pto.mad`, `pto.mad_acc`, `pto.mad_bias`, `pto.mad_mx`, `pto.mad_mx_acc`, `pto.mad_mx_bias`, `pto.mte_l0c_l1`, `pto.mte_l0c_gm`, `pto.mte_l0c_ub` | | 17 | [SIMT Ops](isa/micro-isa/17-simt.md) | SIMT launch, thread/lane queries, vote/shuffle/redux, scalar memory, atomics, scalar math, conversion, entry synchronization, and state preservation | ~65 | `pto.store_vfsimt_info`, `pto.simt_launch`, `pto.get_tid_x`, `pto.get_laneid`, `pto.vote_*`, `pto.shuffle_*`, `pto.redux_*`, `pto.load`, `pto.store`, `pto.atomic_*`, `pto.convert`, `pto.syncthreads`, `pto.keep`, `pto.resume`, etc. | +| 18 | [Special Scalar Operations](isa/micro-isa/18-special-scalar.md) | PTO scalar kernel queries, typed pointer/address calculation, scalar-pipeline memory, and ordinary AICore GM L1-bypass access | 10 | `pto.get_block_idx`, `pto.get_subblock_idx`, `pto.get_block_num`, `pto.get_subblock_num`, `pto.castptr`, `pto.addptr`, `pto.load_scalar`, `pto.store_scalar`, `pto.ld_dev`, `pto.st_dev` | --- @@ -1431,6 +1347,7 @@ This section provides a categorized overview of all PTO micro Instruction operat | Gather | 3 | `pto.vgather2`, `pto.vgatherb` | | Contiguous Store | 3 | `pto.vsts` with `NORM_B8` / `NORM_B16` / `NORM_B32` dist | | Scatter | 3 | `pto.vscatter` | +| Scalar GM access bypassing local L1 data cache | 18 | `pto.ld_dev`, `pto.st_dev` | ### Compute Operations @@ -1462,7 +1379,10 @@ This section provides a categorized overview of all PTO micro Instruction operat ### Scalar & Control Operations -Group 14 covers the full scalar `arith` surface. The rows below list common PTO micro Instruction patterns rather than an exhaustive partition of `arith` ops. +Group 14 covers shared MLIR scalar arithmetic. Group 18 catalogs PTO scalar +queries, pointer/address operations, and scalar-memory operations. SIMT scalar +operations remain in Group 17, while +shared structured-control semantics remain in Group 15. | Operation | Group | Description | |-----------|-------|-------------| @@ -1472,6 +1392,12 @@ Group 14 covers the full scalar `arith` surface. The rows below list common PTO | Scalar Compare & Select | 14 | `arith.cmpi`, `arith.cmpf`, `arith.select` | | Scalar Casts / Width Changes | 14 | `arith.index_cast`, `arith.index_castui`, `arith.extsi`, `arith.extui`, `arith.trunci`, `arith.sitofp`, etc. | | Scalar Bitwise / Shift Ops | 14 | `arith.andi`, `arith.ori`, `arith.xori`, `arith.shli`, `arith.shrsi`, `arith.shrui`, etc. | +| Kernel Execution Queries | 18 | `pto.get_block_idx`, `pto.get_subblock_idx`, `pto.get_block_num`, `pto.get_subblock_num` | +| Typed Pointer / Address Operations | 18 | `pto.castptr`, `pto.addptr` | +| Scalar-Pipeline Memory | 18 | `pto.load_scalar`, `pto.store_scalar` | +| AICore Scalar GM L1-Bypass | 18 | `pto.ld_dev`, `pto.st_dev` | +| SIMT Scalar Memory / Atomics | 17 | `pto.load`, `pto.store`, `pto.ldg`, `pto.stg`, `pto.atomic_*` | +| SIMT Scalar Math / Conversion | 17 | `pto.prmt`, `pto.mulhi`, `pto.sqrt`, `pto.exp`, `pto.fma`, `pto.convert`, etc. | | Counted Loops | 15 | `scf.for` | | Conditional Regions | 15 | `scf.if`, `scf.yield` | | Break-like Structured Loops | 15 | `scf.while`, `scf.condition`, `scf.yield` | diff --git a/include/PTO/IR/VPTOOps.td b/include/PTO/IR/VPTOOps.td index 395807988f..ee23a053b5 100644 --- a/include/PTO/IR/VPTOOps.td +++ b/include/PTO/IR/VPTOOps.td @@ -135,6 +135,8 @@ def PTO_SimtConvertValueType : AnyTypeOf<[I32, I64, F16, BF16, F32, "standard scalar, low-precision scalar payload, or supported vector<2xT> conversion type">; def PTO_AtomicValueType : AnyTypeOf<[I32, I64, F16, BF16, F32, PTO_V2F16Type, PTO_V2BF16Type], "i32, i64, f16, bf16, f32, vector<2xf16> or vector<2xbf16>">; +def PTO_DevScalarValueType : AnyTypeOf<[I8, I16, I32, I64], + "i8, i16, i32 or i64">; class PTO_MicroOp traits = []> : PTO_Op; @@ -258,6 +260,57 @@ def PTOStgOp : PTO_SimtOp<"stg", [ }]; } +def PTOLdDevOp : PTO_MicroOp<"ld_dev", [ + DeclareOpInterfaceMethods + ]> { + let summary = "Load one scalar GM element while bypassing the local L1 data cache."; + let description = [{ + Load one integer scalar element from GM through the AICore device-memory + path while bypassing the local L1 data cache. The offset operand is an + element offset, not a byte offset. The operation does not provide + atomicity, synchronization, memory-ordering, or L2 cache-policy semantics. + }]; + + let arguments = (ins + PTO_BufferType:$ptr, + Index:$offset + ); + + let results = (outs PTO_DevScalarValueType:$value); + + let hasVerifier = 1; + + let assemblyFormat = [{ + $ptr `[` $offset `]` attr-dict `:` type($ptr) `->` type($value) + }]; +} + +def PTOStDevOp : PTO_MicroOp<"st_dev", [ + DeclareOpInterfaceMethods + ]> { + let summary = "Store one scalar GM element while bypassing the local L1 data cache."; + let description = [{ + Store one integer scalar element to GM through the AICore device-memory + path while bypassing the local L1 data cache. The offset operand is an + element offset, not a byte offset. The operation does not provide + atomicity, synchronization, memory-ordering, or L2 cache-policy semantics. + }]; + + let arguments = (ins + PTO_DevScalarValueType:$value, + PTO_BufferType:$ptr, + Index:$offset + ); + + let results = (outs); + + let hasVerifier = 1; + + let assemblyFormat = [{ + $value `,` $ptr `[` $offset `]` attr-dict `:` type($ptr) `,` type($value) + }]; +} + def TensorViewAddrOp : PTO_Op<"tensor_view_addr", [Pure]> { let summary = "Extract address from a tensor view."; let description = [{ diff --git a/lib/PTO/IR/VPTO.cpp b/lib/PTO/IR/VPTO.cpp index a02eac07a8..b3cc857b5f 100644 --- a/lib/PTO/IR/VPTO.cpp +++ b/lib/PTO/IR/VPTO.cpp @@ -546,6 +546,33 @@ static LogicalResult verifyLdgStgAccess(Operation *op, Type ptrType, "packed vector<2/4/8xfp8>, and !pto.hif8x2 value type"; } +static LogicalResult verifyLdStDevAccess(Operation *op, Type ptrType, + Type valueType) { + if (op->hasAttr("l1cache") || op->hasAttr("l2cache")) + return op->emitOpError() + << "does not accept l1cache or l2cache policy attributes"; + + auto ptrTy = dyn_cast(ptrType); + if (!ptrTy) + return op->emitOpError() << "requires !pto.ptr operand"; + if (ptrTy.getMemorySpace().getAddressSpace() != AddressSpace::GM) + return op->emitOpError() << "requires GM pointer"; + + auto intType = dyn_cast(valueType); + if (!intType || (intType.getWidth() != 8 && intType.getWidth() != 16 && + intType.getWidth() != 32 && intType.getWidth() != 64)) + return op->emitOpError() << "supports only i8, i16, i32 or i64 values"; + + if (isInsideSimtExecutionScope(op)) + return op->emitOpError() + << "must be outside pto.simt_entry functions and pto.section.simt"; + auto funcOp = op->getParentOfType(); + if (!funcOp || !pto::isPTOEntryFunction(funcOp)) + return op->emitOpError() + << "requires an enclosing ordinary AICore entry function"; + return success(); +} + LogicalResult PTOLoadOp::verify() { if (failed(verifyVPTOScalarAccessTypes(getOperation(), getPtr().getType(), getValue().getType(), "load"))) @@ -564,16 +591,42 @@ LogicalResult PTOLdgOp::verify() { if (failed(verifyVPTOScalarAccessTypes(getOperation(), getPtr().getType(), getValue().getType(), "ldg"))) return failure(); - return verifyLdgStgAccess(getOperation(), getPtr().getType(), - getValue().getType()); + if (failed(verifyLdgStgAccess(getOperation(), getPtr().getType(), + getValue().getType()))) + return failure(); + if (!isInsideSimtExecutionScope(getOperation())) + return emitOpError() + << "must be inside a pto.simt_entry function or pto.section.simt"; + return success(); } LogicalResult PTOStgOp::verify() { if (failed(verifyVPTOScalarAccessTypes(getOperation(), getPtr().getType(), getValue().getType(), "stg"))) return failure(); - return verifyLdgStgAccess(getOperation(), getPtr().getType(), - getValue().getType()); + if (failed(verifyLdgStgAccess(getOperation(), getPtr().getType(), + getValue().getType()))) + return failure(); + if (!isInsideSimtExecutionScope(getOperation())) + return emitOpError() + << "must be inside a pto.simt_entry function or pto.section.simt"; + return success(); +} + +LogicalResult PTOLdDevOp::verify() { + if (failed(verifyVPTOScalarAccessTypes(getOperation(), getPtr().getType(), + getValue().getType(), "ld_dev"))) + return failure(); + return verifyLdStDevAccess(getOperation(), getPtr().getType(), + getValue().getType()); +} + +LogicalResult PTOStDevOp::verify() { + if (failed(verifyVPTOScalarAccessTypes(getOperation(), getPtr().getType(), + getValue().getType(), "st_dev"))) + return failure(); + return verifyLdStDevAccess(getOperation(), getPtr().getType(), + getValue().getType()); } LogicalResult ShuffleIdxOp::verify() { @@ -706,6 +759,18 @@ void PTOStgOp::getEffects( effects.emplace_back(MemoryEffects::Write::get(), &getPtrMutable()); } +void PTOLdDevOp::getEffects( + SmallVectorImpl> + &effects) { + effects.emplace_back(MemoryEffects::Read::get(), &getPtrMutable()); +} + +void PTOStDevOp::getEffects( + SmallVectorImpl> + &effects) { + effects.emplace_back(MemoryEffects::Write::get(), &getPtrMutable()); +} + template static void getAtomicEffects( OpTy op, diff --git a/lib/PTO/Transforms/VPTOCANN900LLVMEmitter.cpp b/lib/PTO/Transforms/VPTOCANN900LLVMEmitter.cpp index c82a9141d8..df30c8b79e 100644 --- a/lib/PTO/Transforms/VPTOCANN900LLVMEmitter.cpp +++ b/lib/PTO/Transforms/VPTOCANN900LLVMEmitter.cpp @@ -10431,6 +10431,137 @@ class ConvertPtoStgOp final : public OpConversionPattern { LoweringState &state; }; +static std::string buildLdDevCalleeName(unsigned width) { + return "llvm.hivm.LD.DEV.u" + std::to_string(width) + ".GM"; +} + +static std::string buildStDevCalleeName(unsigned width) { + return "llvm.hivm.ST.DEV.u" + std::to_string(width); +} + +class ConvertPtoLdDevOp final : public OpConversionPattern { +public: + ConvertPtoLdDevOp(TypeConverter &typeConverter, MLIRContext *context, + LoweringState &state) + : OpConversionPattern(typeConverter, context), + state(state) {} + + LogicalResult + matchAndRewrite(pto::PTOLdDevOp op, OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override { + auto llvmPtrType = dyn_cast(adaptor.getPtr().getType()); + if (!llvmPtrType) + return rewriter.notifyMatchFailure(op, "expected LLVM pointer operand"); + + auto valueType = dyn_cast(op.getValue().getType()); + if (!valueType) + return rewriter.notifyMatchFailure(op, "expected integer result type"); + + Value offset = adaptor.getOffset(); + if (offset.getType().isIndex()) + offset = rewriter.create(op.getLoc(), + rewriter.getI64Type(), offset); + + Type convertedValueType = + getTypeConverter()->convertType(op.getValue().getType()); + if (!convertedValueType) + return rewriter.notifyMatchFailure(op, + "could not convert ld_dev result type"); + + Value elemPtr = adaptor.getPtr(); + if (!matchPattern(offset, m_Zero())) { + elemPtr = rewriter.create( + op.getLoc(), llvmPtrType, + normalizeGEPElementTypeForLLVMLowering(convertedValueType, rewriter), + adaptor.getPtr(), ValueRange{offset}); + } + + FailureOr gmPtr = reinterpretPointerToAddrSpace( + op, elemPtr, static_cast(pto::AddressSpace::GM)); + if (failed(gmPtr)) + return rewriter.notifyMatchFailure(op, "failed to map ld_dev GM pointer"); + + std::string calleeName = buildLdDevCalleeName(valueType.getWidth()); + Value intrinsicOffset = getI64Constant(rewriter, op.getLoc(), 0); + auto funcType = rewriter.getFunctionType( + TypeRange{gmPtr->getType(), rewriter.getI64Type()}, + TypeRange{rewriter.getI64Type()}); + auto call = rewriter.create( + op.getLoc(), calleeName, TypeRange{rewriter.getI64Type()}, + ValueRange{*gmPtr, intrinsicOffset}); + state.plannedDecls.push_back(PlannedDecl{calleeName, funcType}); + + Value result = call.getResult(0); + if (valueType.getWidth() < 64) + result = rewriter.create(op.getLoc(), convertedValueType, + result); + rewriter.replaceOp(op, result); + return success(); + } + +private: + LoweringState &state; +}; + +class ConvertPtoStDevOp final : public OpConversionPattern { +public: + ConvertPtoStDevOp(TypeConverter &typeConverter, MLIRContext *context, + LoweringState &state) + : OpConversionPattern(typeConverter, context), + state(state) {} + + LogicalResult + matchAndRewrite(pto::PTOStDevOp op, OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override { + auto llvmPtrType = dyn_cast(adaptor.getPtr().getType()); + if (!llvmPtrType) + return rewriter.notifyMatchFailure(op, "expected LLVM pointer operand"); + + auto valueType = dyn_cast(op.getValue().getType()); + if (!valueType) + return rewriter.notifyMatchFailure(op, "expected integer value type"); + + Value offset = adaptor.getOffset(); + if (offset.getType().isIndex()) + offset = rewriter.create(op.getLoc(), + rewriter.getI64Type(), offset); + + Value elemPtr = adaptor.getPtr(); + if (!matchPattern(offset, m_Zero())) { + elemPtr = rewriter.create( + op.getLoc(), llvmPtrType, + normalizeGEPElementTypeForLLVMLowering(adaptor.getValue().getType(), + rewriter), + adaptor.getPtr(), ValueRange{offset}); + } + + FailureOr gmPtr = reinterpretPointerToAddrSpace( + op, elemPtr, static_cast(pto::AddressSpace::GM)); + if (failed(gmPtr)) + return rewriter.notifyMatchFailure(op, "failed to map st_dev GM pointer"); + + Value payload = adaptor.getValue(); + if (valueType.getWidth() < 64) + payload = rewriter.create(op.getLoc(), + rewriter.getI64Type(), payload); + + std::string calleeName = buildStDevCalleeName(valueType.getWidth()); + Value intrinsicOffset = getI64Constant(rewriter, op.getLoc(), 0); + auto funcType = rewriter.getFunctionType( + TypeRange{rewriter.getI64Type(), gmPtr->getType(), + rewriter.getI64Type()}, + TypeRange{}); + rewriter.create(op.getLoc(), calleeName, TypeRange{}, + ValueRange{payload, *gmPtr, intrinsicOffset}); + state.plannedDecls.push_back(PlannedDecl{calleeName, funcType}); + rewriter.eraseOp(op); + return success(); + } + +private: + LoweringState &state; +}; + class ConvertVPTOTypedCarrierOp final : public ConversionPattern { public: ConvertVPTOTypedCarrierOp(TypeConverter &typeConverter, MLIRContext *context) @@ -10874,7 +11005,8 @@ static LogicalResult lowerVPTOTypes(ModuleOp module, llvm::raw_ostream &diagOS) }); target.addIllegalOp(); target.addDynamicallyLegalOp( [&](UnrealizedConversionCastOp op) { @@ -10892,7 +11024,7 @@ static LogicalResult lowerVPTOTypes(ModuleOp module, llvm::raw_ostream &diagOS) ConvertPtoStructGetOp, ConvertPtoStructSetOp>(typeConverter, context); patterns.add( + ConvertPtoStgOp, ConvertPtoLdDevOp, ConvertPtoStDevOp>( typeConverter, context, state); patterns.add(typeConverter, context); patterns.add(typeConverter, context); diff --git a/lib/PTO/Transforms/VPTOLLVMEmitterDispatcher.cpp b/lib/PTO/Transforms/VPTOLLVMEmitterDispatcher.cpp index 7bc5852097..70c3b852eb 100644 --- a/lib/PTO/Transforms/VPTOLLVMEmitterDispatcher.cpp +++ b/lib/PTO/Transforms/VPTOLLVMEmitterDispatcher.cpp @@ -6,6 +6,7 @@ // 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. +#include "PTO/IR/PTO.h" #include "PTO/Transforms/VPTOLLVMEmitter.h" #include "PTO/Support/CANNVersion.h" @@ -18,10 +19,38 @@ static bool usesCANN900Lowering(const VPTOEmissionOptions &options) { options.cannVersion >= CANNVersion::release(9, 0, 0); } +static bool containsLdStDev(ModuleOp module) { + bool found = false; + module.walk([&](Operation *op) { + if (isa(op)) + found = true; + }); + return found; +} + +static LogicalResult verifyLdStDevTarget(ModuleOp module, + const VPTOEmissionOptions &options, + llvm::raw_ostream &diagOS) { + if (!containsLdStDev(module) || usesCANN900Lowering(options)) + return success(); + + const bool isC220 = options.march == "dav-c220-vec" || + options.march == "dav-c220-cube"; + if (isC220) + diagOS << "VPTO LLVM emission failed: pto.ld_dev and pto.st_dev require " + "--pto-arch=a5\n"; + else + diagOS << "VPTO LLVM emission failed: pto.ld_dev and pto.st_dev require " + "CANN 9.0.0 or newer official lowering\n"; + return failure(); +} + LogicalResult lowerVPTOModuleToLLVMModules( ModuleOp module, const VPTOEmissionOptions &options, EmittedLLVMModule &cubeModule, EmittedLLVMModule &vectorModule, llvm::raw_ostream &diagOS) { + if (failed(verifyLdStDevTarget(module, options, diagOS))) + return failure(); if (usesCANN900Lowering(options)) return lowerVPTOModuleToLLVMModulesCANN900(module, options, cubeModule, vectorModule, diagOS); diff --git a/test/lit/vpto/aicore_ld_st_dev_invalid.pto b/test/lit/vpto/aicore_ld_st_dev_invalid.pto new file mode 100644 index 0000000000..fcd403473a --- /dev/null +++ b/test/lit/vpto/aicore_ld_st_dev_invalid.pto @@ -0,0 +1,82 @@ +// 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: split-file %s %t +// RUN: not ptoas --pto-arch=a5 --pto-backend=vpto %t/simt.pto -o %t/simt.o 2>&1 | FileCheck %s --check-prefix=SIMT +// RUN: not ptoas --pto-arch=a5 --pto-backend=vpto %t/ub.pto -o %t/ub.o 2>&1 | FileCheck %s --check-prefix=UB +// RUN: not ptoas --pto-arch=a5 --pto-backend=vpto %t/mismatch.pto -o %t/mismatch.o 2>&1 | FileCheck %s --check-prefix=MISMATCH +// RUN: not ptoas --pto-arch=a5 --pto-backend=vpto %t/helper.pto -o %t/helper.o 2>&1 | FileCheck %s --check-prefix=HELPER +// RUN: not ptoas --pto-arch=a5 --pto-backend=vpto %t/policy.pto -o %t/policy.o 2>&1 | FileCheck %s --check-prefix=POLICY +// RUN: not ptoas --pto-arch=a5 --cann-output-version=9.0.0-beta.1 --pto-backend=vpto --emit-vpto-llvm-ir %t/target.pto -o - 2>&1 | FileCheck %s --check-prefix=BETA +// RUN: not ptoas --pto-arch=a3 --cann-output-version=9.0.0 --pto-backend=vpto --emit-vpto-llvm-ir %t/target.pto -o - 2>&1 | FileCheck %s --check-prefix=A3 + +// SIMT: 'pto.ld_dev' op must be outside pto.simt_entry functions and pto.section.simt +// UB: 'pto.st_dev' op requires GM pointer +// MISMATCH: 'pto.ld_dev' op expects ld_dev value type to match pointer element type +// HELPER: 'pto.ld_dev' op requires an enclosing ordinary AICore entry function +// POLICY: 'pto.ld_dev' op does not accept l1cache or l2cache policy attributes +// BETA: pto.ld_dev and pto.st_dev require CANN 9.0.0 or newer official lowering +// A3: pto.ld_dev and pto.st_dev require --pto-arch=a5 + +//--- simt.pto +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @simt_scope(%gm: !pto.ptr) attributes {pto.simt_entry} { + %c0 = arith.constant 0 : index + %value = pto.ld_dev %gm[%c0] : !pto.ptr -> i32 + pto.st_dev %value, %gm[%c0] : !pto.ptr, i32 + return + } +} + +//--- ub.pto +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @ub_pointer(%ub: !pto.ptr) attributes {pto.aicore} { + %c0 = arith.constant 0 : index + %value = arith.constant 1 : i32 + pto.st_dev %value, %ub[%c0] : !pto.ptr, i32 + return + } +} + +//--- mismatch.pto +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @mismatch(%gm: !pto.ptr) attributes {pto.aicore} { + %c0 = arith.constant 0 : index + %value = pto.ld_dev %gm[%c0] : !pto.ptr -> i16 + return + } +} + +//--- target.pto +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @unsupported_target(%gm: !pto.ptr) attributes {pto.aicore} { + %c0 = arith.constant 0 : index + %value = pto.ld_dev %gm[%c0] : !pto.ptr -> i32 + pto.st_dev %value, %gm[%c0] : !pto.ptr, i32 + return + } +} + +//--- helper.pto +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func private @ordinary_helper(%gm: !pto.ptr) { + %c0 = arith.constant 0 : index + %value = pto.ld_dev %gm[%c0] : !pto.ptr -> i32 + return + } +} + +//--- policy.pto +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @policy(%gm: !pto.ptr) attributes {pto.aicore} { + %c0 = arith.constant 0 : index + %value = pto.ld_dev %gm[%c0] {l1cache = #pto.l1cache} + : !pto.ptr -> i32 + return + } +} diff --git a/test/lit/vpto/aicore_ld_st_dev_vpto_llvm.pto b/test/lit/vpto/aicore_ld_st_dev_vpto_llvm.pto new file mode 100644 index 0000000000..a93bc49712 --- /dev/null +++ b/test/lit/vpto/aicore_ld_st_dev_vpto_llvm.pto @@ -0,0 +1,47 @@ +// 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: ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto --emit-vpto-llvm-ir %s -o - 2>&1 | FileCheck %s + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @aicore_ld_st_dev( + %gm_i8: !pto.ptr, + %gm_i16: !pto.ptr, + %gm_i32: !pto.ptr, + %gm_i64: !pto.ptr, + %dynamic_offset: index) attributes {pto.aicore} { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + + %v8 = pto.ld_dev %gm_i8[%c0] : !pto.ptr -> i8 + %v16 = pto.ld_dev %gm_i16[%c1] : !pto.ptr -> i16 + %v32 = pto.ld_dev %gm_i32[%dynamic_offset] : !pto.ptr -> i32 + %v64 = pto.ld_dev %gm_i64[%c0] : !pto.ptr -> i64 + + pto.st_dev %v8, %gm_i8[%c1] : !pto.ptr, i8 + pto.st_dev %v16, %gm_i16[%dynamic_offset] : !pto.ptr, i16 + pto.st_dev %v32, %gm_i32[%c0] : !pto.ptr, i32 + pto.st_dev %v64, %gm_i64[%c1] : !pto.ptr, i64 + return + } +} + +// CHECK-LABEL: define {{.*}}void @aicore_ld_st_dev_mix_aiv +// CHECK-NOT: @llvm.hivm.ldg. +// CHECK-NOT: @llvm.hivm.stg. +// CHECK-DAG: getelementptr i16, ptr addrspace(1) {{.*}}, i64 1 +// CHECK-DAG: getelementptr i32, ptr addrspace(1) {{.*}} +// CHECK-DAG: call i64 @llvm.hivm.LD.DEV.u8.GM(ptr addrspace(1) {{.*}}, i64 0) +// CHECK-DAG: call i64 @llvm.hivm.LD.DEV.u16.GM(ptr addrspace(1) {{.*}}, i64 0) +// CHECK-DAG: call i64 @llvm.hivm.LD.DEV.u32.GM(ptr addrspace(1) {{.*}}, i64 0) +// CHECK-DAG: call i64 @llvm.hivm.LD.DEV.u64.GM(ptr addrspace(1) {{.*}}, i64 0) +// CHECK-DAG: call void @llvm.hivm.ST.DEV.u8(i64 {{.*}}, ptr addrspace(1) {{.*}}, i64 0) +// CHECK-DAG: call void @llvm.hivm.ST.DEV.u16(i64 {{.*}}, ptr addrspace(1) {{.*}}, i64 0) +// CHECK-DAG: call void @llvm.hivm.ST.DEV.u32(i64 {{.*}}, ptr addrspace(1) {{.*}}, i64 0) +// CHECK-DAG: call void @llvm.hivm.ST.DEV.u64(i64 {{.*}}, ptr addrspace(1) {{.*}}, i64 0) +// CHECK: ret void diff --git a/test/lit/vpto/simt_ldg_stg_scope_invalid.pto b/test/lit/vpto/simt_ldg_stg_scope_invalid.pto new file mode 100644 index 0000000000..777f1cc2c2 --- /dev/null +++ b/test/lit/vpto/simt_ldg_stg_scope_invalid.pto @@ -0,0 +1,34 @@ +// 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: split-file %s %t +// RUN: not ptoas --pto-arch=a5 --pto-backend=vpto %t/ldg.pto -o %t/ldg.o 2>&1 | FileCheck %s --check-prefix=LDG +// RUN: not ptoas --pto-arch=a5 --pto-backend=vpto %t/stg.pto -o %t/stg.o 2>&1 | FileCheck %s --check-prefix=STG + +// LDG: 'pto.ldg' op must be inside a pto.simt_entry function or pto.section.simt +// STG: 'pto.stg' op must be inside a pto.simt_entry function or pto.section.simt + +//--- ldg.pto +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @invalid_ldg(%gm: !pto.ptr) attributes {pto.aicore} { + %c0 = arith.constant 0 : index + %value = pto.ldg %gm[%c0] : !pto.ptr -> i32 + pto.st_dev %value, %gm[%c0] : !pto.ptr, i32 + return + } +} + +//--- stg.pto +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @invalid_stg(%gm: !pto.ptr) attributes {pto.aicore} { + %c0 = arith.constant 0 : index + %value = arith.constant 1 : i32 + pto.stg %value, %gm[%c0] : !pto.ptr, i32 + return + } +} diff --git a/test/lit/vpto/simt_lowlevel_ldst_policy_vpto_llvm.pto b/test/lit/vpto/simt_lowlevel_ldst_policy_vpto_llvm.pto index b2b01574fa..17ee0da096 100644 --- a/test/lit/vpto/simt_lowlevel_ldst_policy_vpto_llvm.pto +++ b/test/lit/vpto/simt_lowlevel_ldst_policy_vpto_llvm.pto @@ -10,6 +10,7 @@ module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { func.func @ldst_policy_kernel(%gm_i8: !pto.ptr, %gm_i16: !pto.ptr, %gm_i32: !pto.ptr, %gm_i64: !pto.ptr, %gm_f16: !pto.ptr, %gm_bf16: !pto.ptr, %gm_f32: !pto.ptr, %gm_f64: !pto.ptr, %dst_i32: !pto.ptr, %dst_i64: !pto.ptr) attributes {pto.aicore} { + pto.section.simt<<<1, 1, 1>>> { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index %c2 = arith.constant 2 : index @@ -45,6 +46,7 @@ module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind, i32 pto.store %load_uncache, %dst_i64[%c0] : !pto.ptr, i64 + } return } } diff --git a/test/lit/vpto/simt_lowlevel_vector_ldg_stg_packed_types_vpto_llvm.pto b/test/lit/vpto/simt_lowlevel_vector_ldg_stg_packed_types_vpto_llvm.pto index b4e9ab5576..9aaf573265 100644 --- a/test/lit/vpto/simt_lowlevel_vector_ldg_stg_packed_types_vpto_llvm.pto +++ b/test/lit/vpto/simt_lowlevel_vector_ldg_stg_packed_types_vpto_llvm.pto @@ -21,6 +21,7 @@ module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind, gm>, %gm_i32: !pto.ptr, gm> ) attributes {pto.aicore} { + pto.section.simt<<<1, 1, 1>>> { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index @@ -58,6 +59,7 @@ module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind, gm> -> vector<2xi32> pto.stg %v_i32, %gm_i32[%c1] : !pto.ptr, gm>, vector<2xi32> + } return } } diff --git a/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/compare.py b/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/compare.py new file mode 100644 index 0000000000..7966705013 --- /dev/null +++ b/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/compare.py @@ -0,0 +1,34 @@ +#!/usr/bin/env 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. + +import os +import sys + +import numpy as np + + +def main() -> None: + strict = os.getenv("COMPARE_STRICT", "1") != "0" + golden = np.fromfile("golden_dst.bin", dtype=np.int32) + output = np.fromfile("dst.bin", dtype=np.int32) + ok = golden.shape == output.shape and np.array_equal(golden, output) + if not ok: + diff = np.nonzero(golden != output)[0] + idx = int(diff[0]) if diff.size else 0 + print( + f"[ERROR] mismatch at idx={idx}, golden={int(golden[idx])}, " + f"out={int(output[idx])}" + ) + if strict: + sys.exit(2) + print("[INFO] compare passed" if ok else "[WARN] compare failed (non-gating)") + + +if __name__ == "__main__": + main() diff --git a/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/golden.py b/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/golden.py new file mode 100644 index 0000000000..b49ae2dcdf --- /dev/null +++ b/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/golden.py @@ -0,0 +1,38 @@ +#!/usr/bin/env 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. + +import argparse +from pathlib import Path + +import numpy as np + + +def generate(output_dir: Path) -> None: + output_dir.mkdir(parents=True, exist_ok=True) + src = np.array( + [0x01020304, -0x1234567, 0x11223344, -7, 99, 100, 101, 102], + dtype=np.int32, + ) + dst = np.full(8, np.int32(0x5A5A5A5A), dtype=np.int32) + golden_dst = dst.copy() + golden_dst[3] = src[1] + src.tofile(output_dir / "src.bin") + dst.tofile(output_dir / "dst.bin") + golden_dst.tofile(output_dir / "golden_dst.bin") + + +def main() -> None: + parser = argparse.ArgumentParser() + parser.add_argument("--output-dir", type=Path, default=Path(".")) + args = parser.parse_args() + generate(args.output_dir) + + +if __name__ == "__main__": + main() diff --git a/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/kernel.pto b/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/kernel.pto new file mode 100644 index 0000000000..3c77fecef6 --- /dev/null +++ b/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/kernel.pto @@ -0,0 +1,17 @@ +// ----------------------------------------------------------------------------- +// case: micro-op/scalar-load-store/ld-st-dev +// family: scalar-load-store +// target_ops: pto.ld_dev, pto.st_dev +// scenarios: core-i32, gm-l1-bypass, nonzero-element-offset +// ----------------------------------------------------------------------------- +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @ld_st_dev_kernel( + %src: !pto.ptr, + %dst: !pto.ptr) attributes {pto.aicore} { + %src_offset = arith.constant 1 : index + %dst_offset = arith.constant 3 : index + %value = pto.ld_dev %src[%src_offset] : !pto.ptr -> i32 + pto.st_dev %value, %dst[%dst_offset] : !pto.ptr, i32 + return + } +} diff --git a/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/launch.cpp b/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/launch.cpp new file mode 100644 index 0000000000..7ae0768f7c --- /dev/null +++ b/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/launch.cpp @@ -0,0 +1,23 @@ +// 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 +#include +#ifndef __CPU_SIM +#include "acl/acl.h" +#endif + +extern "C" __global__ [aicore] void ld_st_dev_kernel(__gm__ int32_t *src, + __gm__ int32_t *dst); + +void LaunchLd_st_dev_kernel(int32_t *src, int32_t *dst, void *stream) { + ld_st_dev_kernel<<<1, nullptr, stream>>>((__gm__ int32_t *)src, + (__gm__ int32_t *)dst); +} diff --git a/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/main.cpp b/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/main.cpp new file mode 100644 index 0000000000..d89cafef69 --- /dev/null +++ b/test/vpto/cases/micro-op/scalar-load-store/ld-st-dev/main.cpp @@ -0,0 +1,81 @@ +// 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. + +#include "test_common.h" +#include "acl/acl.h" +#include +#include +#include + +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) + +void LaunchLd_st_dev_kernel(int32_t *src, int32_t *dst, void *stream); + +int main() { + constexpr size_t elemCount = 8; + size_t fileSize = elemCount * sizeof(int32_t); + int32_t *srcHost = nullptr; + int32_t *dstHost = nullptr; + int32_t *srcDevice = nullptr; + int32_t *dstDevice = nullptr; + int rc = 0; + bool aclInited = false; + bool deviceSet = false; + int deviceId = 0; + aclrtStream stream = nullptr; + + 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 **)(&srcHost), fileSize)); + ACL_CHECK(aclrtMallocHost((void **)(&dstHost), fileSize)); + ACL_CHECK(aclrtMalloc((void **)&srcDevice, fileSize, + ACL_MEM_MALLOC_HUGE_FIRST)); + ACL_CHECK(aclrtMalloc((void **)&dstDevice, fileSize, + ACL_MEM_MALLOC_HUGE_FIRST)); + + ReadFile("./src.bin", fileSize, srcHost, fileSize); + ReadFile("./dst.bin", fileSize, dstHost, fileSize); + ACL_CHECK(aclrtMemcpy(srcDevice, fileSize, srcHost, fileSize, + ACL_MEMCPY_HOST_TO_DEVICE)); + ACL_CHECK(aclrtMemcpy(dstDevice, fileSize, dstHost, fileSize, + ACL_MEMCPY_HOST_TO_DEVICE)); + LaunchLd_st_dev_kernel(srcDevice, dstDevice, stream); + ACL_CHECK(aclrtSynchronizeStream(stream)); + ACL_CHECK(aclrtMemcpy(dstHost, fileSize, dstDevice, fileSize, + ACL_MEMCPY_DEVICE_TO_HOST)); + WriteFile("./dst.bin", dstHost, fileSize); + +cleanup: + aclrtFree(dstDevice); + aclrtFree(srcDevice); + aclrtFreeHost(dstHost); + aclrtFreeHost(srcHost); + if (stream) + aclrtDestroyStream(stream); + if (deviceSet) + aclrtResetDevice(deviceId); + if (aclInited) + aclFinalize(); + return rc; +}