[BUG][ascend] uint16 and uint32 `>>` are lowered as arithmetic shifts, silently corrupting every value with the top bit set

[cheersluvs]

Summary

On the Ascend backend, >> on a tl.uint16 or tl.uint32 value is lowered as
an arithmetic (sign-propagating) shift instead of a logical one. Every input
whose top bit is set comes back sign-extended rather than shifted. There is no
error and no warning -- the result is simply wrong.

tl.uint8 on the same card shifts logically, and every input whose top bit is
clear is correct at all three widths, so the unsigned shift is not missing
outright: it looks like the value is sign-extended into a wider register before
the shift, and the shift then propagates a bit that only exists because of that
extension.

Environment

triton	3.5.1 (/usr/local/lib/python3.11/site-packages/triton)
torch / torch_npu	2.8.0+cpu / 2.8.0.post4.dev20260409
CANN	9.0.0
device	Ascend910B4-1

Reproducer

Self-contained: torch, torch_npu and triton only.

import torch, torch_npu, triton
import triton.language as tl

B, SH = 128, 5

@triton.jit
def k_u8(x_ptr, out_ptr, SH: tl.constexpr, BLOCK: tl.constexpr):
    lane = tl.arange(0, BLOCK)
    bits = tl.load(x_ptr + lane).to(tl.uint8, bitcast=True)
    tl.store(out_ptr + lane, (bits >> SH).to(tl.int32))

@triton.jit
def k_u16(x_ptr, out_ptr, SH: tl.constexpr, BLOCK: tl.constexpr):
    lane = tl.arange(0, BLOCK)
    bits = tl.load(x_ptr + lane).to(tl.uint16, bitcast=True)
    tl.store(out_ptr + lane, (bits >> SH).to(tl.int32))

@triton.jit
def k_u32(x_ptr, out_ptr, SH: tl.constexpr, BLOCK: tl.constexpr):
    lane = tl.arange(0, BLOCK)
    bits = tl.load(x_ptr + lane).to(tl.uint32, bitcast=True)
    tl.store(out_ptr + lane, (bits >> SH).to(tl.int32))

cases = [("uint8 ", k_u8,  torch.int8,  0xFE,       8),
         ("uint16", k_u16, torch.int16, 0xFFFE,     16),
         ("uint32", k_u32, torch.int32, 0xFFFFFFFE, 32)]

for name, kernel, dt, raw, width in cases:
    x = torch.full((B,), -2, dtype=dt, device="npu")     # top bit set
    out = torch.zeros(B, dtype=torch.int32, device="npu")
    kernel[(1,)](x, out, SH=SH, BLOCK=B)
    torch.npu.synchronize()
    print(name, hex(raw), "got", int(out[0]),
          "| logical", raw >> SH, "| arithmetic", (raw - (1 << width)) >> SH)

Actual output

top bit SET -- this is where the two shifts differ
dtype          input          got      logical        arith   verdict
------------------------------------------------------------------------------
uint8           0xfe            7            7           -1   LOGICAL (correct)
uint16        0xfffe           -1         2047           -1   ARITHMETIC (defect)
uint32    0xfffffffe           -1    134217727           -1   ARITHMETIC (defect)

top bit CLEAR -- control, every line must read correct
dtype          input          got     expected   verdict
--------------------------------------------------------------
uint8           0x7e            3            3   correct
uint16        0x7ffe         1023         1023   correct
uint32    0x7ffffffe     67108863     67108863   correct

Expected: got == logical on all three lines of the first block.

Why it matters

The natural way to build a radix bucket out of a float is to bitcast to the
unsigned integer of the same width, flip the non-sign half so the ordering
becomes monotone, and shift. For fp16:

h        = x.to(tl.float16)
bits     = h.to(tl.uint16, bitcast=True)
sign_set = (bits & tl.full(bits.shape, 0x8000, tl.uint16)) != 0
inv      = (~bits) & tl.full(bits.shape, 0x7FFF, tl.uint16)
mapped   = tl.where(sign_set, bits, inv)
bin_idx  = (mapped >> 5).to(tl.uint32)      # <-- destroyed here

For a negative input mapped keeps the raw bits, whose top bit is set by
definition. Measured on this backend, with x = -1.5 (fp16 bits 0xbe00):

fp16 radix bucket of -1.5 (bits 0xbe00) : got -528, expected 1520

0xbe00 read as int16 is -16896, and -16896 >> 5 == -528 exactly. So every
negative value collapses out of its bucket while every positive one is fine --
a top-k selection that is exactly right on half its input and silently wrong on
the other half.

The same construction at 32 bits (bitcast fp32 to uint32, then bits >> 21)
is the standard first radix pass, and it has no natural workaround.

Workaround

Widen first, then mask off the sign extension before shifting:

bin_idx = (mapped.to(tl.int32) & 0xFFFF) >> 5

Verified on this stack: got 2047, expected 2047.

Note

The 32-bit case may be a regression rather than a long-standing defect: we have
a production kernel whose first radix pass is bits >> 21 on a genuine
tl.uint32 with no workaround, and it passes its correctness suite on a
triton 3.2.0 / CANN 8.5.0 Ascend stack -- which it could not do if the shift
were arithmetic there. We have not yet been able to run this reproducer on that
older stack to confirm; we will follow up in a comment when we have.

[cheersluvs]

Correction: not a regression -- the same defect is on the 3.2 line

The "Note" at the end of the issue speculated that the 32-bit case might be a
regression in the 3.5 line. That was wrong, and I am retracting it. I have
now run the same reproducer on the older stack and uint32 is arithmetic there
too:

	triton 3.2.0 / CANN 8.5.0	triton 3.5.1 / CANN 9.0.0
uint8 >>	LOGICAL (correct)	LOGICAL (correct)
uint16 >>	ARITHMETIC (defect)	ARITHMETIC (defect)
uint32 >>	ARITHMETIC (defect)	ARITHMETIC (defect)
top-bit-clear controls, all widths	correct	correct

### host bm-ctyun-wq-910b-64g-0-13
triton 3.2.0  at /usr/local/python3.11.13/lib/python3.11/site-packages/triton
torch  2.6.0+cpu | torch_npu 2.6.0rc1
device Ascend910B4-1

top bit SET -- this is where the two shifts differ
dtype          input          got      logical        arith   verdict
------------------------------------------------------------------------------
uint8           0xfe            7            7           -1   LOGICAL (correct)
uint16        0xfffe           -1         2047           -1   ARITHMETIC (defect)
uint32    0xfffffffe           -1    134217727           -1   ARITHMETIC (defect)

top bit CLEAR -- control, every line must read correct
dtype          input          got     expected   verdict
--------------------------------------------------------------
uint8           0x7e            3            3   correct
uint16        0x7ffe         1023         1023   correct
uint32    0x7ffffffe     67108863     67108863   correct

fp16 radix bucket of -1.5 (bits 0xbe00)  : got -528, expected 1520 -> WRONG

So this is long-standing rather than new, and it spans both the 3.2.0
(triton/backends/ascend) and 3.5.1 (triton/spec/ascend) lines. The
uint8-correct / top-bit-clear-correct pattern holds identically on both,
which keeps the sign-extension reading in the issue body intact.

The reasoning I used to guess "regression" is what failed, and it is worth
naming so nobody repeats it: our kernel's first radix pass is bits >> 21 on a
genuine tl.uint32 with no workaround, and its correctness suite passes on the
3.2.0 stack -- so I inferred the shift must be logical there. It is not. The
suite selects the top 1024 of 129280 randn values, which puts the threshold
near +2.5 sigma, and the mapping only leaves the top bit set for negative
inputs -- so the corrupted half is real but never consulted. A passing test
suite was evidence about the test data, not about the shift.

I am verifying that operator-level reading separately and will not state it as
fact until measured; it does not affect the reproducer or the table above,
which are direct measurements.

[cheersluvs]

What this costs in a real operator

Follow-up to the correction above, with the operator-level measurement I said I
would not state until it was measured.

We hit this in a top-k selection kernel. Its radix pass buckets a float by
bitcasting to the unsigned integer of the same width and flipping the non-sign
half so the ordering becomes monotone:

bits     = x.to(tl.uint32, bitcast=True)
sign_set = (bits & 0x80000000) != 0
inv      = (~bits) & 0x7FFFFFFF
mapped   = tl.where(sign_set, bits, inv)   # bit 31 set iff x was negative
bin_idx  = mapped >> 21                    # <-- arithmetic here

Bit 31 is set for exactly the negative inputs, so on this backend every
negative value is placed in the wrong bucket while every positive one is
correct. Measured on triton 3.2.0 / CANN 8.5.0, Ascend910B4-1, selecting the
top 1024 of 129280 values per row:

input	values in the threshold bucket	result
randn	198	PASS
concentrated, positive	98594	PASS
concentrated, negative	30765	FAIL, max abs error 3.97e-02

The two concentrated cases differ only in sign, which is the signature of this
defect rather than of anything else in the kernel.

Why this is worth fixing rather than documenting

The failure is invisible to an ordinary test suite. Our radix pass only reaches
the 32-bit step when a single coarse bucket holds more than a fixed number of
candidates; with randn inputs the threshold bucket holds 198, so the step
never executes and the suite passed for months. It takes both a concentrated
distribution and negative values to see it -- and a kernel that has those is
not exotic: low-temperature logits, large masked regions, and duplicate scores
all concentrate values into one bucket.

So the practical risk is not "an unusual bit-twiddling idiom is wrong". It is
that a correct-looking kernel silently returns wrong answers on a subset of
real inputs, on a backend where the same source is correct at 8 bits, correct
for positive values at every width, and correct everywhere else.

Workaround, for anyone who lands here first

Shift the low bits, where the shift is measurably correct, and add the top
bit's own contribution back:

sign_set = (bits & 0x80000000) != 0
low31    = bits & 0x7FFFFFFF
bin_idx  = (low31 >> 21) + tl.where(sign_set, 1 << 10, 0)   # 1 << (31 - 21)

Widening to a larger signed type and masking ((bits.to(tl.int64) & 0xFFFFFFFF) >> 21) also works, and is what we use at 16 bits, but it pulls in
64-bit arithmetic. Note that not every shift needs this: a shift whose result
is immediately masked ((bits >> 10) & 0x7FF) keeps only bits the sign fill
never reaches, and a (bits ^ pattern) >> n == 0 test compares against zero,
which both shifts agree on. Only a shift whose result is used directly as a
value is affected.

[zhongsanming]

I've dumped MLIR of all 3 kernels.

ttir

• u8

module attributes {hacc.target = #hacc.target<"Ascend910B4-1">} {
  tt.func public @k_u8(%x_ptr: !tt.ptr<i8> {tt.divisibility = 16 : i32} loc("x_ptr"(#loc)), %out_ptr: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr"(#loc))) attributes {noinline = false} {
    %cst = arith.constant dense<5> : tensor<128xi8> loc(#loc1)
    %lane = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> loc(#loc12)
    %bits = tt.splat %x_ptr : !tt.ptr<i8> -> tensor<128x!tt.ptr<i8>> loc(#loc13)
    %bits_0 = tt.addptr %bits, %lane : tensor<128x!tt.ptr<i8>>, tensor<128xi32> loc(#loc13)
    %bits_1 = tt.load %bits_0 : tensor<128x!tt.ptr<i8>> loc(#loc14)
    %bits_2 = tt.bitcast %bits_1 : tensor<128xi8> -> tensor<128xi8> loc(#loc15)
    %0 = tt.splat %out_ptr : !tt.ptr<i32> -> tensor<128x!tt.ptr<i32>> loc(#loc6)
    %1 = tt.addptr %0, %lane : tensor<128x!tt.ptr<i32>>, tensor<128xi32> loc(#loc6)
    %2 = arith.shrui %bits_2, %cst : tensor<128xi8> loc(#loc1)
    %3 = arith.extui %2 : tensor<128xi8> to tensor<128xi32> loc(#loc7)
    tt.store %1, %3 : tensor<128x!tt.ptr<i32>> loc(#loc8)
    tt.return loc(#loc9)
  } loc(#loc)
} loc(#loc)

• u16

module attributes {hacc.target = #hacc.target<"Ascend910B4-1">} {
  tt.func public @k_u16(%x_ptr: !tt.ptr<i16> {tt.divisibility = 16 : i32} loc("x_ptr"(#loc)), %out_ptr: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr"(#loc))) attributes {noinline = false} {
    %cst = arith.constant dense<5> : tensor<128xi16> loc(#loc1)
    %lane = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> loc(#loc12)
    %bits = tt.splat %x_ptr : !tt.ptr<i16> -> tensor<128x!tt.ptr<i16>> loc(#loc13)
    %bits_0 = tt.addptr %bits, %lane : tensor<128x!tt.ptr<i16>>, tensor<128xi32> loc(#loc13)
    %bits_1 = tt.load %bits_0 : tensor<128x!tt.ptr<i16>> loc(#loc14)
    %bits_2 = tt.bitcast %bits_1 : tensor<128xi16> -> tensor<128xi16> loc(#loc15)
    %0 = tt.splat %out_ptr : !tt.ptr<i32> -> tensor<128x!tt.ptr<i32>> loc(#loc6)
    %1 = tt.addptr %0, %lane : tensor<128x!tt.ptr<i32>>, tensor<128xi32> loc(#loc6)
    %2 = arith.shrui %bits_2, %cst : tensor<128xi16> loc(#loc1)
    %3 = arith.extui %2 : tensor<128xi16> to tensor<128xi32> loc(#loc7)
    tt.store %1, %3 : tensor<128x!tt.ptr<i32>> loc(#loc8)
    tt.return loc(#loc9)
  } loc(#loc)
} loc(#loc)

• u32

module attributes {hacc.target = #hacc.target<"Ascend910B4-1">} {
  tt.func public @k_u32(%x_ptr: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("x_ptr"(#loc)), %out_ptr: !tt.ptr<i32> {tt.divisibility = 16 : i32} loc("out_ptr"(#loc))) attributes {noinline = false} {
    %cst = arith.constant dense<5> : tensor<128xi32> loc(#loc1)
    %lane = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32> loc(#loc11)
    %bits = tt.splat %x_ptr : !tt.ptr<i32> -> tensor<128x!tt.ptr<i32>> loc(#loc12)
    %bits_0 = tt.addptr %bits, %lane : tensor<128x!tt.ptr<i32>>, tensor<128xi32> loc(#loc12)
    %bits_1 = tt.load %bits_0 : tensor<128x!tt.ptr<i32>> loc(#loc13)
    %bits_2 = tt.bitcast %bits_1 : tensor<128xi32> -> tensor<128xi32> loc(#loc14)
    %0 = tt.splat %out_ptr : !tt.ptr<i32> -> tensor<128x!tt.ptr<i32>> loc(#loc6)
    %1 = tt.addptr %0, %lane : tensor<128x!tt.ptr<i32>>, tensor<128xi32> loc(#loc6)
    %2 = arith.shrui %bits_2, %cst : tensor<128xi32> loc(#loc1)
    tt.store %1, %2 : tensor<128x!tt.ptr<i32>> loc(#loc7)
    tt.return loc(#loc8)
  } loc(#loc)
} loc(#loc)

ttadapter

• u8

module attributes {hacc.target = #hacc.target<"Ascend910B4-1">} {
  func.func @k_u8(%arg0: memref<?xi8> loc("/root/flagtree/./dump_ascend_shift_mlir.py":42:0), %arg1: memref<?xi8> loc("/root/flagtree/./dump_ascend_shift_mlir.py":42:0), %x_ptr: memref<?xi8> {tt.divisibility = 16 : i32, tt.tensor_kind = 0 : i32} loc("x_ptr"(#loc)), %out_ptr: memref<?xi32> {tt.divisibility = 16 : i32, tt.tensor_kind = 1 : i32} loc("out_ptr"(#loc)), %arg4: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":42:0), %arg5: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":42:0), %arg6: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":42:0), %arg7: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":42:0), %arg8: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":42:0), %arg9: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":42:0)) attributes {SyncBlockLockArgIdx = 0 : i64, WorkspaceArgIdx = 1 : i64, global_kernel = "local", mix_mode = "aiv", parallel_mode = "simd"} {
    %c5_i8 = arith.constant 5 : i8 loc(#loc1)
    %0 = tensor.empty() : tensor<128xi8> loc(#loc1)
    %1 = linalg.fill ins(%c5_i8 : i8) outs(%0 : tensor<128xi8>) -> tensor<128xi8> loc(#loc1)
    %bits = memref.reinterpret_cast %x_ptr to offset: [0], sizes: [128], strides: [1] : memref<?xi8> to memref<128xi8, strided<[1]>> loc(#loc9)
    %bits_0 = memref.alloc() : memref<128xi8> loc(#loc10)
    memref.copy %bits, %bits_0 : memref<128xi8, strided<[1]>> to memref<128xi8> loc(#loc10)
    %bits_1 = bufferization.to_tensor %bits_0 restrict writable : memref<128xi8> to tensor<128xi8> loc(#loc10)
    %reinterpret_cast = memref.reinterpret_cast %out_ptr to offset: [0], sizes: [128], strides: [1] : memref<?xi32> to memref<128xi32, strided<[1]>> loc(#loc4)
    %2 = arith.shrui %bits_1, %1 : tensor<128xi8> loc(#loc1)
    %3 = arith.extui %2 : tensor<128xi8> to tensor<128xi32> loc(#loc5)
    bufferization.materialize_in_destination %3 in writable %reinterpret_cast : (tensor<128xi32>, memref<128xi32, strided<[1]>>) -> () loc(#loc6)
    return loc(#loc)
  } loc(#loc)
} loc(#loc)

• u16

module attributes {hacc.target = #hacc.target<"Ascend910B4-1">} {
  func.func @k_u16(%arg0: memref<?xi8> loc("/root/flagtree/./dump_ascend_shift_mlir.py":49:0), %arg1: memref<?xi8> loc("/root/flagtree/./dump_ascend_shift_mlir.py":49:0), %x_ptr: memref<?xi16> {tt.divisibility = 16 : i32, tt.tensor_kind = 0 : i32} loc("x_ptr"(#loc)), %out_ptr: memref<?xi32> {tt.divisibility = 16 : i32, tt.tensor_kind = 1 : i32} loc("out_ptr"(#loc)), %arg4: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":49:0), %arg5: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":49:0), %arg6: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":49:0), %arg7: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":49:0), %arg8: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":49:0), %arg9: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":49:0)) attributes {SyncBlockLockArgIdx = 0 : i64, WorkspaceArgIdx = 1 : i64, global_kernel = "local", mix_mode = "aiv", parallel_mode = "simd"} {
    %c5_i16 = arith.constant 5 : i16 loc(#loc1)
    %0 = tensor.empty() : tensor<128xi16> loc(#loc1)
    %1 = linalg.fill ins(%c5_i16 : i16) outs(%0 : tensor<128xi16>) -> tensor<128xi16> loc(#loc1)
    %bits = memref.reinterpret_cast %x_ptr to offset: [0], sizes: [128], strides: [1] : memref<?xi16> to memref<128xi16, strided<[1]>> loc(#loc9)
    %bits_0 = memref.alloc() : memref<128xi16> loc(#loc10)
    memref.copy %bits, %bits_0 : memref<128xi16, strided<[1]>> to memref<128xi16> loc(#loc10)
    %bits_1 = bufferization.to_tensor %bits_0 restrict writable : memref<128xi16> to tensor<128xi16> loc(#loc10)
    %reinterpret_cast = memref.reinterpret_cast %out_ptr to offset: [0], sizes: [128], strides: [1] : memref<?xi32> to memref<128xi32, strided<[1]>> loc(#loc4)
    %2 = arith.shrui %bits_1, %1 : tensor<128xi16> loc(#loc1)
    %3 = arith.extui %2 : tensor<128xi16> to tensor<128xi32> loc(#loc5)
    bufferization.materialize_in_destination %3 in writable %reinterpret_cast : (tensor<128xi32>, memref<128xi32, strided<[1]>>) -> () loc(#loc6)
    return loc(#loc)
  } loc(#loc)
} loc(#loc)

• u32

module attributes {hacc.target = #hacc.target<"Ascend910B4-1">} {
  func.func @k_u32(%arg0: memref<?xi8> loc("/root/flagtree/./dump_ascend_shift_mlir.py":56:0), %arg1: memref<?xi8> loc("/root/flagtree/./dump_ascend_shift_mlir.py":56:0), %x_ptr: memref<?xi32> {tt.divisibility = 16 : i32, tt.tensor_kind = 0 : i32} loc("x_ptr"(#loc)), %out_ptr: memref<?xi32> {tt.divisibility = 16 : i32, tt.tensor_kind = 1 : i32} loc("out_ptr"(#loc)), %arg4: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":56:0), %arg5: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":56:0), %arg6: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":56:0), %arg7: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":56:0), %arg8: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":56:0), %arg9: i32 loc("/root/flagtree/./dump_ascend_shift_mlir.py":56:0)) attributes {SyncBlockLockArgIdx = 0 : i64, WorkspaceArgIdx = 1 : i64, global_kernel = "local", mix_mode = "aiv", parallel_mode = "simd"} {
    %c5_i32 = arith.constant 5 : i32 loc(#loc1)
    %0 = tensor.empty() : tensor<128xi32> loc(#loc1)
    %1 = linalg.fill ins(%c5_i32 : i32) outs(%0 : tensor<128xi32>) -> tensor<128xi32> loc(#loc1)
    %bits = memref.reinterpret_cast %x_ptr to offset: [0], sizes: [128], strides: [1] : memref<?xi32> to memref<128xi32, strided<[1]>> loc(#loc8)
    %bits_0 = memref.alloc() : memref<128xi32> loc(#loc9)
    memref.copy %bits, %bits_0 : memref<128xi32, strided<[1]>> to memref<128xi32> loc(#loc9)
    %bits_1 = bufferization.to_tensor %bits_0 restrict writable : memref<128xi32> to tensor<128xi32> loc(#loc9)
    %reinterpret_cast = memref.reinterpret_cast %out_ptr to offset: [0], sizes: [128], strides: [1] : memref<?xi32> to memref<128xi32, strided<[1]>> loc(#loc4)
    %2 = arith.shrui %bits_1, %1 : tensor<128xi32> loc(#loc1)
    bufferization.materialize_in_destination %2 in writable %reinterpret_cast : (tensor<128xi32>, memref<128xi32, strided<[1]>>) -> () loc(#loc5)
    return loc(#loc)
  } loc(#loc)
} loc(#loc)

All of these kernels correctly uses arith.shrui to perform right shift rather than arith.shrsi.

Following the bishengir-compile pipeline of u16 case, we see the following major changes to IR.

  %3 = arith.shrui %2, %1 : tensor<128xi16>
  %4 = arith.extui %3 : tensor<128xi16> to tensor<128xi32>

pass ConvertArithToHFusion convert it to

    %4 = hfusion.elemwise_binary {fun = #hfusion.binary_fn<shrui>} ins(%2, %1 : tensor<128xi16>, tensor<128xi16>) outs(%3 : tensor<128xi16>) -> tensor<128xi16>
    %6 = hfusion.cast {cast = #hfusion.type_fn<cast_signed>, round_mode = #hfusion.round_mode<rint>} ins(%4 : tensor<128xi16>) outs(%5 : tensor<128xi32>) -> tensor<128xi32>

then pass ConvertHFusionToHIVM further convert it to

    %2 = hivm.hir.vshr ins(%1, %c5_i16 : tensor<128xi16>, i16) outs(%0 : tensor<128xi16>) round : true -> tensor<128xi16>
    %4 = hivm.hir.vcast ins(%2 : tensor<128xi16>) outs(%3 : tensor<128xf32>) -> tensor<128xf32>
    %6 = hivm.hir.vcast ins(%4 : tensor<128xf32>) outs(%5 : tensor<128xi32>) round_mode = <trunc> -> tensor<128xi32>

finally pass ConvertHIVMToStandard convert it to

    call @vshrs_vs_1d_int16_t(%cast_2, %c5_i16, %cast_3) : (memref<?xi16, strided<[?], offset: ?>, #hivm.address_space<ub>>, i16, memref<?xi16, strided<[?], offset: ?>, #hivm.address_space<ub>>) -> ()
    call @vcast_int16_t_to_float_1d_with_mode(%cast_4, %cast_5, %c0_i32) : (memref<?xi16, strided<[?], offset: ?>, #hivm.address_space<ub>>, memref<?xf32, strided<[?], offset: ?>, #hivm.address_space<ub>>, i32) -> ()
    call @vcast_float_to_int32_t_1d_with_mode(%cast_6, %cast_7, %c4_i32) : (memref<?xf32, strided<[?], offset: ?>, #hivm.address_space<ub>>, memref<?xi32, strided<[?], offset: ?>, #hivm.address_space<ub>>, i32) -> ()

From AscendNPU-IR, we don't know exactly whether hivm.hir.vshr support logical shift or not. but from the generated IR perspective, there is no clear sign shows the hivm.hir.vshr is doing logical shift. We need to confirm this from huawei.

Why u8 works

The u8 case is different in pass ConvertHFusionToHIVM

    %2 = hivm.hir.vcast ins(%0 : tensor<128xi8>) outs(%1 : tensor<128xf16>) cast = <cast_unsigned> -> tensor<128xf16>
    %4 = hivm.hir.vcast ins(%2 : tensor<128xf16>) outs(%3 : tensor<128xi16>) round_mode = <trunc> -> tensor<128xi16>
    %5 = hivm.hir.vshr ins(%4, %c5_i16 : tensor<128xi16>, i16) outs(%3 : tensor<128xi16>) round : true -> tensor<128xi16>

The compiler first cast i8 to i16, then performs the shift. We don't have enough information here. but if we assume the cast from i8 to i16 padded the first 8 bit with 0, then it all make sense. Again we need to confirm this from huawei.