Metadata#

class cutlass.operators.metadata.base.OperatorMetadata( operator_name: str, operator_class: type[Operator], supported_targets: list[TargetSm], operands: OperandsMetadata, design: DesignMetadata | None = None, epilogue: EpilogueMetadata | None = None, )#

Bases: object

OperatorMetadata exposes a uniform schema to inspect the properties of an Operator.

It allows users to programmatically and uniformly inspect properties of an Operator without manually inspecting its source code, and without learning a different schema of properties for each Operator.

It describes fixed properties of the Operator in three categories:

OperandsMetadata: properties of the operands passed to the Operator and the constraints it places on them. It describes the conditions for an Operator to functionally support (compute the correct result of) an operation.
DesignMetadata: architectural, performance, or design characteristics of the kernel. Useful for inspecting and comparing otherwise functionally equivalent Operators.
EpilogueMetadata: any custom epilogue computation that was fused into the Operator.

These properties are declared by the Operator itself when it is instantiated, and are fixed after that.

They may be determined by design, or the Operator’s implementation may be parametrized to generate instances supporting different values of a metadata field.

operator_name: str#: A unique human-readable Operator name.

Important

There are no backward-compatibility guarantees for these names, even for patch releases. Users are strongly advised to not use these names for filtering, and use metadata fields instead.

operator_class: type[Operator]#: The derived class of cutlass.operators.Operator that implements interfaces to expose a DSL kernel

supported_targets: list[TargetSm]#

Supported targets that this Operator can be compiled for

Only the earliest portable target in each architecture family is reported, as support for future targets is implied.

Examples

supported_targets=[TargetSm("100f"), TargetSm("120f")] implies support for ["100a", "103a", ... , "120a", "121a", ...] for all members of 100f & 120f arch families
supported_targets=[TargetSm("80")] implies support for ["80","89","90","90a","100", ...]

To check if this Operator is supported on a desired target_sm, use target_sm.supports_operators_from(supported_targets).

operands: OperandsMetadata#

per-operand constraints and any operand-independent choices (e.g. the accumulator type for a GEMM).

Type:: Describes the per-operand requirements of the Operator

design: DesignMetadata | None = None#: Describes characteristics of the design & implementation of the Operator

epilogue: EpilogueMetadata | None = None#: Describes the custom epilogue of the Operator.

property provider: Provider#: The Provider that this Operator’s class is registered with.

supports( args: RuntimeArguments | None, *, target_sm: TargetSm | str | None = None, ) → Status#

Check if the provided args/target_sm satisfy the properties described by this OperatorMetadata.

Parameters:

args (RuntimeArguments | None) – Check if Operator supports these RuntimeArguments. Can be explicitly set to None, in order to check only target_sm.
target_sm (TargetSm | str | None) – Check if Operator can be compiled/ran on given TargetSm architecture, if any.

Returns:

Status indicating whether the args/target_sm are supported, and reason for failure if not.

Return type:

Status

Raises:

ValueError – If both args and target_sm are None.

Operands metadata#

class cutlass.operators.metadata.operands.OperandsMetadata#

Bases: ABC

Describes properties expected of the logical operands passed to an Operator.

It describes the conditions under which an Operator can correctly execute a given RuntimeArguments. It mirrors the fields of a RuntimeArguments subclass implemented by the Operator, and describes requirements placed on each field. For instance, GemmOperandsMetadata mirrors GemmArguments.

RuntimeArguments describes what operation the user wishes to perform, independent of which Operator will perform it.

OperandsMetadata describes conditions under which an Operator can accept and correctly execute RuntimeArguments passed to it.

If an Operator’s OperandsMetadata supports the RuntimeArguments presented by a user, then it produces functionally correct result for those RuntimeArguments.

abstract supports(args: RuntimeArguments) → Status#

Check whether args satisfy the operand requirements described by this metadata.

Parameters:: args (RuntimeArguments) – Runtime arguments to validate.
Returns:: Status indicating whether args satisfy the operand requirements described by this metadata.
Return type:: Status

class cutlass.operators.metadata.operands.GemmOperandsMetadata( A: OperandConstraints, B: OperandConstraints, out: OperandConstraints, accumulator_type: Numeric, )#

Bases: OperandsMetadata

Describes conditions under which an Operation functionally supports GemmArguments.

A: OperandConstraints#: Constraints on operand A of the GEMM

B: OperandConstraints#: Constraints on operand B of the GEMM

out: OperandConstraints#: Constraints on operand out of the GEMM

accumulator_type: Numeric#: The required data type of the accumulator tensor

supports( other: GemmArguments | GemmOperandsMetadata, ) → Status#

Check whether other arguments or metadata are compatible with GEMM operands described by this metadata.

Parameters:: other (GemmArguments | GemmOperandsMetadata) – Either runtime GemmArguments or a peer GemmOperandsMetadata to validate.
Returns:: Status.success() if every operand and the accumulator type match, otherwise Status.fail(...) with the first failing operand’s reason.
Return type:: Status

class cutlass.operators.metadata.operands.GroupedGemmOperandsMetadata( A: OperandConstraints, B: OperandConstraints, out: OperandConstraints, offsets: OperandConstraints, accumulator_type: Numeric, )#

Bases: OperandsMetadata

Describes conditions under which an Operation functionally supports GroupedGemmArguments.

A: OperandConstraints#: Constraints on operand A of the grouped GEMM

B: OperandConstraints#: Constraints on operand B of the grouped GEMM

out: OperandConstraints#: Constraints on operand out of the grouped GEMM

offsets: OperandConstraints#: Constraints on operand offsets of the grouped GEMM.

accumulator_type: Numeric#: The required data type of the accumulator tensor

supports( other: GroupedGemmArguments | GroupedGemmOperandsMetadata, ) → Status#

Check whether other arguments or metadata are compatible with grouped-GEMM operands described by this metadata.

Parameters:: other (GroupedGemmArguments | GroupedGemmOperandsMetadata) – Either runtime GroupedGemmArguments or a peer GroupedGemmOperandsMetadata to validate.
Returns:: Status.success() if every operand, offsets, and the accumulator type match, otherwise Status.fail(...) with the first failure.
Return type:: Status

class cutlass.operators.metadata.operand_constraints.OperandConstraints#

Bases: ABC

Abstract base class describing the constraints an Operator places on an Operand.

An instance of this class describes the requirements an Operator places on one Operand, such as its data type, stride pattern, divisibility, and pointer alignment.

abstract supports( operand: Operand | OperandConstraints, ) → Status#

Check whether operand satisfies these constraints.

Parameters:: operand (Operand | OperandConstraints) – The runtime operand or another constraints object to check against self.
Returns:: Status.success() when operand is accepted, otherwise Status.fail(...) with a human-readable reason.
Return type:: Status

class cutlass.operators.metadata.operand_constraints.DenseTensorConstraints( dtype: Numeric, stride: tuple[int, ...], divisibility: int, ptr_alignment_bytes: int | None = None, )#

Bases: OperandConstraints

Constraints on a single DenseTensor.

Describes the data type, stride pattern, divisibility, and pointer alignment that an Operator requires a DenseTensor to satisfy.

dtype: Numeric#: The data type of the tensor

stride: tuple[int, ...] | None#: The stride of the tensor. If none, any arbitrary stride is allowed.

divisibility: int#: The divisibility requirement on a tensor’s stride and shape elements

ptr_alignment_bytes: int#: The alignment of the tensor’s data pointer, in bytes By default, it matches the number of bytes in stride/shape alignment.

supports( operand: DenseTensor | DenseTensorConstraints, ) → Status#

Check whether operand satisfies the constraints described by this object.

Parameters:: operand (Operand | OperandConstraints) – Operand or peer constraints object to validate.
Returns:: Status.success() if accepted, Status.fail(...) otherwise.
Return type:: Status

class cutlass.operators.metadata.operand_constraints.ScaledOperandConstraints( quantized: DenseTensorConstraints, scale: DenseTensorConstraints, mode: ScaleMode | tuple[int, ...], swizzle: ScaleSwizzleMode, )#

Bases: OperandConstraints

Constraints on a scaled (quantized + scale-factor) operand.

A scaled operand combines a quantized tensor and a scale tensor with a chosen scale mode + scale swizzle. Constraints on each piece are carried in quantized and scale; supports() checks all four.

quantized: DenseTensorConstraints#: Constraints on the narrow-precision tensor holding the quantized values.

scale: DenseTensorConstraints#: Constraints on the scale-factor tensor.

mode: ScaleMode | tuple[int, ...]#: Descriptor of the elements over which each scale factor operates

swizzle: ScaleSwizzleMode#: Swizzling pattern used within scale factor tensor

supports( operand: ScaledOperand | ScaledOperandConstraints, ) → Status#: Check whether operand satisfies these scaled-operand constraints.

Design metadata#

class cutlass.operators.metadata.design.DesignMetadata#

Bases: ABC

Describes architectural, performance, or design characteristics of the kernel separate from its functional contribution.

These properties do not directly alter the logical result of the Operator. Instead, they distinguish the design characteristics of otherwise logically-equivalent Operators, and can be used to guide Operator selection and performance characterization.

This could include properties like:

inherent design choices like which MMA instruction it uses, or use of other architectural features
its performance options, tile scheduling algorithm, etc.
its tile shape, cluster shape, number of shared memory stages, etc.

abstract supports(args: RuntimeArguments) → Status#

Checks whether the provided args satisfy the properties described by this DesignMetadata.

It validates that PerformanceControls within args.performance are supported by the Operator with this DesignMetadata.

DesignMetadata does not directly influence the logical result of the Operator, and thus does not affect what logical operands can be supported in RuntimeArguments.

Parameters:: args (RuntimeArguments) – The arguments to check support for.
Returns:: Status indicating whether the provided args satisfy the properties described by the design in this metadata
Return type:: Status

class cutlass.operators.metadata.design.BLASDesignMetadata( *, mma_instruction_type: type[MmaInstruction], tile_shape: tuple[int, ...], cluster_shape: tuple[int, ...], num_smem_stages: int | None = None, )#

Bases: DesignMetadata

Design metadata for a BLAS (basic-linear algebra subprogram) operation, like GEMM.

It describes fields common to many kernels that perform BLAS-like operations, such as tile/cluster shape, MMA instruction type, etc.

mma_instruction_type: type[MmaInstruction]#: The type of the MMA instruction used by the Operator, if any.

tile_shape: tuple[int, ...]#: Shape of the threadblock-level tile

cluster_shape: tuple[int, ...]#: Shape of the cluster used in terms of threadblocks

num_smem_stages: int | None = None#: Number of shared memory buffers used to pipeline data loading and computation. When None, the Operator determines this internally.

supports(args: RuntimeArguments) → Status#

Checks whether the provided args satisfy the properties described by the design in this metadata.

Parameters:: args (RuntimeArguments) – The arguments to check support for
Returns:: Status indicating whether the provided args satisfy the properties described by the design in this metadata
Return type:: Status

class cutlass.operators.metadata.design.Sm100DesignMetadata( use_2cta_mma: bool, use_tma_store: bool, tile_scheduler: TileSchedulerMetadata | None = None, fallback_cluster_shape: tuple[int, ...] | None = None, *, mma_instruction_type: type[MmaInstruction], tile_shape: tuple[int, ...], cluster_shape: tuple[int, ...], num_smem_stages: int | None = None, )#

Bases: BLASDesignMetadata

Design metadata for operators in the SM100 architecture family.

use_2cta_mma: bool#: Whether to use a 2CTA MMA instruction

use_tma_store: bool#: Whether to use TMA to store the results of the operation

tile_scheduler: TileSchedulerMetadata | None = None#: Describes the tile scheduling strategy used by the Operator

fallback_cluster_shape: tuple[int, ...] | None = None#

Fallback cluster shape for preferred cluster mode.

When None (default), the Operator uses a single fixed cluster shape (BLASDesignMetadata.cluster_shape). When set:

the Operator uses two cluster shapes: cluster_shape as the preferred (larger) shape and fallback_cluster_shape as the smaller fallback for tiles that cannot fill a full preferred cluster.
Preferred cluster depth (Z dimension) must be the same as that of fallback cluster.
Fallback cluster shape must evenly divide the preferred cluster shape.

use_fallback_cluster() → bool#: Return True when this design uses a preferred (larger) and fallback (smaller) cluster shape.

class cutlass.operators.metadata.design.TileSchedulerMetadata#

Bases: ABC

Models an abstract tile scheduler.

abstract classmethod supported_targets( design: DesignMetadata | None = None, operands: OperandsMetadata | None = None, ) → list[TargetSm]#

Returns all the supported targets for the tile scheduler.

Considers optional design and operands metadata when determining which targets are supported.

class cutlass.operators.metadata.design.CLCDynamicPersistentTileSchedulerMetadata#

Bases: TileSchedulerMetadata

Metadata for the CLC dynamic persistent tile scheduler.

classmethod supported_targets( design: DesignMetadata | None = None, operands: OperandsMetadata | None = None, ) → list[TargetSm]#: Returns supported targets for CLC dynamic persistent tile scheduler.

class cutlass.operators.metadata.design.StaticPersistentTileSchedulerMetadata#

Bases: TileSchedulerMetadata

Metadata for the static persistent tile scheduler.

classmethod supported_targets( design: DesignMetadata | None = None, operands: OperandsMetadata | None = None, ) → list[TargetSm]#: Returns supported targets for static persistent tile scheduler.

Epilogue metadata#

class cutlass.operators.metadata.epilogue.EpilogueMetadata(epilogue_args: EpilogueArguments)#

Bases: object

Describes properties related to custom EpilogueArguments that an Operator may support.

classmethod from_args( args: EpilogueArguments, ) → EpilogueMetadata#: Creates an EpilogueMetadata from an EpilogueArguments.

property parameters: list[cute.Tensor | Numeric]#: Returns the list of input and output parameters of the epilogue.

property parameter_names: list[str]#: Returns the list of names of the input and output parameters of the epilogue.

supports(args: RuntimeArguments) → Status#: Checks if the provided EpilogueArguments are supported by the Operator.

MMA instructions#

class cutlass.operators.mma.MmaInstruction#

Bases: object

Models properties of an abstract MMA instruction.

static supported_targets( design: DesignMetadata | None = None, operands: OperandsMetadata | None = None, ) → list[TargetSm]#: Return all the supported targets for the MMA instruction, given optional design and operands metadata.

static shape_k(operands: OperandsMetadata) → int#: Return the K dimension of the shape of matrix multiplication performed by this instruction.

class cutlass.operators.mma.BlackwellTcgen05Mma#

Bases: MmaInstruction

Models the tcgen05.mma instruction for the Blackwell architecture.

static supported_targets( design: DesignMetadata | None = None, operands: OperandsMetadata | None = None, ) → list[TargetSm]#: See MmaInstruction.supported_targets().

static shape_k(operands: OperandsMetadata) → int#: See MmaInstruction.shape_k().

class cutlass.operators.mma.BlackwellTcgen05MmaSparse#

Bases: MmaInstruction

Models the tcgen05.mma.sp instruction for the Blackwell architecture.

static supported_targets( design: DesignMetadata | None = None, operands: OperandsMetadata | None = None, ) → list[TargetSm]#: See MmaInstruction.supported_targets().

class cutlass.operators.mma.HopperWgmma#

Bases: MmaInstruction

Models the wgmma Tensor Core MMA instruction for the Hopper architecture.

static supported_targets( design: DesignMetadata | None = None, operands: OperandsMetadata | None = None, ) → list[TargetSm]#: See MmaInstruction.supported_targets().

class cutlass.operators.mma.AmpereMma#

Bases: MmaInstruction

Models the Tensor Core MMA instruction for the Ampere architecture.

static supported_targets( design: DesignMetadata | None = None, operands: OperandsMetadata | None = None, ) → list[TargetSm]#: See MmaInstruction.supported_targets().