diff --git a/include/infinicore/ops/mul_scalar.hpp b/include/infinicore/ops/mul_scalar.hpp
new file mode 100644
index 000000000..5ee5e0d0b
--- /dev/null
+++ b/include/infinicore/ops/mul_scalar.hpp
@@ -0,0 +1,14 @@
+#pragma once
+
+#include "../device.hpp"
+#include "../graph/graph.hpp"
+#include "common/op.hpp"
+
+namespace infinicore::op {
+
+INFINICORE_GRAPH_OP_CLASS(MulScalar, Tensor, const Tensor &, double);
+
+Tensor mul_scalar(const Tensor &a, double alpha);
+void mul_scalar_(Tensor c, const Tensor &a, double alpha);
+
+} // namespace infinicore::op
diff --git a/include/infiniop.h b/include/infiniop.h
index 7794dfc5c..95bf75a0d 100644
--- a/include/infiniop.h
+++ b/include/infiniop.h
@@ -97,6 +97,7 @@
 #include "infiniop/ops/moe_topk_softmax.h"
 #include "infiniop/ops/mrope.h"
 #include "infiniop/ops/mul.h"
+#include "infiniop/ops/mul_scalar.h"
 #include "infiniop/ops/multi_margin_loss.h"
 #include "infiniop/ops/nrm2.h"
 #include "infiniop/ops/nsa_compress_paged_cache.h"
diff --git a/include/infiniop/handle.h b/include/infiniop/handle.h
index da8689542..207ca2ce9 100644
--- a/include/infiniop/handle.h
+++ b/include/infiniop/handle.h
@@ -7,6 +7,8 @@ struct InfiniopHandle;
 
 typedef struct InfiniopHandle *infiniopHandle_t;
 
+__INFINI_C __export infiniStatus_t infiniopSetRuntimeDevice(infiniDevice_t device, int device_id);
+
 __INFINI_C __export infiniStatus_t infiniopCreateHandle(infiniopHandle_t *handle_ptr);
 
 __INFINI_C __export infiniStatus_t infiniopDestroyHandle(infiniopHandle_t handle);
diff --git a/include/infiniop/ops/mul_scalar.h b/include/infiniop/ops/mul_scalar.h
new file mode 100644
index 000000000..9f90284ee
--- /dev/null
+++ b/include/infiniop/ops/mul_scalar.h
@@ -0,0 +1,25 @@
+#ifndef __INFINIOP_MUL_SCALAR_API_H__
+#define __INFINIOP_MUL_SCALAR_API_H__
+
+#include "../operator_descriptor.h"
+
+typedef struct InfiniopDescriptor *infiniopMulScalarDescriptor_t;
+
+__INFINI_C __export infiniStatus_t infiniopCreateMulScalarDescriptor(infiniopHandle_t handle,
+                                                                     infiniopMulScalarDescriptor_t *desc_ptr,
+                                                                     infiniopTensorDescriptor_t output,
+                                                                     infiniopTensorDescriptor_t input);
+
+__INFINI_C __export infiniStatus_t infiniopGetMulScalarWorkspaceSize(infiniopMulScalarDescriptor_t desc, size_t *size);
+
+__INFINI_C __export infiniStatus_t infiniopMulScalar(infiniopMulScalarDescriptor_t desc,
+                                                     void *workspace,
+                                                     size_t workspace_size,
+                                                     void *output,
+                                                     const void *input,
+                                                     double alpha,
+                                                     void *stream);
+
+__INFINI_C __export infiniStatus_t infiniopDestroyMulScalarDescriptor(infiniopMulScalarDescriptor_t desc);
+
+#endif
diff --git a/python/infinicore/__init__.py b/python/infinicore/__init__.py
index cdefb904b..808995640 100644
--- a/python/infinicore/__init__.py
+++ b/python/infinicore/__init__.py
@@ -114,6 +114,7 @@
 )
 from infinicore.ops.mrope import mrope
 from infinicore.ops.mul import mul
+from infinicore.ops.mul_scalar import mul_scalar
 from infinicore.ops.narrow import narrow
 from infinicore.ops.nrm2 import nrm2
 from infinicore.ops.paged_attention import paged_attention
@@ -234,6 +235,7 @@
     "matmul",
     "equal",
     "mul",
+    "mul_scalar",
     "diff",
     "digamma",
     "dist",
diff --git a/python/infinicore/ops/mul.py b/python/infinicore/ops/mul.py
index e9e1d8c98..8ce703917 100644
--- a/python/infinicore/ops/mul.py
+++ b/python/infinicore/ops/mul.py
@@ -1,8 +1,14 @@
+import numbers
+
 from infinicore.lib import _infinicore
+from infinicore.ops.mul_scalar import mul_scalar
 from infinicore.tensor import Tensor
 
 
 def mul(input, other, *, out=None):
+    if isinstance(other, numbers.Real):
+        return mul_scalar(input, other, out=out)
+
     if out is None:
         return Tensor(_infinicore.mul(input._underlying, other._underlying))
 
diff --git a/python/infinicore/ops/mul_scalar.py b/python/infinicore/ops/mul_scalar.py
new file mode 100644
index 000000000..e7508a570
--- /dev/null
+++ b/python/infinicore/ops/mul_scalar.py
@@ -0,0 +1,12 @@
+from infinicore.lib import _infinicore
+from infinicore.tensor import Tensor
+
+
+def mul_scalar(input, alpha, *, out=None):
+    alpha = float(alpha)
+    if out is None:
+        return Tensor(_infinicore.mul_scalar(input._underlying, alpha))
+
+    _infinicore.mul_scalar_(out._underlying, input._underlying, alpha)
+
+    return out
diff --git a/python/infinicore/tensor.py b/python/infinicore/tensor.py
index 4348887d6..bbe801f93 100644
--- a/python/infinicore/tensor.py
+++ b/python/infinicore/tensor.py
@@ -128,6 +128,9 @@ def __matmul__(self, other):
     def __mul__(self, other):
         return infinicore.mul(self, other)
 
+    def __rmul__(self, other):
+        return infinicore.mul(self, other)
+
     def narrow(self, dim, start, length):
         return infinicore.narrow(self, dim, start, length)
 
diff --git a/src/bridge/infini/rt.hpp b/src/bridge/infini/rt.hpp
new file mode 100644
index 000000000..3b4b350bd
--- /dev/null
+++ b/src/bridge/infini/rt.hpp
@@ -0,0 +1,80 @@
+#pragma once
+
+#include "infinirt.h"
+
+#include <infini/rt.h>
+
+namespace infinicore::bridge::infini::rt {
+
+inline infiniStatus_t translate(::infini::rt::runtime::Error error) {
+    switch (error) {
+    case ::infini::rt::runtime::kSuccess:
+        return INFINI_STATUS_SUCCESS;
+    default:
+        return INFINI_STATUS_INTERNAL_ERROR;
+    }
+}
+
+inline ::infini::rt::Device::Type translate(infiniDevice_t device) {
+    switch (device) {
+    case INFINI_DEVICE_CPU:
+        return ::infini::rt::Device::Type::kCpu;
+    case INFINI_DEVICE_NVIDIA:
+        return ::infini::rt::Device::Type::kNvidia;
+    case INFINI_DEVICE_CAMBRICON:
+        return ::infini::rt::Device::Type::kCambricon;
+    case INFINI_DEVICE_ASCEND:
+        return ::infini::rt::Device::Type::kAscend;
+    case INFINI_DEVICE_METAX:
+        return ::infini::rt::Device::Type::kMetax;
+    case INFINI_DEVICE_MOORE:
+        return ::infini::rt::Device::Type::kMoore;
+    case INFINI_DEVICE_ILUVATAR:
+        return ::infini::rt::Device::Type::kIluvatar;
+    case INFINI_DEVICE_KUNLUN:
+        return ::infini::rt::Device::Type::kKunlun;
+    case INFINI_DEVICE_HYGON:
+        return ::infini::rt::Device::Type::kHygon;
+    case INFINI_DEVICE_QY:
+        return ::infini::rt::Device::Type::kQy;
+    default:
+        return ::infini::rt::Device::Type::kCount;
+    }
+}
+
+inline infiniDevice_t translate(::infini::rt::Device::Type device) {
+    switch (device) {
+    case ::infini::rt::Device::Type::kCpu:
+        return INFINI_DEVICE_CPU;
+    case ::infini::rt::Device::Type::kNvidia:
+        return INFINI_DEVICE_NVIDIA;
+    case ::infini::rt::Device::Type::kCambricon:
+        return INFINI_DEVICE_CAMBRICON;
+    case ::infini::rt::Device::Type::kAscend:
+        return INFINI_DEVICE_ASCEND;
+    case ::infini::rt::Device::Type::kMetax:
+        return INFINI_DEVICE_METAX;
+    case ::infini::rt::Device::Type::kMoore:
+        return INFINI_DEVICE_MOORE;
+    case ::infini::rt::Device::Type::kIluvatar:
+        return INFINI_DEVICE_ILUVATAR;
+    case ::infini::rt::Device::Type::kKunlun:
+        return INFINI_DEVICE_KUNLUN;
+    case ::infini::rt::Device::Type::kHygon:
+        return INFINI_DEVICE_HYGON;
+    case ::infini::rt::Device::Type::kQy:
+        return INFINI_DEVICE_QY;
+    default:
+        return INFINI_DEVICE_TYPE_COUNT;
+    }
+}
+
+inline ::infini::rt::runtime::Stream to_rt_stream(infinirtStream_t stream) {
+    return reinterpret_cast<::infini::rt::runtime::Stream>(stream);
+}
+
+inline infinirtStream_t to_core_stream(::infini::rt::runtime::Stream stream) {
+    return reinterpret_cast<infinirtStream_t>(stream);
+}
+
+} // namespace infinicore::bridge::infini::rt
diff --git a/src/infinicore/context/allocators/device_pinned_allocator.cc b/src/infinicore/context/allocators/device_pinned_allocator.cc
index 36ee599ff..a8764d5f6 100644
--- a/src/infinicore/context/allocators/device_pinned_allocator.cc
+++ b/src/infinicore/context/allocators/device_pinned_allocator.cc
@@ -1,6 +1,6 @@
 #include "device_pinned_allocator.hpp"
 
-#include <infinirt.h>
+#include "../../../bridge/infini/rt.hpp"
 
 #include "../../utils.hpp"
 
@@ -16,7 +16,7 @@ std::byte *DevicePinnedHostAllocator::allocate(size_t size) {
         return nullptr;
     }
     void *ptr;
-    INFINICORE_CHECK_ERROR(infinirtMallocHost(&ptr, size));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::MallocHost(&ptr, size)));
     return (std::byte *)ptr;
 }
 
@@ -25,7 +25,7 @@ void DevicePinnedHostAllocator::deallocate(std::byte *ptr) {
         return;
     }
     if (owner_ == context::getDevice()) {
-        INFINICORE_CHECK_ERROR(infinirtFreeHost(ptr));
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::FreeHost(ptr)));
         gc();
     } else {
         gc_queue_.push(ptr);
@@ -35,7 +35,7 @@ void DevicePinnedHostAllocator::deallocate(std::byte *ptr) {
 void DevicePinnedHostAllocator::gc() {
     while (gc_queue_.empty() == false) {
         std::byte *p = gc_queue_.front();
-        INFINICORE_CHECK_ERROR(infinirtFreeHost(p));
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::FreeHost(p)));
         gc_queue_.pop();
     }
 }
diff --git a/src/infinicore/context/allocators/host_allocator.cc b/src/infinicore/context/allocators/host_allocator.cc
index 41a55b453..0f1cc7072 100644
--- a/src/infinicore/context/allocators/host_allocator.cc
+++ b/src/infinicore/context/allocators/host_allocator.cc
@@ -1,7 +1,5 @@
 #include "host_allocator.hpp"
 
-#include <infinirt.h>
-
 namespace infinicore {
 std::byte *HostAllocator::allocate(size_t size) {
     if (size == 0) {
diff --git a/src/infinicore/context/allocators/pinnable_block_allocator.cc b/src/infinicore/context/allocators/pinnable_block_allocator.cc
index 32e5c5e9b..3c70376f0 100644
--- a/src/infinicore/context/allocators/pinnable_block_allocator.cc
+++ b/src/infinicore/context/allocators/pinnable_block_allocator.cc
@@ -4,8 +4,9 @@
 
 #include "../../utils.hpp"
 
+#include "../../../bridge/infini/rt.hpp"
+
 #include <algorithm>
-#include <infinirt.h>
 #include <stdexcept>
 
 namespace infinicore {
@@ -74,7 +75,7 @@ std::byte *PinnableBlockAllocator::allocate(size_t size) {
             block->in_use = true;
             block->use_count = 1;
 
-            INFINICORE_CHECK_ERROR(infinirtMalloc(&block->ptr, block->size));
+            INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::Malloc(&block->ptr, block->size)));
 
             all_blocks_[block->ptr] = block;
             return reinterpret_cast<std::byte *>(block->ptr);
@@ -101,7 +102,7 @@ std::byte *PinnableBlockAllocator::allocate(size_t size) {
     block->in_use = true;
     block->use_count = 1;
 
-    INFINICORE_CHECK_ERROR(infinirtMalloc(&block->ptr, block->size));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::Malloc(&block->ptr, block->size)));
 
     large_blocks_.push_back(block);
     all_blocks_[block->ptr] = block;
@@ -167,7 +168,7 @@ void PinnableBlockAllocator::trim() {
     for (auto &cls : size_classes_) {
         for (auto it = cls.free_blocks.begin(); it != cls.free_blocks.end();) {
             if (!(*it)->frozen) {
-                INFINICORE_CHECK_ERROR(infinirtFree((*it)->ptr));
+                INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::Free((*it)->ptr)));
                 all_blocks_.erase((*it)->ptr);
                 it = cls.free_blocks.erase(it);
             } else {
@@ -178,7 +179,7 @@ void PinnableBlockAllocator::trim() {
     // Free non-frozen large blocks
     for (auto it = large_blocks_.begin(); it != large_blocks_.end();) {
         if (!(*it)->frozen && !(*it)->in_use) {
-            INFINICORE_CHECK_ERROR(infinirtFree((*it)->ptr));
+            INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::Free((*it)->ptr)));
             all_blocks_.erase((*it)->ptr);
             it = large_blocks_.erase(it);
         } else {
@@ -192,7 +193,7 @@ PinnableBlockAllocator::~PinnableBlockAllocator() {
     std::lock_guard<std::mutex> lock(mutex_);
     for (auto &p : all_blocks_) {
         if (p.second->ptr) {
-            infinirtFree(p.second->ptr);
+            (void)infini::rt::runtime::Free(p.second->ptr);
         }
     }
     all_blocks_.clear();
diff --git a/src/infinicore/context/allocators/stream_ordered_allocator.cc b/src/infinicore/context/allocators/stream_ordered_allocator.cc
index 3bd0aaab1..06ce4a30d 100644
--- a/src/infinicore/context/allocators/stream_ordered_allocator.cc
+++ b/src/infinicore/context/allocators/stream_ordered_allocator.cc
@@ -1,6 +1,6 @@
 #include "stream_ordered_allocator.hpp"
 
-#include <infinirt.h>
+#include "../../../bridge/infini/rt.hpp"
 
 #include "../../utils.hpp"
 
@@ -12,7 +12,14 @@ std::byte *StreamOrderedAllocator::allocate(size_t size) {
         return nullptr;
     }
     void *ptr = nullptr;
-    INFINICORE_CHECK_ERROR(infinirtMallocAsync(&ptr, size, context::getStream()));
+    if (device_.getType() != Device::Type::CPU) {
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::MallocAsync(
+            &ptr,
+            size,
+            bridge::infini::rt::to_rt_stream(context::getStream()))));
+    } else {
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::Malloc(&ptr, size)));
+    }
     return (std::byte *)ptr;
 }
 
@@ -20,6 +27,12 @@ void StreamOrderedAllocator::deallocate(std::byte *ptr) {
     if (ptr == nullptr) {
         return;
     }
-    INFINICORE_CHECK_ERROR(infinirtFreeAsync(ptr, context::getStream()));
+    if (device_.getType() != Device::Type::CPU) {
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::FreeAsync(
+            ptr,
+            bridge::infini::rt::to_rt_stream(context::getStream()))));
+    } else {
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::Free(ptr)));
+    }
 }
 } // namespace infinicore
diff --git a/src/infinicore/context/context_impl.cc b/src/infinicore/context/context_impl.cc
index f126beb1d..64d10dcfc 100644
--- a/src/infinicore/context/context_impl.cc
+++ b/src/infinicore/context/context_impl.cc
@@ -1,4 +1,5 @@
 #include "context_impl.hpp"
+#include "../../bridge/infini/rt.hpp"
 #include "internal.hpp"
 
 #include "../utils.hpp"
@@ -61,26 +62,38 @@ ContextImpl &ContextImpl::singleton() {
 }
 
 ContextImpl::ContextImpl() {
-    std::vector<int> device_counter(static_cast<size_t>(Device::Type::COUNT));
-    INFINICORE_CHECK_ERROR(infinirtGetAllDeviceCount(device_counter.data()));
-
-    // Reserve runtime slot for all devices.
-    runtime_table_[0].resize(device_counter[0]);
-    runtime_table_[0][0] = std::unique_ptr<Runtime>(new Runtime(Device(Device::Type::CPU, 0)));
-
-    // Context will try to use the first non-cpu available device as the default runtime.
-    for (int i = int(Device::Type::COUNT) - 1; i > 0; i--) {
-        if (device_counter[i] > 0) {
-            runtime_table_[i].resize(device_counter[i]);
-            if (current_runtime_ == nullptr) {
-                runtime_table_[i][0] = std::unique_ptr<Runtime>(new Runtime(Device(Device::Type(i), 0)));
-                current_runtime_ = runtime_table_[i][0].get();
-            }
+    std::vector<int> device_counter(static_cast<size_t>(Device::Type::COUNT), 0);
+
+    infini::rt::set_runtime_device_type(bridge::infini::rt::translate(INFINI_DEVICE_CPU));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::GetDeviceCount(&device_counter[static_cast<int>(Device::Type::CPU)])));
+
+    runtime_table_[static_cast<int>(Device::Type::CPU)].resize(device_counter[static_cast<int>(Device::Type::CPU)]);
+    if (device_counter[static_cast<int>(Device::Type::CPU)] > 0) {
+        runtime_table_[static_cast<int>(Device::Type::CPU)][0] = std::unique_ptr<Runtime>(new Runtime(Device(Device::Type::CPU, 0)));
+    }
+
+    if constexpr (infini::rt::DeviceEnabled<infini::rt::Device::Type::kNvidia>::value) {
+        infini::rt::set_runtime_device_type(bridge::infini::rt::translate(INFINI_DEVICE_NVIDIA));
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::GetDeviceCount(&device_counter[static_cast<int>(Device::Type::NVIDIA)])));
+        runtime_table_[static_cast<int>(Device::Type::NVIDIA)].resize(device_counter[static_cast<int>(Device::Type::NVIDIA)]);
+        if (device_counter[static_cast<int>(Device::Type::NVIDIA)] > 0) {
+            runtime_table_[static_cast<int>(Device::Type::NVIDIA)][0] = std::unique_ptr<Runtime>(new Runtime(Device(Device::Type::NVIDIA, 0)));
+            current_runtime_ = runtime_table_[static_cast<int>(Device::Type::NVIDIA)][0].get();
         }
     }
 
-    if (current_runtime_ == nullptr) {
-        current_runtime_ = runtime_table_[0][0].get();
+    if constexpr (infini::rt::DeviceEnabled<infini::rt::Device::Type::kAscend>::value) {
+        infini::rt::set_runtime_device_type(bridge::infini::rt::translate(INFINI_DEVICE_ASCEND));
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::GetDeviceCount(&device_counter[static_cast<int>(Device::Type::ASCEND)])));
+        runtime_table_[static_cast<int>(Device::Type::ASCEND)].resize(device_counter[static_cast<int>(Device::Type::ASCEND)]);
+        if (device_counter[static_cast<int>(Device::Type::ASCEND)] > 0) {
+            runtime_table_[static_cast<int>(Device::Type::ASCEND)][0] = std::unique_ptr<Runtime>(new Runtime(Device(Device::Type::ASCEND, 0)));
+            current_runtime_ = runtime_table_[static_cast<int>(Device::Type::ASCEND)][0].get();
+        }
+    }
+
+    if (current_runtime_ == nullptr && !runtime_table_[static_cast<int>(Device::Type::CPU)].empty()) {
+        current_runtime_ = runtime_table_[static_cast<int>(Device::Type::CPU)][0].get();
     }
 }
 
diff --git a/src/infinicore/context/runtime/runtime.cc b/src/infinicore/context/runtime/runtime.cc
index 92ba0f737..18424d036 100644
--- a/src/infinicore/context/runtime/runtime.cc
+++ b/src/infinicore/context/runtime/runtime.cc
@@ -1,4 +1,5 @@
 #include "runtime.hpp"
+#include "../../../bridge/infini/rt.hpp"
 
 #include "../../utils.hpp"
 
@@ -10,7 +11,9 @@
 namespace infinicore {
 Runtime::Runtime(Device device) : device_(device), graph_manager_(std::make_unique<graph::GraphManager>()) {
     activate();
-    INFINICORE_CHECK_ERROR(infinirtStreamCreate(&stream_));
+    infini::rt::runtime::Stream stream = nullptr;
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::StreamCreate(&stream)));
+    stream_ = bridge::infini::rt::to_core_stream(stream);
     INFINICORE_CHECK_ERROR(infiniopCreateHandle(&infiniop_handle_));
     if (device_.getType() == Device::Type::CPU) {
         device_memory_allocator_ = std::make_unique<PinnableBlockAllocator>(device);
@@ -26,11 +29,14 @@ Runtime::~Runtime() {
     }
     device_memory_allocator_.reset();
     infiniopDestroyHandle(infiniop_handle_);
-    infinirtStreamDestroy(stream_);
+    (void)infini::rt::runtime::StreamDestroy(bridge::infini::rt::to_rt_stream(stream_));
 }
 
 Runtime *Runtime::activate() {
-    INFINICORE_CHECK_ERROR(infinirtSetDevice((infiniDevice_t)device_.getType(), (int)device_.getIndex()));
+    auto rt_device = bridge::infini::rt::translate(static_cast<infiniDevice_t>(device_.getType()));
+    INFINICORE_ASSERT(rt_device != infini::rt::Device::Type::kCount);
+    infini::rt::set_runtime_device_type(rt_device);
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::SetDevice(static_cast<int>(device_.getIndex()))));
     return this;
 }
 
@@ -47,11 +53,11 @@ infiniopHandle_t Runtime::infiniopHandle() const {
 }
 
 void Runtime::syncStream() {
-    INFINICORE_CHECK_ERROR(infinirtStreamSynchronize(stream_));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::StreamSynchronize(bridge::infini::rt::to_rt_stream(stream_))));
 }
 
 void Runtime::syncDevice() {
-    INFINICORE_CHECK_ERROR(infinirtDeviceSynchronize());
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::DeviceSynchronize()));
 }
 
 void Runtime::trimMemory() {
@@ -103,43 +109,47 @@ std::shared_ptr<Memory> Runtime::reinstantiateBlob(std::shared_ptr<Memory> blob)
 }
 
 void Runtime::memcpyH2D(void *dst, const void *src, size_t size, bool async) {
-    if (async) {
-        INFINICORE_CHECK_ERROR(infinirtMemcpyAsync(dst, src, size, INFINIRT_MEMCPY_H2D, stream_));
+    if (async && device_.getType() != Device::Type::CPU) {
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::MemcpyAsync(dst, src, size, infini::rt::runtime::kMemcpyHostToDevice, bridge::infini::rt::to_rt_stream(stream_))));
     } else {
-        INFINICORE_CHECK_ERROR(infinirtMemcpy(dst, src, size, INFINIRT_MEMCPY_H2D));
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::Memcpy(dst, src, size, infini::rt::runtime::kMemcpyHostToDevice)));
     }
 }
 
 void Runtime::memcpyD2H(void *dst, const void *src, size_t size) {
-    INFINICORE_CHECK_ERROR(infinirtMemcpy(dst, src, size, INFINIRT_MEMCPY_D2H));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::Memcpy(dst, src, size, infini::rt::runtime::kMemcpyDeviceToHost)));
 }
 
 void Runtime::memcpyD2D(void *dst, const void *src, size_t size, bool async) {
-    if (async) {
-        INFINICORE_CHECK_ERROR(infinirtMemcpyAsync(dst, src, size, INFINIRT_MEMCPY_D2D, stream_));
+    if (async && device_.getType() != Device::Type::CPU) {
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::MemcpyAsync(dst, src, size, infini::rt::runtime::kMemcpyDeviceToDevice, bridge::infini::rt::to_rt_stream(stream_))));
     } else {
-        INFINICORE_CHECK_ERROR(infinirtMemcpy(dst, src, size, INFINIRT_MEMCPY_D2D));
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::Memcpy(dst, src, size, infini::rt::runtime::kMemcpyDeviceToDevice)));
     }
 }
 
 void Runtime::setDeviceMemory(void *ptr, int value, size_t count) {
-    INFINICORE_CHECK_ERROR(infinirtMemset(ptr, value, count));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::Memset(ptr, value, count)));
 }
 
 void Runtime::setDeviceMemoryAsync(void *ptr, int value, size_t count, infinirtStream_t stream) {
-    INFINICORE_CHECK_ERROR(infinirtMemsetAsync(ptr, value, count, stream));
+    if (device_.getType() != Device::Type::CPU) {
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::MemsetAsync(ptr, value, count, bridge::infini::rt::to_rt_stream(stream))));
+    } else {
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::Memset(ptr, value, count)));
+    }
 }
 
 // Timing method implementations
 infinirtEvent_t Runtime::createEvent() {
     infinirtEvent_t event;
-    INFINICORE_CHECK_ERROR(infinirtEventCreate(&event));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::EventCreate(&event)));
     return event;
 }
 
 infinirtEvent_t Runtime::createEventWithFlags(uint32_t flags) {
     infinirtEvent_t event;
-    INFINICORE_CHECK_ERROR(infinirtEventCreateWithFlags(&event, flags));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::EventCreateWithFlags(&event, flags)));
     return event;
 }
 
@@ -147,26 +157,24 @@ void Runtime::recordEvent(infinirtEvent_t event, infinirtStream_t stream) {
     if (stream == nullptr) {
         stream = stream_;
     }
-    INFINICORE_CHECK_ERROR(infinirtEventRecord(event, stream));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::EventRecord(event, bridge::infini::rt::to_rt_stream(stream))));
 }
 
 bool Runtime::queryEvent(infinirtEvent_t event) {
-    infinirtEventStatus_t status;
-    INFINICORE_CHECK_ERROR(infinirtEventQuery(event, &status));
-    return status == INFINIRT_EVENT_COMPLETE;
+    return bridge::infini::rt::translate(infini::rt::runtime::EventQuery(event)) == INFINI_STATUS_SUCCESS;
 }
 
 void Runtime::synchronizeEvent(infinirtEvent_t event) {
-    INFINICORE_CHECK_ERROR(infinirtEventSynchronize(event));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::EventSynchronize(event)));
 }
 
 void Runtime::destroyEvent(infinirtEvent_t event) {
-    INFINICORE_CHECK_ERROR(infinirtEventDestroy(event));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::EventDestroy(event)));
 }
 
 float Runtime::elapsedTime(infinirtEvent_t start, infinirtEvent_t end) {
     float ms;
-    INFINICORE_CHECK_ERROR(infinirtEventElapsedTime(&ms, start, end));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::EventElapsedTime(&ms, start, end)));
     return ms;
 }
 
@@ -175,7 +183,7 @@ void Runtime::streamWaitEvent(infinirtStream_t stream, infinirtEvent_t event) {
     if (stream == nullptr) {
         stream = stream_;
     }
-    INFINICORE_CHECK_ERROR(infinirtStreamWaitEvent(stream, event));
+    INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(infini::rt::runtime::StreamWaitEvent(bridge::infini::rt::to_rt_stream(stream), event, 0)));
 }
 
 bool Runtime::isGraphRecording() const {
diff --git a/src/infinicore/graph/graph.cc b/src/infinicore/graph/graph.cc
index 3b8fc57e5..441148f8e 100644
--- a/src/infinicore/graph/graph.cc
+++ b/src/infinicore/graph/graph.cc
@@ -1,11 +1,19 @@
 #include "graph_manager.hpp"
 
+#include "../../bridge/infini/rt.hpp"
 #include "../utils.hpp"
 #include "infinicore/context/context.hpp"
-#include <infinirt.h>
+
+#ifdef USE_INFINIRT_GRAPH
+#include <infini/rt.h>
+#endif
 
 namespace infinicore::graph {
 
+#ifdef USE_INFINIRT_GRAPH
+namespace rt_runtime = ::infini::rt::runtime;
+#endif
+
 /* =========================
  * GraphTensor
  * ========================= */
@@ -31,40 +39,49 @@ DispatchableGraphOperator::~DispatchableGraphOperator() {
  * Graph
  * ========================= */
 
+#ifdef USE_INFINIRT_GRAPH
 struct Graph::DeviceGraph {
-    infinirtGraph_t graph;
-    infinirtGraphExec_t exec;
-    infinirtGraphNode_t node;
-    std::vector<char> log_buffer;
-
-    DeviceGraph() {
-        log_buffer.resize(4 * 1024);
-    }
+    rt_runtime::Graph graph = nullptr;
+    rt_runtime::GraphExec exec = nullptr;
+    rt_runtime::Stream stream = nullptr;
+    ::infini::rt::Device::Type device_type = ::infini::rt::Device::Type::kCount;
+    int device_index = 0;
 
     ~DeviceGraph() {
+        if (device_type != ::infini::rt::Device::Type::kCount) {
+            ::infini::rt::set_runtime_device_type(device_type);
+            (void)rt_runtime::SetDevice(device_index);
+        }
         if (exec) {
-            infinirtGraphExecDestroy(exec);
+            (void)rt_runtime::GraphExecDestroy(exec);
         }
         if (graph) {
-            infinirtGraphDestroy(graph);
+            (void)rt_runtime::GraphDestroy(graph);
         }
     }
 
     void launch() {
-        INFINICORE_CHECK_ERROR(infinirtGraphLuanch(exec, context::getStream()));
+        ::infini::rt::set_runtime_device_type(device_type);
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(rt_runtime::SetDevice(device_index)));
+        INFINICORE_CHECK_ERROR(bridge::infini::rt::translate(rt_runtime::GraphLaunch(exec, stream)));
     }
 };
+#else
+struct Graph::DeviceGraph {};
+#endif
 
 Graph::Graph() {
 }
 
 void Graph::run() const {
+#ifdef USE_INFINIRT_GRAPH
     if (device_graph_ != nullptr && device_graph_.get()->exec != nullptr) {
         device_graph_.get()->launch();
-    } else {
-        for (auto &op : op_list_) {
-            op->run();
-        }
+        return;
+    }
+#endif
+    for (auto &op : op_list_) {
+        op->run();
     }
 }
 
@@ -73,8 +90,13 @@ void Graph::add_operator(std::shared_ptr<GraphOperator> op) {
 }
 
 void Graph::instantiate() {
+#ifdef USE_INFINIRT_GRAPH
     // Reset device graph
     device_graph_ = std::make_unique<DeviceGraph>();
+    auto current_device = context::getDevice();
+    device_graph_->device_type = bridge::infini::rt::translate(static_cast<infiniDevice_t>(current_device.getType()));
+    device_graph_->device_index = static_cast<int>(current_device.getIndex());
+    device_graph_->stream = bridge::infini::rt::to_rt_stream(context::getStream());
 
     // warmup
     for (size_t iter = 0; iter < 5; ++iter) {
@@ -82,36 +104,45 @@ void Graph::instantiate() {
     }
     infinicore::context::syncStream();
 
-    if (infinirtStreamBeginCapture(
-            context::getStream(),
-            INFINIRT_STREAM_CAPTURE_MODE_RELAXED)
-        != INFINI_STATUS_SUCCESS) {
+    auto begin_status = bridge::infini::rt::translate(rt_runtime::StreamBeginCapture(
+        device_graph_->stream,
+        rt_runtime::StreamCaptureMode::kStreamCaptureModeRelaxed));
+    if (begin_status != INFINI_STATUS_SUCCESS) {
+        spdlog::warn("Fail to begin device graph capture.");
+        device_graph_.reset();
         return;
     }
 
     // Run and record
     this->run();
 
-    if (infinirtStreamEndCapture(
-            context::getStream(),
-            &device_graph_.get()->graph)
-        != INFINI_STATUS_SUCCESS) {
+    auto end_status = bridge::infini::rt::translate(rt_runtime::StreamEndCapture(
+        device_graph_->stream,
+        &device_graph_->graph));
+    if (end_status != INFINI_STATUS_SUCCESS) {
+        spdlog::warn("Fail to end device graph capture.");
+        device_graph_.reset();
         return;
     }
 
-    if (infinirtGraphInstantiate(
-            &device_graph_.get()->exec,
-            device_graph_.get()->graph,
-            &device_graph_.get()->node,
-            device_graph_.get()->log_buffer.data(),
-            device_graph_.get()->log_buffer.size())
-        != INFINI_STATUS_SUCCESS) {
+    auto instantiate_status = bridge::infini::rt::translate(rt_runtime::GraphInstantiate(
+        &device_graph_->exec,
+        device_graph_->graph));
+    if (instantiate_status != INFINI_STATUS_SUCCESS) {
         static bool warned_once = false;
         if (!warned_once) {
             warned_once = true;
-            spdlog::warn("Fail to instantiate device graph: {}", std::string(device_graph_.get()->log_buffer.data()));
+            spdlog::warn("Fail to instantiate device graph.");
         }
+        device_graph_.reset();
+        return;
     }
+    static bool logged_once = false;
+    if (!logged_once) {
+        logged_once = true;
+        spdlog::info("Using InfiniRT C++ graph runtime API for graph capture and replay.");
+    }
+#endif
 }
 
 Graph::~Graph() = default;
diff --git a/src/infinicore/nn/rope.cc b/src/infinicore/nn/rope.cc
index 115e162e2..c8ab4f02d 100644
--- a/src/infinicore/nn/rope.cc
+++ b/src/infinicore/nn/rope.cc
@@ -1,6 +1,7 @@
 #include "infinicore/nn/rope.hpp"
 #include "../../utils.h"
 #include "../utils.hpp"
+#include "infinicore/context/context.hpp"
 #include "infinicore/ops/mrope.hpp"
 #include "infinicore/ops/rope.hpp"
 #include <algorithm>
@@ -12,6 +13,16 @@
 #include <vector>
 
 namespace infinicore::nn {
+namespace {
+
+void sync_cache_copy(const Device &device) {
+    if (device.getType() == Device::Type::ASCEND) {
+        // Ascend H2D copies are async, and these host buffers leave scope after init.
+        infinicore::context::syncStream();
+    }
+}
+
+} // namespace
 
 RoPE::RoPE(size_t head_dim,
            size_t rotary_dim,
@@ -95,6 +106,7 @@ void RoPE::initialize_cache() {
         auto cos_f32_cpu = Tensor::from_blob(cos_data.data(), {max_seq_len_, cache_dim}, DataType::F32, cpu_device);
         sin_cache_->copy_from(sin_f32_cpu);
         cos_cache_->copy_from(cos_f32_cpu);
+        sync_cache_copy(device_);
     } else if (dtype_ == DataType::BF16) {
         // Convert F32 to BF16 using the same conversion as Python's ml_dtypes.bfloat16
         // This uses round-to-nearest-even (matching _f32_to_bf16 implementation)
@@ -112,6 +124,7 @@ void RoPE::initialize_cache() {
         // copy_from handles cross-device copying to target device
         sin_cache_->copy_from(sin_bf16_cpu);
         cos_cache_->copy_from(cos_bf16_cpu);
+        sync_cache_copy(device_);
     } else if (dtype_ == DataType::F16) {
         // Convert F32 to F16
         std::vector<fp16_t> sin_f16_data(max_seq_len_ * cache_dim);
@@ -127,6 +140,7 @@ void RoPE::initialize_cache() {
 
         sin_cache_->copy_from(sin_f16_cpu);
         cos_cache_->copy_from(cos_f16_cpu);
+        sync_cache_copy(device_);
     } else {
         throw std::runtime_error(
             "RoPE cache dtype conversion not yet supported for dtype: "
diff --git a/src/infinicore/ops/mul_scalar/mul_scalar.cc b/src/infinicore/ops/mul_scalar/mul_scalar.cc
new file mode 100644
index 000000000..b5be25e7e
--- /dev/null
+++ b/src/infinicore/ops/mul_scalar/mul_scalar.cc
@@ -0,0 +1,27 @@
+#include "infinicore/ops/mul_scalar.hpp"
+#include "../../utils.hpp"
+
+namespace infinicore::op {
+
+INFINICORE_GRAPH_OP_DISPATCHERS_IMPL(MulScalar);
+
+MulScalar::MulScalar(Tensor c, const Tensor &a, double alpha) {
+    INFINICORE_ASSERT_TENSORS_SAME_DEVICE(c, a);
+    INFINICORE_GRAPH_OP_DISPATCH(c->device().getType(), c, a, alpha);
+}
+
+void MulScalar::execute(Tensor c, const Tensor &a, double alpha) {
+    INFINICORE_GRAPH_OP_RECORD_OR_RUN(MulScalar, c, a, alpha);
+}
+
+Tensor mul_scalar(const Tensor &a, double alpha) {
+    auto c = Tensor::empty(a->shape(), a->dtype(), a->device());
+    mul_scalar_(c, a, alpha);
+    return c;
+}
+
+void mul_scalar_(Tensor c, const Tensor &a, double alpha) {
+    MulScalar::execute(c, a, alpha);
+}
+
+} // namespace infinicore::op
diff --git a/src/infinicore/ops/mul_scalar/mul_scalar_infiniop.cc b/src/infinicore/ops/mul_scalar/mul_scalar_infiniop.cc
new file mode 100644
index 000000000..13f3c615c
--- /dev/null
+++ b/src/infinicore/ops/mul_scalar/mul_scalar_infiniop.cc
@@ -0,0 +1,52 @@
+#include "infinicore/ops/mul_scalar.hpp"
+
+#include "../infiniop_impl.hpp"
+
+namespace infinicore::op::mul_scalar_impl::infiniop {
+
+INFINIOP_CACHABLE_DESCRIPTOR(Descriptor, MulScalar, 100);
+
+struct PlannedMeta {
+    std::shared_ptr<Descriptor> descriptor;
+    graph::GraphTensor workspace, c, a;
+    double alpha;
+};
+
+void *plan(Tensor c, const Tensor &a, double alpha) {
+    size_t seed = hash_combine(c, a, alpha);
+
+    INFINIOP_CACHABLE_DESCRIPTOR_GET_OR_CREATE(
+        Descriptor, descriptor, MulScalar,
+        seed, c->desc(), a->desc());
+
+    INFINIOP_WORKSPACE_TENSOR(workspace, MulScalar, descriptor);
+
+    return new PlannedMeta{
+        descriptor,
+        graph::GraphTensor(workspace),
+        graph::GraphTensor(c),
+        graph::GraphTensor(a),
+        alpha};
+}
+
+void run(void *planned_meta) {
+    auto planned = reinterpret_cast<PlannedMeta *>(planned_meta);
+
+    INFINICORE_CHECK_ERROR(infiniopMulScalar(
+        planned->descriptor->desc,
+        planned->workspace->data(),
+        planned->workspace->numel(),
+        planned->c->data(),
+        planned->a->data(),
+        planned->alpha,
+        context::getStream()));
+}
+
+void cleanup(void **planned_meta_ptr) {
+    delete *reinterpret_cast<PlannedMeta **>(planned_meta_ptr);
+    *planned_meta_ptr = nullptr;
+}
+
+INFINICORE_GRAPH_OP_REGISTER_ALLDEVICE(MulScalar, &plan, &run, &cleanup);
+
+} // namespace infinicore::op::mul_scalar_impl::infiniop
diff --git a/src/infinicore/pybind11/ops.hpp b/src/infinicore/pybind11/ops.hpp
index 38aceb2fd..087261382 100644
--- a/src/infinicore/pybind11/ops.hpp
+++ b/src/infinicore/pybind11/ops.hpp
@@ -85,6 +85,7 @@
 #include "ops/moe_topk_softmax.hpp"
 #include "ops/mrope.hpp"
 #include "ops/mul.hpp"
+#include "ops/mul_scalar.hpp"
 #include "ops/multi_margin_loss.hpp"
 #include "ops/nrm2.hpp"
 #include "ops/pad.hpp"
@@ -186,6 +187,7 @@ inline void bind(py::module &m) {
     bind_mamba_selective_scan(m);
     bind_kron(m);
     bind_mul(m);
+    bind_mul_scalar(m);
     bind_nrm2(m);
     bind_mha_kvcache(m);
     bind_mha_varlen(m);
diff --git a/src/infinicore/pybind11/ops/mul_scalar.hpp b/src/infinicore/pybind11/ops/mul_scalar.hpp
new file mode 100644
index 000000000..2bce862bc
--- /dev/null
+++ b/src/infinicore/pybind11/ops/mul_scalar.hpp
@@ -0,0 +1,26 @@
+#pragma once
+
+#include <pybind11/pybind11.h>
+
+#include "infinicore/ops/mul_scalar.hpp"
+
+namespace py = pybind11;
+
+namespace infinicore::ops {
+
+inline void bind_mul_scalar(py::module &m) {
+    m.def("mul_scalar",
+          &op::mul_scalar,
+          py::arg("a"),
+          py::arg("alpha"),
+          R"doc(Multiply a tensor by a host scalar.)doc");
+
+    m.def("mul_scalar_",
+          &op::mul_scalar_,
+          py::arg("c"),
+          py::arg("a"),
+          py::arg("alpha"),
+          R"doc(Out-of-place tensor-scalar multiplication into c.)doc");
+}
+
+} // namespace infinicore::ops
diff --git a/src/infiniop/devices/handle.cc b/src/infiniop/devices/handle.cc
index b36d3b6bf..9ded0587b 100644
--- a/src/infiniop/devices/handle.cc
+++ b/src/infiniop/devices/handle.cc
@@ -1,6 +1,6 @@
 #include "infiniop/handle.h"
+#include "../../bridge/infini/rt.hpp"
 #include "../../utils.h"
-#include "infinirt.h"
 
 #ifdef ENABLE_CPU_API
 #include "cpu/cpu_handle.h"
@@ -24,14 +24,29 @@
 #include "metax/metax_handle.h"
 #endif
 
+__INFINI_C infiniStatus_t infiniopSetRuntimeDevice(infiniDevice_t device, int device_id) {
+    auto rt_device = infinicore::bridge::infini::rt::translate(device);
+    if (rt_device == infini::rt::Device::Type::kCount) {
+        return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
+    }
+
+    infini::rt::set_runtime_device_type(rt_device);
+    CHECK_STATUS(infinicore::bridge::infini::rt::translate(infini::rt::runtime::SetDevice(device_id)));
+    return INFINI_STATUS_SUCCESS;
+}
+
 __INFINI_C infiniStatus_t infiniopCreateHandle(infiniopHandle_t *handle_ptr) {
     if (handle_ptr == nullptr) {
         return INFINI_STATUS_NULL_POINTER;
     }
 
-    infiniDevice_t device;
-    int device_id;
-    CHECK_STATUS(infinirtGetDevice(&device, &device_id));
+    infiniDevice_t device = infinicore::bridge::infini::rt::translate(infini::rt::runtime_device_type());
+    if (device == INFINI_DEVICE_TYPE_COUNT) {
+        return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
+    }
+
+    int device_id = 0;
+    CHECK_STATUS(infinicore::bridge::infini::rt::translate(infini::rt::runtime::GetDevice(&device_id)));
 
 #define CREATE(CASE, NAMESPACE) \
     case CASE:                  \
diff --git a/src/infiniop/ops/mul_scalar/cpu/mul_scalar_cpu.cc b/src/infiniop/ops/mul_scalar/cpu/mul_scalar_cpu.cc
new file mode 100644
index 000000000..7049e5b6f
--- /dev/null
+++ b/src/infiniop/ops/mul_scalar/cpu/mul_scalar_cpu.cc
@@ -0,0 +1,98 @@
+#include "mul_scalar_cpu.h"
+#include "../../../devices/cpu/common_cpu.h"
+
+namespace op::mul_scalar::cpu {
+
+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::cpu::Handle *>(handle_);
+    auto result = MulScalarInfo::create(output_desc, input_desc);
+    CHECK_RESULT(result);
+
+    *desc_ptr = new Descriptor(
+        result.take(),
+        nullptr,
+        0,
+        handle->device,
+        handle->device_id);
+
+    return INFINI_STATUS_SUCCESS;
+}
+
+template <typename T>
+inline T mulScalarHost(T value, double alpha) {
+    if constexpr (std::is_same_v<T, fp16_t> || std::is_same_v<T, bf16_t>) {
+        return utils::cast<T>(utils::cast<float>(value) * static_cast<float>(alpha));
+    } else {
+        return value * static_cast<T>(alpha);
+    }
+}
+
+template <typename T>
+infiniStatus_t calculateMulScalar(const MulScalarInfo &info, T *output, const T *input, double alpha) {
+    const auto &elementwise_info = info.elementwise_info;
+    const ptrdiff_t output_size = elementwise_info.getOutputSize();
+
+    if (info.contiguous) {
+#pragma omp parallel for if (output_size > 1024)
+        for (ptrdiff_t i = 0; i < output_size; ++i) {
+            output[i] = mulScalarHost(input[i], alpha);
+        }
+        return INFINI_STATUS_SUCCESS;
+    }
+
+#pragma omp parallel for if (output_size > 1024)
+    for (ptrdiff_t i = 0; i < output_size; ++i) {
+        size_t out_idx = elementwise_info.isOutputContiguous()
+                           ? i
+                           : op::common_cpu::indexToOffset(
+                               i,
+                               elementwise_info.getNdim(),
+                               elementwise_info.getOutputShape(),
+                               elementwise_info.getOutputStrides());
+        size_t in_idx = elementwise_info.getInputContiguous()[0]
+                          ? i
+                          : op::common_cpu::indexToOffset(
+                              i,
+                              elementwise_info.getNdim(),
+                              elementwise_info.getInputShape(0),
+                              elementwise_info.getInputStrides(0));
+        output[out_idx] = mulScalarHost(input[in_idx], alpha);
+    }
+
+    return INFINI_STATUS_SUCCESS;
+}
+
+infiniStatus_t Descriptor::calculate(
+    void *workspace,
+    size_t workspace_size,
+    void *output,
+    const void *input,
+    double alpha,
+    void *stream) const {
+
+    (void)workspace;
+    (void)workspace_size;
+    (void)stream;
+
+    switch (_info.data_type) {
+    case INFINI_DTYPE_F16:
+        return calculateMulScalar(_info, static_cast<fp16_t *>(output), static_cast<const fp16_t *>(input), alpha);
+    case INFINI_DTYPE_BF16:
+        return calculateMulScalar(_info, static_cast<bf16_t *>(output), static_cast<const bf16_t *>(input), alpha);
+    case INFINI_DTYPE_F32:
+        return calculateMulScalar(_info, static_cast<float *>(output), static_cast<const float *>(input), alpha);
+    case INFINI_DTYPE_F64:
+        return calculateMulScalar(_info, static_cast<double *>(output), static_cast<const double *>(input), alpha);
+    default:
+        return INFINI_STATUS_BAD_TENSOR_DTYPE;
+    }
+}
+
+} // namespace op::mul_scalar::cpu
diff --git a/src/infiniop/ops/mul_scalar/cpu/mul_scalar_cpu.h b/src/infiniop/ops/mul_scalar/cpu/mul_scalar_cpu.h
new file mode 100644
index 000000000..19217a35c
--- /dev/null
+++ b/src/infiniop/ops/mul_scalar/cpu/mul_scalar_cpu.h
@@ -0,0 +1,9 @@
+#ifndef __MUL_SCALAR_CPU_H__
+#define __MUL_SCALAR_CPU_H__
+
+#include "../../../elementwise/cpu/elementwise_cpu.h"
+#include "../mul_scalar.h"
+
+MUL_SCALAR_DESCRIPTOR(cpu)
+
+#endif // __MUL_SCALAR_CPU_H__
diff --git a/src/infiniop/ops/mul_scalar/info.h b/src/infiniop/ops/mul_scalar/info.h
new file mode 100644
index 000000000..68ad91d31
--- /dev/null
+++ b/src/infiniop/ops/mul_scalar/info.h
@@ -0,0 +1,47 @@
+#ifndef __MUL_SCALAR_INFO_H__
+#define __MUL_SCALAR_INFO_H__
+
+#include "../../../utils.h"
+#include "../../elementwise/elementwise.h"
+#include "../../tensor.h"
+
+class MulScalarInfo {
+private:
+    MulScalarInfo(
+        op::elementwise::ElementwiseInfo elementwise_info_,
+        infiniDtype_t data_type_,
+        bool contiguous_)
+        : elementwise_info(std::move(elementwise_info_)),
+          data_type(data_type_),
+          contiguous(contiguous_) {}
+
+public:
+    op::elementwise::ElementwiseInfo elementwise_info;
+    infiniDtype_t data_type;
+    bool contiguous;
+
+    size_t numel() const { return elementwise_info.getOutputSize(); }
+
+    static utils::Result<MulScalarInfo> create(
+        infiniopTensorDescriptor_t output_desc,
+        infiniopTensorDescriptor_t input_desc) {
+
+        CHECK_OR_RETURN(output_desc != nullptr, INFINI_STATUS_NULL_POINTER);
+        CHECK_OR_RETURN(input_desc != nullptr, INFINI_STATUS_NULL_POINTER);
+
+        auto data_type = output_desc->dtype();
+        CHECK_OR_RETURN(input_desc->dtype() == data_type, INFINI_STATUS_BAD_TENSOR_DTYPE);
+        CHECK_DTYPE(data_type, INFINI_DTYPE_F16, INFINI_DTYPE_BF16, INFINI_DTYPE_F32, INFINI_DTYPE_F64);
+        CHECK_SAME_SHAPE(output_desc->shape(), input_desc->shape());
+
+        auto elementwise_info_result = op::elementwise::ElementwiseInfo::create(output_desc, {input_desc});
+        CHECK_RESULT(elementwise_info_result);
+
+        return utils::Result<MulScalarInfo>(MulScalarInfo(
+            elementwise_info_result.take(),
+            data_type,
+            output_desc->isContiguous() && input_desc->isContiguous()));
+    }
+};
+
+#endif // __MUL_SCALAR_INFO_H__
diff --git a/src/infiniop/ops/mul_scalar/mul_scalar.h b/src/infiniop/ops/mul_scalar/mul_scalar.h
new file mode 100644
index 000000000..c0ca68e18
--- /dev/null
+++ b/src/infiniop/ops/mul_scalar/mul_scalar.h
@@ -0,0 +1,49 @@
+#ifndef __MUL_SCALAR_H__
+#define __MUL_SCALAR_H__
+
+#include "../../operator.h"
+#include "info.h"
+
+#include <memory>
+
+#define MUL_SCALAR_DESCRIPTOR(NAMESPACE)                                      \
+                                                                              \
+    namespace op::mul_scalar::NAMESPACE {                                     \
+    class Descriptor final : public InfiniopDescriptor {                      \
+        MulScalarInfo _info;                                                  \
+        std::unique_ptr<op::elementwise::NAMESPACE::DeviceImpl> _device_info; \
+        size_t _workspace_size;                                               \
+                                                                              \
+        Descriptor(                                                           \
+            MulScalarInfo info,                                               \
+            op::elementwise::NAMESPACE::DeviceImpl *device_info,              \
+            size_t workspace_size_,                                           \
+            infiniDevice_t device_type,                                       \
+            int device_id)                                                    \
+            : InfiniopDescriptor{device_type, device_id},                     \
+              _info(std::move(info)),                                         \
+              _device_info(std::move(device_info)),                           \
+              _workspace_size(workspace_size_) {}                             \
+                                                                              \
+    public:                                                                   \
+        ~Descriptor();                                                        \
+                                                                              \
+        size_t workspaceSize() const { return _workspace_size; }              \
+                                                                              \
+        static infiniStatus_t create(                                         \
+            infiniopHandle_t handle,                                          \
+            Descriptor **desc_ptr,                                            \
+            infiniopTensorDescriptor_t output_desc,                           \
+            infiniopTensorDescriptor_t input_desc);                           \
+                                                                              \
+        infiniStatus_t calculate(                                             \
+            void *workspace,                                                  \
+            size_t workspace_size,                                            \
+            void *output,                                                     \
+            const void *input,                                                \
+            double alpha,                                                     \
+            void *stream) const;                                              \
+    };                                                                        \
+    }
+
+#endif // __MUL_SCALAR_H__
diff --git a/src/infiniop/ops/mul_scalar/nvidia/mul_scalar_nvidia.cu b/src/infiniop/ops/mul_scalar/nvidia/mul_scalar_nvidia.cu
new file mode 100644
index 000000000..06ca80553
--- /dev/null
+++ b/src/infiniop/ops/mul_scalar/nvidia/mul_scalar_nvidia.cu
@@ -0,0 +1,144 @@
+#include "mul_scalar_nvidia.cuh"
+
+#include "../../../devices/nvidia/nvidia_common.cuh"
+#include "../../../devices/nvidia/nvidia_kernel_common.cuh"
+#include "../../../elementwise/nvidia/elementwise_nvidia.cuh"
+
+#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<cudaStream_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 = cudaGetLastError();
+    return err == cudaSuccess ? 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::nvidia::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::nvidia {
+
+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::nvidia::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::nvidia::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::nvidia
diff --git a/src/infiniop/ops/mul_scalar/nvidia/mul_scalar_nvidia.cuh b/src/infiniop/ops/mul_scalar/nvidia/mul_scalar_nvidia.cuh
new file mode 100644
index 000000000..12f3553b6
--- /dev/null
+++ b/src/infiniop/ops/mul_scalar/nvidia/mul_scalar_nvidia.cuh
@@ -0,0 +1,9 @@
+#ifndef __MUL_SCALAR_NVIDIA_H__
+#define __MUL_SCALAR_NVIDIA_H__
+
+#include "../../../elementwise/nvidia/elementwise_nvidia_api.cuh"
+#include "../mul_scalar.h"
+
+MUL_SCALAR_DESCRIPTOR(nvidia)
+
+#endif // __MUL_SCALAR_NVIDIA_H__
diff --git a/src/infiniop/ops/mul_scalar/operator.cc b/src/infiniop/ops/mul_scalar/operator.cc
new file mode 100644
index 000000000..a8f50496f
--- /dev/null
+++ b/src/infiniop/ops/mul_scalar/operator.cc
@@ -0,0 +1,155 @@
+#include "../../operator.h"
+#include "../../handle.h"
+#include "infiniop/ops/mul_scalar.h"
+
+#ifdef ENABLE_CPU_API
+#include "cpu/mul_scalar_cpu.h"
+#endif
+#if defined(ENABLE_NVIDIA_API) || defined(ENABLE_ILUVATAR_API) || defined(ENABLE_QY_API) || defined(ENABLE_ALI_API) || defined(ENABLE_HYGON_API)
+#include "nvidia/mul_scalar_nvidia.cuh"
+#endif
+
+__INFINI_C infiniStatus_t infiniopCreateMulScalarDescriptor(
+    infiniopHandle_t handle,
+    infiniopMulScalarDescriptor_t *desc_ptr,
+    infiniopTensorDescriptor_t output_desc,
+    infiniopTensorDescriptor_t input_desc) {
+
+#define CREATE(CASE, NAMESPACE)                                                   \
+    case CASE:                                                                    \
+        return op::mul_scalar::NAMESPACE::Descriptor::create(                     \
+            handle,                                                               \
+            reinterpret_cast<op::mul_scalar::NAMESPACE::Descriptor **>(desc_ptr), \
+            output_desc,                                                          \
+            input_desc)
+
+    switch (handle->device) {
+#ifdef ENABLE_CPU_API
+        CREATE(INFINI_DEVICE_CPU, cpu);
+#endif
+#ifdef ENABLE_NVIDIA_API
+        CREATE(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_ILUVATAR_API
+        CREATE(INFINI_DEVICE_ILUVATAR, nvidia);
+#endif
+#ifdef ENABLE_QY_API
+        CREATE(INFINI_DEVICE_QY, nvidia);
+#endif
+#ifdef ENABLE_ALI_API
+        CREATE(INFINI_DEVICE_ALI, nvidia);
+#endif
+#ifdef ENABLE_HYGON_API
+        CREATE(INFINI_DEVICE_HYGON, nvidia);
+#endif
+    default:
+        return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
+    }
+
+#undef CREATE
+}
+
+__INFINI_C infiniStatus_t infiniopGetMulScalarWorkspaceSize(infiniopMulScalarDescriptor_t desc, size_t *size) {
+
+#define GET(CASE, NAMESPACE)                                                                      \
+    case CASE:                                                                                    \
+        *size = reinterpret_cast<op::mul_scalar::NAMESPACE::Descriptor *>(desc)->workspaceSize(); \
+        return INFINI_STATUS_SUCCESS
+
+    switch (desc->device_type) {
+#ifdef ENABLE_CPU_API
+        GET(INFINI_DEVICE_CPU, cpu);
+#endif
+#ifdef ENABLE_NVIDIA_API
+        GET(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_ILUVATAR_API
+        GET(INFINI_DEVICE_ILUVATAR, nvidia);
+#endif
+#ifdef ENABLE_QY_API
+        GET(INFINI_DEVICE_QY, nvidia);
+#endif
+#ifdef ENABLE_ALI_API
+        GET(INFINI_DEVICE_ALI, nvidia);
+#endif
+#ifdef ENABLE_HYGON_API
+        GET(INFINI_DEVICE_HYGON, nvidia);
+#endif
+    default:
+        return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
+    }
+#undef GET
+}
+
+__INFINI_C infiniStatus_t infiniopMulScalar(
+    infiniopMulScalarDescriptor_t desc,
+    void *workspace,
+    size_t workspace_size,
+    void *output,
+    const void *input,
+    double alpha,
+    void *stream) {
+
+#define CALCULATE(CASE, NAMESPACE)                                                   \
+    case CASE:                                                                       \
+        return reinterpret_cast<const op::mul_scalar::NAMESPACE::Descriptor *>(desc) \
+            ->calculate(workspace, workspace_size, output, input, alpha, stream)
+
+    switch (desc->device_type) {
+#ifdef ENABLE_CPU_API
+        CALCULATE(INFINI_DEVICE_CPU, cpu);
+#endif
+#ifdef ENABLE_NVIDIA_API
+        CALCULATE(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_ILUVATAR_API
+        CALCULATE(INFINI_DEVICE_ILUVATAR, nvidia);
+#endif
+#ifdef ENABLE_QY_API
+        CALCULATE(INFINI_DEVICE_QY, nvidia);
+#endif
+#ifdef ENABLE_ALI_API
+        CALCULATE(INFINI_DEVICE_ALI, nvidia);
+#endif
+#ifdef ENABLE_HYGON_API
+        CALCULATE(INFINI_DEVICE_HYGON, nvidia);
+#endif
+    default:
+        return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
+    }
+
+#undef CALCULATE
+}
+
+__INFINI_C infiniStatus_t infiniopDestroyMulScalarDescriptor(infiniopMulScalarDescriptor_t desc) {
+
+#define DELETE(CASE, NAMESPACE)                                                       \
+    case CASE:                                                                        \
+        delete reinterpret_cast<const op::mul_scalar::NAMESPACE::Descriptor *>(desc); \
+        return INFINI_STATUS_SUCCESS
+
+    switch (desc->device_type) {
+#ifdef ENABLE_CPU_API
+        DELETE(INFINI_DEVICE_CPU, cpu);
+#endif
+#ifdef ENABLE_NVIDIA_API
+        DELETE(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_ILUVATAR_API
+        DELETE(INFINI_DEVICE_ILUVATAR, nvidia);
+#endif
+#ifdef ENABLE_QY_API
+        DELETE(INFINI_DEVICE_QY, nvidia);
+#endif
+#ifdef ENABLE_ALI_API
+        DELETE(INFINI_DEVICE_ALI, nvidia);
+#endif
+#ifdef ENABLE_HYGON_API
+        DELETE(INFINI_DEVICE_HYGON, nvidia);
+#endif
+    default:
+        return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
+    }
+
+#undef DELETE
+}
diff --git a/src/infiniop/ops/paged_attention/ascend/paged_attention_ascend_kernel.cpp b/src/infiniop/ops/paged_attention/ascend/paged_attention_ascend_kernel.cpp
index 0f47c2ff5..e6df4f50b 100644
--- a/src/infiniop/ops/paged_attention/ascend/paged_attention_ascend_kernel.cpp
+++ b/src/infiniop/ops/paged_attention/ascend/paged_attention_ascend_kernel.cpp
@@ -164,22 +164,22 @@ class PagedAttentionKernel {
 
 private:
     __aicore__ inline ptrdiff_t keyBase(size_t seq_idx, size_t kv_head_idx, size_t token_idx) {
-        const size_t block_idx = token_idx / _page_block_size;
+        const size_t page_block_idx = token_idx / _page_block_size;
         const size_t token_offset = token_idx % _page_block_size;
         const int64_t physical_block = loadIndex(
             _block_tables_gm,
-            static_cast<ptrdiff_t>(seq_idx) * _block_table_batch_stride + static_cast<ptrdiff_t>(block_idx));
+            static_cast<ptrdiff_t>(seq_idx) * _block_table_batch_stride + static_cast<ptrdiff_t>(page_block_idx));
         return static_cast<ptrdiff_t>(physical_block) * _k_batch_stride
              + static_cast<ptrdiff_t>(kv_head_idx) * _k_head_stride
              + static_cast<ptrdiff_t>(token_offset) * _k_row_stride;
     }
 
     __aicore__ inline ptrdiff_t valueBase(size_t seq_idx, size_t kv_head_idx, size_t token_idx) {
-        const size_t block_idx = token_idx / _page_block_size;
+        const size_t page_block_idx = token_idx / _page_block_size;
         const size_t token_offset = token_idx % _page_block_size;
         const int64_t physical_block = loadIndex(
             _block_tables_gm,
-            static_cast<ptrdiff_t>(seq_idx) * _block_table_batch_stride + static_cast<ptrdiff_t>(block_idx));
+            static_cast<ptrdiff_t>(seq_idx) * _block_table_batch_stride + static_cast<ptrdiff_t>(page_block_idx));
         return static_cast<ptrdiff_t>(physical_block) * _v_batch_stride
              + static_cast<ptrdiff_t>(kv_head_idx) * _v_head_stride
              + static_cast<ptrdiff_t>(token_offset) * _v_row_stride;
diff --git a/src/infiniop/ops/paged_attention_prefill/ascend/paged_attention_prefill_ascend_kernel.cpp b/src/infiniop/ops/paged_attention_prefill/ascend/paged_attention_prefill_ascend_kernel.cpp
index e1fce2785..db90b2056 100644
--- a/src/infiniop/ops/paged_attention_prefill/ascend/paged_attention_prefill_ascend_kernel.cpp
+++ b/src/infiniop/ops/paged_attention_prefill/ascend/paged_attention_prefill_ascend_kernel.cpp
@@ -197,22 +197,22 @@ class PagedAttentionPrefillKernel {
     }
 
     __aicore__ inline ptrdiff_t keyBase(size_t seq_idx, size_t kv_head_idx, size_t token_idx) {
-        const size_t block_idx = token_idx / _page_block_size;
+        const size_t page_block_idx = token_idx / _page_block_size;
         const size_t token_offset = token_idx % _page_block_size;
         const int64_t physical_block = loadPrefillIndex(
             _block_tables_gm,
-            static_cast<ptrdiff_t>(seq_idx) * _block_table_batch_stride + static_cast<ptrdiff_t>(block_idx));
+            static_cast<ptrdiff_t>(seq_idx) * _block_table_batch_stride + static_cast<ptrdiff_t>(page_block_idx));
         return static_cast<ptrdiff_t>(physical_block) * _k_batch_stride
              + static_cast<ptrdiff_t>(kv_head_idx) * _k_head_stride
              + static_cast<ptrdiff_t>(token_offset) * _k_row_stride;
     }
 
     __aicore__ inline ptrdiff_t valueBase(size_t seq_idx, size_t kv_head_idx, size_t token_idx) {
-        const size_t block_idx = token_idx / _page_block_size;
+        const size_t page_block_idx = token_idx / _page_block_size;
         const size_t token_offset = token_idx % _page_block_size;
         const int64_t physical_block = loadPrefillIndex(
             _block_tables_gm,
-            static_cast<ptrdiff_t>(seq_idx) * _block_table_batch_stride + static_cast<ptrdiff_t>(block_idx));
+            static_cast<ptrdiff_t>(seq_idx) * _block_table_batch_stride + static_cast<ptrdiff_t>(page_block_idx));
         return static_cast<ptrdiff_t>(physical_block) * _v_batch_stride
              + static_cast<ptrdiff_t>(kv_head_idx) * _v_head_stride
              + static_cast<ptrdiff_t>(token_offset) * _v_row_stride;
diff --git a/src/infiniop/ops/quant/per_channel_quant_int8/operator.cc b/src/infiniop/ops/quant/per_channel_quant_int8/operator.cc
index 3c5e7596c..273cf991b 100644
--- a/src/infiniop/ops/quant/per_channel_quant_int8/operator.cc
+++ b/src/infiniop/ops/quant/per_channel_quant_int8/operator.cc
@@ -2,7 +2,7 @@
 #include "../../../handle.h"
 #include "infiniop/ops/quant/per_channel_quant_int8.h"
 
-#if defined(ENABLE_NVIDIA_API) || defined(ENABLE_QY_API)
+#if defined(ENABLE_NVIDIA_API) || defined(ENABLE_QY_API) || (defined(ENABLE_HYGON_API) && !defined(INFINIOP_DISABLE_HYGON_PER_CHANNEL_QUANT_INT8))
 #include "nvidia/per_channel_quant_int8_nvidia.cuh"
 #endif
 #if defined(ENABLE_MOORE_API)
@@ -31,6 +31,9 @@ __INFINI_C infiniStatus_t infiniopCreatePerChannelQuantI8Descriptor(infiniopHand
 #ifdef ENABLE_QY_API
         CREATE(INFINI_DEVICE_QY, nvidia)
 #endif
+#if defined(ENABLE_HYGON_API) && !defined(INFINIOP_DISABLE_HYGON_PER_CHANNEL_QUANT_INT8)
+        CREATE(INFINI_DEVICE_HYGON, nvidia)
+#endif
 #ifdef ENABLE_MOORE_API
         CREATE(INFINI_DEVICE_MOORE, moore)
 #endif
@@ -52,6 +55,9 @@ __INFINI_C infiniStatus_t infiniopGetPerChannelQuantI8WorkspaceSize(infiniopPerC
 #ifdef ENABLE_QY_API
         GET(INFINI_DEVICE_QY, nvidia)
 #endif
+#if defined(ENABLE_HYGON_API) && !defined(INFINIOP_DISABLE_HYGON_PER_CHANNEL_QUANT_INT8)
+        GET(INFINI_DEVICE_HYGON, nvidia)
+#endif
 #ifdef ENABLE_MOORE_API
         GET(INFINI_DEVICE_MOORE, moore)
 #endif
@@ -81,6 +87,9 @@ __INFINI_C infiniStatus_t infiniopPerChannelQuantI8(infiniopPerChannelQuantI8Des
 #ifdef ENABLE_QY_API
         QUANT(INFINI_DEVICE_QY, nvidia)
 #endif
+#if defined(ENABLE_HYGON_API) && !defined(INFINIOP_DISABLE_HYGON_PER_CHANNEL_QUANT_INT8)
+        QUANT(INFINI_DEVICE_HYGON, nvidia)
+#endif
 #ifdef ENABLE_MOORE_API
         QUANT(INFINI_DEVICE_MOORE, moore)
 #endif
@@ -103,6 +112,9 @@ __INFINI_C infiniStatus_t infiniopDestroyPerChannelQuantI8Descriptor(infiniopPer
 #ifdef ENABLE_QY_API
         DESTROY(INFINI_DEVICE_QY, nvidia)
 #endif
+#if defined(ENABLE_HYGON_API) && !defined(INFINIOP_DISABLE_HYGON_PER_CHANNEL_QUANT_INT8)
+        DESTROY(INFINI_DEVICE_HYGON, nvidia)
+#endif
 #ifdef ENABLE_MOORE_API
         DESTROY(INFINI_DEVICE_MOORE, moore)
 #endif
diff --git a/src/infiniop/ops/scaled_mm/nvidia/int8_gemm_nvidia.cu b/src/infiniop/ops/scaled_mm/nvidia/int8_gemm_nvidia.cu
index 36351cf42..7532ac93e 100644
--- a/src/infiniop/ops/scaled_mm/nvidia/int8_gemm_nvidia.cu
+++ b/src/infiniop/ops/scaled_mm/nvidia/int8_gemm_nvidia.cu
@@ -64,7 +64,7 @@ infiniStatus_t Descriptor::create(
     return INFINI_STATUS_SUCCESS;
 }
 
-#ifdef ENABLE_QY_API
+#if defined(ENABLE_QY_API) || defined(ENABLE_HYGON_API)
 template <unsigned int BLOCK_SIZE, typename Tdata>
 infiniStatus_t Descriptor::launchKernel(const I8GemmInfo &info, Tdata *y, const Tdata *bias, const int8_t *x_packed, const float *x_scale, const int8_t *w_packed, const float *w_scale, void *stream_, void *workspace) const {
     cudaStream_t stream = (cudaStream_t)stream_;
@@ -179,7 +179,7 @@ infiniStatus_t Descriptor::calculate(
     } else {
         return INFINI_STATUS_NOT_IMPLEMENTED;
     }
-#elif defined ENABLE_QY_API
+#elif defined(ENABLE_QY_API) || defined(ENABLE_HYGON_API)
 #define CALCULATE_LINEAR(BLOCK_SIZE, TDATA) \
     launchKernel<BLOCK_SIZE, TDATA>(_info, (TDATA *)out, (const TDATA *)bias, (const int8_t *)a, (const float *)a_scale, (const int8_t *)b, (const float *)b_scale, stream, workspace)
 #define CALCULATE_LINEAR_WITH_BLOCK_SIZE(BLOCK_SIZE)            \
diff --git a/src/infiniop/ops/scaled_mm/operator.cc b/src/infiniop/ops/scaled_mm/operator.cc
index b852ab6c0..4b24330b6 100644
--- a/src/infiniop/ops/scaled_mm/operator.cc
+++ b/src/infiniop/ops/scaled_mm/operator.cc
@@ -2,7 +2,7 @@
 #include "../../handle.h"
 #include "infiniop/ops/int8_gemm.h"
 
-#if defined(ENABLE_NVIDIA_API) || defined(ENABLE_QY_API)
+#if defined(ENABLE_NVIDIA_API) || defined(ENABLE_QY_API) || defined(ENABLE_HYGON_API)
 #include "nvidia/int8_gemm_nvidia.cuh"
 #endif
 
@@ -36,6 +36,9 @@ __INFINI_C infiniStatus_t infiniopCreateI8GemmDescriptor(infiniopHandle_t handle
 #if defined(ENABLE_QY_API)
         CREATE(INFINI_DEVICE_QY, nvidia)
 #endif
+#if defined(ENABLE_HYGON_API)
+        CREATE(INFINI_DEVICE_HYGON, nvidia)
+#endif
 #if defined(ENABLE_MOORE_API)
         CREATE(INFINI_DEVICE_MOORE, moore)
 #endif
@@ -57,6 +60,9 @@ __INFINI_C infiniStatus_t infiniopGetI8GemmWorkspaceSize(infiniopI8GemmDescripto
 #if defined(ENABLE_QY_API)
         GET(INFINI_DEVICE_QY, nvidia)
 #endif
+#if defined(ENABLE_HYGON_API)
+        GET(INFINI_DEVICE_HYGON, nvidia)
+#endif
 #if defined(ENABLE_MOORE_API)
         GET(INFINI_DEVICE_MOORE, moore)
 #endif
@@ -87,6 +93,9 @@ __INFINI_C infiniStatus_t infiniopI8Gemm(infiniopI8GemmDescriptor_t desc,
 #if defined(ENABLE_QY_API)
         CACULATE(INFINI_DEVICE_QY, nvidia)
 #endif
+#if defined(ENABLE_HYGON_API)
+        CACULATE(INFINI_DEVICE_HYGON, nvidia)
+#endif
 #if defined(ENABLE_MOORE_API)
         CACULATE(INFINI_DEVICE_MOORE, moore)
 #endif
@@ -108,6 +117,9 @@ __INFINI_C infiniStatus_t infiniopDestroyI8GemmDescriptor(infiniopI8GemmDescript
 #if defined(ENABLE_QY_API)
         DESTROY(INFINI_DEVICE_QY, nvidia)
 #endif
+#if defined(ENABLE_HYGON_API)
+        DESTROY(INFINI_DEVICE_HYGON, nvidia)
+#endif
 #if defined(ENABLE_MOORE_API)
         DESTROY(INFINI_DEVICE_MOORE, moore)
 #endif
diff --git a/test/infiniop/libinfiniop/liboperators.py b/test/infiniop/libinfiniop/liboperators.py
index ceffd5141..48ce04fe4 100644
--- a/test/infiniop/libinfiniop/liboperators.py
+++ b/test/infiniop/libinfiniop/liboperators.py
@@ -1,11 +1,12 @@
+import ctypes
 import os
 import platform
-import ctypes
-from ctypes import c_int, c_int64, c_uint64, POINTER
+from ctypes import POINTER, c_int, c_int64, c_uint64
+from pathlib import Path
+
 from .datatypes import *
 from .devices import *
 from .op_register import OpRegister
-from pathlib import Path
 from .structs import *
 
 INFINI_ROOT = os.getenv("INFINI_ROOT") or str(Path.home() / ".infini")
@@ -45,12 +46,12 @@ def find_library_in_ld_path(subdir, library_name):
         libop_path = find_library_in_ld_path("lib", "libinfiniop.so")
         librt_path = find_library_in_ld_path("lib", "libinfinirt.so")
 
-    assert (
-        libop_path is not None
-    ), f"Cannot find infiniop.dll or libinfiniop.so. Check if INFINI_ROOT is set correctly."
-    assert (
-        librt_path is not None
-    ), f"Cannot find infinirt.dll or libinfinirt.so. Check if INFINI_ROOT is set correctly."
+    assert libop_path is not None, (
+        "Cannot find infiniop.dll or libinfiniop.so. Check if INFINI_ROOT is set correctly."
+    )
+    assert librt_path is not None, (
+        "Cannot find infinirt.dll or libinfinirt.so. Check if INFINI_ROOT is set correctly."
+    )
 
     librt = ctypes.CDLL(librt_path)
     libop = ctypes.CDLL(libop_path)
@@ -65,13 +66,12 @@ def find_library_in_ld_path(subdir, library_name):
     lib.infiniopCreateTensorDescriptor.restype = c_int
     lib.infiniopDestroyTensorDescriptor.argtypes = [infiniopTensorDescriptor_t]
     lib.infiniopDestroyTensorDescriptor.restype = c_int
+    lib.infiniopSetRuntimeDevice.argtypes = [c_int, c_int]
+    lib.infiniopSetRuntimeDevice.restype = c_int
     lib.infiniopCreateHandle.argtypes = [POINTER(infiniopHandle_t)]
     lib.infiniopCreateHandle.restype = c_int
     lib.infiniopDestroyHandle.argtypes = [infiniopHandle_t]
     lib.infiniopDestroyHandle.restype = c_int
-    lib.infinirtSetDevice.argtypes = [c_int, c_int]
-    lib.infinirtSetDevice.restype = c_int
-
     OpRegister.register_lib(lib)
 
     return lib
diff --git a/test/infiniop/libinfiniop/utils.py b/test/infiniop/libinfiniop/utils.py
index c0560dbde..8b0b1c173 100644
--- a/test/infiniop/libinfiniop/utils.py
+++ b/test/infiniop/libinfiniop/utils.py
@@ -1,10 +1,12 @@
-from typing import Sequence
-import torch
 import ctypes
+from typing import Sequence
+
 import numpy as np
+import torch
+
 from .datatypes import *
 from .devices import *
-from .liboperators import infiniopTensorDescriptor_t, LIBINFINIOP, infiniopHandle_t
+from .liboperators import LIBINFINIOP, infiniopHandle_t, infiniopTensorDescriptor_t
 
 
 def check_error(status):
@@ -602,9 +604,10 @@ def print_discrepancy(
     if actual.shape != expected.shape:
         raise ValueError("Tensors must have the same shape to compare.")
 
-    import torch
     import sys
 
+    import torch
+
     is_terminal = sys.stdout.isatty()
 
     actual = actual.to("cpu")
@@ -730,7 +733,7 @@ def test_operator(device, test_func, test_cases, tensor_dtypes):
         to be passed to `test_func`.
     - tensor_dtypes (list): A list of tensor data types (e.g., `torch.float32`) to test.
     """
-    LIBINFINIOP.infinirtSetDevice(device, ctypes.c_int(0))
+    check_error(LIBINFINIOP.infiniopSetRuntimeDevice(device, ctypes.c_int(0)))
     handle = create_handle()
     tensor_dtypes = filter_tensor_dtypes_by_device(device, tensor_dtypes)
     try:
@@ -768,12 +771,9 @@ def get_test_devices(args):
     if args.qy:
         devices_to_test.append(InfiniDeviceEnum.QY)
     if args.cambricon:
-        import torch_mlu
-
         devices_to_test.append(InfiniDeviceEnum.CAMBRICON)
     if args.ascend:
         import torch
-        import torch_npu
 
         torch.npu.set_device(0)  # Ascend NPU needs explicit device initialization
         devices_to_test.append(InfiniDeviceEnum.ASCEND)
@@ -783,12 +783,9 @@ def get_test_devices(args):
         devices_to_test.append(InfiniDeviceEnum.METAX)
     if args.moore:
         import torch
-        import torch_musa
 
         devices_to_test.append(InfiniDeviceEnum.MOORE)
     if args.kunlun:
-        import torch_xmlir
-
         devices_to_test.append(InfiniDeviceEnum.KUNLUN)
     if args.hygon:
         import torch
diff --git a/xmake.lua b/xmake.lua
index fcbcf423d..960ad5dcd 100644
--- a/xmake.lua
+++ b/xmake.lua
@@ -292,10 +292,10 @@ option_end()
 option("infinirt-root")
     set_default("")
     set_showmenu(true)
-    set_description("Path to an installed standalone InfiniRT prefix used by --infiniops")
+    set_description("Path to an installed InfiniRT prefix")
 option_end()
 
-local function get_standalone_infinirt_root()
+local function get_infinirt_root()
     local infinirt_root = get_config("infinirt-root")
     if not infinirt_root or infinirt_root == "" then
         infinirt_root = os.getenv("INFINI_RT_ROOT")
@@ -327,6 +327,60 @@ end
 
 local infiniops_external_built = false
 
+local function find_infinirt_library(infinirt_root, xmake_os)
+    local infinirt_library = path.join(infinirt_root, "lib", "libinfinirt.so")
+    if not xmake_os.isfile(infinirt_library) then
+        infinirt_library = path.join(infinirt_root, "lib64", "libinfinirt.so")
+    end
+    if not xmake_os.isfile(infinirt_library) then
+        raise("InfiniRT library not found under: " .. infinirt_root)
+    end
+    return infinirt_library
+end
+
+local function add_external_infinirt()
+    local INFINI_ROOT = os.getenv("INFINI_ROOT") or (os.getenv(is_host("windows") and "HOMEPATH" or "HOME") .. "/.infini")
+    local infinirt_root = get_infinirt_root()
+    if infinirt_root and infinirt_root ~= "" then
+        add_includedirs(path.join(infinirt_root, "include"), { public = true })
+        add_linkdirs(path.join(infinirt_root, "lib"), path.join(infinirt_root, "lib64"))
+        add_rpathdirs(path.join(infinirt_root, "lib"), path.join(infinirt_root, "lib64"))
+    else
+        add_includedirs(INFINI_ROOT .. "/include", { public = true })
+        add_linkdirs(INFINI_ROOT .. "/lib")
+        add_rpathdirs(INFINI_ROOT .. "/lib")
+    end
+    if has_config("ascend-npu") then
+        local ascend_home = os.getenv("ASCEND_HOME") or os.getenv("ASCEND_TOOLKIT_HOME") or os.getenv("ASCEND_HOME_PATH")
+        if ascend_home and ascend_home ~= "" then
+            for _, include_dir in ipairs({
+                path.join(ascend_home, "include"),
+                path.join(ascend_home, "include/aclnn"),
+                path.join(ascend_home, "aarch64-linux/include"),
+                path.join(ascend_home, "aarch64-linux/include/aclnn"),
+            }) do
+                if os.isdir(include_dir) then
+                    add_includedirs(include_dir, { public = true })
+                end
+            end
+        end
+    end
+    add_links("infinirt")
+end
+
+local function patch_infiniops_private_infinirt(xmake_os, infiniops_lib, infinirt_library, private_infinirt)
+    local private_soname = path.filename(private_infinirt)
+    xmake_os.cp(infinirt_library, private_infinirt)
+    xmake_os.execv("patchelf", {"--set-soname", private_soname, private_infinirt})
+    xmake_os.execv("patchelf", {"--replace-needed", infinirt_library, private_soname, infiniops_lib})
+    xmake_os.execv("patchelf", {"--replace-needed", "libinfinirt.so", private_soname, infiniops_lib})
+
+    local needed = xmake_os.iorunv("patchelf", {"--print-needed", infiniops_lib})
+    if not needed:find(private_soname, 1, true) then
+        xmake_os.execv("patchelf", {"--add-needed", private_soname, infiniops_lib})
+    end
+end
+
 local function filter_infiniops_ops_for_backend(infiniops_ops)
     if not infiniops_ops or #infiniops_ops == 0 then
         return infiniops_ops
@@ -376,7 +430,7 @@ local function build_infiniops_external(xmake_os)
     local infiniops_root = path.absolute(get_config("infiniops-root") or "submodules/InfiniOps", os.projectdir())
     local infiniops_builddir = path.join(infiniops_root, "build")
     local INFINI_ROOT = os.getenv("INFINI_ROOT") or (os.getenv(is_host("windows") and "HOMEPATH" or "HOME") .. "/.infini")
-    local infinirt_root = get_standalone_infinirt_root()
+    local infinirt_root = get_infinirt_root()
     local cmake_config_args = {
         "-S", infiniops_root,
         "-B", infiniops_builddir,
@@ -408,20 +462,10 @@ local function build_infiniops_external(xmake_os)
     xmake_os.execv("cmake", {"--build", infiniops_builddir, "--target", "infiniops"})
     xmake_os.execv("cmake", {"--install", infiniops_builddir, "--prefix", INFINI_ROOT})
     if infinirt_root and infinirt_root ~= "" then
-        local standalone_infinirt = path.join(infinirt_root, "lib", "libinfinirt.so")
-        if not xmake_os.isfile(standalone_infinirt) then
-            standalone_infinirt = path.join(infinirt_root, "lib64", "libinfinirt.so")
-        end
-        if not xmake_os.isfile(standalone_infinirt) then
-            raise("Standalone InfiniRT library not found under: " .. infinirt_root)
-        end
+        local infinirt_library = find_infinirt_library(infinirt_root, xmake_os)
         local infiniops_lib = path.join(INFINI_ROOT, "lib", "libinfiniops.so")
-        local private_soname = "libinfiniops_infinirt.so"
-        local private_infinirt = path.join(INFINI_ROOT, "lib", private_soname)
-        xmake_os.cp(standalone_infinirt, private_infinirt)
-        xmake_os.execv("patchelf", {"--set-soname", private_soname, private_infinirt})
-        xmake_os.execv("patchelf", {"--replace-needed", standalone_infinirt, private_soname, infiniops_lib})
-        xmake_os.execv("patchelf", {"--replace-needed", "libinfinirt.so", private_soname, infiniops_lib})
+        local private_infinirt = path.join(INFINI_ROOT, "lib", "libinfiniops_infinirt.so")
+        patch_infiniops_private_infinirt(xmake_os, infiniops_lib, infinirt_library, private_infinirt)
     end
     infiniops_external_built = true
 end
@@ -461,57 +505,9 @@ target("infini-utils")
     add_files("src/utils/*.cc")
 target_end()
 
-target("infinirt")
-    set_kind("shared")
-
-    if has_config("cpu") then
-        add_deps("infinirt-cpu")
-    end
-    if has_config("nv-gpu") then
-        add_deps("infinirt-nvidia")
-    end
-    if has_config("cambricon-mlu") then
-        add_deps("infinirt-cambricon")
-    end
-    if has_config("ascend-npu") then
-        add_deps("infinirt-ascend")
-    end
-    if has_config("metax-gpu") then
-        add_deps("infinirt-metax")
-    end
-    if has_config("moore-gpu") then
-        add_deps("infinirt-moore")
-    end
-    if has_config("iluvatar-gpu") then
-        add_deps("infinirt-iluvatar")
-    end
-    if has_config("ali-ppu") then
-        add_deps("infinirt-ali")
-    end
-    if has_config("qy-gpu") then
-        add_deps("infinirt-qy")
-        add_files("build/.objs/infinirt-qy/rules/qy.cuda/src/infinirt/cuda/*.cu.o", {public = true})
-    end
-    if has_config("kunlun-xpu") then
-        add_deps("infinirt-kunlun")
-    end
-    if has_config("hygon-dcu") then
-        add_deps("infinirt-hygon")
-    end
-    set_languages("cxx17")
-    if not is_plat("windows") then
-        add_cxflags("-fPIC")
-        add_cxxflags("-fPIC")
-        add_ldflags("-fPIC", {force = true})
-    end
-    set_installdir(os.getenv("INFINI_ROOT") or (os.getenv(is_host("windows") and "HOMEPATH" or "HOME") .. "/.infini"))
-    add_files("src/infinirt/*.cc")
-    add_installfiles("include/infinirt.h", {prefixdir = "include"})
-target_end()
-
 target("infiniop")
     set_kind("shared")
-    add_deps("infinirt")
+    add_external_infinirt()
 
     local public_cuda_root = get_config("cuda") or os.getenv("CUDA_HOME") or os.getenv("CUDA_PATH")
     if public_cuda_root and public_cuda_root ~= "" then
@@ -585,11 +581,11 @@ target("infiniop")
     add_installfiles("include/infiniop/*.h", {prefixdir = "include/infiniop"})
     add_installfiles("include/infiniop.h", {prefixdir = "include"})
     add_installfiles("include/infinicore.h", {prefixdir = "include"})
+    add_installfiles("include/infinirt.h", {prefixdir = "include"})
 target_end()
 
 target("infiniccl")
     set_kind("shared")
-    add_deps("infinirt")
 
     if has_config("nv-gpu") then
         add_deps("infiniccl-nvidia")
@@ -635,7 +631,7 @@ target_end()
 
 target("infinicore_c_api")
     set_kind("phony")
-    add_deps("infiniop", "infinirt", "infiniccl")
+    add_deps("infiniop", "infiniccl")
     after_build(function (target) print(YELLOW .. "[Congratulations!] Now you can install the libraries with \"xmake install\"" .. NC) end)
 target_end()
 
@@ -650,7 +646,7 @@ target_end()
 
 target("infinicore_cpp_api")
     set_kind("shared")
-    add_deps("infiniop", "infinirt", "infiniccl")
+    add_deps("infiniop", "infiniccl")
     set_languages("cxx17")
     set_symbols("visibility")
 
@@ -669,6 +665,7 @@ target("infinicore_cpp_api")
         add_defines("CHAR_BIT=8", "INT_MIN=(-2147483647 - 1)", "INT_MAX=2147483647", "UINT_MAX=4294967295U")
     end
     add_includedirs(INFINI_ROOT.."/include", { public = true })
+    add_external_infinirt()
     if has_config("nv-gpu") then
         local cuda_root = os.getenv("CUDA_HOME") or os.getenv("CUDA_PATH") or get_config("cuda") or "/usr/local/cuda"
         add_includedirs(cuda_root .. "/include")
@@ -679,7 +676,7 @@ target("infinicore_cpp_api")
             raise("InfiniOps root not found: " .. infiniops_root)
         end
         get_infiniops_backend_cmake_arg()
-        local infinirt_root = get_standalone_infinirt_root()
+        local infinirt_root = get_infinirt_root()
         if infinirt_root and infinirt_root ~= "" then
             add_includedirs(infinirt_root .. "/include")
             add_linkdirs(infinirt_root .. "/lib", infinirt_root .. "/lib64")
@@ -694,27 +691,18 @@ target("infinicore_cpp_api")
         end)
         after_install(function (target)
             local INFINI_ROOT = os.getenv("INFINI_ROOT") or (os.getenv(is_host("windows") and "HOMEPATH" or "HOME") .. "/.infini")
-            local infinirt_root = get_standalone_infinirt_root()
+            local infinirt_root = get_infinirt_root()
             if infinirt_root and infinirt_root ~= "" then
-                local standalone_infinirt = path.join(infinirt_root, "lib", "libinfinirt.so")
-                if not os.isfile(standalone_infinirt) then
-                    standalone_infinirt = path.join(infinirt_root, "lib64", "libinfinirt.so")
-                end
-                if not os.isfile(standalone_infinirt) then
-                    raise("Standalone InfiniRT library not found under: " .. infinirt_root)
-                end
+                local infinirt_library = find_infinirt_library(infinirt_root, os)
                 local infiniops_lib = path.join(INFINI_ROOT, "lib", "libinfiniops.so")
-                local private_soname = "libinfiniops_infinirt.so"
-                local private_infinirt = path.join(INFINI_ROOT, "lib", private_soname)
-                os.cp(standalone_infinirt, private_infinirt)
-                os.execv("patchelf", {"--set-soname", private_soname, private_infinirt})
-                os.execv("patchelf", {"--replace-needed", "libinfinirt.so", private_soname, infiniops_lib})
+                local private_infinirt = path.join(INFINI_ROOT, "lib", "libinfiniops_infinirt.so")
+                patch_infiniops_private_infinirt(os, infiniops_lib, infinirt_library, private_infinirt)
             end
         end)
     end
 
     add_linkdirs(INFINI_ROOT.."/lib")
-    add_links("infiniop", "infinirt", "infiniccl")
+    add_links("infiniop", "infiniccl")
 
     if get_config("flash-attn") and get_config("flash-attn") ~= "" then
         add_installfiles("(builddir)/$(plat)/$(arch)/$(mode)/flash-attn*.so", {prefixdir = "lib"})
@@ -943,14 +931,15 @@ target("_infinicore")
     set_kind("shared")
     local INFINI_ROOT = os.getenv("INFINI_ROOT") or (os.getenv(is_host("windows") and "HOMEPATH" or "HOME") .. "/.infini")
     add_includedirs(INFINI_ROOT.."/include", { public = true })
+    add_external_infinirt()
 
     add_linkdirs(INFINI_ROOT.."/lib")
-    add_links("infiniop", "infinirt", "infiniccl")
+    add_links("infiniop", "infiniccl")
 
     add_files("src/infinicore/pybind11/**.cc")
 
     if has_config("infiniops") then
-        local infinirt_root = get_standalone_infinirt_root()
+        local infinirt_root = get_infinirt_root()
         if infinirt_root and infinirt_root ~= "" then
             add_includedirs(infinirt_root .. "/include")
             add_linkdirs(infinirt_root .. "/lib", infinirt_root .. "/lib64")
@@ -965,21 +954,11 @@ target("_infinicore")
                 infiniops_lib = path.join(infiniops_root, "build", "src", "libinfiniops.so")
                 infiniops_lib_installed = false
             end
-            local infinirt_root = get_standalone_infinirt_root()
+            local infinirt_root = get_infinirt_root()
             if infinirt_root and infinirt_root ~= "" then
-                local standalone_infinirt = path.join(infinirt_root, "lib", "libinfinirt.so")
-                if not os.isfile(standalone_infinirt) then
-                    standalone_infinirt = path.join(infinirt_root, "lib64", "libinfinirt.so")
-                end
-                if not os.isfile(standalone_infinirt) then
-                    raise("Standalone InfiniRT library not found under: " .. infinirt_root)
-                end
-                local private_soname = "libinfiniops_infinirt.so"
-                local private_infinirt = path.join(INFINI_ROOT, "lib", private_soname)
-                os.cp(standalone_infinirt, private_infinirt)
-                os.execv("patchelf", {"--set-soname", private_soname, private_infinirt})
-                os.execv("patchelf", {"--replace-needed", standalone_infinirt, private_soname, infiniops_lib})
-                os.execv("patchelf", {"--replace-needed", "libinfinirt.so", private_soname, infiniops_lib})
+                local infinirt_library = find_infinirt_library(infinirt_root, os)
+                local private_infinirt = path.join(INFINI_ROOT, "lib", "libinfiniops_infinirt.so")
+                patch_infiniops_private_infinirt(os, infiniops_lib, infinirt_library, private_infinirt)
             end
             os.mkdir(path.join(INFINI_ROOT, "lib"))
             if not infiniops_lib_installed then
diff --git a/xmake/ali.lua b/xmake/ali.lua
index 940650d67..550f7e0fc 100644
--- a/xmake/ali.lua
+++ b/xmake/ali.lua
@@ -104,7 +104,6 @@ target_end()
 
 target("infiniccl-ali")
     set_kind("static")
-    add_deps("infinirt")
     on_install(function (target) end)
     if has_config("ccl") then
         set_policy("build.cuda.devlink", true)
diff --git a/xmake/ascend.lua b/xmake/ascend.lua
index 6179cac13..f8b2fe9ac 100644
--- a/xmake/ascend.lua
+++ b/xmake/ascend.lua
@@ -1,16 +1,38 @@
 add_defines("ENABLE_ASCEND_API")
-local ASCEND_HOME = os.getenv("ASCEND_HOME") or os.getenv("ASCEND_TOOLKIT_HOME")
+local ASCEND_HOME = os.getenv("ASCEND_HOME") or os.getenv("ASCEND_TOOLKIT_HOME") or os.getenv("ASCEND_HOME_PATH")
 local SOC_VERSION = os.getenv("SOC_VERSION")
 
 -- Add include dirs
-add_includedirs(ASCEND_HOME .. "/include")
-add_includedirs(ASCEND_HOME .. "/include/aclnn")
-add_linkdirs(ASCEND_HOME .. "/lib64")
+for _, include_dir in ipairs({
+    path.join(ASCEND_HOME, "include"),
+    path.join(ASCEND_HOME, "include/aclnn"),
+    path.join(ASCEND_HOME, "aarch64-linux/include"),
+    path.join(ASCEND_HOME, "aarch64-linux/include/aclnn"),
+}) do
+    if os.isdir(include_dir) then
+        add_includedirs(include_dir)
+    end
+end
+for _, lib_dir in ipairs({
+    path.join(ASCEND_HOME, "lib64"),
+    path.join(ASCEND_HOME, "aarch64-linux/lib64"),
+}) do
+    if os.isdir(lib_dir) then
+        add_linkdirs(lib_dir)
+    end
+end
 add_links("libascendcl.so")
 add_links("libnnopbase.so")
 add_links("libopapi.so")
 add_links("libruntime.so")
-add_linkdirs(ASCEND_HOME .. "/../../driver/lib64/driver")
+for _, driver_dir in ipairs({
+    path.join(ASCEND_HOME, "../driver/lib64/driver"),
+    path.join(ASCEND_HOME, "../../driver/lib64/driver"),
+}) do
+    if os.isdir(driver_dir) then
+        add_linkdirs(driver_dir)
+    end
+end
 add_links("libascend_hal.so")
 local builddir = string.format(
         "%s/build/%s/%s/%s",
@@ -69,7 +91,6 @@ target_end()
 
 target("infiniccl-ascend")
     set_kind("static")
-    add_deps("infinirt")
     add_deps("infini-utils")
     set_warnings("all", "error")
     set_languages("cxx17")
diff --git a/xmake/bang.lua b/xmake/bang.lua
index ffa85ef6d..0be267c94 100644
--- a/xmake/bang.lua
+++ b/xmake/bang.lua
@@ -65,7 +65,6 @@ target_end()
 
 target("infiniccl-cambricon")
     set_kind("static")
-    add_deps("infinirt")
     add_deps("infini-utils")
     set_warnings("all", "error")
     set_languages("cxx17")
diff --git a/xmake/cpu.lua b/xmake/cpu.lua
index e192fbbbd..a1974b823 100644
--- a/xmake/cpu.lua
+++ b/xmake/cpu.lua
@@ -26,6 +26,7 @@ target_end()
 
 target("infinirt-cpu")
     set_kind("static")
+    set_default(false)
     add_deps("infini-utils")
     on_install(function (target) end)
 
diff --git a/xmake/hygon.lua b/xmake/hygon.lua
index cf4e970f8..e5c86b863 100644
--- a/xmake/hygon.lua
+++ b/xmake/hygon.lua
@@ -101,7 +101,12 @@ local function resolve_hygon_arch()
 end
 
 local HYGON_ARCH = resolve_hygon_arch()
+local HYGON_SKIP_PER_CHANNEL_QUANT_INT8 = HYGON_ARCH == "gfx906"
 print("编译海光DCU架构: " .. HYGON_ARCH)
+if HYGON_SKIP_PER_CHANNEL_QUANT_INT8 then
+    print("Hygon gfx906: skip per_channel_quant_int8 operator")
+    add_defines("INFINIOP_DISABLE_HYGON_PER_CHANNEL_QUANT_INT8")
+end
 
 target("infiniop-hygon")
     set_kind("static")
@@ -119,7 +124,7 @@ target("infiniop-hygon")
     add_hygon_dtk_paths()
 
     set_warnings("all", "error")
-    add_cuflags("-Wno-error=unused-private-field")
+    add_cuflags("-Wno-error=unused-private-field", {force = true})
     add_cuflags("-Wno-return-type", {force = true})  -- 抑制return语句警告
     add_cuflags("-Wno-error=macro-redefined", {force = true})
     add_cuflags("-Wno-error=ignored-attributes", {force = true})
@@ -140,6 +145,9 @@ target("infiniop-hygon")
 
     -- 复用NVIDIA的CUDA实现，通过HIP兼容层
     add_files("../src/infiniop/devices/nvidia/*.cu", "../src/infiniop/ops/*/nvidia/*.cu")
+    if not HYGON_SKIP_PER_CHANNEL_QUANT_INT8 then
+        add_files("../src/infiniop/ops/quant/per_channel_quant_int8/nvidia/*.cu")
+    end
 
     -- Keep platform-specific or currently unregistered NVIDIA sources out of the Hygon target.
     remove_files("../src/infiniop/ops/avg_pool3d/nvidia/*.cu")
@@ -148,8 +156,7 @@ target("infiniop-hygon")
     remove_files("../src/infiniop/ops/dist/nvidia/*.cu")
     remove_files("../src/infiniop/ops/gptq_qyblas_gemm/nvidia/*.cu")
     remove_files("../src/infiniop/ops/histc/nvidia/*.cu")
-    remove_files("../src/infiniop/ops/quant*/nvidia/*.cu")
-    remove_files("../src/infiniop/ops/scaled_mm/nvidia/*.cu")
+    remove_files("../src/infiniop/ops/quant/per_tensor_quant_int8/nvidia/*.cu")
 
     if has_config("ninetoothed") then
         add_files("../build/ninetoothed/*.c", "../build/ninetoothed/*.cpp", {cxxflags = {"-Wno-return-type"}})
@@ -190,7 +197,6 @@ target_end()
 
 target("infiniccl-hygon")
     set_kind("static")
-    add_deps("infinirt")
     on_install(function (target) end)
 
     if has_config("ccl") then
diff --git a/xmake/iluvatar.lua b/xmake/iluvatar.lua
index 34a913a11..5f69fd56e 100644
--- a/xmake/iluvatar.lua
+++ b/xmake/iluvatar.lua
@@ -88,7 +88,6 @@ target_end()
 
 target("infiniccl-iluvatar")
     set_kind("static")
-    add_deps("infinirt")
     on_install(function (target) end)
 
     if has_config("ccl") then
diff --git a/xmake/kunlun.lua b/xmake/kunlun.lua
index 84ba14082..4f3a40db2 100644
--- a/xmake/kunlun.lua
+++ b/xmake/kunlun.lua
@@ -108,7 +108,6 @@ target_end()
 
 target("infiniccl-kunlun")
     set_kind("static")
-    add_deps("infinirt")
     add_deps("infini-utils")
     set_warnings("all", "error")
     set_languages("cxx17")
diff --git a/xmake/metax.lua b/xmake/metax.lua
index 85407ed1b..3e700fff2 100644
--- a/xmake/metax.lua
+++ b/xmake/metax.lua
@@ -158,7 +158,6 @@ target_end()
 
 target("infiniccl-metax")
     set_kind("static")
-    add_deps("infinirt")
     on_install(function (target) end)
     set_warnings("all", "error")
     if not is_plat("windows") then
diff --git a/xmake/moore.lua b/xmake/moore.lua
index d34dd24f7..6e8800be8 100644
--- a/xmake/moore.lua
+++ b/xmake/moore.lua
@@ -81,7 +81,6 @@ target_end()
 
 target("infiniccl-moore")
     set_kind("static")
-    add_deps("infinirt")
     on_install(function (target) end)
     set_warnings("all", "error")
     if not is_plat("windows") then
diff --git a/xmake/nvidia.lua b/xmake/nvidia.lua
index 282f5876f..aca0dad66 100644
--- a/xmake/nvidia.lua
+++ b/xmake/nvidia.lua
@@ -246,7 +246,6 @@ target_end()
 
 target("infiniccl-nvidia")
     set_kind("static")
-    add_deps("infinirt")
     on_install(function (target) end)
     if has_config("ccl") then
         set_policy("build.cuda.devlink", true)
diff --git a/xmake/qy.lua b/xmake/qy.lua
index 35d53d6ef..82cf92d6b 100644
--- a/xmake/qy.lua
+++ b/xmake/qy.lua
@@ -174,7 +174,6 @@ target_end()
 
 target("infiniccl-qy")
     set_kind("static")
-    add_deps("infinirt")
     on_install(function (target) end)
     if has_config("ccl") then
         add_rules("qy.cuda", {override = true})
diff --git a/xmake/test.lua b/xmake/test.lua
index 64533163b..d4ef4abe7 100644
--- a/xmake/test.lua
+++ b/xmake/test.lua
@@ -54,14 +54,20 @@ target_end()
 
 target("infinirt-test")
     set_kind("binary")
-    add_deps("infinirt")
+    set_default(false)
     on_install(function (target) end)
 
     set_languages("cxx17")
     set_warnings("all", "error")
 
+    local INFINI_ROOT = os.getenv("INFINI_ROOT") or (os.getenv(is_host("windows") and "HOMEPATH" or "HOME") .. "/.infini")
+    add_includedirs(INFINI_ROOT.."/include")
+    add_linkdirs(INFINI_ROOT.."/lib")
+    add_links("infinirt")
+    add_rpathdirs(INFINI_ROOT.."/lib")
+
     add_files(os.projectdir().."/src/infinirt-test/*.cc")
-    set_installdir(os.getenv("INFINI_ROOT") or (os.getenv(is_host("windows") and "HOMEPATH" or "HOME") .. "/.infini"))
+    set_installdir(INFINI_ROOT)
 target_end()
 
 target("infinicore-test")