For the complete documentation index, see llms.txt. Markdown versions of all pages are available by appending .md to any URL (e.g. /max/get-started.md).

Mojo struct

DecodeKVConsumer

struct DecodeKVConsumer[dtype: DType, config: MLA_SM100_Decode_Config, num_producer: Int = Int(1)]

Consumer side of the decode KV pipeline that waits for and releases KV stages.

Parameters

dtype (DType): Element type of the KV tiles stored in SMEM.
config (MLA_SM100_Decode_Config): Decode config supplying num_kv_stages, BN_QK, and q_depth for pipeline stage sizing.
num_producer (Int): Number of producer threads arriving on each producer mbarrier (defaults to 1).

Fields

pipe (DecodeKVConsumer[dtype, config, num_producer].KVPipeType):
smem (Pointer[Scalar[dtype], MutAnyOrigin, address_space=AddressSpace.SHARED, _safe=False]):

Implemented traits

AnyType, Copyable, ImplicitlyCopyable, ImplicitlyDeletable, Movable, RegisterPassable, TrivialRegisterPassable

`comptime` members

`kv_stage_elems`

comptime kv_stage_elems = (config * config)

`KVPipeType`

comptime KVPipeType = KVPipelineGeneric[config.num_kv_stages, Int(1), num_producer, Int(2)]

Methods

`init`

def __init__(pipe: KVPipelineGeneric[config.num_kv_stages, Int(1), num_producer, Int(2)], smem: Pointer[Scalar[dtype], MutAnyOrigin, address_space=AddressSpace.SHARED, _safe=False]) -> Self

`stage_base_ptr`

def stage_base_ptr[*, qk_stage: Int = Int(0)](self) -> Pointer[Scalar[dtype], MutAnyOrigin, address_space=AddressSpace.SHARED, _safe=False]

Returns:

Pointer[Scalar[dtype], MutAnyOrigin, address_space=AddressSpace.SHARED, _safe=False]

`stage_index`

def stage_index[*, qk_stage: Int = Int(0)](self) -> UInt32

Returns:

UInt32

`wait`

def wait[*, qk_stage: Int = Int(0)](self)

`release`

def release[*, qk_stage: Int = Int(0)](mut self, e: Int32)

Parameters​

Fields​

Implemented traits​

comptime members​

kv_stage_elems​

KVPipeType​

Methods​

__init__​

stage_base_ptr​

stage_index​

wait​

release​