bug(vpto): non-zero pto.vci / llvm.hivm.vci base ignored on A5 camodel

[peanutchan]

Summary

On Ascend950PR camodel (Bisheng + llvm.hivm.vci), a non-zero scalar base passed to pto.vci is ignored. The hardware/backend always behaves as vci(0) and emits lane ids [0..N), even when LLVM IR correctly carries the base immediate.

This breaks absolute index generation such as vci(i * 64) for multi-vreg expert indexing (e.g. TileLang topk_gate).

Spec expectation

For ASC order:

result[i] = base + i

So pto.vci %c64 on !pto.vreg<64xi32> should yield [64, 65, ..., 127].

Observed

Camodel returns [0, 1, ..., 63].

End-to-end symptom in TileLang topk_gate (E=384): top-1 expert 144 was reported as 16 (144 % 64), and 292 as 36, until indices were materialized as vci(0) + vadds(base).

Minimal repro (legacy VPTO)

// kernel.pto — expected [64..127], observed [0..63]
module attributes {pto.target_arch = "a5", pto.backend = "vpto"} {
  func.func @vci_nonzero_base(%arg0: !pto.ptr<i32, gm>)
      attributes {pto.entry, pto.kernel_kind = #pto.kernel_kind<vector>} {
    %c0_i64 = arith.constant 0 : i64
    %ub = pto.castptr %c0_i64 : i64 -> !pto.ptr<i32, ub>
    %c0 = arith.constant 0 : index
    %c64_i32 = arith.constant 64 : i32
    pto.vecscope {
      %mask = pto.pset_b32 "PAT_ALL" : !pto.mask<b32>
      %idx = pto.vci %c64_i32 {order = "ASC"} : i32 -> !pto.vreg<64xi32>
      pto.vsts %idx, %ub[%c0], %mask
        : !pto.vreg<64xi32>, !pto.ptr<i32, ub>, !pto.mask<b32>
    }
    pto.set_flag["PIPE_V", "PIPE_MTE3", "EVENT_ID0"]
    pto.wait_flag["PIPE_V", "PIPE_MTE3", "EVENT_ID0"]
    %c1 = arith.constant 1 : i64
    %c64 = arith.constant 64 : i64
    %c0b = arith.constant 0 : i64
    pto.set_loop_size_ubtoout %c1, %c1 : i64, i64
    pto.copy_ubuf_to_gm %ub, %arg0, %c0b, %c1, %c64, %c0b, %c64, %c64
      : !pto.ptr<i32, ub>, !pto.ptr<i32, gm>, i64, i64, i64, i64, i64, i64
    pto.barrier #pto.pipe<PIPE_ALL>
    return
  }
}

PTOAS still emits the base into LLVM correctly:

declare <64 x i32> @llvm.hivm.vci.v64s32(i32, i32)
%7 = call <64 x i32> @llvm.hivm.vci.v64s32(i32 64, i32 0)

So this looks like a Bisheng / HIVM intrinsic / camodel bug (or undocumented encoding) rather than PTOAS dropping the operand in IR.

VMI form

pto.vmi.vci %c64 (size 64, contiguous) currently lowers through contiguous iota to a single pto.vci %c64, which hits the same backend bug.

Workaround

Materialize absolute indices as:

lanes = vci(0)
result = vadds(lanes, base)   # or ASC: vci(0)+vadds; DESC: vdup(base)-vci(0)

TileLang ASC already used vci(0) + vadds(i*64) for this reason. A local PTOAS workaround rewrites contiguous VMI iota the same way; ideally the backend honors the intrinsic base so that rewrite can be removed.

Environment

• Soc: Ascend950PR_9599 camodel via msprof op simulator
• PTOAS built against llvm-vpto (HIVM) + CANN 9.0.0
• Reproduced via TileLang examples/ascend/vmi/topk_gate before the vci(0)+vadds rewrite

Ask

1. Confirm whether llvm.hivm.vci.*(base, order) is supposed to honor a non-zero base.
2. If yes, fix Bisheng/HIVM lowering / camodel so vci(64) yields [64..127].
3. If the intrinsic is lane-only by design, please document that and keep the vci(0)+vadds(base) materialization as the supported pattern (and update ISA docs / gym cases that pass non-zero bases).

[peanutchan]

Root cause update (not Bisheng/camodel): PTOAS default/9.0.0-beta.1 A5 emission used llvm.hivm.vci.v64s32 (signedness-aware sN fragment). AscendC/CCE and CANN 9.0.0 emit the real HIVM overload llvm.hivm.vci.v64i32. hivmc-a5 has v64i32 and not v64s32; with the wrong name, camodel ignored a non-zero base even though LLVM IR carried i32 64.

Fix PR: will link once created — branch peanutchan:fix/vci-nonzero-base-hivm.

Camodel repro after fix: test/vpto/cases/micro-op/dsa-sfu/vci (non-zero bases at 64/128/…) compare passed.

[peanutchan]

Root cause (PTOAS emission, not Bisheng): default/9.0.0-beta.1 A5 lowering emitted llvm.hivm.vci.v64s32. AscendC/CCE and CANN 9.0.0 use llvm.hivm.vci.v64i32. hivmc-a5 recognizes v64i32 only; the wrong name made camodel ignore a non-zero base (IR still had i32 64).

Fix: #1093

Evidence after fix:

• LLVM IR: call @llvm.hivm.vci.v64i32(i32 64, i32 0)
• Camodel: test/vpto/cases/micro-op/dsa-sfu/vci compare passed (arange with bases 0/64/128/…)

[peanutchan]

Follow-up from TileLang topk_gate (AABB): #1093 signless v64i32 is necessary but not sufficient for unrolled constant bases.

Path	LLVM call	Camodel
Gym micro-op/dsa-sfu/vci (loop IV → sreg)	vci.v64i32(i32 %reg, i32 0)	PASS
TileLang vci(i*64) immediates after #1093	vci.v64i32(i32 64, i32 0)	still drops base (144→16); EXU shows VCI×6
TileLang + VMI vci(0)+vadds materialize	VCI×1 + VADDS×5	PASS

Tried volatile alloca→load to force a register operand: IR looks right (load volatile + vci(%reg)), but bisheng VF outline fails (Could not outline function!).

So for TileLang/unrolled constants we still need either (a) keep vci(0)+vadds, or (b) a non-stack way to materialize imm→sreg that hivmc accepts inside VF.

[peanutchan]

E2E gap / toolchain context (please don’t treat gym-only as done)

Prior fix #1093 was validated mainly on:

• Lit IR: llvm.hivm.vci.v64i32
• Gym camodel: test/vpto/cases/micro-op/dsa-sfu/vci under cann_9rel/cann-9.0.0 — this path uses register bases (vci(%reg) from loop IV) and PASSes

It was not e2e-validated on the TileLang consumer that originally filed the bug:

# Host has multiple CANNs; perf docs often use 9.1b3:
#   /usr/local/Ascend/cann_91b3/cann-9.1.0-beta.3   (camodel dav_3510)
#   /usr/local/Ascend/cann_9rel/cann-9.0.0            (what gym/topk ran on)
#   /usr/local/Ascend/cann_90b1/cann-9.0.0-beta.1
#   /usr/local/Ascend/cann_9b2/cann-9.0.0-beta.2

cd tilelang-deepseek/examples/ascend/vmi/topk_gate
source $ASCEND_HOME_PATH/set_env.sh
export PTOAS_ROOT=…/PTOAS-main
export PTOAS_KEEP_MSPROF_STAGING=1
export TILELANG_DISABLE_CACHE=1
# sync rebuilt libPTOASCompiler.so + _core.so into the venv site-packages
./run.sh sim --kernel vmi --schedule aabb --tokens 4 --experts 384 --topk 9 --token-tile 4
# Expect expert 144 (not 16). Symptom 144→16 means base dropped (144%64).

Findings after applying #1093 on TileLang

Case	LLVM	Camodel
Gym (sreg base)	vci.v64i32(%reg, 0)	PASS on 9.0.0
TileLang unrolled vci(i*64) imm	vci.v64i32(i32 64, 0)	FAIL 144→16 even with v64i32; EXU VCI×6
Volatile alloca→load to force reg	IR OK	bisheng Could not outline function!
VMI vci(0)+vadds workaround	VCI×1+VADDS×5	PASS

Assumption to verify: immediate-base encoding may still be broken (or version-dependent) in hivmc/camodel; register-base works. Fix may need imm→sreg materialization that is VF-outline-safe, and/or validation on cann_91b3 (9.1.0-beta.3) which is what our dual-issue / hand-CCE rooftop hosts use.

Please treat TileLang topk_gate VMI aabb N=4 E=384 K=9 as the e2e gate, not gym alone.

[peanutchan]

E2E fix landed on #1093 (second commit): TileLang unrolled imm vci(i*64) now PASSes camodel on both cann_9rel 9.0.0 and cann_91b3 9.1.0-beta.3 with VCI×6 (no vadds workaround). Root cause of remaining imm drop after signless rename: HIVM llvm.hivm.vci.*(i32 <imm>) encodes base 0; outline-safe fix is alloca outside aivector carrier + volatile load → sreg operand.