Skip to main content

Mojo function

blackwell_tma_umma_warp_specialized_kernel

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_tensor_layout: Layout, c_desc_shape: IndexList[c_tile_rank], block_tile_shape: IndexList[3], mma_shape: IndexList[3], cluster_shape: StaticTuple[Int32, 3], num_pipeline_stages: UInt, num_accum_pipeline_stages: Int, num_output_stages: Int = 2, output_tile_shape: IndexList[2] = Index(128, 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 = 2, elementwise_lambda_fn: Optional[elementwise_epilogue_type] = None, transpose_c: Bool = False](num_active_experts: Int, a_tma_op: TMATensorTile[a_type, a_tile_rank, a_tile_shape, a_desc_shape], expert_ids: NDBuffer[DType.int32, ImmutAnyOrigin, DimList.create_unknown[1](), DimList.create_unknown[1]()], b_tma_op: TMATensorTile[b_type, b_tile_rank, b_tile_shape, b_desc_shape], b_offsets: NDBuffer[DType.uint32, ImmutAnyOrigin, DimList.create_unknown[1](), DimList.create_unknown[1]()], c_tma_op: TMATensorTile[c_type, c_tile_rank, c_tile_shape_param, c_desc_shape], c: LayoutTensor[c_type, c_tensor_layout, MutAnyOrigin], mnk: StaticTuple[UInt32, 3])

Was this page helpful?