abs / neg

abs and neg operate independently on local register values. Their opcode selects the numeric type and applicable qualifiers. Operand order is dst,src; each scalar tensor operand ends in (1,) after the group and repeat axes. An ordinary dtype view does not change the numeric type of the requested operation.

Example: an explicit two-instruction scalar expression

from tvm.script import tirx as T


@T.function
def negative_magnitude(A: T.Tensor((4, 1), "float32"), Out: T.Tensor((4, 1), "float32")):
    T.device_entry(launch=T.cuda.LaunchConfig(grid=1, block=1))
    r = T.alloc_local((4, 1), "float32", layout="default")
    T.cuda.tile.apply(T.ptx["ld.global.f32"], r, A)
    T.cuda.tile.apply(T.ptx["abs.f32"], r, r)
    T.cuda.tile.apply(T.ptx["neg.f32"], r, r)
    T.cuda.tile.apply(T.ptx["st.global.f32"], Out, r)

For finite nonzero inputs this computes -abs(A). The tile expansion has four abs calls followed by four neg calls. This order comes from the two explicit operations, not from a hidden unary-op combination. Later compiler optimizations may simplify the expression. For a single caller-defined scalar expression with no exact-instruction requirement, use ewise: caller-defined scalar compositions.

NaN, signed-zero, and integer boundary behavior belongs to the selected PTX type/form, not to a generic Python abs reference. Destination storage may be gapped but must be distinct across repeats. Register ownership cannot change between source and destination. Predication follows mov.

Implementation: python/tvm/backend/cuda/tile/patterns/register.py. No temporary memory, numeric cast, barrier, or packed form is inserted.