diff --git a/src/infiniop/ops/dequantize_gptq/metax/dequantize_w42f16_metax.h b/src/infiniop/ops/dequantize_gptq/metax/dequantize_w42f16_metax.h
new file mode 100644
index 000000000..255f61241
--- /dev/null
+++ b/src/infiniop/ops/dequantize_gptq/metax/dequantize_w42f16_metax.h
@@ -0,0 +1,8 @@
+#ifndef __DEQUANTIZE_GPTQ_METAX_H__
+#define __DEQUANTIZE_GPTQ_METAX_H__
+
+#include "../dequantize_gptq.h"
+
+DESCRIPTOR(metax)
+
+#endif // __DEQUANTIZE_GPTQ_METAX_H__
diff --git a/src/infiniop/ops/dequantize_gptq/metax/dequantize_w42f16_metax.maca b/src/infiniop/ops/dequantize_gptq/metax/dequantize_w42f16_metax.maca
new file mode 100644
index 000000000..5a0830894
--- /dev/null
+++ b/src/infiniop/ops/dequantize_gptq/metax/dequantize_w42f16_metax.maca
@@ -0,0 +1,136 @@
+#ifdef ENABLE_METAX_MC_API
+#include <mc_runtime.h>
+#else
+#include <hc_runtime.h>
+#endif
+
+#if defined(ENABLE_METAX_API) || defined(ENABLE_QY_API)
+
+#include "../../../devices/metax/metax_common.h"
+#include "../../../devices/metax/metax_handle.h"
+#include "../../../devices/metax/metax_kernel_common.h"
+#include "../dequantize_gptq.h"
+#include "dequantize_w42f16_metax.h"
+
+namespace op::dequantize_gptq::metax {
+
+struct Descriptor::Opaque {
+    std::shared_ptr<device::metax::Handle::Internal> internal;
+};
+
+Descriptor::~Descriptor() { delete _opaque; }
+
+infiniStatus_t Descriptor::create(
+    infiniopHandle_t handle_,
+    Descriptor **desc_ptr,
+    infiniopTensorDescriptor_t out_desc,
+    infiniopTensorDescriptor_t qweight_desc,
+    infiniopTensorDescriptor_t scales_desc,
+    infiniopTensorDescriptor_t zeros_desc,
+    infiniopTensorDescriptor_t g_idx_desc) {
+
+    auto handle = reinterpret_cast<device::metax::Handle *>(handle_);
+    auto result = DequantizeGPTQInfo::create(out_desc, qweight_desc, scales_desc, zeros_desc, g_idx_desc);
+
+    *desc_ptr = new Descriptor(
+        0,
+        new Opaque{handle->internal()},
+        result.take(),
+        handle->device, handle->device_id);
+    return INFINI_STATUS_SUCCESS;
+}
+
+// 仅保留这一版内核（支持 g_idx）
+// qweight: [in_packed, out_features] packing 8 input channels per word
+// zeros: [num_groups, out_packed] packing 8 output channels per word
+// scales: [num_groups, out_features], g_idx: [in_features]
+__global__ void __launch_bounds__(128)
+    dequantize_weights_gptq(const uint32_t *__restrict__ qweight,
+                            const half *__restrict__ scales,
+                            const uint32_t *__restrict__ zeros,
+                            const int *__restrict__ g_idx,
+                            half *__restrict__ out,
+                            int in_features,
+                            int out_features,
+                            int out_packed, // ceil(out_features / 8)
+                            int num_groups) {
+    // Each thread handles one packed output column (8 real output cols).
+    const int col_pack = blockIdx.x * blockDim.x + threadIdx.x; // packed output column
+    const int row = blockIdx.y * blockDim.y + threadIdx.y;      // real input row
+    if (col_pack >= out_packed || row >= in_features) {
+        return;
+    }
+
+    // Clamp gid to valid range
+    const int gid_raw = g_idx ? g_idx[row] : 0;
+    const int gid = ((gid_raw % num_groups) + num_groups) % num_groups;
+
+    const int pack_row = row >> 3; // packed input row
+
+    const int zero_idx = gid * out_packed + col_pack; // zeros layout: [num_groups, out_packed]
+    const uint32_t zeros_loaded = zeros[zero_idx];
+
+    const int q_shift = (row & 7) * 4;  // qweight packs 8 input rows
+    const int col_base = col_pack << 3; // 8 real cols per pack
+    const int scale_base = gid * out_features + col_base;
+
+#pragma unroll
+    for (int j = 0; j < 8; ++j) {
+        const int col = col_base + j;
+        if (col >= out_features) {
+            break;
+        }
+
+        const uint32_t q_loaded = qweight[pack_row * out_features + col];
+        const int q_nib = (q_loaded >> q_shift) & 0xF;
+
+        const int z_nib = (zeros_loaded >> (j * 4)) & 0xF;
+        const half scale = scales[scale_base + j];
+
+        // GPTQ quirk: The stored zero point is usually offset by 1.
+        // Standard formula: (q - (z + 1)) * s
+        const float v = float(q_nib - (z_nib + 1)) * __half2float(scale);
+        out[row * out_features + col] = __float2half(v);
+    }
+}
+
+infiniStatus_t
+Descriptor::calculate(
+    void *workspace, size_t workspace_size,
+    void *out,
+    const void *qweight,
+    const void *scales,
+    const void *zeros,
+    const void *g_idx,
+    void *stream) const {
+
+    const int in_features = _info.in_features();
+    const int out_features = _info.out_features();
+    const int out_packed = _info.out_packed();
+    const int in_packed = _info.in_packed();
+    const int num_groups = _info.num_groups();
+
+    if (num_groups <= 0 || in_features <= 0 || out_features <= 0 || out_packed <= 0 || in_packed <= 0) {
+        return INFINI_STATUS_BAD_PARAM;
+    }
+
+    constexpr int BLOCK_X = 16; // packed columns
+    constexpr int BLOCK_Y = 4;  // rows
+    dim3 threads(BLOCK_X, BLOCK_Y);
+    dim3 blocks((out_packed + BLOCK_X - 1) / BLOCK_X,
+                (in_features + BLOCK_Y - 1) / BLOCK_Y);
+
+    dequantize_weights_gptq<<<blocks, threads, 0,
+                              reinterpret_cast<hcStream_t>(stream)>>>(
+        reinterpret_cast<const uint32_t *>(qweight),
+        reinterpret_cast<const half *>(scales),
+        reinterpret_cast<const uint32_t *>(zeros),
+        reinterpret_cast<const int *>(g_idx),
+        reinterpret_cast<half *>(out),
+        in_features, out_features, out_packed, num_groups);
+    return INFINI_STATUS_SUCCESS;
+}
+
+} // namespace op::dequantize_gptq::metax
+
+#endif
diff --git a/src/infiniop/ops/dequantize_gptq/operator.cc b/src/infiniop/ops/dequantize_gptq/operator.cc
index d48430445..047f37b65 100644
--- a/src/infiniop/ops/dequantize_gptq/operator.cc
+++ b/src/infiniop/ops/dequantize_gptq/operator.cc
@@ -5,6 +5,10 @@
 #if defined(ENABLE_NVIDIA_API) || defined(ENABLE_QY_API)
 #include "nvidia/dequantize_w42f16_nvidia.cuh"
 #endif
+
+#ifdef ENABLE_METAX_API
+#include "metax/dequantize_w42f16_metax.h"
+#endif
 #ifdef ENABLE_MOORE_API
 #include "moore/dequantize_w42f16_moore.h"
 #endif
@@ -34,6 +38,9 @@ __INFINI_C infiniStatus_t infiniopCreateDequantizeGPTQDescriptor(
 #ifdef ENABLE_NVIDIA_API
         CREATE(INFINI_DEVICE_NVIDIA, nvidia);
 #endif
+#ifdef ENABLE_METAX_API
+        CREATE(INFINI_DEVICE_METAX, metax);
+#endif
 #ifdef ENABLE_MOORE_API
         CREATE(INFINI_DEVICE_MOORE, moore);
 #endif
@@ -61,6 +68,9 @@ __INFINI_C infiniStatus_t infiniopGetDequantizeGPTQWorkspaceSize(infiniopDequant
 #ifdef ENABLE_NVIDIA_API
         GET(INFINI_DEVICE_NVIDIA, nvidia);
 #endif
+#ifdef ENABLE_METAX_API
+        GET(INFINI_DEVICE_METAX, metax);
+#endif
 #ifdef ENABLE_MOORE_API
         GET(INFINI_DEVICE_MOORE, moore);
 #endif
@@ -96,6 +106,9 @@ __INFINI_C infiniStatus_t infiniopDequantizeGPTQ(
 #ifdef ENABLE_NVIDIA_API
         CALCULATE(INFINI_DEVICE_NVIDIA, nvidia);
 #endif
+#ifdef ENABLE_METAX_API
+        CALCULATE(INFINI_DEVICE_METAX, metax);
+#endif
 #ifdef ENABLE_MOORE_API
         CALCULATE(INFINI_DEVICE_MOORE, moore);
 #endif
@@ -124,6 +137,9 @@ infiniopDestroyDequantizeGPTQDescriptor(infiniopDequantizeGPTQDescriptor_t desc)
 #ifdef ENABLE_NVIDIA_API
         DELETE(INFINI_DEVICE_NVIDIA, nvidia);
 #endif
+#ifdef ENABLE_METAX_API
+        DELETE(INFINI_DEVICE_METAX, metax);
+#endif
 #ifdef ENABLE_MOORE_API
         DELETE(INFINI_DEVICE_MOORE, moore);
 #endif
diff --git a/src/infiniop/ops/mamba_selective_scan/metax/mamba_selective_scan_metax.h b/src/infiniop/ops/mamba_selective_scan/metax/mamba_selective_scan_metax.h
new file mode 100644
index 000000000..0a690d942
--- /dev/null
+++ b/src/infiniop/ops/mamba_selective_scan/metax/mamba_selective_scan_metax.h
@@ -0,0 +1,5 @@
+#ifndef __MAMBA_SELECTIVE_SCAN_METAX_H__
+#define __MAMBA_SELECTIVE_SCAN_METAX_H__
+#include "../mamba_selective_scan.h"
+DESCRIPTOR(metax)
+#endif
diff --git a/src/infiniop/ops/mamba_selective_scan/metax/mamba_selective_scan_metax.maca b/src/infiniop/ops/mamba_selective_scan/metax/mamba_selective_scan_metax.maca
new file mode 100644
index 000000000..3d8cad013
--- /dev/null
+++ b/src/infiniop/ops/mamba_selective_scan/metax/mamba_selective_scan_metax.maca
@@ -0,0 +1,118 @@
+#include "../../../devices/metax/metax_common.h"
+#include "../../../devices/metax/metax_kernel_common.h"
+#include "mamba_selective_scan_metax.h"
+#ifdef ENABLE_METAX_MC_API
+#include <mc_runtime.h>
+#else
+#include <hc_runtime.h>
+#endif
+
+namespace op::mamba_selective_scan::metax {
+
+struct Descriptor::Opaque {
+    std::shared_ptr<device::metax::Handle::Internal> internal;
+};
+Descriptor::~Descriptor() { delete _opaque; }
+
+namespace {
+template <typename T>
+__device__ float to_float(T v) { return static_cast<float>(v); }
+template <>
+__device__ float to_float<half>(half v) { return __half2float(v); }
+template <>
+__device__ float to_float<__nv_bfloat16>(__nv_bfloat16 v) { return __bfloat162float(v); }
+template <typename T>
+__device__ T from_float(float v) { return static_cast<T>(v); }
+template <>
+__device__ half from_float<half>(float v) { return __float2half_rn(v); }
+template <>
+__device__ __nv_bfloat16 from_float<__nv_bfloat16>(float v) { return __float2bfloat16_rn(v); }
+__device__ float softplusf_stable(float x) { return x > 20.0f ? x : log1pf(expf(x)); }
+__device__ float siluf(float x) { return x / (1.0f + expf(-x)); }
+
+template <typename T>
+__global__ void mamba_selective_scan_kernel(
+    T *out, const T *x, const T *dt, const T *b, const T *c, const T *a_log,
+    const T *d, const T *gate, const T *dt_bias, float *state,
+    size_t batch, size_t seq_len, size_t intermediate, size_t state_size) {
+    size_t linear = blockIdx.x * blockDim.x + threadIdx.x;
+    size_t total = batch * intermediate;
+    if (linear >= total) {
+        return;
+    }
+    size_t ch = linear % intermediate;
+    size_t batch_idx = linear / intermediate;
+    size_t state_base = (batch_idx * intermediate + ch) * state_size;
+    for (size_t t = 0; t < seq_len; ++t) {
+        size_t x_idx = (batch_idx * seq_len + t) * intermediate + ch;
+        float xt = to_float<T>(x[x_idx]);
+        float dtv = softplusf_stable(to_float<T>(dt[x_idx]) + to_float<T>(dt_bias[ch]));
+        float y = 0.0f;
+        for (size_t n = 0; n < state_size; ++n) {
+            float a = -expf(to_float<T>(a_log[ch * state_size + n]));
+            float discrete_a = expf(a * dtv);
+            float bn = to_float<T>(b[(batch_idx * seq_len + t) * state_size + n]);
+            float cn = to_float<T>(c[(batch_idx * seq_len + t) * state_size + n]);
+            float s = discrete_a * state[state_base + n] + dtv * bn * xt;
+            state[state_base + n] = s;
+            y += s * cn;
+        }
+        y += xt * to_float<T>(d[ch]);
+        y *= siluf(to_float<T>(gate[x_idx]));
+        out[x_idx] = from_float<T>(y);
+    }
+}
+
+template <typename T>
+infiniStatus_t launch_typed(const MambaSelectiveScanInfo &info, void *out, const void *x,
+                            const void *dt, const void *b, const void *c, const void *a_log,
+                            const void *d, const void *gate, const void *dt_bias, void *state,
+                            hcStream_t stream) {
+    constexpr int threads = 256;
+    size_t total = info.batch * info.intermediate;
+    dim3 blocks((total + threads - 1) / threads);
+    mamba_selective_scan_kernel<T><<<blocks, threads, 0, stream>>>(
+        reinterpret_cast<T *>(out), reinterpret_cast<const T *>(x), reinterpret_cast<const T *>(dt),
+        reinterpret_cast<const T *>(b), reinterpret_cast<const T *>(c), reinterpret_cast<const T *>(a_log),
+        reinterpret_cast<const T *>(d), reinterpret_cast<const T *>(gate), reinterpret_cast<const T *>(dt_bias),
+        reinterpret_cast<float *>(state), info.batch, info.seq_len, info.intermediate, info.state_size);
+    CHECK_METAX(hcGetLastError());
+    return INFINI_STATUS_SUCCESS;
+}
+} // namespace
+
+infiniStatus_t Descriptor::create(infiniopHandle_t handle, Descriptor **desc_ptr,
+                                  infiniopTensorDescriptor_t out_desc, infiniopTensorDescriptor_t x_desc,
+                                  infiniopTensorDescriptor_t dt_desc, infiniopTensorDescriptor_t b_desc,
+                                  infiniopTensorDescriptor_t c_desc, infiniopTensorDescriptor_t a_log_desc,
+                                  infiniopTensorDescriptor_t d_desc, infiniopTensorDescriptor_t gate_desc,
+                                  infiniopTensorDescriptor_t dt_bias_desc, infiniopTensorDescriptor_t state_desc) {
+    auto result = MambaSelectiveScanInfo::create(out_desc, x_desc, dt_desc, b_desc, c_desc, a_log_desc, d_desc, gate_desc, dt_bias_desc, state_desc);
+    CHECK_RESULT(result);
+    auto info = result.take();
+    *desc_ptr = new Descriptor(new Opaque{reinterpret_cast<device::metax::Handle *>(handle)->internal()}, info, 0, handle->device, handle->device_id);
+    return INFINI_STATUS_SUCCESS;
+}
+
+infiniStatus_t Descriptor::calculate(void *workspace, size_t workspace_size, void *out,
+                                     const void *x, const void *dt, const void *b, const void *c,
+                                     const void *a_log, const void *d, const void *gate,
+                                     const void *dt_bias, void *state, void *stream_) const {
+    (void)workspace;
+    if (workspace_size < _workspace_size) {
+        return INFINI_STATUS_INSUFFICIENT_WORKSPACE;
+    }
+    auto stream = reinterpret_cast<hcStream_t>(stream_);
+    if (_info.dtype == INFINI_DTYPE_F16) {
+        return launch_typed<half>(_info, out, x, dt, b, c, a_log, d, gate, dt_bias, state, stream);
+    }
+    if (_info.dtype == INFINI_DTYPE_BF16) {
+        return launch_typed<__nv_bfloat16>(_info, out, x, dt, b, c, a_log, d, gate, dt_bias, state, stream);
+    }
+    if (_info.dtype == INFINI_DTYPE_F32) {
+        return launch_typed<float>(_info, out, x, dt, b, c, a_log, d, gate, dt_bias, state, stream);
+    }
+    return INFINI_STATUS_BAD_TENSOR_DTYPE;
+}
+
+} // namespace op::mamba_selective_scan::metax
diff --git a/src/infiniop/ops/mamba_selective_scan/operator.cc b/src/infiniop/ops/mamba_selective_scan/operator.cc
index c7ea98736..ce31335f9 100644
--- a/src/infiniop/ops/mamba_selective_scan/operator.cc
+++ b/src/infiniop/ops/mamba_selective_scan/operator.cc
@@ -5,6 +5,10 @@
 #include "nvidia/mamba_selective_scan_nvidia.cuh"
 #endif
 
+#ifdef ENABLE_METAX_API
+#include "metax/mamba_selective_scan_metax.h"
+#endif
+
 __INFINI_C infiniStatus_t infiniopCreateMambaSelectiveScanDescriptor(
     infiniopHandle_t handle, infiniopMambaSelectiveScanDescriptor_t *desc_ptr,
     infiniopTensorDescriptor_t out_desc, infiniopTensorDescriptor_t x_desc,
@@ -30,6 +34,9 @@ __INFINI_C infiniStatus_t infiniopCreateMambaSelectiveScanDescriptor(
 #endif
 #ifdef ENABLE_HYGON_API
         CREATE(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        CREATE(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -57,6 +64,9 @@ __INFINI_C infiniStatus_t infiniopGetMambaSelectiveScanWorkspaceSize(infiniopMam
 #endif
 #ifdef ENABLE_HYGON_API
         GET(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        GET(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -87,6 +97,9 @@ __INFINI_C infiniStatus_t infiniopMambaSelectiveScan(
 #endif
 #ifdef ENABLE_HYGON_API
         CALC(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        CALC(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -114,6 +127,9 @@ __INFINI_C infiniStatus_t infiniopDestroyMambaSelectiveScanDescriptor(infiniopMa
 #endif
 #ifdef ENABLE_HYGON_API
         DESTROY(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        DESTROY(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
diff --git a/src/infiniop/ops/moe_sum/metax/moe_sum_metax.h b/src/infiniop/ops/moe_sum/metax/moe_sum_metax.h
new file mode 100644
index 000000000..c5512af7f
--- /dev/null
+++ b/src/infiniop/ops/moe_sum/metax/moe_sum_metax.h
@@ -0,0 +1,8 @@
+#ifndef __MOE_SUM_METAX_H__
+#define __MOE_SUM_METAX_H__
+
+#include "../moe_sum.h"
+
+DESCRIPTOR(metax)
+
+#endif
diff --git a/src/infiniop/ops/moe_sum/metax/moe_sum_metax.maca b/src/infiniop/ops/moe_sum/metax/moe_sum_metax.maca
new file mode 100644
index 000000000..a1c74314c
--- /dev/null
+++ b/src/infiniop/ops/moe_sum/metax/moe_sum_metax.maca
@@ -0,0 +1,127 @@
+#ifdef ENABLE_METAX_MC_API
+#include <mc_runtime.h>
+#else
+#include <hc_runtime.h>
+#endif
+
+#ifdef ENABLE_METAX_API
+
+#include "moe_sum_metax.h"
+
+#include "../../../devices/metax/metax_common.h"
+#include "../../../devices/metax/metax_kernel_common.h"
+
+#include <algorithm>
+#include <memory>
+#include <utility>
+
+namespace op::moe_sum::metax {
+
+struct Descriptor::Opaque {
+    std::shared_ptr<device::metax::Handle::Internal> internal;
+};
+
+Descriptor::~Descriptor() {
+    delete _opaque;
+}
+
+infiniStatus_t Descriptor::create(
+    infiniopHandle_t handle,
+    Descriptor **desc_ptr,
+    infiniopTensorDescriptor_t output_desc,
+    infiniopTensorDescriptor_t input_desc) {
+    auto result = MoeSumInfo::create(output_desc, input_desc);
+    CHECK_RESULT(result);
+    auto info = result.take();
+
+    *desc_ptr = new Descriptor(
+        new Opaque{reinterpret_cast<device::metax::Handle *>(handle)->internal()},
+        std::move(info),
+        0,
+        handle->device,
+        handle->device_id);
+    return INFINI_STATUS_SUCCESS;
+}
+
+namespace {
+
+template <typename scalar_t, int TOPK>
+__global__ void moe_sum_kernel(
+    scalar_t *__restrict__ output,
+    const scalar_t *__restrict__ input,
+    size_t hidden_size) {
+    const size_t token_idx = blockIdx.x;
+    for (size_t idx = threadIdx.x; idx < hidden_size; idx += blockDim.x) {
+        scalar_t x = static_cast<scalar_t>(0.0f);
+#pragma unroll
+        for (int k = 0; k < TOPK; ++k) {
+            x += input[token_idx * TOPK * hidden_size + k * hidden_size + idx];
+        }
+        output[token_idx * hidden_size + idx] = x;
+    }
+}
+
+template <typename scalar_t>
+infiniStatus_t launch_moe_sum(
+    void *output,
+    const void *input,
+    size_t num_tokens,
+    size_t topk,
+    size_t hidden_size,
+    hcStream_t stream) {
+    dim3 grid(static_cast<unsigned int>(num_tokens));
+    dim3 block(static_cast<unsigned int>(std::min<size_t>(hidden_size, 1024)));
+
+    switch (topk) {
+    case 2:
+        moe_sum_kernel<scalar_t, 2><<<grid, block, 0, stream>>>(
+            static_cast<scalar_t *>(output),
+            static_cast<const scalar_t *>(input),
+            hidden_size);
+        return INFINI_STATUS_SUCCESS;
+    case 3:
+        moe_sum_kernel<scalar_t, 3><<<grid, block, 0, stream>>>(
+            static_cast<scalar_t *>(output),
+            static_cast<const scalar_t *>(input),
+            hidden_size);
+        return INFINI_STATUS_SUCCESS;
+    case 4:
+        moe_sum_kernel<scalar_t, 4><<<grid, block, 0, stream>>>(
+            static_cast<scalar_t *>(output),
+            static_cast<const scalar_t *>(input),
+            hidden_size);
+        return INFINI_STATUS_SUCCESS;
+    default:
+        return INFINI_STATUS_BAD_PARAM;
+    }
+}
+
+} // namespace
+
+infiniStatus_t Descriptor::calculate(
+    void *workspace,
+    size_t workspace_size,
+    void *output,
+    const void *input,
+    void *stream) const {
+    (void)workspace;
+    if (workspace_size < _workspace_size) {
+        return INFINI_STATUS_INSUFFICIENT_WORKSPACE;
+    }
+
+    auto cuda_stream = reinterpret_cast<hcStream_t>(stream);
+    switch (_info.dtype) {
+    case INFINI_DTYPE_F16:
+        return launch_moe_sum<half>(output, input, _info.num_tokens, _info.topk, _info.hidden_size, cuda_stream);
+    case INFINI_DTYPE_BF16:
+        return launch_moe_sum<__nv_bfloat16>(output, input, _info.num_tokens, _info.topk, _info.hidden_size, cuda_stream);
+    case INFINI_DTYPE_F32:
+        return launch_moe_sum<float>(output, input, _info.num_tokens, _info.topk, _info.hidden_size, cuda_stream);
+    default:
+        return INFINI_STATUS_BAD_TENSOR_DTYPE;
+    }
+}
+
+} // namespace op::moe_sum::metax
+
+#endif // ENABLE_METAX_API
diff --git a/src/infiniop/ops/moe_sum/operator.cc b/src/infiniop/ops/moe_sum/operator.cc
index 04e75434a..ce8bcc23c 100644
--- a/src/infiniop/ops/moe_sum/operator.cc
+++ b/src/infiniop/ops/moe_sum/operator.cc
@@ -6,6 +6,10 @@
 #include "nvidia/moe_sum_nvidia.cuh"
 #endif
 
+#ifdef ENABLE_METAX_API
+#include "metax/moe_sum_metax.h"
+#endif
+
 __INFINI_C infiniStatus_t infiniopCreateMoeSumDescriptor(
     infiniopHandle_t handle,
     infiniopMoeSumDescriptor_t *desc_ptr,
@@ -23,6 +27,9 @@ __INFINI_C infiniStatus_t infiniopCreateMoeSumDescriptor(
     switch (handle->device) {
 #ifdef ENABLE_NVIDIA_API
         CREATE(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        CREATE(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -43,6 +50,9 @@ __INFINI_C infiniStatus_t infiniopGetMoeSumWorkspaceSize(
     switch (desc->device_type) {
 #ifdef ENABLE_NVIDIA_API
         GET(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        GET(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -67,6 +77,9 @@ __INFINI_C infiniStatus_t infiniopMoeSum(
     switch (desc->device_type) {
 #ifdef ENABLE_NVIDIA_API
         CALCULATE(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        CALCULATE(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -86,6 +99,9 @@ __INFINI_C infiniStatus_t infiniopDestroyMoeSumDescriptor(
     switch (desc->device_type) {
 #ifdef ENABLE_NVIDIA_API
         DESTROY(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        DESTROY(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
diff --git a/src/infiniop/ops/mul_scalar/metax/mul_scalar_metax.h b/src/infiniop/ops/mul_scalar/metax/mul_scalar_metax.h
new file mode 100644
index 000000000..d2efae78a
--- /dev/null
+++ b/src/infiniop/ops/mul_scalar/metax/mul_scalar_metax.h
@@ -0,0 +1,9 @@
+#ifndef __MUL_SCALAR_METAX_H__
+#define __MUL_SCALAR_METAX_H__
+
+#include "../../../elementwise/metax/elementwise_metax_api.h"
+#include "../mul_scalar.h"
+
+MUL_SCALAR_DESCRIPTOR(metax)
+
+#endif // __MUL_SCALAR_METAX_H__
diff --git a/src/infiniop/ops/mul_scalar/metax/mul_scalar_metax.maca b/src/infiniop/ops/mul_scalar/metax/mul_scalar_metax.maca
new file mode 100644
index 000000000..b97a9e72f
--- /dev/null
+++ b/src/infiniop/ops/mul_scalar/metax/mul_scalar_metax.maca
@@ -0,0 +1,150 @@
+#ifdef ENABLE_METAX_MC_API
+#include <mc_runtime.h>
+#else
+#include <hc_runtime.h>
+#endif
+
+#include "mul_scalar_metax.h"
+
+#include "../../../devices/metax/metax_common.h"
+#include "../../../devices/metax/metax_kernel_common.h"
+#include "../../../elementwise/metax/elementwise_metax.h"
+
+#include <algorithm>
+
+namespace {
+
+template <typename T>
+__device__ __forceinline__ T mulScalarValue(T value, double alpha) {
+    return value * static_cast<T>(alpha);
+}
+
+template <>
+__device__ __forceinline__ half mulScalarValue<half>(half value, double alpha) {
+    return __hmul(value, __float2half(static_cast<float>(alpha)));
+}
+
+template <>
+__device__ __forceinline__ cuda_bfloat16 mulScalarValue<cuda_bfloat16>(cuda_bfloat16 value, double alpha) {
+    return __hmul(value, __float2bfloat16_rn(static_cast<float>(alpha)));
+}
+
+template <typename T>
+__global__ void mulScalarKernel(size_t n, T *output, const T *input, double alpha) {
+    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+    size_t stride = blockDim.x * gridDim.x;
+    for (; idx < n; idx += stride) {
+        output[idx] = mulScalarValue(input[idx], alpha);
+    }
+}
+
+template <typename T>
+infiniStatus_t launchMulScalarKernel(size_t numel, void *output, const void *input, double alpha, void *stream) {
+    if (numel == 0) {
+        return INFINI_STATUS_SUCCESS;
+    }
+
+    constexpr int block = 256;
+    int grid = static_cast<int>((numel + block - 1) / block);
+    grid = std::min(grid, 65535);
+    auto cuda_stream = reinterpret_cast<hcStream_t>(stream);
+    mulScalarKernel<T><<<grid, block, 0, cuda_stream>>>(
+        numel,
+        reinterpret_cast<T *>(output),
+        reinterpret_cast<const T *>(input),
+        static_cast<double>(alpha));
+    auto err = hcGetLastError();
+    return err == hcSuccess ? INFINI_STATUS_SUCCESS : INFINI_STATUS_INTERNAL_ERROR;
+}
+
+struct MulScalarOp {
+    static constexpr size_t num_inputs = 1;
+
+    template <typename T>
+    __device__ __forceinline__ T operator()(const T &value, double alpha) const {
+        return mulScalarValue(value, alpha);
+    }
+};
+
+template <typename T>
+infiniStatus_t calculateMulScalar(
+    const MulScalarInfo &info,
+    op::elementwise::metax::DeviceImpl *device_info,
+    size_t workspace_size,
+    void *workspace,
+    void *output,
+    const void *input,
+    double alpha,
+    void *stream) {
+
+    if (info.contiguous) {
+        return launchMulScalarKernel<T>(info.numel(), output, input, alpha, stream);
+    }
+
+    if (workspace_size < info.elementwise_info.getMetaMemSize() + sizeof(void *)) {
+        return INFINI_STATUS_INSUFFICIENT_WORKSPACE;
+    }
+
+    return device_info->calculate<256, MulScalarOp, T>(
+        info.elementwise_info,
+        workspace,
+        output,
+        {input},
+        stream,
+        static_cast<double>(alpha));
+}
+
+} // namespace
+
+namespace op::mul_scalar::metax {
+
+Descriptor::~Descriptor() = default;
+
+infiniStatus_t Descriptor::create(
+    infiniopHandle_t handle_,
+    Descriptor **desc_ptr,
+    infiniopTensorDescriptor_t output_desc,
+    infiniopTensorDescriptor_t input_desc) {
+
+    auto handle = reinterpret_cast<device::metax::Handle *>(handle_);
+    auto result = MulScalarInfo::create(output_desc, input_desc);
+    CHECK_RESULT(result);
+
+    auto info = result.take();
+    auto workspace_size = info.elementwise_info.getMetaMemSize() + sizeof(void *);
+    auto device_impl_result = op::elementwise::metax::DeviceImpl::create(handle->internal());
+    CHECK_RESULT(device_impl_result);
+
+    *desc_ptr = new Descriptor(
+        std::move(info),
+        device_impl_result.take(),
+        workspace_size,
+        handle->device,
+        handle->device_id);
+
+    return INFINI_STATUS_SUCCESS;
+}
+
+infiniStatus_t Descriptor::calculate(
+    void *workspace,
+    size_t workspace_size,
+    void *output,
+    const void *input,
+    double alpha,
+    void *stream) const {
+
+    switch (_info.data_type) {
+    case INFINI_DTYPE_F16:
+        return calculateMulScalar<half>(_info, _device_info.get(), workspace_size, workspace, output, input, alpha, stream);
+    case INFINI_DTYPE_BF16:
+        return calculateMulScalar<cuda_bfloat16>(_info, _device_info.get(), workspace_size, workspace, output, input, alpha, stream);
+    case INFINI_DTYPE_F32:
+        return calculateMulScalar<float>(_info, _device_info.get(), workspace_size, workspace, output, input, alpha, stream);
+    case INFINI_DTYPE_F64:
+        return calculateMulScalar<double>(_info, _device_info.get(), workspace_size, workspace, output, input, alpha, stream);
+    default:
+        return INFINI_STATUS_BAD_TENSOR_DTYPE;
+    }
+}
+
+} // namespace op::mul_scalar::metax
diff --git a/src/infiniop/ops/mul_scalar/operator.cc b/src/infiniop/ops/mul_scalar/operator.cc
index a8f50496f..cd0aeb656 100644
--- a/src/infiniop/ops/mul_scalar/operator.cc
+++ b/src/infiniop/ops/mul_scalar/operator.cc
@@ -9,6 +9,10 @@
 #include "nvidia/mul_scalar_nvidia.cuh"
 #endif
 
+#ifdef ENABLE_METAX_API
+#include "metax/mul_scalar_metax.h"
+#endif
+
 __INFINI_C infiniStatus_t infiniopCreateMulScalarDescriptor(
     infiniopHandle_t handle,
     infiniopMulScalarDescriptor_t *desc_ptr,
@@ -41,6 +45,9 @@ __INFINI_C infiniStatus_t infiniopCreateMulScalarDescriptor(
 #endif
 #ifdef ENABLE_HYGON_API
         CREATE(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        CREATE(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -74,6 +81,9 @@ __INFINI_C infiniStatus_t infiniopGetMulScalarWorkspaceSize(infiniopMulScalarDes
 #endif
 #ifdef ENABLE_HYGON_API
         GET(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        GET(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -113,6 +123,9 @@ __INFINI_C infiniStatus_t infiniopMulScalar(
 #endif
 #ifdef ENABLE_HYGON_API
         CALCULATE(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        CALCULATE(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -146,6 +159,9 @@ __INFINI_C infiniStatus_t infiniopDestroyMulScalarDescriptor(infiniopMulScalarDe
 #endif
 #ifdef ENABLE_HYGON_API
         DELETE(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        DELETE(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
diff --git a/src/infiniop/ops/prepare_moe_input/metax/prepare_moe_input_metax.h b/src/infiniop/ops/prepare_moe_input/metax/prepare_moe_input_metax.h
new file mode 100644
index 000000000..88b9bb00b
--- /dev/null
+++ b/src/infiniop/ops/prepare_moe_input/metax/prepare_moe_input_metax.h
@@ -0,0 +1,8 @@
+#ifndef __PREPARE_MOE_INPUT_METAX_H__
+#define __PREPARE_MOE_INPUT_METAX_H__
+
+#include "../prepare_moe_input.h"
+
+DESCRIPTOR(metax)
+
+#endif
diff --git a/src/infiniop/ops/prepare_moe_input/metax/prepare_moe_input_metax.maca b/src/infiniop/ops/prepare_moe_input/metax/prepare_moe_input_metax.maca
new file mode 100644
index 000000000..ad31abf34
--- /dev/null
+++ b/src/infiniop/ops/prepare_moe_input/metax/prepare_moe_input_metax.maca
@@ -0,0 +1,217 @@
+#ifdef ENABLE_METAX_MC_API
+#include <mc_runtime.h>
+#else
+#include <hc_runtime.h>
+#endif
+
+#ifdef ENABLE_METAX_API
+
+#include "prepare_moe_input_metax.h"
+
+#include "../../../devices/metax/metax_common.h"
+#include "../../../devices/metax/metax_kernel_common.h"
+
+#include <algorithm>
+#include <cstdint>
+#include <memory>
+#include <utility>
+
+namespace op::prepare_moe_input::metax {
+
+struct Descriptor::Opaque {
+    std::shared_ptr<device::metax::Handle::Internal> internal;
+};
+
+Descriptor::~Descriptor() {
+    delete _opaque;
+}
+
+infiniStatus_t Descriptor::create(
+    infiniopHandle_t handle,
+    Descriptor **desc_ptr,
+    infiniopTensorDescriptor_t expert_offsets_desc,
+    infiniopTensorDescriptor_t blockscale_offsets_desc,
+    infiniopTensorDescriptor_t problem_sizes1_desc,
+    infiniopTensorDescriptor_t problem_sizes2_desc,
+    infiniopTensorDescriptor_t input_permutation_desc,
+    infiniopTensorDescriptor_t output_permutation_desc,
+    infiniopTensorDescriptor_t topk_ids_desc,
+    size_t num_experts,
+    size_t n,
+    size_t k) {
+    auto result = PrepareMoeInputInfo::create(
+        expert_offsets_desc,
+        blockscale_offsets_desc,
+        problem_sizes1_desc,
+        problem_sizes2_desc,
+        input_permutation_desc,
+        output_permutation_desc,
+        topk_ids_desc,
+        num_experts,
+        n,
+        k);
+    CHECK_RESULT(result);
+    auto info = result.take();
+    const size_t workspace_size = info.num_experts * sizeof(int32_t);
+
+    *desc_ptr = new Descriptor(
+        new Opaque{reinterpret_cast<device::metax::Handle *>(handle)->internal()},
+        std::move(info),
+        workspace_size,
+        handle->device,
+        handle->device_id);
+    return INFINI_STATUS_SUCCESS;
+}
+
+namespace {
+
+constexpr uint32_t THREADS_PER_EXPERT = 512;
+constexpr int32_t BLOCKSCALE_ALIGNMENT = 128;
+
+__global__ void compute_problem_sizes_kernel(
+    const int32_t *__restrict__ topk_ids,
+    int32_t *__restrict__ problem_sizes1,
+    int32_t *__restrict__ problem_sizes2,
+    int32_t *__restrict__ atomic_buffer,
+    size_t topk_length,
+    int32_t n,
+    int32_t k) {
+    const int expert_id = blockIdx.x;
+
+    int32_t occurrences = 0;
+    for (size_t i = threadIdx.x; i < topk_length; i += THREADS_PER_EXPERT) {
+        occurrences += (topk_ids[i] == expert_id);
+    }
+    atomicAdd(&atomic_buffer[expert_id], occurrences);
+    __syncthreads();
+
+    if (threadIdx.x == 0) {
+        const int32_t final_occurrences = atomic_buffer[expert_id];
+        problem_sizes1[expert_id * 3] = final_occurrences;
+        problem_sizes1[expert_id * 3 + 1] = 2 * n;
+        problem_sizes1[expert_id * 3 + 2] = k;
+        problem_sizes2[expert_id * 3] = final_occurrences;
+        problem_sizes2[expert_id * 3 + 1] = k;
+        problem_sizes2[expert_id * 3 + 2] = n;
+    }
+}
+
+__global__ void compute_expert_offsets_kernel(
+    const int32_t *__restrict__ problem_sizes1,
+    int32_t *__restrict__ expert_offsets,
+    int32_t *__restrict__ atomic_buffer,
+    size_t num_experts) {
+    int32_t total_offset = 0;
+    expert_offsets[0] = 0;
+    for (size_t expert = 0; expert < num_experts; ++expert) {
+        atomic_buffer[expert] = total_offset;
+        total_offset += problem_sizes1[expert * 3];
+        expert_offsets[expert + 1] = total_offset;
+    }
+}
+
+__global__ void compute_expert_blockscale_offsets_kernel(
+    const int32_t *__restrict__ problem_sizes1,
+    int32_t *__restrict__ expert_offsets,
+    int32_t *__restrict__ blockscale_offsets,
+    int32_t *__restrict__ atomic_buffer,
+    size_t num_experts) {
+    int32_t total_offset = 0;
+    int32_t total_rounded_offset = 0;
+    expert_offsets[0] = 0;
+    blockscale_offsets[0] = 0;
+    for (size_t expert = 0; expert < num_experts; ++expert) {
+        atomic_buffer[expert] = total_offset;
+        const int32_t num_tokens = problem_sizes1[expert * 3];
+        const int32_t rounded_num_tokens = ((num_tokens + BLOCKSCALE_ALIGNMENT - 1) / BLOCKSCALE_ALIGNMENT) * BLOCKSCALE_ALIGNMENT;
+        total_offset += num_tokens;
+        total_rounded_offset += rounded_num_tokens;
+        expert_offsets[expert + 1] = total_offset;
+        blockscale_offsets[expert + 1] = total_rounded_offset;
+    }
+}
+
+__global__ void compute_arg_sorts_kernel(
+    const int32_t *__restrict__ topk_ids,
+    int32_t *__restrict__ input_permutation,
+    int32_t *__restrict__ output_permutation,
+    int32_t *__restrict__ atomic_buffer,
+    size_t topk_length,
+    size_t topk) {
+    const int expert_id = blockIdx.x;
+
+    for (size_t i = threadIdx.x; i < topk_length; i += THREADS_PER_EXPERT) {
+        if (topk_ids[i] == expert_id) {
+            const int32_t start = atomicAdd(&atomic_buffer[expert_id], 1);
+            input_permutation[start] = static_cast<int32_t>(i / topk);
+            output_permutation[i] = start;
+        }
+    }
+}
+
+} // namespace
+
+infiniStatus_t Descriptor::calculate(
+    void *workspace,
+    size_t workspace_size,
+    void *expert_offsets,
+    void *blockscale_offsets,
+    void *problem_sizes1,
+    void *problem_sizes2,
+    void *input_permutation,
+    void *output_permutation,
+    const void *topk_ids,
+    void *stream) const {
+    if (workspace_size < _workspace_size) {
+        return INFINI_STATUS_INSUFFICIENT_WORKSPACE;
+    }
+    if (_info.has_blockscale_offsets && blockscale_offsets == nullptr) {
+        return INFINI_STATUS_BAD_PARAM;
+    }
+
+    auto cuda_stream = reinterpret_cast<hcStream_t>(stream);
+    auto *atomic_buffer = static_cast<int32_t *>(workspace);
+    hcMemsetAsync(atomic_buffer, 0, _workspace_size, cuda_stream);
+
+    const uint32_t threads = static_cast<uint32_t>(std::min<size_t>(THREADS_PER_EXPERT, _info.topk_length));
+    const uint32_t safe_threads = std::max<uint32_t>(threads, 1);
+    const uint32_t blocks = static_cast<uint32_t>(_info.num_experts);
+
+    compute_problem_sizes_kernel<<<blocks, safe_threads, 0, cuda_stream>>>(
+        static_cast<const int32_t *>(topk_ids),
+        static_cast<int32_t *>(problem_sizes1),
+        static_cast<int32_t *>(problem_sizes2),
+        atomic_buffer,
+        _info.topk_length,
+        static_cast<int32_t>(_info.n),
+        static_cast<int32_t>(_info.k));
+
+    if (_info.has_blockscale_offsets) {
+        compute_expert_blockscale_offsets_kernel<<<1, 1, 0, cuda_stream>>>(
+            static_cast<const int32_t *>(problem_sizes1),
+            static_cast<int32_t *>(expert_offsets),
+            static_cast<int32_t *>(blockscale_offsets),
+            atomic_buffer,
+            _info.num_experts);
+    } else {
+        compute_expert_offsets_kernel<<<1, 1, 0, cuda_stream>>>(
+            static_cast<const int32_t *>(problem_sizes1),
+            static_cast<int32_t *>(expert_offsets),
+            atomic_buffer,
+            _info.num_experts);
+    }
+
+    compute_arg_sorts_kernel<<<blocks, safe_threads, 0, cuda_stream>>>(
+        static_cast<const int32_t *>(topk_ids),
+        static_cast<int32_t *>(input_permutation),
+        static_cast<int32_t *>(output_permutation),
+        atomic_buffer,
+        _info.topk_length,
+        _info.topk);
+
+    return INFINI_STATUS_SUCCESS;
+}
+
+} // namespace op::prepare_moe_input::metax
+
+#endif // ENABLE_METAX_API
diff --git a/src/infiniop/ops/prepare_moe_input/operator.cc b/src/infiniop/ops/prepare_moe_input/operator.cc
index b5ea0d81d..aba7cd0a3 100644
--- a/src/infiniop/ops/prepare_moe_input/operator.cc
+++ b/src/infiniop/ops/prepare_moe_input/operator.cc
@@ -6,6 +6,10 @@
 #include "nvidia/prepare_moe_input_nvidia.cuh"
 #endif
 
+#ifdef ENABLE_METAX_API
+#include "metax/prepare_moe_input_metax.h"
+#endif
+
 __INFINI_C infiniStatus_t infiniopCreatePrepareMoeInputDescriptor(
     infiniopHandle_t handle,
     infiniopPrepareMoeInputDescriptor_t *desc_ptr,
@@ -39,6 +43,9 @@ __INFINI_C infiniStatus_t infiniopCreatePrepareMoeInputDescriptor(
     switch (handle->device) {
 #ifdef ENABLE_NVIDIA_API
         CREATE(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        CREATE(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -59,6 +66,9 @@ __INFINI_C infiniStatus_t infiniopGetPrepareMoeInputWorkspaceSize(
     switch (desc->device_type) {
 #ifdef ENABLE_NVIDIA_API
         GET(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        GET(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -89,6 +99,9 @@ __INFINI_C infiniStatus_t infiniopPrepareMoeInput(
     switch (desc->device_type) {
 #ifdef ENABLE_NVIDIA_API
         CALCULATE(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        CALCULATE(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -108,6 +121,9 @@ __INFINI_C infiniStatus_t infiniopDestroyPrepareMoeInputDescriptor(
     switch (desc->device_type) {
 #ifdef ENABLE_NVIDIA_API
         DESTROY(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        DESTROY(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
diff --git a/src/infiniop/ops/rwkv5_wkv/metax/rwkv5_wkv_metax.h b/src/infiniop/ops/rwkv5_wkv/metax/rwkv5_wkv_metax.h
new file mode 100644
index 000000000..64bfd83c1
--- /dev/null
+++ b/src/infiniop/ops/rwkv5_wkv/metax/rwkv5_wkv_metax.h
@@ -0,0 +1,8 @@
+#ifndef __RWKV5_WKV_METAX_H__
+#define __RWKV5_WKV_METAX_H__
+
+#include "../rwkv5_wkv.h"
+
+DESCRIPTOR(metax)
+
+#endif
diff --git a/src/infiniop/ops/rwkv5_wkv/metax/rwkv5_wkv_metax.maca b/src/infiniop/ops/rwkv5_wkv/metax/rwkv5_wkv_metax.maca
new file mode 100644
index 000000000..5913d8530
--- /dev/null
+++ b/src/infiniop/ops/rwkv5_wkv/metax/rwkv5_wkv_metax.maca
@@ -0,0 +1,193 @@
+#ifdef ENABLE_METAX_MC_API
+#include <mc_runtime.h>
+#else
+#include <hc_runtime.h>
+#endif
+
+#include "../../../devices/metax/metax_common.h"
+#include "../../../devices/metax/metax_kernel_common.h"
+#include "rwkv5_wkv_metax.h"
+
+namespace op::rwkv5_wkv::metax {
+
+struct Descriptor::Opaque {
+    std::shared_ptr<device::metax::Handle::Internal> internal;
+};
+
+Descriptor::~Descriptor() {
+    delete _opaque;
+}
+
+namespace {
+
+template <typename T>
+__device__ float to_float(T value) {
+    return static_cast<float>(value);
+}
+
+template <>
+__device__ float to_float<half>(half value) {
+    return __half2float(value);
+}
+
+template <>
+__device__ float to_float<__nv_bfloat16>(__nv_bfloat16 value) {
+    return __bfloat162float(value);
+}
+
+template <typename T>
+__device__ T from_float(float value) {
+    return static_cast<T>(value);
+}
+
+template <>
+__device__ half from_float<half>(float value) {
+    return __float2half_rn(value);
+}
+
+template <>
+__device__ __nv_bfloat16 from_float<__nv_bfloat16>(float value) {
+    return __float2bfloat16_rn(value);
+}
+
+template <typename T>
+__global__ void rwkv5_wkv_kernel(
+    T *out,
+    const T *receptance,
+    const T *key,
+    const T *value,
+    const T *time_decay,
+    const T *time_faaaa,
+    float *state,
+    size_t batch,
+    size_t seq_len,
+    size_t hidden_size,
+    size_t num_heads,
+    size_t head_size) {
+
+    const size_t linear = blockIdx.x * blockDim.x + threadIdx.x;
+    const size_t total = batch * num_heads * head_size;
+    if (linear >= total) {
+        return;
+    }
+
+    const size_t s_out = linear % head_size;
+    const size_t tmp = linear / head_size;
+    const size_t h = tmp % num_heads;
+    const size_t b = tmp / num_heads;
+    const size_t channel_base = h * head_size;
+    const size_t state_col_base = ((b * num_heads + h) * head_size) * head_size + s_out;
+
+    for (size_t t = 0; t < seq_len; ++t) {
+        const size_t token_base = (b * seq_len + t) * hidden_size + channel_base;
+        const float vj = to_float<T>(value[token_base + s_out]);
+        float acc = 0.0f;
+        for (size_t i = 0; i < head_size; ++i) {
+            const float r = to_float<T>(receptance[token_base + i]);
+            const float k = to_float<T>(key[token_base + i]);
+            const float time_first = to_float<T>(time_faaaa[channel_base + i]);
+            const float decay = __expf(-__expf(to_float<T>(time_decay[channel_base + i])));
+            const size_t state_idx = state_col_base + i * head_size;
+            const float att = k * vj;
+            acc += r * (time_first * att + state[state_idx]);
+            state[state_idx] = att + decay * state[state_idx];
+        }
+        out[token_base + s_out] = from_float<T>(acc);
+    }
+}
+
+template <typename T>
+infiniStatus_t launch_typed(
+    const Rwkv5WkvInfo &info,
+    void *out,
+    const void *receptance,
+    const void *key,
+    const void *value,
+    const void *time_decay,
+    const void *time_faaaa,
+    void *state,
+    hcStream_t stream) {
+
+    constexpr int threads = 256;
+    const size_t total = info.batch * info.num_heads * info.head_size;
+    const dim3 blocks(static_cast<unsigned int>((total + threads - 1) / threads));
+    rwkv5_wkv_kernel<T><<<blocks, threads, 0, stream>>>(
+        reinterpret_cast<T *>(out),
+        reinterpret_cast<const T *>(receptance),
+        reinterpret_cast<const T *>(key),
+        reinterpret_cast<const T *>(value),
+        reinterpret_cast<const T *>(time_decay),
+        reinterpret_cast<const T *>(time_faaaa),
+        reinterpret_cast<float *>(state),
+        info.batch,
+        info.seq_len,
+        info.hidden_size,
+        info.num_heads,
+        info.head_size);
+    CHECK_METAX(hcGetLastError());
+    return INFINI_STATUS_SUCCESS;
+}
+
+} // namespace
+
+infiniStatus_t Descriptor::create(
+    infiniopHandle_t handle,
+    Descriptor **desc_ptr,
+    infiniopTensorDescriptor_t out_desc,
+    infiniopTensorDescriptor_t receptance_desc,
+    infiniopTensorDescriptor_t key_desc,
+    infiniopTensorDescriptor_t value_desc,
+    infiniopTensorDescriptor_t time_decay_desc,
+    infiniopTensorDescriptor_t time_faaaa_desc,
+    infiniopTensorDescriptor_t state_desc) {
+
+    auto result = Rwkv5WkvInfo::create(
+        out_desc,
+        receptance_desc,
+        key_desc,
+        value_desc,
+        time_decay_desc,
+        time_faaaa_desc,
+        state_desc);
+    CHECK_RESULT(result);
+    auto info = result.take();
+
+    *desc_ptr = new Descriptor(
+        new Opaque{reinterpret_cast<device::metax::Handle *>(handle)->internal()},
+        info,
+        0,
+        handle->device,
+        handle->device_id);
+    return INFINI_STATUS_SUCCESS;
+}
+
+infiniStatus_t Descriptor::calculate(
+    void *workspace,
+    size_t workspace_size,
+    void *out,
+    const void *receptance,
+    const void *key,
+    const void *value,
+    const void *time_decay,
+    const void *time_faaaa,
+    void *state,
+    void *stream_) const {
+
+    (void)workspace;
+    if (workspace_size < _workspace_size) {
+        return INFINI_STATUS_INSUFFICIENT_WORKSPACE;
+    }
+    auto stream = reinterpret_cast<hcStream_t>(stream_);
+    if (_info.dtype == INFINI_DTYPE_F16) {
+        return launch_typed<half>(_info, out, receptance, key, value, time_decay, time_faaaa, state, stream);
+    }
+    if (_info.dtype == INFINI_DTYPE_BF16) {
+        return launch_typed<__nv_bfloat16>(_info, out, receptance, key, value, time_decay, time_faaaa, state, stream);
+    }
+    if (_info.dtype == INFINI_DTYPE_F32) {
+        return launch_typed<float>(_info, out, receptance, key, value, time_decay, time_faaaa, state, stream);
+    }
+    return INFINI_STATUS_BAD_TENSOR_DTYPE;
+}
+
+} // namespace op::rwkv5_wkv::metax
diff --git a/src/infiniop/ops/rwkv5_wkv/operator.cc b/src/infiniop/ops/rwkv5_wkv/operator.cc
index 75d7f6c00..6cac196f7 100644
--- a/src/infiniop/ops/rwkv5_wkv/operator.cc
+++ b/src/infiniop/ops/rwkv5_wkv/operator.cc
@@ -6,6 +6,10 @@
 #include "nvidia/rwkv5_wkv_nvidia.cuh"
 #endif
 
+#ifdef ENABLE_METAX_API
+#include "metax/rwkv5_wkv_metax.h"
+#endif
+
 __INFINI_C infiniStatus_t infiniopCreateRwkv5WkvDescriptor(
     infiniopHandle_t handle,
     infiniopRwkv5WkvDescriptor_t *desc_ptr,
@@ -39,6 +43,9 @@ __INFINI_C infiniStatus_t infiniopCreateRwkv5WkvDescriptor(
 #endif
 #ifdef ENABLE_HYGON_API
         CREATE(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        CREATE(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -70,6 +77,9 @@ __INFINI_C infiniStatus_t infiniopGetRwkv5WkvWorkspaceSize(
 #endif
 #ifdef ENABLE_HYGON_API
         GET(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        GET(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -111,6 +121,9 @@ __INFINI_C infiniStatus_t infiniopRwkv5Wkv(
 #endif
 #ifdef ENABLE_HYGON_API
         CALCULATE(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        CALCULATE(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -141,6 +154,9 @@ __INFINI_C infiniStatus_t infiniopDestroyRwkv5WkvDescriptor(
 #endif
 #ifdef ENABLE_HYGON_API
         DESTROY(INFINI_DEVICE_HYGON, nvidia);
+#endif
+#ifdef ENABLE_METAX_API
+        DESTROY(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;