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 function

blackwell_tma_umma_warp_specialized_kernel

def blackwell_tma_umma_warp_specialized_kernel[a_type: DType, b_type: DType, c_type: DType, expert_m: Int, a_tile_rank: Int, a_tile_shape: IndexList[a_tile_rank], a_desc_shape: IndexList[a_tile_rank], b_tile_rank: Int, b_tile_shape: IndexList[b_tile_rank], b_desc_shape: IndexList[b_tile_rank], c_tile_rank: Int, c_tile_shape_param: IndexList[c_tile_rank], c_desc_shape: IndexList[c_tile_rank], block_tile_shape: IndexList[Int(3)], mma_shape: IndexList[Int(3)], cluster_shape: StaticTuple[Int32, Int(3)], num_pipeline_stages: Int, num_accum_pipeline_stages: Int, num_output_stages: Int = Int(2), output_tile_shape: IndexList[Int(2)] = Index[Int, Int](Int(128), Int(32)), transpose_b: Bool = True, a_swizzle: TensorMapSwizzle = TensorMapSwizzle.SWIZZLE_128B, b_swizzle: TensorMapSwizzle = TensorMapSwizzle.SWIZZLE_128B, c_swizzle: TensorMapSwizzle = TensorMapSwizzle.SWIZZLE_128B, cta_group: Int = Int(2), elementwise_lambda_fn: Optional[def[dtype: DType, width: SIMDSize, *, alignment: Int = Int(1)](IndexList[Int(2)], SIMD[dtype, width]) capturing -> None] = None, transpose_c: Bool = False, use_tma: Bool = True, K_actual: Int = Int(0), a_gmem_layout: Layout = Layout.row_major(Int(1), Int(1)), b_gmem_layout: Layout = Layout.row_major(Int(1), Int(1))](expert_usage_stats: UnsafePointer[UInt32, ImmutAnyOrigin], a_tma_op: TMATensorTile[a_type, a_tile_rank, a_tile_shape, a_desc_shape], expert_ids: UnsafePointer[Int32, ImmutAnyOrigin], b_tma_op: TMATensorTile[b_type, b_tile_rank, b_tile_shape, b_desc_shape], b_offsets: UnsafePointer[UInt32, ImmutAnyOrigin], c_tma_op: TMATensorTile[c_type, c_tile_rank, c_tile_shape_param, c_desc_shape], c_ptr: UnsafePointer[Scalar[c_type], MutAnyOrigin], mnk: StaticTuple[UInt32, Int(3)], a_gmem: LayoutTensor[a_type, a_gmem_layout, ImmutAnyOrigin], b_gmem: LayoutTensor[b_type, b_gmem_layout, ImmutAnyOrigin])