From b002b8f9b24ea04a4704af38660cfe691b6d25b7 Mon Sep 17 00:00:00 2001 From: zhushuang Date: Thu, 23 Jul 2026 10:26:11 +0800 Subject: [PATCH] feat: add Thead PPU runtime backend --- CMakeLists.txt | 52 +++++++- cmake/InfiniRTConfig.cmake.in | 1 + scripts/compare_performance_results.py | 2 +- scripts/generate_public_headers.py | 8 +- scripts/run_performance_tests.py | 6 +- src/CMakeLists.txt | 14 ++ src/device.h | 5 +- src/native/cuda/thead/data_type_.h | 30 +++++ src/native/cuda/thead/device_.h | 13 ++ src/native/cuda/thead/runtime_.h | 175 +++++++++++++++++++++++++ tests/CMakeLists.txt | 19 +++ tests/install_consumer/CMakeLists.txt | 2 +- tests/install_consumer_smoke.cc | 4 + tests/test_runtime_dispatch.cc | 5 + 14 files changed, 323 insertions(+), 13 deletions(-) create mode 100644 src/native/cuda/thead/data_type_.h create mode 100644 src/native/cuda/thead/device_.h create mode 100644 src/native/cuda/thead/runtime_.h diff --git a/CMakeLists.txt b/CMakeLists.txt index 65e63ec..8fe090d 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_THEAD "Enable T-Head 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_THEAD 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_THEAD 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_THEAD`, `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_THEAD) + add_compile_definitions(WITH_THEAD=1) + + set(THEAD_CUDA_ROOT "") + foreach(_thead_cuda_env THEAD_CUDA_ROOT PPU_CUDA_ROOT) + if(NOT THEAD_CUDA_ROOT AND DEFINED ENV{${_thead_cuda_env}} AND + NOT "$ENV{${_thead_cuda_env}}" STREQUAL "") + set(THEAD_CUDA_ROOT "$ENV{${_thead_cuda_env}}") + endif() + endforeach() + if(NOT THEAD_CUDA_ROOT AND DEFINED ENV{PPU_SDK_ROOT} AND + NOT "$ENV{PPU_SDK_ROOT}" STREQUAL "") + set(THEAD_CUDA_ROOT "$ENV{PPU_SDK_ROOT}/CUDA_SDK") + endif() + if(NOT THEAD_CUDA_ROOT AND EXISTS "/usr/local/PPU_SDK/CUDA_SDK") + set(THEAD_CUDA_ROOT "/usr/local/PPU_SDK/CUDA_SDK") + endif() + + if(NOT THEAD_CUDA_ROOT OR NOT EXISTS "${THEAD_CUDA_ROOT}/bin/nvcc") + message(FATAL_ERROR + "`WITH_THEAD` is `ON` but the Thead PPU CUDA SDK was not found. " + "Set `THEAD_CUDA_ROOT`, `PPU_CUDA_ROOT`, or `PPU_SDK_ROOT`.") + endif() + + set(CMAKE_CUDA_COMPILER "${THEAD_CUDA_ROOT}/bin/nvcc" + CACHE FILEPATH "Thead PPU CUDA compiler" FORCE) + set(CUDAToolkit_ROOT "${THEAD_CUDA_ROOT}" + CACHE PATH "Thead PPU CUDA toolkit root" FORCE) + if(NOT CMAKE_CUDA_ARCHITECTURES) + set(CMAKE_CUDA_ARCHITECTURES 80 CACHE STRING + "Thead PPU CUDA architecture" FORCE) + endif() + + message(STATUS "Thead PPU: CUDA toolkit ${THEAD_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_THEAD 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_THEAD) + list(APPEND INFINI_RT_PUBLIC_HEADER_DEVICES thead) +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_THEAD) set(INFINI_RT_PACKAGE_NEEDS_CUDATOOLKIT ON) endif() diff --git a/cmake/InfiniRTConfig.cmake.in b/cmake/InfiniRTConfig.cmake.in index 09bce5a..f4de493 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_THEAD @WITH_THEAD@) set(InfiniRT_WITH_METAX @WITH_METAX@) set(InfiniRT_WITH_MOORE @WITH_MOORE@) set(InfiniRT_WITH_CAMBRICON @WITH_CAMBRICON@) diff --git a/scripts/compare_performance_results.py b/scripts/compare_performance_results.py index 7d8bc04..245758d 100644 --- a/scripts/compare_performance_results.py +++ b/scripts/compare_performance_results.py @@ -8,7 +8,7 @@ def _load_results(path): if isinstance(data, dict) and "results" in data: data = data["results"] if not isinstance(data, list): - raise ValueError(f"{path} must contain a JSON array or an object with results") + raise TypeError(f"{path} must contain a JSON array or an object with results") return data diff --git a/scripts/generate_public_headers.py b/scripts/generate_public_headers.py index 3363741..e6a6f41 100644 --- a/scripts/generate_public_headers.py +++ b/scripts/generate_public_headers.py @@ -3,7 +3,6 @@ import pathlib import re - _DETAIL_PREFIX = "infini/rt/detail" _DEVICE_HEADERS = { @@ -27,6 +26,11 @@ ("hygon", "device_.h", "native/cuda/hygon/device_.h"), ("hygon", "runtime_.h", "native/cuda/hygon/runtime_.h"), ), + "thead": ( + ("thead", "data_type_.h", "native/cuda/thead/data_type_.h"), + ("thead", "device_.h", "native/cuda/thead/device_.h"), + ("thead", "runtime_.h", "native/cuda/thead/runtime_.h"), + ), "metax": ( ("metax", "data_type_.h", "native/cuda/metax/data_type_.h"), ("metax", "device_.h", "native/cuda/metax/device_.h"), @@ -54,6 +58,7 @@ "nvidia": "Device::Type::kNvidia", "iluvatar": "Device::Type::kIluvatar", "hygon": "Device::Type::kHygon", + "thead": "Device::Type::kThead", "metax": "Device::Type::kMetax", "moore": "Device::Type::kMoore", "cambricon": "Device::Type::kCambricon", @@ -64,6 +69,7 @@ "nvidia", "iluvatar", "hygon", + "thead", "metax", "moore", "cambricon", diff --git a/scripts/run_performance_tests.py b/scripts/run_performance_tests.py index 4e08c0d..54f28b3 100644 --- a/scripts/run_performance_tests.py +++ b/scripts/run_performance_tests.py @@ -6,7 +6,6 @@ import subprocess import sys - _BACKEND_OPTIONS = ( ("WITH_NVIDIA", "nvidia"), ("WITH_ILUVATAR", "iluvatar"), @@ -39,7 +38,7 @@ def _read_cmake_cache(build_dir): return values for line in cache_path.read_text(encoding="utf-8", errors="replace").splitlines(): - if line.startswith("//") or line.startswith("#") or "=" not in line: + if line.startswith(("//", "#")) or "=" not in line: continue key_type, value = line.split("=", 1) key = key_type.split(":", 1)[0] @@ -105,8 +104,7 @@ def _run_executable(path): [str(path)], cwd=path.parent, text=True, - stdout=subprocess.PIPE, - stderr=subprocess.PIPE, + capture_output=True, check=False, ) if completed.stderr: diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 2091963..b6acd58 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -63,6 +63,20 @@ if(WITH_HYGON) ) endif() +if(WITH_THEAD) + enable_language(CUDA) + + target_compile_definitions(infinirt PUBLIC WITH_THEAD=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..db3a20d 100644 --- a/src/device.h +++ b/src/device.h @@ -21,6 +21,7 @@ class Device { kMoore = 5, kIluvatar = 6, kHygon = 7, + kThead = 8, kCount }; @@ -64,6 +65,7 @@ class Device { {Type::kMoore, "moore"}, {Type::kIluvatar, "iluvatar"}, {Type::kHygon, "hygon"}, + {Type::kThead, "thead"}, }}}; static constexpr ConstexprMap; + Device::Type::kIluvatar, Device::Type::kHygon, Device::Type::kThead>; // 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/thead/data_type_.h b/src/native/cuda/thead/data_type_.h new file mode 100644 index 0000000..745e455 --- /dev/null +++ b/src/native/cuda/thead/data_type_.h @@ -0,0 +1,30 @@ +#ifndef INFINI_RT_THEAD_DATA_TYPE__H_ +#define INFINI_RT_THEAD_DATA_TYPE__H_ + +// clang-format off +#include +#include +// clang-format on + +#include "data_type.h" +#include "native/cuda/thead/device_.h" + +namespace infini::rt { + +using cuda_bfloat16 = nv_bfloat16; + +using cuda_bfloat162 = nv_bfloat162; + +template <> +struct TypeMap { + using type = half; +}; + +template <> +struct TypeMap { + using type = __nv_bfloat16; +}; + +} // namespace infini::rt + +#endif diff --git a/src/native/cuda/thead/device_.h b/src/native/cuda/thead/device_.h new file mode 100644 index 0000000..af2bbeb --- /dev/null +++ b/src/native/cuda/thead/device_.h @@ -0,0 +1,13 @@ +#ifndef INFINI_RT_THEAD_DEVICE__H_ +#define INFINI_RT_THEAD_DEVICE__H_ + +#include "device.h" + +namespace infini::rt { + +template <> +struct DeviceEnabled : std::true_type {}; + +} // namespace infini::rt + +#endif diff --git a/src/native/cuda/thead/runtime_.h b/src/native/cuda/thead/runtime_.h new file mode 100644 index 0000000..d96be70 --- /dev/null +++ b/src/native/cuda/thead/runtime_.h @@ -0,0 +1,175 @@ +#ifndef INFINI_RT_THEAD_RUNTIME__H_ +#define INFINI_RT_THEAD_RUNTIME__H_ + +#include +#include + +// clang-format off +#include +// clang-format on + +#include "native/cuda/runtime_.h" +#include "native/cuda/thead/device_.h" + +namespace infini::rt::runtime { + +template <> +struct Runtime + : GraphRuntime, + CudaRuntime>> { + 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::kThead; + + 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(args)...); + }; + + static constexpr auto MallocHost = [](auto&&... args) { + return cudaMallocHost(std::forward(args)...); + }; + + static constexpr auto MallocAsync = [](auto&&... args) { + return cudaMallocAsync(std::forward(args)...); + }; + + static constexpr auto Free = cudaFree; + + static constexpr auto FreeHost = [](auto&&... args) { + return cudaFreeHost(std::forward(args)...); + }; + + static constexpr auto FreeAsync = [](auto&&... args) { + return cudaFreeAsync(std::forward(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(args)...); + }; + + static constexpr auto StreamCreate = cudaStreamCreate; + + static constexpr auto StreamDestroy = [](auto&&... args) { + return cudaStreamDestroy(std::forward(args)...); + }; + + static constexpr auto StreamSynchronize = [](auto&&... args) { + return cudaStreamSynchronize(std::forward(args)...); + }; + + static constexpr auto StreamWaitEvent = [](auto&&... args) { + return cudaStreamWaitEvent(std::forward(args)...); + }; + + static constexpr auto EventCreate = [](auto&&... args) { + return cudaEventCreate(std::forward(args)...); + }; + + static constexpr auto EventCreateWithFlags = [](auto&&... args) { + return cudaEventCreateWithFlags(std::forward(args)...); + }; + + static constexpr auto EventRecord = [](auto&&... args) { + return cudaEventRecord(std::forward(args)...); + }; + + static constexpr auto EventQuery = [](auto&&... args) { + return cudaEventQuery(std::forward(args)...); + }; + + static constexpr auto EventSynchronize = [](auto&&... args) { + return cudaEventSynchronize(std::forward(args)...); + }; + + static constexpr auto EventDestroy = [](auto&&... args) { + return cudaEventDestroy(std::forward(args)...); + }; + + static constexpr auto EventElapsedTime = [](auto&&... args) { + return cudaEventElapsedTime(std::forward(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(args)...); + }; + + static constexpr auto StreamEndCapture = [](auto&&... args) { + return cudaStreamEndCapture(std::forward(args)...); + }; + + static constexpr auto GraphDestroy = [](auto&&... args) { + return cudaGraphDestroy(std::forward(args)...); + }; + + static constexpr auto GraphInstantiate = [](auto&&... args) { + return cudaGraphInstantiate(std::forward(args)...); + }; + + static constexpr auto GraphExecDestroy = [](auto&&... args) { + return cudaGraphExecDestroy(std::forward(args)...); + }; + + static constexpr auto GraphLaunch = [](auto&&... args) { + return cudaGraphLaunch(std::forward(args)...); + }; +}; + +static_assert(Runtime::Validate()); + +} // namespace infini::rt::runtime + +#endif diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 8ffa4b3..2172c73 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_THEAD) + set(INFINI_RT_TEST_HAS_RUNTIME_BACKEND ON) + add_infini_rt_backend_runtime_test( + THEAD infini::rt::Device::Type::kThead + infini/rt/thead/runtime_.h + 1 1 1 1 1 1 1 1) + add_infini_rt_backend_graph_test( + THEAD infini::rt::Device::Type::kThead 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_THEAD) + target_compile_definitions(test_runtime_dispatch + PRIVATE INFINI_RT_TEST_WITH_THEAD=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_THEAD) + set(INFINI_RT_TEST_CONSUMER_BACKEND THEAD) + if(NOT INFINI_RT_TEST_CUDATOOLKIT_ROOT) + set(INFINI_RT_TEST_CUDATOOLKIT_ROOT "${THEAD_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..4030592 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 THEAD 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..748d55d 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_THEAD) || \ 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_THEAD) + constexpr auto kExpectedDeviceType = infini::rt::Device::Type::kThead; + 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..368f6dc 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_THEAD) + TestDispatch(&context, infini::rt::Device::Type::kThead, "THEAD", + {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});