cvt

cvt#

cvt performs a numeric conversion selected by the opcode. A dtype view only reinterprets bits; it is used here to expose the native carrier layout, not to perform the conversion itself. Destination and source widths, operand group lengths, rounding, saturation and packing come from the PTX form.

Scalar conversion uses tensor suffix (1,) on both sides. Packed cvt.rn.f16x2.f32 has three operands: one b32 destination, a float32 source for its high halfword, and a float32 source for its low halfword. All three have suffix (1,). The pair is not an inferred vectorization of a scalar cast.

Example: eight float32 values to four packed half pairs#

from tvm.script import tirx as T


@T.function
def narrow_pairs(A: T.Tensor((4, 2), "float32"), Out: T.Tensor((4, 1), "uint32")):
    T.device_entry(launch=T.cuda.LaunchConfig(grid=1, block=1))
    a = T.alloc_local((4, 2), "float32", layout="default")
    d = T.alloc_local((4, 1), "uint32", layout="default")
    T.cuda.tile.apply(T.ptx["ld.global.v2.f32"], a, A)
    T.cuda.tile.apply(T.ptx["cvt.rn.f16x2.f32"], d, a[:, 1:2], a[:, 0:1])
    T.cuda.tile.apply(T.ptx["st.global.b32"], Out, d)

For each input row [x0,x1], the output word contains converted x0 in its low half and x1 in its high half. Reversing the two source operands would exchange the output elements. The expansion emits four packed cvt instructions, one per row.

import numpy as np
import tvm

x = np.linspace(-2, 2, 8, dtype="float32").reshape(4, 2)
compiled = tvm.tirx.compile(narrow_pairs, target="cuda", pipeline="tirx")
out = tvm.runtime.tensor(np.zeros((4, 1), dtype="uint32"), tvm.cuda())
compiled(tvm.runtime.tensor(x, tvm.cuda()), out)
np.testing.assert_array_equal(out.numpy().view("float16"), x.astype("float16"))

The inverse operation is two scalar cvt.f32.f16 instructions per packed word; there is no corresponding cvt.f32x2.f16x2 spelling in this API. Expose the packed source with a uint16 view, add a singleton register-group axis per half, and apply the scalar opcode. For a composite conversion expression, including explicit intermediate rounding, use ewise: caller-defined scalar compositions.

Register ownership must agree on both sides even when their storage widths differ. A wrong carrier width, missing group axis, or overlapping destination is rejected. There is no implicit staging through shared memory.

Implementation: python/tvm/backend/cuda/tile/patterns/register.py.

Contents