Skip to main content

Mojo function

grouped_matmul_sm100_blockwise_scaled_fp8

grouped_matmul_sm100_blockwise_scaled_fp8[a_layout: Layout, b_layout: Layout, c_layout: Layout, a_scales_layout: Layout, b_scales_layout: Layout, a_offsets_layout: Layout, expert_ids_layout: Layout, c_type: DType, a_type: DType, b_type: DType, a_scales_type: DType, b_scales_type: DType, a_offsets_type: DType, expert_ids_type: DType, //, *, transpose_b: Bool, umma_shape: IndexList[3], block_tile_shape: IndexList[3], a_swizzle: TensorMapSwizzle = 3, b_swizzle: TensorMapSwizzle = 3, elementwise_lambda_fn: OptionalReg[fn[dtype: DType, width: Int, *, alignment: Int = 1](IndexList[2], SIMD[dtype, width]) capturing -> None] = None](c: LayoutTensor[c_type, c_layout, origin, address_space=address_space, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], a: LayoutTensor[a_type, a_layout, origin, address_space=address_space, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], b: LayoutTensor[b_type, b_layout, origin, address_space=address_space, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], a_scales: LayoutTensor[a_scales_type, a_scales_layout, origin, address_space=address_space, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], b_scales: LayoutTensor[b_scales_type, b_scales_layout, origin, address_space=address_space, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], a_offsets: LayoutTensor[a_offsets_type, a_offsets_layout, origin, address_space=address_space, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], expert_ids: LayoutTensor[expert_ids_type, expert_ids_layout, origin, address_space=address_space, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], max_num_tokens_per_expert: Int, num_active_experts: Int, ctx: DeviceContext)

Was this page helpful?