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_layout: Layout, b_layout: Layout, c_layout: Layout, c_tensor_layout: Layout, a_desc_layout: Layout, b_desc_layout: Layout, c_desc_layout: Layout, block_tile_shape: IndexList[3], mma_shape: IndexList[3], cluster_shape: StaticTuple[Int32, 3], num_pipeline_stages: UInt, num_accum_pipeline_stages: UInt, num_output_stages: UInt = 2, output_tile_shape: IndexList[2] = Index(128, 32), transpose_b: Bool = True, a_swizzle: TensorMapSwizzle = 3, b_swizzle: TensorMapSwizzle = 3, c_swizzle: TensorMapSwizzle = 3, cta_group: Int = 2, elementwise_compute_lambda_fn: OptionalReg[fn[dtype: DType, width: Int, *, alignment: Int = 1](IndexList[2], SIMD[dtype, width]) capturing -> SIMD[dtype, width]] = None, register_based_epilogue: Bool = True, transpose_c: Bool = False](num_active_experts: Int, a_tma_op: TMATensorTile[a_type, a_layout, a_desc_layout], expert_ids: NDBuffer[DType.int32, 1, MutableAnyOrigin], b_tma_op: TMATensorTile[b_type, b_layout, b_desc_layout], b_offsets: NDBuffer[DType.uint32, 1, MutableAnyOrigin], c_tma_op: TMATensorTile[c_type, c_layout, c_desc_layout], c: LayoutTensor[c_type, c_tensor_layout, MutableAnyOrigin], mnk: StaticTuple[UInt32, 3])
Was this page helpful?
Thank you! We'll create more content like this.
Thank you for helping us improve!