mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-03-20 07:07:43 +00:00
2312eef6c3
[CK_TILE][FMHA] Support microscaling (mxfp8 and mxfp4) on gfx950 (#4368) ## Motivation Microscaling types (mxfp8 and mxfp4) for fwd qr pipeline ## Technical Details The microscaling is used when quant scale mode is `BlockAttentionQuantScaleEnum::MX` and `Q/K/P/VDataType` are fp8/bf8/fp4. Supported features: * only "qr" pipeline is implemented * hdim 128 and 256 (smaller hdim are not possible due to restrictions of "qr" pipeline, but they can be computed using instances with padding) * both 32x32x64 and 16x16x128 scale MFMAs are supported * Q and K scales are applied in hdim, V scales - in seqlen dimension * column-major V only * batch and group mode * bias, Alibi (tested but no instances by default, just like fp8) * masking etc. Aiter PR with new API args: https://github.com/ROCm/aiter/pull/2008 ## Test Plan ``` ninja test_ck_tile_fmha_fwd_mxfp8 && bin/test_ck_tile_fmha_fwd_mxfp8 ninja test_ck_tile_fmha_fwd_mxfp4 && bin/test_ck_tile_fmha_fwd_mxfp4 ``` ## Test Result The tests must pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
55 lines
2.8 KiB
C++
55 lines
2.8 KiB
C++
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
|
|
// SPDX-License-Identifier: MIT
|
|
#pragma once
|
|
|
|
#include "ck_tile/host/arg_parser.hpp"
|
|
#include "ck_tile/host/check_err.hpp"
|
|
#include "ck_tile/host/concat.hpp"
|
|
#include "ck_tile/host/convolution_host_tensor_descriptor_helper.hpp"
|
|
#include "ck_tile/host/convolution_parameter.hpp"
|
|
#include "ck_tile/host/device_memory.hpp"
|
|
#include "ck_tile/host/device_prop.hpp"
|
|
#include "ck_tile/host/fill.hpp"
|
|
#include "ck_tile/host/flush_icache.hpp"
|
|
#include "ck_tile/host/high_res_cpu_clock.hpp"
|
|
#include "ck_tile/host/hip_check_error.hpp"
|
|
#include "ck_tile/host/host_tensor.hpp"
|
|
#include "ck_tile/host/joinable_thread.hpp"
|
|
#include "ck_tile/host/kernel_launch.hpp"
|
|
#include "ck_tile/host/permute_pk_int4.hpp"
|
|
#include "ck_tile/host/ranges.hpp"
|
|
#include "ck_tile/host/reference/reference_batched_contraction.hpp"
|
|
#include "ck_tile/host/reference/reference_batched_dropout.hpp"
|
|
#include "ck_tile/host/reference/reference_batched_dropout_randval.hpp"
|
|
#include "ck_tile/host/reference/reference_batched_elementwise.hpp"
|
|
#include "ck_tile/host/reference/reference_batched_gemm.hpp"
|
|
#include "ck_tile/host/reference/reference_batched_masking.hpp"
|
|
#include "ck_tile/host/reference/reference_batched_mx_descale.hpp"
|
|
#include "ck_tile/host/reference/reference_batched_rotary_position_embedding.hpp"
|
|
#include "ck_tile/host/reference/reference_batched_softmax.hpp"
|
|
#include "ck_tile/host/reference/reference_batched_transpose.hpp"
|
|
#include "ck_tile/host/reference/reference_blocked_attention.hpp"
|
|
#include "ck_tile/host/reference/reference_elementwise.hpp"
|
|
#include "ck_tile/host/reference/reference_fused_moe.hpp"
|
|
#include "ck_tile/host/reference/reference_gemm.hpp"
|
|
#include "ck_tile/host/reference/reference_grouped_conv_bwd_data.hpp"
|
|
#include "ck_tile/host/reference/reference_grouped_conv_bwd_weight.hpp"
|
|
#include "ck_tile/host/reference/reference_grouped_conv_fwd.hpp"
|
|
#include "ck_tile/host/reference/reference_im2col.hpp"
|
|
#include "ck_tile/host/reference/reference_layernorm2d_fwd.hpp"
|
|
#include "ck_tile/host/reference/reference_moe_gemm.hpp"
|
|
#include "ck_tile/host/reference/reference_moe_sorting.hpp"
|
|
#include "ck_tile/host/reference/reference_permute.hpp"
|
|
#include "ck_tile/host/reference/reference_pool.hpp"
|
|
#include "ck_tile/host/reference/reference_reduce.hpp"
|
|
#include "ck_tile/host/reference/reference_rmsnorm2d_fwd.hpp"
|
|
#include "ck_tile/host/reference/reference_rowwise_quantization2d.hpp"
|
|
#include "ck_tile/host/reference/reference_softmax.hpp"
|
|
#include "ck_tile/host/reference/reference_topk.hpp"
|
|
#include "ck_tile/host/reference/reference_transpose.hpp"
|
|
#include "ck_tile/host/rotating_buffers.hpp"
|
|
#include "ck_tile/host/stream_config.hpp"
|
|
#include "ck_tile/host/stream_utils.hpp"
|
|
#include "ck_tile/host/tensor_shuffle_utils.hpp"
|
|
#include "ck_tile/host/timer.hpp"
|