Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion 3rdparty/composable_kernel
Submodule composable_kernel updated 83 files
+13 −0 CMakeLists.txt
+82 −0 dispatcher/bindings/ctypes/CMakeLists.txt
+422 −0 dispatcher/bindings/ctypes/grouped_gemm_abquant_ctypes_lib.cpp
+362 −0 dispatcher/bindings/ctypes/grouped_gemm_aquant_ctypes_lib.cpp
+137 −0 dispatcher/codegen/codegen_common.py
+2 −0 dispatcher/codegen/unified_gemm_codegen.py
+664 −0 dispatcher/codegen/unified_grouped_gemm_abquant_codegen.py
+632 −0 dispatcher/codegen/unified_grouped_gemm_aquant_codegen.py
+312 −0 dispatcher/examples/gemm/python/14_grouped_gemm_aquant.py
+348 −0 dispatcher/examples/gemm/python/15_grouped_gemm_abquant.py
+57 −44 dispatcher/python/gemm_utils.py
+859 −0 dispatcher/python/grouped_gemm_abquant_utils.py
+787 −0 dispatcher/python/grouped_gemm_aquant_utils.py
+8 −6 dispatcher/python/grouped_gemm_bquant_utils.py
+14 −0 dispatcher/tests/CMakeLists.txt
+49 −0 dispatcher/tests/test_gemm_utils.py
+363 −0 dispatcher/tests/test_grouped_gemm_abquant_utils.py
+410 −0 dispatcher/tests/test_grouped_gemm_aquant_utils.py
+6 −0 example/65_gemm_multiply_multiply/CMakeLists.txt
+10 −3 example/65_gemm_multiply_multiply/moe_gemm1_xdl_fp8.cpp
+3 −1 example/65_gemm_multiply_multiply/moe_gemm1_xdl_fp8_blockscale.cpp
+6 −4 example/65_gemm_multiply_multiply/moe_gemm1_xdl_fp8_blockscale_splitk.cpp
+5 −4 example/67_gemm_microscaling/moe_gemm1_xdl_mx_fp4.cpp
+5 −4 example/67_gemm_microscaling/moe_gemm1_xdl_mx_fp4_bns.cpp
+5 −4 example/67_gemm_microscaling/moe_gemm1_xdl_mx_fp4_bpreshuffle.cpp
+2 −6 example/ck_tile/01_fmha/codegen/ops/fmha_fwd.py
+5 −0 example/ck_tile/18_flatmm/CMakeLists.txt
+596 −0 example/ck_tile/18_flatmm/mxgemm/mx_flatmm_data_cache_prefetch.cpp
+17 −0 include/ck/ck.hpp
+2 −1 include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp
+1 −1 include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4r2.hpp
+78 −0 include/ck/tensor_operation/gpu/grid/gridwise_moe_gemm.hpp
+46 −6 include/ck/tensor_operation/gpu/grid/gridwise_moe_gemm_blockscale.hpp
+37 −0 include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm.hpp
+36 −0 include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm_bns.hpp
+38 −0 include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm_bpreshuffle.hpp
+22 −3 include/ck_tile/core/arch/arch.hpp
+37 −0 include/ck_tile/core/config.hpp
+11 −12 include/ck_tile/core/tensor/tile_window.hpp
+2 −2 include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_tdm_v1.hpp
+14 −1 library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm.hpp
+12 −1 library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm1_blockscale.hpp
+7 −1 library/include/ck/library/reference_tensor_operation/cpu/reference_moe_mx_gemm1.hpp
+9 −1 library/src/tensor_operation_instance/gpu/CMakeLists.txt
+1 −1 script/tools/ck-docker
+8 −2 script/tools/ck-exec
+1 −1 script/tools/ck-rocprof.md
+5 −2 script/tools/ck-shell
+1 −1 script/tools/ck-start
+1 −1 script/tools/ck-status
+37 −0 script/tools/common.sh
+119 −0 script/tools/test_ck_smi_helpers.sh
+8 −8 test/CMakeLists.txt
+1 −1 test/ck_tile/add_rmsnorm2d_rdquant/CMakeLists.txt
+1 −2 test/ck_tile/batched_gemm/CMakeLists.txt
+1 −1 test/ck_tile/container/CMakeLists.txt
+1 −1 test/ck_tile/core/arch/CMakeLists.txt
+8 −8 test/ck_tile/core/arch/mma/CMakeLists.txt
+1 −0 test/ck_tile/core/arch/mma/get_cmake_targets_helper.hpp
+23 −0 test/ck_tile/core/arch/test_arch.cpp
+1 −1 test/ck_tile/data_type/CMakeLists.txt
+1 −1 test/ck_tile/elementwise/CMakeLists.txt
+1 −1 test/ck_tile/grouped_conv/CMakeLists.txt
+1 −1 test/ck_tile/grouped_gemm/CMakeLists.txt
+1 −1 test/ck_tile/image_to_column/CMakeLists.txt
+1 −1 test/ck_tile/layernorm2d/CMakeLists.txt
+1 −1 test/ck_tile/moe_smoothquant/CMakeLists.txt
+1 −1 test/ck_tile/moe_sorting/CMakeLists.txt
+1 −1 test/ck_tile/permute/CMakeLists.txt
+1 −1 test/ck_tile/pooling/CMakeLists.txt
+1 −1 test/ck_tile/reduce/CMakeLists.txt
+1 −1 test/ck_tile/rmsnorm2d/CMakeLists.txt
+1 −1 test/ck_tile/smoothquant/CMakeLists.txt
+1 −1 test/ck_tile/topk_softmax/CMakeLists.txt
+4 −4 test/prefetch_op/prefetch_op_util.hpp
+15 −11 test_data/generate_test_dataset.sh
+66 −0 tile_engine/ops/common/smi_cli.py
+224 −0 tile_engine/ops/common/smi_utils.py
+95 −0 tile_engine/ops/common/test_smi_utils.py
+4 −25 tile_engine/ops/gemm/block_scale_gemm/gemm_bquant/gemm_bquant_full_benchmark.py
+14 −33 tile_engine/ops/gemm/gemm_full_benchmark.py
+4 −27 tile_engine/ops/gemm/gemm_multi_d_full_benchmark.py
+4 −27 tile_engine/ops/gemm/streamk_gemm_full_benchmark.py
1 change: 1 addition & 0 deletions aiter/ops/quant.py
Original file line number Diff line number Diff line change
Expand Up @@ -682,6 +682,7 @@ def get_torch_act(aType):
ActivationType.No: lambda *a, **k: a[0],
ActivationType.Silu: F.silu,
ActivationType.Gelu: F.gelu,
ActivationType.GeluTanh: lambda x: F.gelu(x, approximate="tanh"),
}
return tmp.get(aType, NotImplementedError)

Expand Down
8 changes: 7 additions & 1 deletion aiter/utility/dtypes.py
Original file line number Diff line number Diff line change
Expand Up @@ -163,5 +163,11 @@ def _convert(s):


def str2ActivationType(s):
"""Convert string to ActivationType."""
members = getattr(ActivationType, "__members__", None)
if members is not None:
s_lower = s.lower()
for name, member in members.items():
if name.lower() == s_lower:
return member
raise argparse.ArgumentTypeError(f"invalid activation type: {s}")
return getattr(ActivationType, s.capitalize())
Comment on lines 165 to 173
5 changes: 4 additions & 1 deletion csrc/ck_gemm_moe_2stages_codegen/gemm_moe_ck2stages.cu
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,8 @@
using MoeKernelMap = std::unordered_map<std::string, MoeKernel>;

// Map aiter ActivationType (stage1) to the CK stage1 activation op value.
// CK act values: gelu=0, silu=1, swiglu(OAI swiglu_oai)=3. Value 2 is reserved.
// CK act values: gelu=0, silu=1, swiglu(OAI swiglu_oai)=3, gelu_tanh_and_mul=4.
// Value 2 is reserved.
static inline int map_activation_to_ck_stage1(int activation)
{
switch (static_cast<ActivationType>(activation))
Expand All @@ -25,6 +26,8 @@ static inline int map_activation_to_ck_stage1(int activation)
return 0;
case ActivationType::Swiglu:
return 3;
case ActivationType::GeluTanh:
return 4;
default:
TORCH_CHECK(false, "Unsupported activation for ck_moe_stage1: ", activation);
return -1;
Expand Down
5 changes: 3 additions & 2 deletions csrc/ck_gemm_moe_2stages_codegen/gemm_moe_ck2stages_common.py
Original file line number Diff line number Diff line change
Expand Up @@ -18,8 +18,9 @@
from chip_info import get_gfx

# CK stage1 activation op values. swiglu (OAI swiglu_oai) maps to CK act value 3;
# value 2 is intentionally skipped (reserved for a future activation variant).
ACT_OP_MAP = {"gelu": 0, "silu": 1, "swiglu": 3}
# gelutanh (gelu_tanh_and_mul) maps to CK act value 4; value 2 is intentionally
# skipped (reserved for a future activation variant).
ACT_OP_MAP = {"gelu": 0, "silu": 1, "swiglu": 3, "gelutanh": 4}
Comment on lines 20 to +23
ACT_OP_NAME = {v: k for k, v in ACT_OP_MAP.items()}


Expand Down
37 changes: 36 additions & 1 deletion csrc/ck_gemm_moe_2stages_codegen/gen_instances.py
Original file line number Diff line number Diff line change
Expand Up @@ -1074,7 +1074,7 @@ def generate_instance_and_lookUpTable(self):
default="silu",
required=False,
type=str,
choices=["silu", "gelu", "swiglu"],
choices=["silu", "gelu", "swiglu", "gelutanh"],
help="select activation",
)

Expand Down Expand Up @@ -1243,6 +1243,41 @@ def generate_instance_and_lookUpTable(self):
False, # splitk
)
codegen.generate_instance_and_lookUpTable()

# gelutanh (gelu_tanh_and_mul) plain-f8 quant moe (per_tensor / per_token = PTPC).
# Same rationale as the swiglu loop above: the general quant loop uses
# acts=["silu","gelu"] and so ships 0 gelutanh x plain-f8 instances in the
# AOT/wheel prebuild. Runtime JIT (Path B, gen_func passes -b f8) already
# generates them on demand, so this loop only extends the AOT prebuild set
# and removes the ~first-run cold compile -- it is NOT required for
# correctness. Mirrors the general loop's plain-f8 path (f8 x f8, tag=a8w8,
# CDEElementOp=MulABScale, plain gridwise_moe_gemm.hpp). act="gelutanh" ->
# CK ActOP=4. preshuffle forced True for plain f8. Targets Gemma-family MoE
# experts that use the tanh-approx GELU. Independent loop so existing
# silu/gelu/swiglu/no-quant coverage is untouched.
gelutanh_plain_c_dtypes = ["f16", "b16"]
gelutanh_plain_quant_l = ["per_tensor", "per_token"]
for (
c_dtype,
routed_weight,
quant,
) in itertools.product(
gelutanh_plain_c_dtypes,
routed_weight_l,
gelutanh_plain_quant_l,
):
codegen = ck_moe_2stage_gemm_codegen(
args.working_path,
"f8", # a_dtype
"f8", # b_dtype
c_dtype,
quant_dict[quant],
"gelutanh",
routed_weight,
True, # preshuffle (plain f8 forced True, mirrors general loop)
False, # splitk
)
codegen.generate_instance_and_lookUpTable()
else:
for b_dtype in args.b_dtype:
a_dtype = b_dtype if b_dtype != "i4" else "f8"
Expand Down
11 changes: 6 additions & 5 deletions csrc/include/aiter_enum.h
Original file line number Diff line number Diff line change
Expand Up @@ -6,11 +6,12 @@

enum class ActivationType : int
{
No = -1,
Silu = 0,
Gelu = 1,
Swiglu = 2,
Situv2 = 3,
No = -1,
Silu = 0,
Gelu = 1,
Swiglu = 2,
Situv2 = 3,
GeluTanh = 4,
};

enum class QuantType : int
Expand Down
1 change: 1 addition & 0 deletions csrc/include/rocm_ops.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,7 @@ namespace py = pybind11;
.value("Gelu", ActivationType::Gelu) \
.value("Swiglu", ActivationType::Swiglu) \
.value("Situv2", ActivationType::Situv2) \
.value("GeluTanh", ActivationType::GeluTanh) \
.export_values(); \
pybind11::enum_<MlaVersion>(m, "MlaVersion") \
.value("V32", MlaVersion::V32) \
Expand Down
Loading