ROCm · valarLip · Mar 18, 2025 · Mar 18, 2025 · Mar 20, 2025 · Mar 20, 2025
diff --git a/3rdparty/composable_kernel b/3rdparty/composable_kernel
diff --git a/aiter/__init__.py b/aiter/__init__.py
@@ -4,14 +4,40 @@
 import torch
 import os
 import logging
+
+
 logger = logging.getLogger("aiter")
+
+
+def getLogger():
+    global logger
+    if not logger.handlers:
+        logger.setLevel(logging.DEBUG)
+
+        console_handler = logging.StreamHandler()
+        if int(os.environ.get("AITER_LOG_MORE", 0)):
+            formatter = logging.Formatter(
+                fmt="[%(name)s %(levelname)s] %(asctime)s.%(msecs)03d - %(processName)s:%(process)d - %(pathname)s:%(lineno)d - %(funcName)s\n%(message)s",
+                datefmt="%Y-%m-%d %H:%M:%S",
+            )
+        else:
+            formatter = logging.Formatter(
+                fmt="[%(name)s] %(message)s",
+            )
+        console_handler.setFormatter(formatter)
+        console_handler.setLevel(logging.INFO)
+        logger.addHandler(console_handler)
+
+    return logger
+
+
+logger = getLogger()
+
 import importlib.util
-if importlib.util.find_spec('aiter_') is not None:
+
+if importlib.util.find_spec("aiter_") is not None:
     from aiter_ import *
-# if importlib.util.find_spec('hipbsolidxgemm_') is not None:
-#     from hipbsolidxgemm_ import *
-# if importlib.util.find_spec('rocsolidxgemm_') is not None:
-#     from rocsolidxgemm_ import *
+from .jit import core
 from .ops.norm import *
 from .ops.quant import *
 from .ops.gemm_op_a8w8 import *
@@ -33,23 +59,3 @@
 from .ops.mha import *
 from .ops.gradlib import *
 from . import mla
-
-def getLogger():
-    global logger
-    if not logger.handlers:
-        logger.setLevel(logging.DEBUG)
-
-        console_handler = logging.StreamHandler()
-        if int(os.environ.get('AITER_LOG_MORE', 0)):
-            formatter = logging.Formatter(
-                fmt="[%(name)s %(levelname)s] %(asctime)s.%(msecs)03d - %(process)d:%(processName)s - %(pathname)s:%(lineno)d - %(funcName)s\n%(message)s",
-                datefmt="%Y-%m-%d %H:%M:%S",
-            )
-            console_handler.setFormatter(formatter)
-        console_handler.setLevel(logging.INFO)
-        logger.addHandler(console_handler)
-
-    return logger
-
-
-logger = getLogger()
diff --git a/aiter/configs/tuned_fmoe.csv b/aiter/configs/tuned_fmoe.csv
@@ -0,0 +1,18 @@
+token,model_dim,inter_dim,expert,topk,dtype,q_dtype,q_type,use_g1u1,us,tag,err
+1,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,41.02714000000011,fmoe_stage1_bf16_pertokenFp8_g1u1_32x128,0.1%
+2,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,59.21461000000003,ck_32,0.0%
+4,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,97.44835000000003,ck_32,0.0%
+8,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,96.15475000000002,ck_32,0.0%
+16,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,135.1984999999999,ck_32,0.0%
+32,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,137.15732999999992,ck_32,0.0%
+64,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,145.41974000000016,ck_32,0.0%
+128,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,169.0600200000003,fmoe_stage1_bf16_pertokenFp8_g1u1_48x128,0.1%
+256,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,261.0999800000003,fmoe_stage1_bf16_pertokenFp8_g1u1_48x128,0.1%
+512,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,381.25242999999966,fmoe_stage1_bf16_pertokenFp8_g1u1_48x128,0.1%
+1024,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,674.98728,fmoe_stage1_bf16_pertokenFp8_g1u1_128x128,0.1%
+1536,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,925.0758300000005,fmoe_stage1_bf16_pertokenFp8_g1u1_128x128,0.1%
+2048,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,1258.0477299999995,fmoe_stage1_bf16_pertokenFp8_g1u1_128x128,0.1%
+3072,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,1759.5858300000023,fmoe_stage1_bf16_pertokenFp8_g1u1_128x128,0.1%
+4096,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,2345.5597700000003,fmoe_stage1_bf16_pertokenFp8_g1u1_128x128,0.1%
+6144,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,3356.622479999999,fmoe_stage1_bf16_pertokenFp8_g1u1_128x128,0.1%
+8192,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1,4446.33572,fmoe_stage1_bf16_pertokenFp8_g1u1_128x128,0.1%
diff --git a/aiter/configs/untuned_fmoe.csv b/aiter/configs/untuned_fmoe.csv
@@ -0,0 +1,18 @@
+token,model_dim,inter_dim,expert,topk,dtype,q_dtype,q_type,use_g1u1
+1,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+2,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+4,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+8,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+16,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+32,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+64,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+128,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+256,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+512,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+1024,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+1536,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+2048,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+3072,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+4096,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+6144,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+8192,6144,4096,8,2,torch.bfloat16,torch.float8_e4m3fnuz,QuantType.per_Token,1
+3 −3		CHANGELOG.md
+52 −7		Jenkinsfile
+3 −0		client_example/10_grouped_convnd_bwd_data/CMakeLists.txt
+3 −3		client_example/10_grouped_convnd_bwd_data/README.md
+205 −0		client_example/10_grouped_convnd_bwd_data/grouped_conv2d_bwd_data_ngchw.cpp
+3 −3		client_example/11_grouped_conv_bwd_weight/README.md
+6 −6		example/01_gemm/CMakeLists.txt
+10 −1		example/09_convnd_fwd/CMakeLists.txt
+18 −2		example/15_grouped_gemm/run_grouped_gemm_example.inc
+2 −2		example/65_gemm_multiply_multiply/CMakeLists.txt
+3 −3		example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py
+1 −1		example/ck_tile/01_fmha/codegen/ops/fmha_fwd.py
+7 −2		example/ck_tile/01_fmha/codegen/ops/fmha_fwd_splitkv.py
+0 −0		example/ck_tile/03_gemm/gemm_basic.cpp
+14 −0		example/ck_tile/03_gemm/script/benchmark_basic_bf16.sh
+14 −0		example/ck_tile/03_gemm/script/benchmark_basic_bf8.sh
+3 −0		example/ck_tile/03_gemm/script/run_full_test.sh
+2 −8		example/ck_tile/15_fused_moe/instances/fused_moe_api.cpp
+250 −53		example/ck_tile/16_batched_gemm/batched_gemm.cpp
+32 −8		example/ck_tile/16_batched_gemm/batched_gemm.hpp
+0 −1		example/ck_tile/16_batched_gemm/run_batched_gemm_example.inc
+253 −101		example/ck_tile/17_grouped_gemm/grouped_gemm.cpp
+26 −2		example/ck_tile/17_grouped_gemm/grouped_gemm.hpp
+6 −3		example/ck_tile/17_grouped_gemm/run_grouped_gemm_example.inc
+397 −40		include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp
+7 −0		include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp
+19 −0		include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp
+25 −2		include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
+26 −3		include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp
+61 −9		include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_dl.hpp
+49 −8		include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_xdl.hpp
+51 −9		include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_xdl_splitk_cshuffle.hpp
+4 −3		include/ck/tensor_operation/gpu/grid/gridwise_elementwise_2d.hpp
+7 −10		include/ck/tensor_operation/gpu/grid/gridwise_moe_gemm.hpp
+4 −4		include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v7r3_scatter.hpp
+19 −13		include/ck/tensor_operation/operator_transform/transform_conv_ngchw_to_nhwgc.hpp
+1 −1		include/ck/utility/amd_ck_fp8.hpp
+4 −0		include/ck/utility/data_type.hpp
+1 −2		include/ck_tile/host/kernel_launch.hpp
+1 −0		include/ck_tile/ops/common/utils.hpp
+7 −7		include/ck_tile/ops/fmha/kernel/fmha_bwd_kernel.hpp
+3 −3		include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp
+3 −3		include/ck_tile/ops/fmha/kernel/fmha_fwd_splitkv_combine_kernel.hpp
+4 −3		include/ck_tile/ops/fmha/kernel/fmha_fwd_splitkv_kernel.hpp
+2 −2		include/ck_tile/ops/gemm/kernel/batched_gemm_kernel.hpp
+14 −14		include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp
+12 −33		include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp
+23 −12		include/ck_tile/ops/gemm/pipeline/gemm_pipeline_agmem_bgmem_creg_v1.hpp
+2 −2		include/ck_tile/ops/gemm/pipeline/gemm_pipeline_agmem_bgmem_creg_v1_default_policy.hpp
+144 −0		...tensor_operation_instance/gpu/grouped_conv_bwd_data/device_grouped_conv_bwd_data_transpose_xdl_instance.hpp
+2 −0		...or_operation_instance/gpu/grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_two_stage_xdl_instance.hpp
+61 −1		library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_backward_data.hpp
+91 −0		library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_backward_data_xdl.inc
+3 −0		library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/CMakeLists.txt
+48 −0		...ance/gpu/grouped_conv2d_bwd_data/xdl/device_grouped_conv2d_bwd_data_xdl_ngchw_gkyxc_ngkhw_bf16_instance.cpp
+48 −0		...tance/gpu/grouped_conv2d_bwd_data/xdl/device_grouped_conv2d_bwd_data_xdl_ngchw_gkyxc_ngkhw_f16_instance.cpp
+48 −0		...tance/gpu/grouped_conv2d_bwd_data/xdl/device_grouped_conv2d_bwd_data_xdl_ngchw_gkyxc_ngkhw_f32_instance.cpp
+3 −0		library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/CMakeLists.txt
+49 −0		...e/gpu/grouped_conv3d_bwd_data/xdl/device_grouped_conv3d_bwd_data_xdl_ngcdhw_gkzyxc_ngkdhw_bf16_instance.cpp
+49 −0		...ce/gpu/grouped_conv3d_bwd_data/xdl/device_grouped_conv3d_bwd_data_xdl_ngcdhw_gkzyxc_ngkdhw_f16_instance.cpp
+49 −0		...ce/gpu/grouped_conv3d_bwd_data/xdl/device_grouped_conv3d_bwd_data_xdl_ngcdhw_gkzyxc_ngkdhw_f32_instance.cpp
+6 −1		profiler/include/profiler/profile_grouped_conv_bwd_data_impl.hpp
+33 −1		profiler/src/profile_grouped_conv_bwd_data.cpp
+2 −1		script/convert_miopen_driver_to_profiler.py
+2 −2		test/ck_tile/batched_gemm/test_batched_gemm_ut_cases.inc
+128 −64		test/ck_tile/batched_gemm/test_batched_gemm_util.hpp
+3 −3		test/ck_tile/grouped_gemm/test_grouped_gemm_ut_cases.inc
+135 −83		test/ck_tile/grouped_gemm/test_grouped_gemm_util.hpp
+7 −1		test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_xdl.cpp