triton.language.max

triton.language.max = <function max>

Returns the maximum 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

• return_indices (bool) -- if true, return index corresponding to the maximum value

• return_indices_tie_break_left (bool) -- if true, in case of a tie (i.e., multiple elements have the same maximum 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.max(...) instead of max(x, ...). .. rubric:: Example

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


@triton.jit
def triton_max_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.max(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_max():
    M, N, dim = 4, 8, 1
    x = torch.randn(M, N, dtype=torch.float32).npu()
    out = torch.empty(M, dtype=torch.float32).npu()
    triton_max_2d[1, 1, 1](x, out, dim, M, N, M, N)
    ref = torch.max(x, dim=dim)[0]
    assert torch.allclose(out.cpu(), ref.cpu()), "max result mismatch"
    print("max result:", out)


if __name__ == "__main__":
    test_max()

Special Restrictions

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

• return_indices: return_indices=True is not supported when axis=None.

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