tilelang.language.dtypes¶

Attributes¶

`AnyDType`
`dtype_name`
`int_`
`__all__`

Classes¶

`dtype`	Abstract base class for generic types.
`bool`	Abstract base class for generic types.

Functions¶

`__dtype_call__`(self, *args[, is_size_var])
`__dtype_as_torch__`(self)	Convert TileLang dtype to PyTorch dtype.
`__dtype_new__`(cls, value)
`__dtype_bytes__`(self)	Return the number of bytes for this dtype.
`__dtype_is_float4_e2m1fn__`(self)	Packed 4-bit FP4 E2M1 (CUTLASS float_e2m1_t).
`__dtype_is_float4_e2m1_unpacked__`(self)	8-bit FP4 E2M1 unpacked shared-memory storage.
`__dtype_is_float4__`(self)	Whether this is any FP4 E2M1 logical variant (packed or unpacked storage).
`is_float4_e2m1fn`(value)	Whether value is the packed 4-bit FP4 E2M1 variant.
`is_float4_e2m1_unpacked`(value)	Whether value is the 8-bit FP4 E2M1 unpacked shared-memory variant.
`is_float4`(value)	Whether value is any FP4 E2M1 logical variant (packed or unpacked).
`is_f8f6f4_family`(value)	Whether value is a tcgen05 `kind::f8f6f4` / `mxf8f6f4` operand dtype.
`validate_gemm_ab_dtypes`(a_dtype, b_dtype, *[, ...])	Validate mixed A/B dtypes for GEMM tile ops.
`get_tvm_dtype`(value)

Module Contents¶

class tilelang.language.dtypes.dtype¶

Bases: Generic[_T]

Abstract base class for generic types.

A generic type is typically declared by inheriting from this class parameterized with one or more type variables. For example, a generic mapping type might be defined as:

class Mapping(Generic[KT, VT]):
    def __getitem__(self, key: KT) -> VT:
        ...
    # Etc.

This class can then be used as follows:

def lookup_name(mapping: Mapping[KT, VT], key: KT, default: VT) -> VT:
    try:
        return mapping[key]
    except KeyError:
        return default

property bits: int¶

Return type:: int

property bytes: int¶

Return type:: int

as_torch()¶

Return type:: torch.dtype

is_float4_e2m1fn()¶

Return type:: bool

is_float4_e2m1_unpacked()¶

Return type:: bool

is_float4()¶

Return type:: bool

tilelang.language.dtypes.AnyDType¶

tilelang.language.dtypes.dtype_name = 'uint16'¶

tilelang.language.dtypes.int_¶

tilelang.language.dtypes.__dtype_call__(self, *args, is_size_var=False)¶

Parameters:

self (dtype)
is_size_var (bool)

Return type:

tvm.tirx.Var

tilelang.language.dtypes.__dtype_as_torch__(self)¶

Convert TileLang dtype to PyTorch dtype.

Parameters:: self (dtype)
Return type:: torch.dtype

tilelang.language.dtypes.__dtype_new__(cls, value)¶

Parameters:: value (AnyDType)
Return type:: dtype

tilelang.language.dtypes.__dtype_bytes__(self)¶

Return the number of bytes for this dtype.

Parameters:: self (dtype)
Return type:: int

tilelang.language.dtypes.__dtype_is_float4_e2m1fn__(self)¶

Packed 4-bit FP4 E2M1 (CUTLASS float_e2m1_t).

Use for pure FP4 workloads on tcgen05 mxf4 / mxf4nvf4 (packed SMEM, half the footprint). Global tensors and packed SMEM roundtrips use this dtype. TMA: CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B.

Parameters:: self (dtype)
Return type:: bool

tilelang.language.dtypes.__dtype_is_float4_e2m1_unpacked__(self)¶

8-bit FP4 E2M1 unpacked shared-memory storage.

Only for tcgen05 kind::f8f6f4 and block-scaled kind::mxf8f6f4 mixed-precision paths. It models the SMEM-side unpacked layout, one FP4 value per byte slot, and pairs with TMA CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B. Not for mxf4 / mxf4nvf4 packed FP4 kernels or general SIMT copy/cast paths.

Parameters:: self (dtype)
Return type:: bool

tilelang.language.dtypes.__dtype_is_float4__(self)¶

Whether this is any FP4 E2M1 logical variant (packed or unpacked storage).

Parameters:: self (dtype)
Return type:: bool

tilelang.language.dtypes.is_float4_e2m1fn(value)¶

Whether value is the packed 4-bit FP4 E2M1 variant.

Parameters:: value (AnyDType)
Return type:: bool

tilelang.language.dtypes.is_float4_e2m1_unpacked(value)¶

Whether value is the 8-bit FP4 E2M1 unpacked shared-memory variant.

Parameters:: value (AnyDType)
Return type:: bool

tilelang.language.dtypes.is_float4(value)¶

Whether value is any FP4 E2M1 logical variant (packed or unpacked).

Parameters:: value (AnyDType)
Return type:: bool

tilelang.language.dtypes.is_f8f6f4_family(value)¶

Whether value is a tcgen05 kind::f8f6f4 / mxf8f6f4 operand dtype.

Parameters:: value (AnyDType)
Return type:: bool

tilelang.language.dtypes.validate_gemm_ab_dtypes(a_dtype, b_dtype, *, a_in_tmem=False, allow_f8f6f4_mixed=False)¶

Validate mixed A/B dtypes for GEMM tile ops.

TS variants (A in TMEM) skip validation because A/B dtypes may legitimately differ. Only TCGEN05 f8f6f4-family paths should opt in to mixed A/B dtypes.

Parameters:

a_dtype (AnyDType)
b_dtype (AnyDType)
a_in_tmem (bool)
allow_f8f6f4_mixed (bool)

Return type:

None

tilelang.language.dtypes.get_tvm_dtype(value)¶

Parameters:: value (AnyDType)
Return type:: dtype

class tilelang.language.dtypes.bool¶

Bases: dtype

Abstract base class for generic types.

A generic type is typically declared by inheriting from this class parameterized with one or more type variables. For example, a generic mapping type might be defined as:

class Mapping(Generic[KT, VT]):
    def __getitem__(self, key: KT) -> VT:
        ...
    # Etc.

This class can then be used as follows:

def lookup_name(mapping: Mapping[KT, VT], key: KT, default: VT) -> VT:
    try:
        return mapping[key]
    except KeyError:
        return default

tilelang.language.dtypes.__all__¶