triton.language.argmax

triton.language.argmax = <function argmax>

Returns the maximum index of all elements in the input tensor along the provided axis

参数:

• input (Tensor) -- the input values

• axis (int) -- the dimension along which the reduction should be done. If None, reduce all dimensions

• keep_dims (bool) -- if true, keep the reduced dimensions with length 1

• tie_break_left (bool) -- if true, in case of a tie (i.e., multiple elements have the same maximum index value), return the left-most index for values that aren't NaN

This function can also be called as a member function on tensor, as x.argmax(...) instead of argmax(x, ...). .. rubric:: Example

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


@triton.jit
def triton_argmax_2d(in_ptr0, out_ptr0, dim: tl.constexpr, M: tl.constexpr, N: tl.constexpr, MNUMEL: tl.constexpr,
                     NNUMEL: tl.constexpr):
    mblk_idx = tl.arange(0, MNUMEL)
    nblk_idx = tl.arange(0, NNUMEL)
    mmask = mblk_idx < M
    nmask = nblk_idx < N
    mask = (mmask[:, None]) & (nmask[None, :])
    idx = mblk_idx[:, None] * N + nblk_idx[None, :]
    x = tl.load(in_ptr0 + idx, mask=mask, other=-float('inf'))
    tmp4 = tl.argmax(x, dim)
    if dim == 0:
        tl.store(out_ptr0 + tl.arange(0, N), tmp4, None)
    else:
        tl.store(out_ptr0 + tl.arange(0, M), tmp4, None)


def test_argmax():
    M, N, dim = 4, 8, 1
    x = torch.randn(M, N, dtype=torch.float32).npu()
    out = torch.empty(M, dtype=torch.int32).npu()
    triton_argmax_2d[1, 1, 1](x, out, dim, M, N, M, N)
    ref = torch.argmax(x, dim=dim).to(torch.int32)
    assert torch.equal(out.cpu(), ref.cpu()), "argmax result mismatch"
    print("argmax result:", out)


if __name__ == "__main__":
    test_argmax()

Special Restrictions

• DataType: Ascend A2/A3 does not support fp64, uint16, uint32, uint64; Ascend 950 does not fp64, fp8e4(E4M3), fp8e5(E5M2) (hardware limitation).

• keep_dims=True requires more test coverage; currently verified for 3D tensor with dim=2.