Skip to main content

Mojo function

matmul_sm100_grouped_blockwise_scaled_fp8_1d2d_kernel

matmul_sm100_grouped_blockwise_scaled_fp8_1d2d_kernel[a_type: DType, b_type: DType, c_type: DType, a_scales_type: DType, b_scales_type: DType, accum_type: DType, a_layout: Layout, b_layout: Layout, a_offsets_layout: Layout, expert_ids_layout: Layout, a_scales_layout: Layout, b_scales_layout: Layout, c_static_N: 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], block_tile_shape: IndexList[3], mma_shape: IndexList[3], transpose_b: Bool = True, cluster_shape: StaticTuple[Int32, 3] = StaticTuple(SIMD(1), SIMD(1), SIMD(1)), a_swizzle: TensorMapSwizzle = TensorMapSwizzle.SWIZZLE_128B, b_swizzle: TensorMapSwizzle = TensorMapSwizzle.SWIZZLE_128B, num_threads: Int = 128, elementwise_lambda_fn: Optional[def[dtype: DType, width: Int, *, alignment: Int = 1](IndexList[2], SIMD[dtype, width]) capturing -> None] = None](a_tma_op: TMATensorTile[a_type, a_tile_rank, a_tile_shape, a_desc_shape], b_tma_op: TMATensorTile[b_type, b_tile_rank, b_tile_shape, b_desc_shape], a_offsets: LayoutTensor[DType.uint32, a_offsets_layout, MutAnyOrigin], expert_ids: LayoutTensor[DType.int32, expert_ids_layout, MutAnyOrigin], c_ptr: UnsafePointer[Scalar[c_type], MutAnyOrigin], a_scales: LayoutTensor[a_scales_type, a_scales_layout, MutAnyOrigin], b_scales: LayoutTensor[b_scales_type, b_scales_layout, MutAnyOrigin], num_iters: Int)

Was this page helpful?