diff --git a/CMakeLists.txt b/CMakeLists.txt
index 65e63ec..d47705a 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -33,6 +33,7 @@ option(WITH_CPU "Enable CPU backend" OFF)
 option(WITH_NVIDIA "Enable CUDA backend" OFF)
 option(WITH_ILUVATAR "Enable Iluvatar GPU backend" OFF)
 option(WITH_HYGON "Enable Hygon GPU backend" OFF)
+option(WITH_ALI "Enable Ali PPU backend" OFF)
 option(WITH_METAX "Enable MetaX backend" OFF)
 option(WITH_CAMBRICON "Enable Cambricon backend" OFF)
 option(WITH_MOORE "Enable Moore backend" OFF)
@@ -118,7 +119,7 @@ if(AUTO_DETECT_DEVICES)
 
     if(WITH_NVIDIA)
         set(_non_nvidia_gpu_detected FALSE)
-        foreach(_gpu_backend WITH_ILUVATAR WITH_HYGON WITH_METAX WITH_CAMBRICON WITH_MOORE WITH_ASCEND)
+        foreach(_gpu_backend WITH_ILUVATAR WITH_HYGON WITH_ALI WITH_METAX WITH_CAMBRICON WITH_MOORE WITH_ASCEND)
             if(${_gpu_backend})
                 set(_non_nvidia_gpu_detected TRUE)
             endif()
@@ -135,14 +136,14 @@ include_directories(${CMAKE_CURRENT_SOURCE_DIR}/src)
 
 # Only one CUDA-like GPU backend can be enabled at a time.
 set(_gpu_backend_count 0)
-foreach(_gpu_backend WITH_NVIDIA WITH_ILUVATAR WITH_HYGON WITH_METAX WITH_MOORE WITH_ASCEND)
+foreach(_gpu_backend WITH_NVIDIA WITH_ILUVATAR WITH_HYGON WITH_ALI WITH_METAX WITH_MOORE WITH_ASCEND)
     if(${_gpu_backend})
         math(EXPR _gpu_backend_count "${_gpu_backend_count} + 1")
     endif()
 endforeach()
 
 if(_gpu_backend_count GREATER 1)
-    message(FATAL_ERROR "`WITH_NVIDIA`, `WITH_ILUVATAR`, `WITH_HYGON`, `WITH_METAX`, `WITH_MOORE`, and `WITH_ASCEND` are mutually exclusive. Build one GPU backend at a time.")
+    message(FATAL_ERROR "`WITH_NVIDIA`, `WITH_ILUVATAR`, `WITH_HYGON`, `WITH_ALI`, `WITH_METAX`, `WITH_MOORE`, and `WITH_ASCEND` are mutually exclusive. Build one GPU backend at a time.")
 endif()
 
 if(WITH_NVIDIA)
@@ -217,6 +218,44 @@ if(WITH_HYGON)
     find_package(CUDAToolkit REQUIRED)
 endif()
 
+if(WITH_ALI)
+    add_compile_definitions(WITH_ALI=1)
+
+    set(ALI_CUDA_ROOT "")
+    foreach(_ali_cuda_env ALI_CUDA_ROOT PPU_CUDA_ROOT)
+        if(NOT ALI_CUDA_ROOT AND DEFINED ENV{${_ali_cuda_env}} AND
+           NOT "$ENV{${_ali_cuda_env}}" STREQUAL "")
+            set(ALI_CUDA_ROOT "$ENV{${_ali_cuda_env}}")
+        endif()
+    endforeach()
+    if(NOT ALI_CUDA_ROOT AND DEFINED ENV{PPU_SDK_ROOT} AND
+       NOT "$ENV{PPU_SDK_ROOT}" STREQUAL "")
+        set(ALI_CUDA_ROOT "$ENV{PPU_SDK_ROOT}/CUDA_SDK")
+    endif()
+    if(NOT ALI_CUDA_ROOT AND EXISTS "/usr/local/PPU_SDK/CUDA_SDK")
+        set(ALI_CUDA_ROOT "/usr/local/PPU_SDK/CUDA_SDK")
+    endif()
+
+    if(NOT ALI_CUDA_ROOT OR NOT EXISTS "${ALI_CUDA_ROOT}/bin/nvcc")
+        message(FATAL_ERROR
+            "`WITH_ALI` is `ON` but the Ali PPU CUDA SDK was not found. "
+            "Set `ALI_CUDA_ROOT`, `PPU_CUDA_ROOT`, or `PPU_SDK_ROOT`.")
+    endif()
+
+    set(CMAKE_CUDA_COMPILER "${ALI_CUDA_ROOT}/bin/nvcc"
+        CACHE FILEPATH "Ali PPU CUDA compiler" FORCE)
+    set(CUDAToolkit_ROOT "${ALI_CUDA_ROOT}"
+        CACHE PATH "Ali PPU CUDA toolkit root" FORCE)
+    if(NOT CMAKE_CUDA_ARCHITECTURES)
+        set(CMAKE_CUDA_ARCHITECTURES 80 CACHE STRING
+            "Ali PPU CUDA architecture" FORCE)
+    endif()
+
+    message(STATUS "Ali PPU: CUDA toolkit ${ALI_CUDA_ROOT}")
+    enable_language(CUDA)
+    find_package(CUDAToolkit REQUIRED)
+endif()
+
 if(WITH_METAX)
     add_compile_definitions(WITH_METAX=1)
 
@@ -294,7 +333,7 @@ if(WITH_ASCEND)
 endif()
 
 # If all other platforms are not enabled, CPU is enabled by default.
-if(NOT WITH_NVIDIA AND NOT WITH_ILUVATAR AND NOT WITH_HYGON AND NOT WITH_METAX AND NOT WITH_MOORE AND NOT WITH_CAMBRICON AND NOT WITH_ASCEND)
+if(NOT WITH_NVIDIA AND NOT WITH_ILUVATAR AND NOT WITH_HYGON AND NOT WITH_ALI AND NOT WITH_METAX AND NOT WITH_MOORE AND NOT WITH_CAMBRICON AND NOT WITH_ASCEND)
     set(WITH_CPU ON)
     add_compile_definitions(WITH_CPU=1)
 endif()
@@ -312,6 +351,9 @@ endif()
 if(WITH_HYGON)
     list(APPEND INFINI_RT_PUBLIC_HEADER_DEVICES hygon)
 endif()
+if(WITH_ALI)
+    list(APPEND INFINI_RT_PUBLIC_HEADER_DEVICES ali)
+endif()
 if(WITH_METAX)
     list(APPEND INFINI_RT_PUBLIC_HEADER_DEVICES metax)
 endif()
@@ -342,7 +384,7 @@ endif()
 add_subdirectory(src)
 
 set(INFINI_RT_PACKAGE_NEEDS_CUDATOOLKIT OFF)
-if(WITH_NVIDIA OR WITH_ILUVATAR OR WITH_HYGON)
+if(WITH_NVIDIA OR WITH_ILUVATAR OR WITH_HYGON OR WITH_ALI)
     set(INFINI_RT_PACKAGE_NEEDS_CUDATOOLKIT ON)
 endif()
 
diff --git a/cmake/InfiniRTConfig.cmake.in b/cmake/InfiniRTConfig.cmake.in
index 09bce5a..f27581f 100644
--- a/cmake/InfiniRTConfig.cmake.in
+++ b/cmake/InfiniRTConfig.cmake.in
@@ -9,6 +9,7 @@ set(InfiniRT_WITH_CPU @WITH_CPU@)
 set(InfiniRT_WITH_NVIDIA @WITH_NVIDIA@)
 set(InfiniRT_WITH_ILUVATAR @WITH_ILUVATAR@)
 set(InfiniRT_WITH_HYGON @WITH_HYGON@)
+set(InfiniRT_WITH_ALI @WITH_ALI@)
 set(InfiniRT_WITH_METAX @WITH_METAX@)
 set(InfiniRT_WITH_MOORE @WITH_MOORE@)
 set(InfiniRT_WITH_CAMBRICON @WITH_CAMBRICON@)
diff --git a/scripts/generate_public_headers.py b/scripts/generate_public_headers.py
index 3363741..7c3b5c5 100644
--- a/scripts/generate_public_headers.py
+++ b/scripts/generate_public_headers.py
@@ -27,6 +27,11 @@
         ("hygon", "device_.h", "native/cuda/hygon/device_.h"),
         ("hygon", "runtime_.h", "native/cuda/hygon/runtime_.h"),
     ),
+    "ali": (
+        ("ali", "data_type_.h", "native/cuda/ali/data_type_.h"),
+        ("ali", "device_.h", "native/cuda/ali/device_.h"),
+        ("ali", "runtime_.h", "native/cuda/ali/runtime_.h"),
+    ),
     "metax": (
         ("metax", "data_type_.h", "native/cuda/metax/data_type_.h"),
         ("metax", "device_.h", "native/cuda/metax/device_.h"),
@@ -54,6 +59,7 @@
     "nvidia": "Device::Type::kNvidia",
     "iluvatar": "Device::Type::kIluvatar",
     "hygon": "Device::Type::kHygon",
+    "ali": "Device::Type::kAli",
     "metax": "Device::Type::kMetax",
     "moore": "Device::Type::kMoore",
     "cambricon": "Device::Type::kCambricon",
@@ -64,6 +70,7 @@
     "nvidia",
     "iluvatar",
     "hygon",
+    "ali",
     "metax",
     "moore",
     "cambricon",
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 2091963..463114e 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -63,6 +63,20 @@ if(WITH_HYGON)
     )
 endif()
 
+if(WITH_ALI)
+    enable_language(CUDA)
+
+    target_compile_definitions(infinirt PUBLIC WITH_ALI=1)
+
+    find_package(CUDAToolkit REQUIRED)
+    target_link_libraries(infinirt PUBLIC CUDA::cudart)
+
+    set_target_properties(infinirt PROPERTIES
+        CUDA_STANDARD 17
+        CUDA_STANDARD_REQUIRED ON
+    )
+endif()
+
 if(WITH_METAX)
     target_compile_definitions(infinirt PRIVATE WITH_METAX=1)
     _infinirt_use_backend_runtime(
diff --git a/src/device.h b/src/device.h
index 482b2fb..cd52361 100644
--- a/src/device.h
+++ b/src/device.h
@@ -21,6 +21,7 @@ class Device {
     kMoore = 5,
     kIluvatar = 6,
     kHygon = 7,
+    kAli = 8,
     kCount
   };
 
@@ -64,6 +65,7 @@ class Device {
           {Type::kMoore, "moore"},
           {Type::kIluvatar, "iluvatar"},
           {Type::kHygon, "hygon"},
+          {Type::kAli, "ali"},
       }}};
 
   static constexpr ConstexprMap<std::string_view, Device::Type,
@@ -77,6 +79,7 @@ class Device {
           {"moore", Type::kMoore},
           {"iluvatar", Type::kIluvatar},
           {"hygon", Type::kHygon},
+          {"ali", Type::kAli},
       }}};
 
   int index_{0};
@@ -91,7 +94,7 @@ struct DeviceEnabled : std::false_type {};
 using AllDeviceTypes =
     List<Device::Type::kCpu, Device::Type::kNvidia, Device::Type::kCambricon,
          Device::Type::kAscend, Device::Type::kMetax, Device::Type::kMoore,
-         Device::Type::kIluvatar, Device::Type::kHygon>;
+         Device::Type::kIluvatar, Device::Type::kHygon, Device::Type::kAli>;
 
 // Deferred computation of active devices. The `Filter` and `FilterList`
 // evaluation are nested inside a class template so that `DeviceEnabled`
diff --git a/src/native/cuda/ali/data_type_.h b/src/native/cuda/ali/data_type_.h
new file mode 100644
index 0000000..5c88b89
--- /dev/null
+++ b/src/native/cuda/ali/data_type_.h
@@ -0,0 +1,30 @@
+#ifndef INFINI_RT_ALI_DATA_TYPE__H_
+#define INFINI_RT_ALI_DATA_TYPE__H_
+
+// clang-format off
+#include <cuda_bf16.h>
+#include <cuda_fp16.h>
+// clang-format on
+
+#include "data_type.h"
+#include "native/cuda/ali/device_.h"
+
+namespace infini::rt {
+
+using cuda_bfloat16 = nv_bfloat16;
+
+using cuda_bfloat162 = nv_bfloat162;
+
+template <>
+struct TypeMap<Device::Type::kAli, DataType::kFloat16> {
+  using type = half;
+};
+
+template <>
+struct TypeMap<Device::Type::kAli, DataType::kBFloat16> {
+  using type = __nv_bfloat16;
+};
+
+}  // namespace infini::rt
+
+#endif
diff --git a/src/native/cuda/ali/device_.h b/src/native/cuda/ali/device_.h
new file mode 100644
index 0000000..6ced8f0
--- /dev/null
+++ b/src/native/cuda/ali/device_.h
@@ -0,0 +1,13 @@
+#ifndef INFINI_RT_ALI_DEVICE__H_
+#define INFINI_RT_ALI_DEVICE__H_
+
+#include "device.h"
+
+namespace infini::rt {
+
+template <>
+struct DeviceEnabled<Device::Type::kAli> : std::true_type {};
+
+}  // namespace infini::rt
+
+#endif
diff --git a/src/native/cuda/ali/runtime_.h b/src/native/cuda/ali/runtime_.h
new file mode 100644
index 0000000..37725e2
--- /dev/null
+++ b/src/native/cuda/ali/runtime_.h
@@ -0,0 +1,175 @@
+#ifndef INFINI_RT_ALI_RUNTIME__H_
+#define INFINI_RT_ALI_RUNTIME__H_
+
+#include <cstddef>
+#include <utility>
+
+// clang-format off
+#include <cuda_runtime.h>
+// clang-format on
+
+#include "native/cuda/ali/device_.h"
+#include "native/cuda/runtime_.h"
+
+namespace infini::rt::runtime {
+
+template <>
+struct Runtime<Device::Type::kAli>
+    : GraphRuntime<Runtime<Device::Type::kAli>,
+                   CudaRuntime<Runtime<Device::Type::kAli>>> {
+  using Error = cudaError_t;
+
+  using Stream = cudaStream_t;
+
+  using Graph = cudaGraph_t;
+
+  using GraphExec = cudaGraphExec_t;
+
+  using Event = cudaEvent_t;
+
+  using StreamCaptureMode = cudaStreamCaptureMode;
+
+  static constexpr Device::Type kDeviceType = Device::Type::kAli;
+
+  static constexpr Error kSuccess = cudaSuccess;
+
+  static constexpr auto SetDevice = cudaSetDevice;
+
+  static constexpr auto GetDevice = cudaGetDevice;
+
+  static constexpr auto GetDeviceCount = cudaGetDeviceCount;
+
+  static constexpr auto DeviceSynchronize = cudaDeviceSynchronize;
+
+  static constexpr auto Malloc = [](auto&&... args) {
+    return cudaMalloc(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto MallocHost = [](auto&&... args) {
+    return cudaMallocHost(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto MallocAsync = [](auto&&... args) {
+    return cudaMallocAsync(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto Free = cudaFree;
+
+  static constexpr auto FreeHost = [](auto&&... args) {
+    return cudaFreeHost(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto FreeAsync = [](auto&&... args) {
+    return cudaFreeAsync(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto MemGetInfo = [](std::size_t* free, std::size_t* total) {
+    const auto status = cudaMemGetInfo(free, total);
+    // On PPU-SMI 1.22, Driver Version: 1.6.3-8ee7e7
+    // cudaMemGetInfo may return cudaSuccess while reporting free memory greater
+    // than total memory. Clamp free to total as a temporary workaround
+    // until the vendor SDK/driver fixes memory accounting.
+    if (status == cudaSuccess && *free > *total) {
+      *free = *total;
+    }
+    return status;
+  };
+
+  static constexpr auto Memcpy = cudaMemcpy;
+
+  static constexpr auto MemcpyAsync = cudaMemcpyAsync;
+
+  static constexpr auto kMemcpyHostToHost = cudaMemcpyHostToHost;
+
+  static constexpr auto kMemcpyHostToDevice = cudaMemcpyHostToDevice;
+
+  static constexpr auto kMemcpyDeviceToHost = cudaMemcpyDeviceToHost;
+
+  static constexpr auto kMemcpyDeviceToDevice = cudaMemcpyDeviceToDevice;
+
+  static constexpr auto Memset = cudaMemset;
+
+  static constexpr auto MemsetAsync = [](auto&&... args) {
+    return cudaMemsetAsync(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto StreamCreate = cudaStreamCreate;
+
+  static constexpr auto StreamDestroy = [](auto&&... args) {
+    return cudaStreamDestroy(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto StreamSynchronize = [](auto&&... args) {
+    return cudaStreamSynchronize(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto StreamWaitEvent = [](auto&&... args) {
+    return cudaStreamWaitEvent(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto EventCreate = [](auto&&... args) {
+    return cudaEventCreate(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto EventCreateWithFlags = [](auto&&... args) {
+    return cudaEventCreateWithFlags(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto EventRecord = [](auto&&... args) {
+    return cudaEventRecord(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto EventQuery = [](auto&&... args) {
+    return cudaEventQuery(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto EventSynchronize = [](auto&&... args) {
+    return cudaEventSynchronize(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto EventDestroy = [](auto&&... args) {
+    return cudaEventDestroy(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto EventElapsedTime = [](auto&&... args) {
+    return cudaEventElapsedTime(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto kStreamCaptureModeGlobal = cudaStreamCaptureModeGlobal;
+
+  static constexpr auto kStreamCaptureModeThreadLocal =
+      cudaStreamCaptureModeThreadLocal;
+
+  static constexpr auto kStreamCaptureModeRelaxed =
+      cudaStreamCaptureModeRelaxed;
+
+  static constexpr auto StreamBeginCapture = [](auto&&... args) {
+    return cudaStreamBeginCapture(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto StreamEndCapture = [](auto&&... args) {
+    return cudaStreamEndCapture(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto GraphDestroy = [](auto&&... args) {
+    return cudaGraphDestroy(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto GraphInstantiate = [](auto&&... args) {
+    return cudaGraphInstantiate(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto GraphExecDestroy = [](auto&&... args) {
+    return cudaGraphExecDestroy(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto GraphLaunch = [](auto&&... args) {
+    return cudaGraphLaunch(std::forward<decltype(args)>(args)...);
+  };
+};
+
+static_assert(Runtime<Device::Type::kAli>::Validate());
+
+}  // namespace infini::rt::runtime
+
+#endif
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index 8ffa4b3..8b05314 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -81,6 +81,16 @@ if(WITH_HYGON)
         HYGON infini::rt::Device::Type::kHygon 1)
 endif()
 
+if(WITH_ALI)
+    set(INFINI_RT_TEST_HAS_RUNTIME_BACKEND ON)
+    add_infini_rt_backend_runtime_test(
+        ALI infini::rt::Device::Type::kAli
+        infini/rt/ali/runtime_.h
+        1 1 1 1 1 1 1 1)
+    add_infini_rt_backend_graph_test(
+        ALI infini::rt::Device::Type::kAli 1)
+endif()
+
 if(WITH_METAX)
     set(INFINI_RT_TEST_HAS_RUNTIME_BACKEND ON)
     add_infini_rt_backend_runtime_test(
@@ -132,6 +142,10 @@ if(INFINI_RT_TEST_HAS_RUNTIME_BACKEND)
         target_compile_definitions(test_runtime_dispatch
                                    PRIVATE INFINI_RT_TEST_WITH_HYGON=1)
     endif()
+    if(WITH_ALI)
+        target_compile_definitions(test_runtime_dispatch
+                                   PRIVATE INFINI_RT_TEST_WITH_ALI=1)
+    endif()
     if(WITH_METAX)
         target_compile_definitions(test_runtime_dispatch
                                    PRIVATE INFINI_RT_TEST_WITH_METAX=1)
@@ -174,6 +188,11 @@ elseif(WITH_HYGON)
     if(NOT INFINI_RT_TEST_CUDATOOLKIT_ROOT)
         set(INFINI_RT_TEST_CUDATOOLKIT_ROOT "${HYGON_CUDA_ROOT}")
     endif()
+elseif(WITH_ALI)
+    set(INFINI_RT_TEST_CONSUMER_BACKEND ALI)
+    if(NOT INFINI_RT_TEST_CUDATOOLKIT_ROOT)
+        set(INFINI_RT_TEST_CUDATOOLKIT_ROOT "${ALI_CUDA_ROOT}")
+    endif()
 elseif(WITH_METAX)
     set(INFINI_RT_TEST_CONSUMER_BACKEND METAX)
     set(INFINI_RT_TEST_BACKEND_ROOT_VARIABLE InfiniRT_METAX_ROOT)
diff --git a/tests/install_consumer/CMakeLists.txt b/tests/install_consumer/CMakeLists.txt
index c1a3c33..90e4229 100644
--- a/tests/install_consumer/CMakeLists.txt
+++ b/tests/install_consumer/CMakeLists.txt
@@ -24,7 +24,7 @@ if(NOT "${InfiniRT_ENABLED_BACKENDS}" STREQUAL
             "InfiniRT reported enabled backends '${InfiniRT_ENABLED_BACKENDS}', expected '${_infinirt_expected_backends}'.")
 endif()
 
-foreach(_infinirt_backend CPU NVIDIA ILUVATAR HYGON METAX MOORE CAMBRICON ASCEND)
+foreach(_infinirt_backend CPU NVIDIA ILUVATAR HYGON ALI METAX MOORE CAMBRICON ASCEND)
     string(TOLOWER "${_infinirt_backend}" _infinirt_backend_lower)
     list(FIND _infinirt_expected_backends "${_infinirt_backend_lower}"
          _infinirt_backend_index)
diff --git a/tests/install_consumer_smoke.cc b/tests/install_consumer_smoke.cc
index 6b72130..622b42c 100644
--- a/tests/install_consumer_smoke.cc
+++ b/tests/install_consumer_smoke.cc
@@ -23,6 +23,7 @@ int main() {
     defined(INFINI_RT_CONSUMER_BACKEND_NVIDIA) ||    \
     defined(INFINI_RT_CONSUMER_BACKEND_ILUVATAR) ||  \
     defined(INFINI_RT_CONSUMER_BACKEND_HYGON) ||     \
+    defined(INFINI_RT_CONSUMER_BACKEND_ALI) ||       \
     defined(INFINI_RT_CONSUMER_BACKEND_METAX) ||     \
     defined(INFINI_RT_CONSUMER_BACKEND_MOORE) ||     \
     defined(INFINI_RT_CONSUMER_BACKEND_CAMBRICON) || \
@@ -40,6 +41,9 @@ int main() {
 #elif defined(INFINI_RT_CONSUMER_BACKEND_HYGON)
   constexpr auto kExpectedDeviceType = infini::rt::Device::Type::kHygon;
   constexpr bool kExpectAsyncMemcpySuccess = true;
+#elif defined(INFINI_RT_CONSUMER_BACKEND_ALI)
+  constexpr auto kExpectedDeviceType = infini::rt::Device::Type::kAli;
+  constexpr bool kExpectAsyncMemcpySuccess = true;
 #elif defined(INFINI_RT_CONSUMER_BACKEND_METAX)
   constexpr auto kExpectedDeviceType = infini::rt::Device::Type::kMetax;
   constexpr bool kExpectAsyncMemcpySuccess = true;
diff --git a/tests/test_runtime_dispatch.cc b/tests/test_runtime_dispatch.cc
index be49b3b..06bb042 100644
--- a/tests/test_runtime_dispatch.cc
+++ b/tests/test_runtime_dispatch.cc
@@ -420,6 +420,11 @@ int main() {
                {true, true, true, true, true, true, true, true});
 #endif
 
+#if defined(INFINI_RT_TEST_WITH_ALI)
+  TestDispatch(&context, infini::rt::Device::Type::kAli, "ALI",
+               {true, true, true, true, true, true, true, true});
+#endif
+
 #if defined(INFINI_RT_TEST_WITH_METAX)
   TestDispatch(&context, infini::rt::Device::Type::kMetax, "METAX",
                {true, true, true, true, true, true, true, true});