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_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[Int, Int](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[def[dtype: DType, width: Int, *, alignment: Int = 1](IndexList[2], SIMD[dtype, width]) capturing -> None] = None, transpose_c: Bool = False, use_tma: Bool = True, K_actual: Int = 0, a_gmem_layout: Layout = Layout.row_major(1, 1), b_gmem_layout: Layout = Layout.row_major(1, 1)](num_active_experts: Int, 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, 3], a_gmem: LayoutTensor[a_type, a_gmem_layout, ImmutAnyOrigin], b_gmem: LayoutTensor[b_type, b_gmem_layout, ImmutAnyOrigin])

Was this page helpful?