ROCm · anhminhnguyenhoang · Jan 16, 2026 · Jan 19, 2026 · Jan 20, 2026 · Jan 20, 2026
diff --git a/3rdparty/composable_kernel b/3rdparty/composable_kernel
diff --git a/aiter/ops/triton/_triton_kernels/fusions/__init__.py b/aiter/ops/triton/_triton_kernels/fusions/__init__.py
@@ -0,0 +1,20 @@
+# SPDX-License-Identifier: MIT
+# Copyright (C) 2024-2026, Advanced Micro Devices, Inc. All rights reserved.
+
+from aiter.ops.triton._triton_kernels.fusions.mhc import (
+    _mhc_fused_kernel,
+    _mhc_fused_split_kernel,
+    _mhc_fused_reduce_kernel,
+    _mhc_lite_fused_split_kernel,
+    _mhc_lite_fused_reduce_kernel,
+    _sinkhorn_knopp_log_domain_kernel,
+)
+
+__all__ = [
+    "_mhc_fused_kernel",
+    "_mhc_fused_split_kernel",
+    "_mhc_fused_reduce_kernel",
+    "_mhc_lite_fused_split_kernel",
+    "_mhc_lite_fused_reduce_kernel",
+    "_sinkhorn_knopp_log_domain_kernel",
+]
+6 −6		.github/CODEOWNERS
+13 −3		CMakeLists.txt
+6 −34		Jenkinsfile
+16 −0		README.md
+30 −0		example/CMakeLists.txt
+1 −1		example/ck_tile/01_fmha/codegen/ops/fmha_batch_prefill.py
+12 −3		example/ck_tile/38_block_scale_gemm/CMakeLists.txt
+6 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_bf8.cpp
+6 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_bf8i4.cpp
+6 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_fp8.cpp
+6 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_fp8i4.cpp
+0 −222		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb.cpp
+53 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_bf8.cpp
+57 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_bf8i4.cpp
+53 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_fp8.cpp
+57 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_fp8i4.cpp
+0 −62		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_preshufflequant.cpp
+50 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_preshufflequant_bf8.cpp
+52 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_preshufflequant_bf8i4.cpp
+50 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_preshufflequant_fp8.cpp
+52 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_preshufflequant_fp8i4.cpp
+0 −270		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshufflequant.cpp
+55 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshufflequant_bf8.cpp
+59 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshufflequant_bf8i4.cpp
+55 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshufflequant_fp8.cpp
+59 −0		example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshufflequant_fp8i4.cpp
+33 −6		example/ck_tile/38_block_scale_gemm/gemm_quant.cpp
+35 −26		experimental/builder/include/ck_tile/builder/reflect/conv_describe.hpp
+86 −641		experimental/builder/include/ck_tile/builder/reflect/conv_traits.hpp
+84 −0		...l/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
+84 −0		...uilder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp
+84 −0		...nclude/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp
+739 −0		experimental/builder/include/ck_tile/builder/reflect/conv_traits_helpers.hpp
+8 −0		experimental/builder/include/ck_tile/builder/reflect/instance_to_conv_traits.hpp
+8 −0		...ilder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
+8 −0		...er/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp
+8 −0		...de/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp
+2 −1		experimental/builder/test/CMakeLists.txt
+78 −78		experimental/builder/test/conv/ck/test_conv_traits.cpp
+0 −1,127		experimental/builder/test/conv/ck/unit_instance_to_conv_traits.cpp
+800 −0		experimental/builder/test/conv/ck/unit_instance_to_conv_traits_features.cpp
+262 −0		experimental/builder/test/conv/ck/unit_instance_to_conv_traits_instances.cpp
+175 −63		include/ck/library/utility/gpu_verification.hpp
+4 −0		include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_wmma_cshuffle_v3.hpp
+8 −0		include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp
+4 −0		include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_wmma_cshuffle_v3.hpp
+4 −0		include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_wmma_cshuffle_v3.hpp
+1 −1		include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
+1 −1		include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp
+59 −17		include/ck_tile/ops/fmha/pipeline/block_fmha_batch_prefill_pipeline_qr_ks_vs_async.hpp
+2 −10		include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_problem.hpp
+17 −1		include/ck_tile/ops/gemm_quant/block/block_universal_gemm_as_bs_bquant_cr.hpp
+50 −27		include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp
+5 −4		include/ck_tile/ops/gemm_quant/pipeline/gemm_bquant_pipeline_ag_bg_cr_base.hpp
+3 −1		include/ck_tile/ops/gemm_quant/pipeline/gemm_bquant_pipeline_ag_bg_cr_policy.hpp
+3 −1		include/ck_tile/ops/gemm_quant/pipeline/gemm_bquant_pipeline_ag_bg_cr_v3.hpp
+48 −26		include/ck_tile/ops/gemm_quant/pipeline/gemm_group_quant_utils.hpp
+0 −2		include/ck_tile/ops/gemm_quant/pipeline/gemm_quant_pipeline_problem.hpp
+22 −13		include/ck_tile/ops/gemm_quant/pipeline/gemm_wp_bquant_pipeline_ag_bg_cr_v2.hpp
+54 −56		profiler/include/profiler/profile_grouped_conv_bwd_data_impl.hpp
+44 −62		profiler/include/profiler/profile_grouped_conv_bwd_weight_impl.hpp
+44 −30		profiler/include/profiler/profile_grouped_conv_fwd_impl.hpp
+254 −0		script/run-tests.ps1
+26 −0		test/ck_tile/gemm_block_scale/CMakeLists.txt
+39 −0		test/ck_tile/gemm_block_scale/test_gemm_quant_bquant_preshuffleQuant_decode_1d.cpp
+54 −0		test/ck_tile/gemm_block_scale/test_gemm_quant_bquant_preshuffleQuant_decode_2d.cpp
+41 −0		test/ck_tile/gemm_block_scale/test_gemm_quant_bquant_preshuffleQuant_prefill_1d.cpp
+63 −0		test/ck_tile/gemm_block_scale/test_gemm_quant_bquant_preshuffleQuant_prefill_2d.cpp
+8 −5		test/ck_tile/gemm_block_scale/test_gemm_quant_bquant_preshuffle_decode_2d.cpp
+10 −5		test/ck_tile/gemm_block_scale/test_gemm_quant_bquant_preshuffle_prefill_2d.cpp
+22 −29		test/ck_tile/gemm_block_scale/test_gemm_quant_fixtures.hpp
+4 −3		test/ck_tile/gemm_streamk/CMakeLists.txt
+1 −1		test/gpu_verification/test_gpu_verification.cpp
+21 −17		test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight.cpp
+4 −4		test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_bilinear.cpp
+4 −4		test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_scale.cpp