diff --git a/python/freetoken/moe/fused.py b/python/freetoken/moe/fused.py index a90aa4c79..5c34fb33c 100644 --- a/python/freetoken/moe/fused.py +++ b/python/freetoken/moe/fused.py @@ -84,7 +84,18 @@ def moe_align_block_size( - The padding ensures that the total number of tokens is now divisible by block_size for proper block matrix operations. """ - from freetoken.kernel.backend import is_sgl_kernel_installed + from freetoken.kernel.backend import is_rocm, is_sgl_kernel_installed + + if is_rocm(): + import triton + from packaging.version import Version + + # FIXME: remove once Triton 3.8 is unsupported. Its AMD backend folds compares on tl.histogram results to false. + # https://github.com/triton-lang/triton/pull/11246 + if Version(triton.__version__).release[:2] == (3, 8): + from freetoken.kernel import moe_align_block_size_triton + + return moe_align_block_size_triton(topk_ids, block_size, num_experts) if not is_sgl_kernel_installed(): from freetoken.kernel.triton.moe_align import (