ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-03 13:11:25 +00:00
Code Issues Packages Projects Releases Wiki Activity
Files
7998ae89693dbc24793334bdb5e12568fa30fe2b
composable_kernel/include/ck/tensor_operation/gpu/block
History
Mingtao Gu 7998ae8969 [CK] Mxfp4 moe blockscale buf2lds version support (#2455) 
* change cshuffle size

* added mxfp4 moe async buffer loading without B preshuffle

* added mx moe B shuffling + scale shuffling (async loads)

* minor fix

---------

Co-authored-by: mtgu0705 <mtgu@amd.com>
2025-07-06 15:42:00 +08:00
..
blockwise_gemm_dl_v2r3.hpp
Implement DPP8 based GEMM for Navi21 (#826)
2023-08-14 15:46:27 -05:00
blockwise_gemm_dlops_v2r2.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
blockwise_gemm_dlops_v3.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
blockwise_gemm_dpp.hpp
Fix cmake warnings (#1342)
2024-06-21 09:47:58 +02:00
blockwise_gemm_mx_pipeline_xdlops_base.hpp
Do not use warpSize as compile time constant as it is removed (#2320)
2025-06-17 11:54:30 -07:00
blockwise_gemm_pipeline_wmma_selector.hpp
WMMA GEMM universal pipeline v1, mixed precision and paddings, examples (#2230)
2025-06-04 12:22:33 +06:00
blockwise_gemm_pipeline_wmmaops_base.hpp
WMMA GEMM universal pipeline v1, mixed precision and paddings, examples (#2230)
2025-06-04 12:22:33 +06:00
blockwise_gemm_pipeline_wmmaops_v1.hpp
WMMA GEMM universal pipeline v1, mixed precision and paddings, examples (#2230)
2025-06-04 12:22:33 +06:00
blockwise_gemm_pipeline_wmmaops_v3.hpp
WMMA GEMM universal pipeline v1, mixed precision and paddings, examples (#2230)
2025-06-04 12:22:33 +06:00
blockwise_gemm_pipeline_wmmaops.hpp
DeviceGemm_Wmma_CShuffleV3 with BlockGemmPipelineVersion::v3 (#2096)
2025-04-28 10:14:21 +05:00
blockwise_gemm_pipeline_xdlops_ab_scale_selector.hpp
[A8W8 GEMM] Optimized weight-preshuffled implementation & add quantization datatype for CK TILE rms_norm (#1862)
2025-02-20 14:00:27 -08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v1.hpp
Ck int4 moe develop (#1949)
2025-03-10 11:16:44 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v3.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_gufusion_dequant_v1.hpp
fix moe i4 bug from aiter (#2339)
2025-06-24 14:51:29 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_gufusion_v1.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_gufusion_v3.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_gufusion_v3.hpp
[CK] Mxfp4 moe blockscale buf2lds version support (#2455)
2025-07-06 15:42:00 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_selector.hpp
[CK] Mxfp4 moe blockscale buf2lds version support (#2455)
2025-07-06 15:42:00 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_v3.hpp
[CK] Mxfp4 moe blockscale buf2lds version support (#2455)
2025-07-06 15:42:00 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_selector.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_v1.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_v2.hpp
Add the instances for small sized GEMM in preshuffle and improve CMake Flag (#2212)
2025-05-20 15:05:08 -07:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_v3.hpp
Narrowing error fix for codegen compilation (#2194)
2025-05-16 11:11:54 -07:00
blockwise_gemm_pipeline_xdlops_b_scale_selector.hpp
[A8W8 GEMM] Optimized weight-preshuffled implementation & add quantization datatype for CK TILE rms_norm (#1862)
2025-02-20 14:00:27 -08:00
blockwise_gemm_pipeline_xdlops_base.hpp
Do not use warpSize as compile time constant as it is removed (#2320)
2025-06-17 11:54:30 -07:00
blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_selector.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v1.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v3.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_gufusion_v1.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_gufusion_v3.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_selector.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_v1.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_v3.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_mx_bpreshuffle_selector.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
blockwise_gemm_pipeline_xdlops_mx_moe_gufusion_v3.hpp
[CK] Mxfp4 moe blockscale buf2lds version support (#2455)
2025-07-06 15:42:00 +08:00
blockwise_gemm_pipeline_xdlops_mx_moe_nbs_gufusion_v3.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_mx_moe_nbs_selector.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_mx_moe_nbs_v1.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_mx_moe_nbs_v3.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_mx_moe_selector.hpp
[CK] Mxfp4 moe blockscale buf2lds version support (#2455)
2025-07-06 15:42:00 +08:00
blockwise_gemm_pipeline_xdlops_mx_moe_v3.hpp
[CK] Mxfp4 moe blockscale buf2lds version support (#2455)
2025-07-06 15:42:00 +08:00
blockwise_gemm_pipeline_xdlops_mx_selector.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
blockwise_gemm_pipeline_xdlops_selector.hpp
[A8W8 GEMM] Optimized weight-preshuffled implementation & add quantization datatype for CK TILE rms_norm (#1862)
2025-02-20 14:00:27 -08:00
blockwise_gemm_pipeline_xdlops_v1_ab_scale.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
blockwise_gemm_pipeline_xdlops_v1_b_scale.hpp
Implement the fp16xint4 scale weight only kernel for Ali (#1786)
2025-01-03 18:35:21 +08:00
blockwise_gemm_pipeline_xdlops_v1_mx.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
blockwise_gemm_pipeline_xdlops_v1.hpp
Fix cmake warnings (#1342)
2024-06-21 09:47:58 +02:00
blockwise_gemm_pipeline_xdlops_v2_ab_scale.hpp
Do not use warpSize as compile time constant as it is removed (#2320)
2025-06-17 11:54:30 -07:00
blockwise_gemm_pipeline_xdlops_v2_b_scale.hpp
Do not use warpSize as compile time constant as it is removed (#2320)
2025-06-17 11:54:30 -07:00
blockwise_gemm_pipeline_xdlops_v2.hpp
Do not use warpSize as compile time constant as it is removed (#2320)
2025-06-17 11:54:30 -07:00
blockwise_gemm_pipeline_xdlops_v3_ab_scale.hpp
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
blockwise_gemm_pipeline_xdlops_v3_b_scale.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
blockwise_gemm_pipeline_xdlops_v3_mx_bpreshuffle.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
blockwise_gemm_pipeline_xdlops_v3_mx.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
blockwise_gemm_pipeline_xdlops_v3.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
blockwise_gemm_pipeline_xdlops_v4_b_scale.hpp
Implement the fp16xint4 scale weight only kernel for Ali (#1786)
2025-01-03 18:35:21 +08:00
blockwise_gemm_pipeline_xdlops_v4.hpp
[A8W8 GEMM] Optimized weight-preshuffled implementation & add quantization datatype for CK TILE rms_norm (#1862)
2025-02-20 14:00:27 -08:00
blockwise_gemm_pipeline_xdlops_v5.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
blockwise_gemm_pipeline_xdlops.hpp
Fix compilation errors with Clang20.0. (#1533)
2024-09-25 13:45:38 -07:00
blockwise_gemm_smfmac_xdlops.hpp
Added structural sparsity blockwise gemm (#1435)
2024-09-11 15:19:42 +02:00
blockwise_gemm_wmma.hpp
fix compilation errors for gfx12 with clang20 (#1606)
2024-10-28 19:02:48 -07:00
blockwise_gemm_xdlops_skip_b_lds.hpp
Fix cmake warnings (#1342)
2024-06-21 09:47:58 +02:00
blockwise_gemm_xdlops.hpp
Merging the gfx12 code into public repo. (#1362)
2024-06-27 00:33:34 -07:00
blockwise_softmax.hpp
[HotFix] add config and version files to pass on build info (#856)
2023-08-23 11:36:17 -07:00
blockwise_tensor_slice_transfer_v5r1.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
blockwise_welford.hpp
Batchnorm splitk single kernel (#771)
2023-07-06 10:58:55 -05:00
reduction_functions_blockwise.hpp
Batchnorm splitk single kernel (#771)
2023-07-06 10:58:55 -05:00
thread_group_tensor_slice_transfer_direct_load.hpp
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
thread_group_tensor_slice_transfer_gather_direct_load.hpp
[CK] Mxfp4 moe blockscale buf2lds version support (#2455)
2025-07-06 15:42:00 +08:00
thread_group_tensor_slice_transfer_v4r1_dequant.hpp
Navi3 rel (#1176)
2024-03-08 17:11:51 -08:00
thread_group_tensor_slice_transfer_v4r1_gather.hpp
Moe gemm activation (#2026)
2025-04-23 10:35:34 +08:00
thread_group_tensor_slice_transfer_v4r1.hpp
[A8W8 GEMM] Optimized weight-preshuffled implementation & add quantization datatype for CK TILE rms_norm (#1862)
2025-02-20 14:00:27 -08:00
thread_group_tensor_slice_transfer_v4r2.hpp
Add elementwise with dynamic vector dim (#1198)
2024-03-22 10:40:43 +01:00
thread_group_tensor_slice_transfer_v6r1.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
thread_group_tensor_slice_transfer_v6r1r2.hpp
initial stream-k implementation with example (#699)
2023-07-26 14:18:15 -05:00
thread_group_tensor_slice_transfer_v6r2.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
thread_group_tensor_slice_transfer_v6r3.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
thread_group_tensor_slice_transfer_v7.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
thread_group_tensor_slice_transfer_v7r2.hpp
Codegen hipRTC compilation (#1579)
2025-01-31 09:48:39 -08:00
thread_group_tensor_slice_transfer_v7r3_scatter.hpp
Moe gemm activation (#2026)
2025-04-23 10:35:34 +08:00
thread_group_tensor_slice_transfer_v7r3.hpp
add f8 gemm multiD with both row/col wise scale (#1300)
2024-05-28 12:04:22 -05:00