ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-15 10:37:44 +00:00
Code Issues Packages Projects Releases Wiki Activity
Files
a8dbac647029057abbd5d7fcd1473d5f9b33cded
composable_kernel/include/ck/tensor_operation/gpu/block
History
Xudong Yuan a8dbac6470 Ck moe mxfp4 blockm32 (#3098) 
* block_m = 32

* ck block_m = 32

* aiter/3rdparty/composable_kernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_v3.hpp format

* mxfp4_moe v1 pipe

* update format

---------

Co-authored-by: zhimding <zhimding@amd.com>
Co-authored-by: lalala-sh <Jiaxing.Wen@amd.com>
Co-authored-by: felix <felix.li@amd.com>

[ROCm/composable_kernel commit: d04eba4ae3]
2025-11-07 08:45:41 +08:00
..
blockwise_gemm_dl_v2r3.hpp
…
blockwise_gemm_dlops_v2r2.hpp
…
blockwise_gemm_dlops_v3.hpp
…
blockwise_gemm_dpp.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_mx_pipeline_xdlops_base.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_wmma_selector.hpp
Implement batched gemm gemm for RDNA (3 and 4) (#2612)
2025-09-04 14:10:24 -07:00
blockwise_gemm_pipeline_wmmaops_base.hpp
WMMA gemm_add_relu_add_layernorm (#2989)
2025-10-31 11:19:26 -07:00
blockwise_gemm_pipeline_wmmaops_v1.hpp
Kabraham/fix block gemm v1 b scale (#3129)
2025-10-31 07:19:01 -07:00
blockwise_gemm_pipeline_wmmaops_v3.hpp
Implement batched gemm gemm for RDNA (3 and 4) (#2612)
2025-09-04 14:10:24 -07: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
…
blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v1.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v3.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_gufusion_dequant_v1.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_gufusion_v1.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_gufusion_v3.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_gufusion_v3.hpp
Ck moe mxfp4 blockm32 (#3098)
2025-11-07 08:45:41 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_selector.hpp
Ck moe mxfp4 blockm32 (#3098)
2025-11-07 08:45:41 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_v1.hpp
Ck moe mxfp4 blockm32 (#3098)
2025-11-07 08:45:41 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_v3.hpp
Ck moe mxfp4 blockm32 (#3098)
2025-11-07 08:45:41 +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
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_v2.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_b_preshuffle_v3.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_b_scale_selector.hpp
…
blockwise_gemm_pipeline_xdlops_base.hpp
Conv:TF32: add more instances - 1 (#2867)
2025-09-25 09:27:18 +08: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
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v3.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_gufusion_v1.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_gufusion_v3.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +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
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_v3.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +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
Grouped conv fwd with direct load (#3082)
2025-10-29 09:54:42 +01:00
blockwise_gemm_pipeline_xdlops_v1_ab_scale.hpp
Conv:TF32: add more instances - 1 (#2867)
2025-09-25 09:27:18 +08:00
blockwise_gemm_pipeline_xdlops_v1_b_scale.hpp
Conv:TF32: add more instances - 1 (#2867)
2025-09-25 09:27:18 +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
Grouped conv fwd with direct load (#3082)
2025-10-29 09:54:42 +01:00
blockwise_gemm_pipeline_xdlops_v2_ab_scale.hpp
Conv:TF32: add more instances - 1 (#2867)
2025-09-25 09:27:18 +08:00
blockwise_gemm_pipeline_xdlops_v2_b_scale.hpp
Conv:TF32: add more instances - 1 (#2867)
2025-09-25 09:27:18 +08:00
blockwise_gemm_pipeline_xdlops_v2.hpp
Conv:TF32: add more instances - 1 (#2867)
2025-09-25 09:27:18 +08:00
blockwise_gemm_pipeline_xdlops_v3_ab_scale.hpp
Conv:TF32: add more instances - 1 (#2867)
2025-09-25 09:27:18 +08:00
blockwise_gemm_pipeline_xdlops_v3_b_scale.hpp
Conv:TF32: add more instances - 1 (#2867)
2025-09-25 09:27:18 +08:00
blockwise_gemm_pipeline_xdlops_v3_mx_bpreshuffle.hpp
[CK] Fixed MPerBlock=32 build issue for MXFP4 GEMM decode (#2512)
2025-07-18 14:35:54 +08: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
Conv:TF32: add more instances - 1 (#2867)
2025-09-25 09:27:18 +08:00
blockwise_gemm_pipeline_xdlops_v4_b_scale.hpp
Conv:TF32: add more instances - 1 (#2867)
2025-09-25 09:27:18 +08:00
blockwise_gemm_pipeline_xdlops_v4.hpp
Grouped conv fwd with direct load (#3082)
2025-10-29 09:54:42 +01:00
blockwise_gemm_pipeline_xdlops_v5.hpp
Conv:TF32: add more instances - 1 (#2867)
2025-09-25 09:27:18 +08:00
blockwise_gemm_pipeline_xdlops.hpp
Extend XDL kernel to Support RDNA3/4 - Part 3 (#2723)
2025-09-09 11:22:36 +08:00
blockwise_gemm_smfmac_xdlops.hpp
Extend XDL kernel to Support RDNA3/4 - Part 3 (#2723)
2025-09-09 11:22:36 +08:00
blockwise_gemm_wmma.hpp
…
blockwise_gemm_xdlops_skip_b_lds.hpp
Extend XDL kernel to Support RDNA3/4 - Part 2 (#2722)
2025-09-04 08:33:40 +08:00
blockwise_gemm_xdlops.hpp
TF32 POC in Conv3d on MI30x platform #2763 (second attempt) (#2852)
2025-09-17 14:50:15 -07:00
blockwise_softmax.hpp
…
blockwise_tensor_slice_transfer_v5r1.hpp
…
blockwise_welford.hpp
…
reduction_functions_blockwise.hpp
…
thread_group_tensor_slice_transfer_direct_load.hpp
upgrade from clang-format-12 to clang-format-18 (#2568)
2025-07-28 11:34:07 -07:00
thread_group_tensor_slice_transfer_gather_direct_load.hpp
upgrade from clang-format-12 to clang-format-18 (#2568)
2025-07-28 11:34:07 -07:00
thread_group_tensor_slice_transfer_global.hpp
Wave Tile Transfer supporting global load with transpose (#3027)
2025-10-16 11:33:56 -07:00
thread_group_tensor_slice_transfer_v4r1_dequant.hpp
…
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
…
thread_group_tensor_slice_transfer_v4r2.hpp
…
thread_group_tensor_slice_transfer_v6r1.hpp
…
thread_group_tensor_slice_transfer_v6r1r2.hpp
…
thread_group_tensor_slice_transfer_v6r2.hpp
…
thread_group_tensor_slice_transfer_v6r3.hpp
…
thread_group_tensor_slice_transfer_v7.hpp
…
thread_group_tensor_slice_transfer_v7r2.hpp
…
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
WMMA gemm_add_relu_add_layernorm (#2989)
2025-10-31 11:19:26 -07:00