From 69e4a92c4750d1409f1e394ef73807ea6feccef6 Mon Sep 17 00:00:00 2001 From: Zimin Li Date: Thu, 10 Sep 2026 15:38:49 +0800 Subject: [PATCH 1/9] feat: support Ali PPU as a valid device/platform, including build, runtime, and CCL --- CMakeLists.txt | 148 ++++++++++++++++++++++++++++++-- examples/CMakeLists.txt | 4 + src/CMakeLists.txt | 23 +++++ src/backend_device_map.h | 4 + src/backends/ccl/nccl/ali/api.h | 24 ++++++ src/device.h | 10 ++- src/devices/ali/caster_.cuh | 64 ++++++++++++++ src/devices/ali/data_type_.h | 26 ++++++ src/devices/ali/device_.h | 32 +++++++ src/devices/ali/runtime_.h | 59 +++++++++++++ 10 files changed, 384 insertions(+), 10 deletions(-) create mode 100644 src/backends/ccl/nccl/ali/api.h create mode 100644 src/devices/ali/caster_.cuh create mode 100644 src/devices/ali/data_type_.h create mode 100644 src/devices/ali/device_.h create mode 100644 src/devices/ali/runtime_.h diff --git a/CMakeLists.txt b/CMakeLists.txt index 48181fc..5abf69d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -8,6 +8,7 @@ set(CMAKE_CXX_STANDARD_REQUIRED ON) # --- DEVICE OPTIONS (Hardware Runtimes) --- # ========================================================= option(WITH_NVIDIA "Enable NVIDIA GPU support" OFF) +option(WITH_ALI "Enable Ali PPU support" OFF) option(WITH_ILUVATAR "Enable ILUVATAR GPU support" OFF) option(WITH_METAX "Enable MetaX GPU support" OFF) option(WITH_MOORE "Enable Moore GPU support" OFF) @@ -38,6 +39,57 @@ option(BUILD_EXAMPLES "Build internal examples" ON) if(AUTO_DETECT_DEVICES) message(STATUS "Auto-detecting available devices...") + # Ali PPU + set(ALI_FOUND FALSE) + + if(WITH_ALI) + set(ALI_FOUND TRUE) + else() + set(_ALI_SDK_HINTS) + foreach(_ali_env PPU_SDK PPU_HOME PPU_PATH) + if(DEFINED ENV{${_ali_env}} AND NOT "$ENV{${_ali_env}}" STREQUAL "") + list(APPEND _ALI_SDK_HINTS "$ENV{${_ali_env}}") + endif() + endforeach() + foreach(_ali_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) + if(DEFINED ENV{${_ali_cuda_env}} AND NOT "$ENV{${_ali_cuda_env}}" STREQUAL "") + list(APPEND _ALI_SDK_HINTS "$ENV{${_ali_cuda_env}}" "$ENV{${_ali_cuda_env}}/..") + endif() + endforeach() + list(APPEND _ALI_SDK_HINTS /usr/local/PPU_SDK /usr/local/PPU_SDK/CUDA_SDK) + + foreach(_ali_hint ${_ALI_SDK_HINTS}) + if(EXISTS "${_ali_hint}/CUDA_SDK" OR EXISTS "${_ali_hint}/bin/ppu-smi" OR EXISTS "${_ali_hint}/ppu-smi/bin/ppu-smi") + set(ALI_FOUND TRUE) + endif() + endforeach() + + if(NOT ALI_FOUND) + find_program(ALI_SMI_PATH + NAMES ppu-smi nvidia-smi + HINTS /usr/local/PPU_SDK/bin /usr/local/PPU_SDK/ppu-smi/bin + ) + if(ALI_SMI_PATH) + execute_process( + COMMAND ${ALI_SMI_PATH} -L + RESULT_VARIABLE SMI_RESULT + OUTPUT_VARIABLE SMI_OUTPUT + ERROR_QUIET + ) + if(SMI_RESULT EQUAL 0 AND SMI_OUTPUT MATCHES "PPU") + set(ALI_FOUND TRUE) + endif() + endif() + endif() + endif() + + if(ALI_FOUND) + set(WITH_ALI ON) + message(STATUS "Auto-detected Ali PPU environment.") + else() + message(STATUS "Ali PPU environment not detected.") + endif() + # NVIDIA set(NVIDIA_FOUND FALSE) @@ -60,6 +112,11 @@ if(AUTO_DETECT_DEVICES) endif() endif() + if(NVIDIA_FOUND AND ALI_FOUND AND NOT WITH_NVIDIA) + set(NVIDIA_FOUND FALSE) + message(STATUS "Ali PPU environment detected; skipping NVIDIA auto-detection.") + endif() + if(NVIDIA_FOUND) set(WITH_NVIDIA ON) message(STATUS "Auto-detected NVIDIA environment.") @@ -251,12 +308,26 @@ if(AUTO_DETECT_BACKENDS) endif() # Detect NCCL Dependencies - if(WITH_NVIDIA OR WITH_ILUVATAR OR WITH_HYGON) + if(WITH_NVIDIA OR WITH_ALI OR WITH_ILUVATAR OR WITH_HYGON) set(_NCCL_HINTS) set(_NCCL_HEADER_NAMES nccl.h) set(_NCCL_LIBRARY_NAMES nccl) - if (WITH_ILUVATAR) + if(WITH_ALI) + foreach(_ali_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) + if(DEFINED ENV{${_ali_cuda_env}} AND NOT "$ENV{${_ali_cuda_env}}" STREQUAL "") + list(APPEND _NCCL_HINTS "$ENV{${_ali_cuda_env}}") + endif() + endforeach() + foreach(_ali_env PPU_SDK PPU_HOME PPU_PATH) + if(DEFINED ENV{${_ali_env}} AND NOT "$ENV{${_ali_env}}" STREQUAL "") + list(APPEND _NCCL_HINTS "$ENV{${_ali_env}}/CUDA_SDK" "$ENV{${_ali_env}}") + endif() + endforeach() + list(APPEND _NCCL_HINTS /usr/local/PPU_SDK/CUDA_SDK /usr/local/PPU_SDK) + list(REMOVE_DUPLICATES _NCCL_HINTS) + message(STATUS "Ali PPU detected. Searching for NCCL in ${_NCCL_HINTS}") + elseif(WITH_ILUVATAR) set(_NCCL_HINTS /usr/local/corex) message(STATUS "Iluvatar detected. Searching for NCCL in ${_NCCL_HINTS}") elseif(WITH_HYGON) @@ -277,8 +348,14 @@ if(AUTO_DETECT_BACKENDS) message(STATUS "Hygon detected. Searching for NCCL-compatible RCCL in ${_NCCL_HINTS}") endif() - find_path(AUTO_NCCL_INC NAMES ${_NCCL_HEADER_NAMES} HINTS ${_NCCL_HINTS} PATH_SUFFIXES include QUIET) - find_library(AUTO_NCCL_LIB NAMES ${_NCCL_LIBRARY_NAMES} HINTS ${_NCCL_HINTS} PATH_SUFFIXES lib lib64 rccl/lib QUIET) + find_path(AUTO_NCCL_INC NAMES ${_NCCL_HEADER_NAMES} + HINTS ${_NCCL_HINTS} + PATH_SUFFIXES include targets/x86_64-linux/include + QUIET) + find_library(AUTO_NCCL_LIB NAMES ${_NCCL_LIBRARY_NAMES} + HINTS ${_NCCL_HINTS} + PATH_SUFFIXES lib lib64 targets/x86_64-linux/lib rccl/lib + QUIET) if(AUTO_NCCL_INC AND AUTO_NCCL_LIB) set(WITH_NCCL ON) @@ -352,6 +429,40 @@ endif() # ========================================================= # --- DEPENDENCY DISCOVERY --- # ========================================================= +if(WITH_ALI) + set(_ALI_CUDA_HINTS) + foreach(_ali_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) + if(DEFINED ENV{${_ali_cuda_env}} AND NOT "$ENV{${_ali_cuda_env}}" STREQUAL "") + list(APPEND _ALI_CUDA_HINTS "$ENV{${_ali_cuda_env}}") + endif() + endforeach() + foreach(_ali_env PPU_SDK PPU_HOME PPU_PATH) + if(DEFINED ENV{${_ali_env}} AND NOT "$ENV{${_ali_env}}" STREQUAL "") + list(APPEND _ALI_CUDA_HINTS "$ENV{${_ali_env}}/CUDA_SDK" "$ENV{${_ali_env}}") + endif() + endforeach() + list(APPEND _ALI_CUDA_HINTS /usr/local/PPU_SDK/CUDA_SDK /usr/local/PPU_SDK /usr/local/cuda) + list(REMOVE_DUPLICATES _ALI_CUDA_HINTS) + + set(ALI_CUDA_ROOT "") + foreach(_ali_cuda_hint ${_ALI_CUDA_HINTS}) + if(NOT ALI_CUDA_ROOT AND EXISTS "${_ali_cuda_hint}/include/cuda_runtime.h") + set(ALI_CUDA_ROOT "${_ali_cuda_hint}") + endif() + endforeach() + + if(NOT ALI_CUDA_ROOT) + message(FATAL_ERROR "`WITH_ALI` is `ON` but the Ali PPU CUDA-compatible SDK was not found. Set `PPU_SDK`, `PPU_HOME`, `CUDA_SDK`, `CUDA_HOME`, or `CUDA_PATH`.") + endif() + + set(CUDAToolkit_ROOT "${ALI_CUDA_ROOT}" CACHE PATH "Ali PPU CUDA-compatible toolkit root" FORCE) + list(APPEND _NCCL_HINTS ${_ALI_CUDA_HINTS}) + list(REMOVE_DUPLICATES _NCCL_HINTS) + + enable_language(CUDA) + find_package(CUDAToolkit REQUIRED) +endif() + if(WITH_NVIDIA) enable_language(CUDA) find_package(CUDAToolkit REQUIRED) @@ -504,23 +615,42 @@ if(WITH_OMPI OR WITH_MPICH) endif() if(WITH_NCCL) - if (NOT WITH_NVIDIA AND NOT WITH_ILUVATAR AND NOT WITH_HYGON) - message(FATAL_ERROR "NCCL backend requires NVIDIA, Iluvatar, or Hygon GPU support. Please enable `WITH_NVIDIA`, `WITH_ILUVATAR`, or `WITH_HYGON`.") + if (NOT WITH_NVIDIA AND NOT WITH_ALI AND NOT WITH_ILUVATAR AND NOT WITH_HYGON) + message(FATAL_ERROR "NCCL backend requires NVIDIA, Ali PPU, Iluvatar, or Hygon GPU support. Please enable `WITH_NVIDIA`, `WITH_ALI`, `WITH_ILUVATAR`, or `WITH_HYGON`.") endif() set(_NCCL_HEADER_NAMES nccl.h) set(_NCCL_LIBRARY_NAMES nccl) set(NCCL_COMPILE_DEFINITIONS) - if(WITH_HYGON) + if(WITH_ALI) + foreach(_ali_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) + if(DEFINED ENV{${_ali_cuda_env}} AND NOT "$ENV{${_ali_cuda_env}}" STREQUAL "") + list(APPEND _NCCL_HINTS "$ENV{${_ali_cuda_env}}") + endif() + endforeach() + foreach(_ali_env PPU_SDK PPU_HOME PPU_PATH) + if(DEFINED ENV{${_ali_env}} AND NOT "$ENV{${_ali_env}}" STREQUAL "") + list(APPEND _NCCL_HINTS "$ENV{${_ali_env}}/CUDA_SDK" "$ENV{${_ali_env}}") + endif() + endforeach() + list(APPEND _NCCL_HINTS /usr/local/PPU_SDK/CUDA_SDK /usr/local/PPU_SDK) + elseif(WITH_HYGON) list(APPEND _NCCL_HINTS "${HYGON_DTK_ROOT}" "${HYGON_DTK_ROOT}/rccl") set(_NCCL_HEADER_NAMES rccl/rccl.h rccl.h) set(_NCCL_LIBRARY_NAMES rccl) set(NCCL_COMPILE_DEFINITIONS INFINI_CCL_USE_RCCL) endif() - find_library(NCCL_LIB NAMES ${_NCCL_LIBRARY_NAMES} HINTS ${_NCCL_HINTS} PATH_SUFFIXES lib lib64 rccl/lib REQUIRED) - find_path(NCCL_INC NAMES ${_NCCL_HEADER_NAMES} HINTS ${_NCCL_HINTS} PATH_SUFFIXES include REQUIRED) + list(REMOVE_DUPLICATES _NCCL_HINTS) + find_library(NCCL_LIB NAMES ${_NCCL_LIBRARY_NAMES} + HINTS ${_NCCL_HINTS} + PATH_SUFFIXES lib lib64 targets/x86_64-linux/lib rccl/lib + REQUIRED) + find_path(NCCL_INC NAMES ${_NCCL_HEADER_NAMES} + HINTS ${_NCCL_HINTS} + PATH_SUFFIXES include targets/x86_64-linux/include + REQUIRED) include_directories(${NCCL_INC}) endif() diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 2a39876..29741d9 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -26,6 +26,10 @@ foreach(source_file ${EXAMPLE_SOURCES}) target_link_libraries(${target_name} PRIVATE CUDA::cudart) endif() + if(WITH_ALI) + target_link_libraries(${target_name} PRIVATE CUDA::cudart) + endif() + if(WITH_ILUVATAR) set_source_files_properties(${source_file} PROPERTIES LANGUAGE CXX) set_target_properties(${target_name} PROPERTIES diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index c191b68..bfda536 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -68,6 +68,29 @@ if(WITH_NVIDIA) ) endif() +# Ali PPU +if(WITH_ALI) + list(APPEND DEVICE_LIST "ali") + + set(ALI_PATTERNS + "devices/cuda/*.cc" + "devices/cuda/*.cpp" + "devices/cuda/*.cu" + "devices/ali/*.cc" + "devices/ali/*.cpp" + "devices/ali/*.cu" + ) + file(GLOB ALI_SOURCES ${ALI_PATTERNS}) + + target_sources(infiniccl PRIVATE ${ALI_SOURCES}) + target_link_libraries(infiniccl PRIVATE CUDA::cudart CUDA::cuda_driver) + + set_target_properties(infiniccl PROPERTIES + CUDA_STANDARD 17 + CUDA_STANDARD_REQUIRED ON + ) +endif() + # Iluvatar if(WITH_ILUVATAR) list(APPEND DEVICE_LIST "iluvatar") diff --git a/src/backend_device_map.h b/src/backend_device_map.h index 8c65d82..51ac898 100644 --- a/src/backend_device_map.h +++ b/src/backend_device_map.h @@ -17,6 +17,10 @@ template <> struct IsSupportedCombination : std::true_type {}; +template <> +struct IsSupportedCombination + : std::true_type {}; + template <> struct IsSupportedCombination : std::true_type {}; diff --git a/src/backends/ccl/nccl/ali/api.h b/src/backends/ccl/nccl/ali/api.h new file mode 100644 index 0000000..92cfbbc --- /dev/null +++ b/src/backends/ccl/nccl/ali/api.h @@ -0,0 +1,24 @@ +#ifndef INFINI_CCL_BACKENDS_CCL_NCCL_ALI_API_H_ +#define INFINI_CCL_BACKENDS_CCL_NCCL_ALI_API_H_ + +#include "backends/ccl/nccl/api.h" +#include "devices/ali/runtime_.h" + +namespace infini::ccl { + +template <> +struct NcclDataTypeTraits { +#if defined(__CUDA_BF16_TYPES_EXIST__) + static constexpr ncclDataType_t kBFloat16 = ncclBfloat16; +#else + static constexpr ncclDataType_t kBFloat16 = ncclNumTypes; +#endif +}; + +template <> +struct CclApi + : NcclApi {}; + +} // namespace infini::ccl + +#endif // INFINI_CCL_BACKENDS_CCL_NCCL_ALI_API_H_ diff --git a/src/device.h b/src/device.h index 4809420..4dbf0de 100644 --- a/src/device.h +++ b/src/device.h @@ -22,6 +22,7 @@ class Device { kKunlun = 7, kHygon = 8, kQy = 9, + kAli = 10, kCount }; @@ -67,6 +68,7 @@ class Device { {Type::kKunlun, "kunlun"}, {Type::kHygon, "hygon"}, {Type::kQy, "qy"}, + {Type::kAli, "ali"}, }}}; static constexpr ConstexprMap; + Device::Type::kAli, Device::Type::kQy>; // Deferred computation of active devices. The `Filter` and `FilterList` // evaluation are nested inside a class template so that `DeviceEnabled` @@ -161,6 +164,11 @@ struct DevicePriority { static constexpr int value = 5; }; +template <> +struct DevicePriority { + static constexpr int value = 5; +}; + enum class MemorySpace : std::uint8_t { kHost = 0, kDevice = 1, kUnknown }; template diff --git a/src/devices/ali/caster_.cuh b/src/devices/ali/caster_.cuh new file mode 100644 index 0000000..18e8cd7 --- /dev/null +++ b/src/devices/ali/caster_.cuh @@ -0,0 +1,64 @@ +#ifndef INFINI_CCL_DEVICES_ALI_CASTER_CUH_ +#define INFINI_CCL_DEVICES_ALI_CASTER_CUH_ + +#include "caster.h" +#include "data_type_.h" + +namespace infini::ccl { + +template <> +struct HardwareCastImpl { + __host__ __device__ static float Apply(half x) { return __half2float(x); } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static half Apply(float x) { return __float2half(x); } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static float Apply(__nv_bfloat16 x) { + return __bfloat162float(x); + } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static __nv_bfloat16 Apply(float x) { + return __float2bfloat16(x); + } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static __nv_bfloat16 Apply(int x) { + return __int2bfloat16_rn(x); + } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static half Apply(int x) { return __int2half_rn(x); } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static __nv_bfloat16 Apply(double x) { + return __double2bfloat16(x); + } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static half Apply(double x) { return __double2half(x); } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static half Apply(__nv_bfloat16 x) { return __half(x); } +}; + +} // namespace infini::ccl + +#endif // INFINI_CCL_DEVICES_ALI_CASTER_CUH_ diff --git a/src/devices/ali/data_type_.h b/src/devices/ali/data_type_.h new file mode 100644 index 0000000..b4600f5 --- /dev/null +++ b/src/devices/ali/data_type_.h @@ -0,0 +1,26 @@ +#ifndef INFINI_CCL_DEVICES_ALI_DATA_TYPE_H_ +#define INFINI_CCL_DEVICES_ALI_DATA_TYPE_H_ + +// clang-format off +#include +#include +// clang-format on + +#include "data_type_impl.h" +#include "devices/ali/device_.h" + +namespace infini::ccl { + +template <> +struct TypeMap { + using type = half; +}; + +template <> +struct TypeMap { + using type = __nv_bfloat16; +}; + +} // namespace infini::ccl + +#endif // INFINI_CCL_DEVICES_ALI_DATA_TYPE_H_ diff --git a/src/devices/ali/device_.h b/src/devices/ali/device_.h new file mode 100644 index 0000000..86a5519 --- /dev/null +++ b/src/devices/ali/device_.h @@ -0,0 +1,32 @@ +#ifndef INFINI_CCL_DEVICES_ALI_DEVICE_H_ +#define INFINI_CCL_DEVICES_ALI_DEVICE_H_ + +#include + +#include "device.h" +#include "devices/cuda/checks.h" + +namespace infini::ccl { + +template <> +struct DeviceEnabled : std::true_type {}; + +template <> +MemorySpace GetMemorySpace(const void* ptr) { + if (!ptr) { + return MemorySpace::kHost; + } + + cudaPointerAttributes attr; + INFINI_CHECK_CUDA(cudaPointerGetAttributes(&attr, ptr)); + + if (attr.type == cudaMemoryTypeDevice || attr.type == cudaMemoryTypeManaged) { + return MemorySpace::kDevice; + } + + return MemorySpace::kHost; +} + +} // namespace infini::ccl + +#endif // INFINI_CCL_DEVICES_ALI_DEVICE_H_ diff --git a/src/devices/ali/runtime_.h b/src/devices/ali/runtime_.h new file mode 100644 index 0000000..c12066c --- /dev/null +++ b/src/devices/ali/runtime_.h @@ -0,0 +1,59 @@ +#ifndef INFINI_CCL_DEVICES_ALI_RUNTIME_H_ +#define INFINI_CCL_DEVICES_ALI_RUNTIME_H_ + +#include + +// clang-format off +#include +// clang-format on + +#include "devices/ali/device_.h" +#include "devices/cuda/runtime_.h" +#include "logging.h" +#include "return_status_impl.h" + +namespace infini::ccl { + +template <> +struct Runtime : CudaRuntime> { + using Stream = cudaStream_t; + + static constexpr Device::Type kDeviceType = Device::Type::kAli; + + static constexpr auto Check = + [](auto status, ReturnStatus err_code = ReturnStatus::kSystemError) { + if (status != cudaSuccess) { + LOG(cudaGetErrorString(static_cast(status))); + return err_code; + } + return ReturnStatus::kSuccess; + }; + + static constexpr auto Malloc = [](auto &&...args) { + return cudaMalloc(std::forward(args)...); + }; + + static constexpr auto Memcpy = cudaMemcpy; + + static constexpr auto Free = cudaFree; + + static constexpr auto MemcpyHostToDevice = cudaMemcpyHostToDevice; + + static constexpr auto MemcpyDeviceToHost = cudaMemcpyDeviceToHost; + + static constexpr auto Memset = cudaMemset; + + static constexpr auto GetDevice = cudaGetDevice; + + static constexpr auto SetDevice = cudaSetDevice; + + static constexpr auto DeviceSynchronize = cudaDeviceSynchronize; + + static constexpr auto StreamSynchronize = cudaStreamSynchronize; +}; + +static_assert(Runtime::Validate()); + +} // namespace infini::ccl + +#endif // INFINI_CCL_DEVICES_ALI_RUNTIME_H_ From 114b02cf3dd6df2e0603dbe0d0d06b86244bc283 Mon Sep 17 00:00:00 2001 From: Zimin Li Date: Thu, 10 Sep 2026 15:56:37 +0800 Subject: [PATCH 2/9] fix: handle local aliases in icclrun --- scripts/icclrun_logic.py | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/scripts/icclrun_logic.py b/scripts/icclrun_logic.py index c0f0d63..065cdf6 100644 --- a/scripts/icclrun_logic.py +++ b/scripts/icclrun_logic.py @@ -75,10 +75,15 @@ def __init__(self, manual_config_path): def _is_local(self, ip): """Check if the IP/hostname refers to the local machine.""" try: - local_ips = {"127.0.0.1", "localhost"} + local_ips = {"127.0.0.1", "localhost", "::1", socket.gethostname()} + if ip in local_ips: + return True # Add hostname-resolved IPs. - local_ips.update(socket.gethostbyname_ex(socket.gethostname())[2]) + try: + local_ips.update(socket.gethostbyname_ex(socket.gethostname())[2]) + except Exception: + pass # Add all interface IPs via `hostname -I`. try: From 18d690ffa4948062afa955c132ee02972eba9f57 Mon Sep 17 00:00:00 2001 From: Zimin Li Date: Thu, 24 Sep 2026 17:38:44 +0800 Subject: [PATCH 3/9] fix: rename and align T-Head platform - Rename public Ali identifiers to T-Head - Prevent false NVIDIA auto-detection while preserving explicit NVIDIA selection. --- CMakeLists.txt | 203 ++++++++++++------------ examples/CMakeLists.txt | 2 +- scripts/gen_bridge.py | 10 +- src/CMakeLists.txt | 18 +-- src/backend_device_map.h | 2 +- src/backends/ccl/nccl/ali/api.h | 24 --- src/backends/ccl/nccl/t-head/api.h | 24 +++ src/device.h | 10 +- src/devices/ali/data_type_.h | 26 --- src/devices/{ali => t-head}/caster_.cuh | 24 +-- src/devices/t-head/data_type_.h | 26 +++ src/devices/{ali => t-head}/device_.h | 10 +- src/devices/{ali => t-head}/runtime_.h | 14 +- 13 files changed, 201 insertions(+), 192 deletions(-) delete mode 100644 src/backends/ccl/nccl/ali/api.h create mode 100644 src/backends/ccl/nccl/t-head/api.h delete mode 100644 src/devices/ali/data_type_.h rename src/devices/{ali => t-head}/caster_.cuh (58%) create mode 100644 src/devices/t-head/data_type_.h rename src/devices/{ali => t-head}/device_.h (62%) rename src/devices/{ali => t-head}/runtime_.h (75%) diff --git a/CMakeLists.txt b/CMakeLists.txt index 5abf69d..c4a393f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -8,12 +8,12 @@ set(CMAKE_CXX_STANDARD_REQUIRED ON) # --- DEVICE OPTIONS (Hardware Runtimes) --- # ========================================================= option(WITH_NVIDIA "Enable NVIDIA GPU support" OFF) -option(WITH_ALI "Enable Ali PPU support" OFF) option(WITH_ILUVATAR "Enable ILUVATAR GPU support" OFF) option(WITH_METAX "Enable MetaX GPU support" OFF) option(WITH_MOORE "Enable Moore GPU support" OFF) option(WITH_CAMBRICON "Enable Cambricon MLU support" OFF) option(WITH_HYGON "Enable Hygon DCU support" OFF) +option(WITH_THEAD "Enable T-Head PPU support" OFF) set(WITH_CPU ON CACHE INTERNAL "CPU backend is always enabled") @@ -39,58 +39,8 @@ option(BUILD_EXAMPLES "Build internal examples" ON) if(AUTO_DETECT_DEVICES) message(STATUS "Auto-detecting available devices...") - # Ali PPU - set(ALI_FOUND FALSE) - - if(WITH_ALI) - set(ALI_FOUND TRUE) - else() - set(_ALI_SDK_HINTS) - foreach(_ali_env PPU_SDK PPU_HOME PPU_PATH) - if(DEFINED ENV{${_ali_env}} AND NOT "$ENV{${_ali_env}}" STREQUAL "") - list(APPEND _ALI_SDK_HINTS "$ENV{${_ali_env}}") - endif() - endforeach() - foreach(_ali_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) - if(DEFINED ENV{${_ali_cuda_env}} AND NOT "$ENV{${_ali_cuda_env}}" STREQUAL "") - list(APPEND _ALI_SDK_HINTS "$ENV{${_ali_cuda_env}}" "$ENV{${_ali_cuda_env}}/..") - endif() - endforeach() - list(APPEND _ALI_SDK_HINTS /usr/local/PPU_SDK /usr/local/PPU_SDK/CUDA_SDK) - - foreach(_ali_hint ${_ALI_SDK_HINTS}) - if(EXISTS "${_ali_hint}/CUDA_SDK" OR EXISTS "${_ali_hint}/bin/ppu-smi" OR EXISTS "${_ali_hint}/ppu-smi/bin/ppu-smi") - set(ALI_FOUND TRUE) - endif() - endforeach() - - if(NOT ALI_FOUND) - find_program(ALI_SMI_PATH - NAMES ppu-smi nvidia-smi - HINTS /usr/local/PPU_SDK/bin /usr/local/PPU_SDK/ppu-smi/bin - ) - if(ALI_SMI_PATH) - execute_process( - COMMAND ${ALI_SMI_PATH} -L - RESULT_VARIABLE SMI_RESULT - OUTPUT_VARIABLE SMI_OUTPUT - ERROR_QUIET - ) - if(SMI_RESULT EQUAL 0 AND SMI_OUTPUT MATCHES "PPU") - set(ALI_FOUND TRUE) - endif() - endif() - endif() - endif() - - if(ALI_FOUND) - set(WITH_ALI ON) - message(STATUS "Auto-detected Ali PPU environment.") - else() - message(STATUS "Ali PPU environment not detected.") - endif() - # NVIDIA + set(NVIDIA_REQUESTED ${WITH_NVIDIA}) set(NVIDIA_FOUND FALSE) file(GLOB NVIDIA_DEV_FILES "/dev/nvidia0") @@ -112,18 +62,6 @@ if(AUTO_DETECT_DEVICES) endif() endif() - if(NVIDIA_FOUND AND ALI_FOUND AND NOT WITH_NVIDIA) - set(NVIDIA_FOUND FALSE) - message(STATUS "Ali PPU environment detected; skipping NVIDIA auto-detection.") - endif() - - if(NVIDIA_FOUND) - set(WITH_NVIDIA ON) - message(STATUS "Auto-detected NVIDIA environment.") - else() - message(STATUS "NVIDIA environment not detected.") - endif() - # Iluvatar set(ILUVATAR_FOUND FALSE) @@ -273,6 +211,71 @@ if(AUTO_DETECT_DEVICES) else() message(STATUS "Hygon environment not detected.") endif() + + # T-Head + set(THEAD_FOUND FALSE) + + if(WITH_THEAD) + set(THEAD_FOUND TRUE) + else() + set(_THEAD_SDK_HINTS) + foreach(_thead_env PPU_SDK PPU_HOME PPU_PATH) + if(DEFINED ENV{${_thead_env}} AND NOT "$ENV{${_thead_env}}" STREQUAL "") + list(APPEND _THEAD_SDK_HINTS "$ENV{${_thead_env}}") + endif() + endforeach() + foreach(_thead_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) + if(DEFINED ENV{${_thead_cuda_env}} AND NOT "$ENV{${_thead_cuda_env}}" STREQUAL "") + list(APPEND _THEAD_SDK_HINTS "$ENV{${_thead_cuda_env}}" "$ENV{${_thead_cuda_env}}/..") + endif() + endforeach() + list(APPEND _THEAD_SDK_HINTS /usr/local/PPU_SDK /usr/local/PPU_SDK/CUDA_SDK) + + foreach(_thead_hint ${_THEAD_SDK_HINTS}) + if(EXISTS "${_thead_hint}/CUDA_SDK" OR EXISTS "${_thead_hint}/bin/ppu-smi" OR EXISTS "${_thead_hint}/ppu-smi/bin/ppu-smi") + set(THEAD_FOUND TRUE) + endif() + endforeach() + + if(NOT THEAD_FOUND) + find_program(THEAD_SMI_PATH + NAMES ppu-smi nvidia-smi + HINTS /usr/local/PPU_SDK/bin /usr/local/PPU_SDK/ppu-smi/bin + ) + if(THEAD_SMI_PATH) + execute_process( + COMMAND ${THEAD_SMI_PATH} -L + RESULT_VARIABLE SMI_RESULT + OUTPUT_VARIABLE SMI_OUTPUT + ERROR_QUIET + ) + if(SMI_RESULT EQUAL 0 AND SMI_OUTPUT MATCHES "PPU") + set(THEAD_FOUND TRUE) + endif() + endif() + endif() + endif() + + if(THEAD_FOUND) + set(WITH_THEAD ON) + message(STATUS "Auto-detected T-Head environment.") + else() + message(STATUS "T-Head environment not detected.") + endif() + + if(THEAD_FOUND AND NVIDIA_FOUND AND NOT NVIDIA_REQUESTED) + set(NVIDIA_FOUND FALSE) + set(WITH_NVIDIA OFF CACHE BOOL "Enable NVIDIA GPU support" FORCE) + message(STATUS "T-Head environment detected; ignoring NVIDIA auto-detection.") + endif() + + if(NVIDIA_FOUND) + set(WITH_NVIDIA ON) + message(STATUS "Auto-detected NVIDIA environment.") + else() + message(STATUS "NVIDIA environment not detected.") + endif() + endif() # ========================================================= @@ -308,25 +311,25 @@ if(AUTO_DETECT_BACKENDS) endif() # Detect NCCL Dependencies - if(WITH_NVIDIA OR WITH_ALI OR WITH_ILUVATAR OR WITH_HYGON) + if(WITH_NVIDIA OR WITH_THEAD OR WITH_ILUVATAR OR WITH_HYGON) set(_NCCL_HINTS) set(_NCCL_HEADER_NAMES nccl.h) set(_NCCL_LIBRARY_NAMES nccl) - if(WITH_ALI) - foreach(_ali_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) - if(DEFINED ENV{${_ali_cuda_env}} AND NOT "$ENV{${_ali_cuda_env}}" STREQUAL "") - list(APPEND _NCCL_HINTS "$ENV{${_ali_cuda_env}}") + if(WITH_THEAD) + foreach(_thead_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) + if(DEFINED ENV{${_thead_cuda_env}} AND NOT "$ENV{${_thead_cuda_env}}" STREQUAL "") + list(APPEND _NCCL_HINTS "$ENV{${_thead_cuda_env}}") endif() endforeach() - foreach(_ali_env PPU_SDK PPU_HOME PPU_PATH) - if(DEFINED ENV{${_ali_env}} AND NOT "$ENV{${_ali_env}}" STREQUAL "") - list(APPEND _NCCL_HINTS "$ENV{${_ali_env}}/CUDA_SDK" "$ENV{${_ali_env}}") + foreach(_thead_env PPU_SDK PPU_HOME PPU_PATH) + if(DEFINED ENV{${_thead_env}} AND NOT "$ENV{${_thead_env}}" STREQUAL "") + list(APPEND _NCCL_HINTS "$ENV{${_thead_env}}/CUDA_SDK" "$ENV{${_thead_env}}") endif() endforeach() list(APPEND _NCCL_HINTS /usr/local/PPU_SDK/CUDA_SDK /usr/local/PPU_SDK) list(REMOVE_DUPLICATES _NCCL_HINTS) - message(STATUS "Ali PPU detected. Searching for NCCL in ${_NCCL_HINTS}") + message(STATUS "T-Head detected. Searching for NCCL in ${_NCCL_HINTS}") elseif(WITH_ILUVATAR) set(_NCCL_HINTS /usr/local/corex) message(STATUS "Iluvatar detected. Searching for NCCL in ${_NCCL_HINTS}") @@ -429,34 +432,34 @@ endif() # ========================================================= # --- DEPENDENCY DISCOVERY --- # ========================================================= -if(WITH_ALI) - set(_ALI_CUDA_HINTS) - foreach(_ali_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) - if(DEFINED ENV{${_ali_cuda_env}} AND NOT "$ENV{${_ali_cuda_env}}" STREQUAL "") - list(APPEND _ALI_CUDA_HINTS "$ENV{${_ali_cuda_env}}") +if(WITH_THEAD) + set(_THEAD_CUDA_HINTS) + foreach(_thead_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) + if(DEFINED ENV{${_thead_cuda_env}} AND NOT "$ENV{${_thead_cuda_env}}" STREQUAL "") + list(APPEND _THEAD_CUDA_HINTS "$ENV{${_thead_cuda_env}}") endif() endforeach() - foreach(_ali_env PPU_SDK PPU_HOME PPU_PATH) - if(DEFINED ENV{${_ali_env}} AND NOT "$ENV{${_ali_env}}" STREQUAL "") - list(APPEND _ALI_CUDA_HINTS "$ENV{${_ali_env}}/CUDA_SDK" "$ENV{${_ali_env}}") + foreach(_thead_env PPU_SDK PPU_HOME PPU_PATH) + if(DEFINED ENV{${_thead_env}} AND NOT "$ENV{${_thead_env}}" STREQUAL "") + list(APPEND _THEAD_CUDA_HINTS "$ENV{${_thead_env}}/CUDA_SDK" "$ENV{${_thead_env}}") endif() endforeach() - list(APPEND _ALI_CUDA_HINTS /usr/local/PPU_SDK/CUDA_SDK /usr/local/PPU_SDK /usr/local/cuda) - list(REMOVE_DUPLICATES _ALI_CUDA_HINTS) + list(APPEND _THEAD_CUDA_HINTS /usr/local/PPU_SDK/CUDA_SDK /usr/local/PPU_SDK /usr/local/cuda) + list(REMOVE_DUPLICATES _THEAD_CUDA_HINTS) - set(ALI_CUDA_ROOT "") - foreach(_ali_cuda_hint ${_ALI_CUDA_HINTS}) - if(NOT ALI_CUDA_ROOT AND EXISTS "${_ali_cuda_hint}/include/cuda_runtime.h") - set(ALI_CUDA_ROOT "${_ali_cuda_hint}") + set(THEAD_CUDA_ROOT "") + foreach(_thead_cuda_hint ${_THEAD_CUDA_HINTS}) + if(NOT THEAD_CUDA_ROOT AND EXISTS "${_thead_cuda_hint}/include/cuda_runtime.h") + set(THEAD_CUDA_ROOT "${_thead_cuda_hint}") endif() endforeach() - if(NOT ALI_CUDA_ROOT) - message(FATAL_ERROR "`WITH_ALI` is `ON` but the Ali PPU CUDA-compatible SDK was not found. Set `PPU_SDK`, `PPU_HOME`, `CUDA_SDK`, `CUDA_HOME`, or `CUDA_PATH`.") + if(NOT THEAD_CUDA_ROOT) + message(FATAL_ERROR "`WITH_THEAD` is `ON` but the T-Head CUDA-compatible SDK was not found. Set `PPU_SDK`, `PPU_HOME`, `CUDA_SDK`, `CUDA_HOME`, or `CUDA_PATH`.") endif() - set(CUDAToolkit_ROOT "${ALI_CUDA_ROOT}" CACHE PATH "Ali PPU CUDA-compatible toolkit root" FORCE) - list(APPEND _NCCL_HINTS ${_ALI_CUDA_HINTS}) + set(CUDAToolkit_ROOT "${THEAD_CUDA_ROOT}" CACHE PATH "T-Head CUDA-compatible toolkit root" FORCE) + list(APPEND _NCCL_HINTS ${_THEAD_CUDA_HINTS}) list(REMOVE_DUPLICATES _NCCL_HINTS) enable_language(CUDA) @@ -615,23 +618,23 @@ if(WITH_OMPI OR WITH_MPICH) endif() if(WITH_NCCL) - if (NOT WITH_NVIDIA AND NOT WITH_ALI AND NOT WITH_ILUVATAR AND NOT WITH_HYGON) - message(FATAL_ERROR "NCCL backend requires NVIDIA, Ali PPU, Iluvatar, or Hygon GPU support. Please enable `WITH_NVIDIA`, `WITH_ALI`, `WITH_ILUVATAR`, or `WITH_HYGON`.") + if (NOT WITH_NVIDIA AND NOT WITH_THEAD AND NOT WITH_ILUVATAR AND NOT WITH_HYGON) + message(FATAL_ERROR "NCCL backend requires NVIDIA, T-Head, Iluvatar, or Hygon GPU support. Please enable `WITH_NVIDIA`, `WITH_THEAD`, `WITH_ILUVATAR`, or `WITH_HYGON`.") endif() set(_NCCL_HEADER_NAMES nccl.h) set(_NCCL_LIBRARY_NAMES nccl) set(NCCL_COMPILE_DEFINITIONS) - if(WITH_ALI) - foreach(_ali_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) - if(DEFINED ENV{${_ali_cuda_env}} AND NOT "$ENV{${_ali_cuda_env}}" STREQUAL "") - list(APPEND _NCCL_HINTS "$ENV{${_ali_cuda_env}}") + if(WITH_THEAD) + foreach(_thead_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) + if(DEFINED ENV{${_thead_cuda_env}} AND NOT "$ENV{${_thead_cuda_env}}" STREQUAL "") + list(APPEND _NCCL_HINTS "$ENV{${_thead_cuda_env}}") endif() endforeach() - foreach(_ali_env PPU_SDK PPU_HOME PPU_PATH) - if(DEFINED ENV{${_ali_env}} AND NOT "$ENV{${_ali_env}}" STREQUAL "") - list(APPEND _NCCL_HINTS "$ENV{${_ali_env}}/CUDA_SDK" "$ENV{${_ali_env}}") + foreach(_thead_env PPU_SDK PPU_HOME PPU_PATH) + if(DEFINED ENV{${_thead_env}} AND NOT "$ENV{${_thead_env}}" STREQUAL "") + list(APPEND _NCCL_HINTS "$ENV{${_thead_env}}/CUDA_SDK" "$ENV{${_thead_env}}") endif() endforeach() list(APPEND _NCCL_HINTS /usr/local/PPU_SDK/CUDA_SDK /usr/local/PPU_SDK) diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 29741d9..266a548 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -26,7 +26,7 @@ foreach(source_file ${EXAMPLE_SOURCES}) target_link_libraries(${target_name} PRIVATE CUDA::cudart) endif() - if(WITH_ALI) + if(WITH_THEAD) target_link_libraries(${target_name} PRIVATE CUDA::cudart) endif() diff --git a/scripts/gen_bridge.py b/scripts/gen_bridge.py index 59d4dd2..4ea917e 100644 --- a/scripts/gen_bridge.py +++ b/scripts/gen_bridge.py @@ -47,6 +47,11 @@ "cncl": "backends/ccl/cncl", } +# Map public device names to their C++ enum names. +DEVICE_ENUM_MAP = { + "t-head": "THead", +} + # ================================================================= # LOGIC # ================================================================= @@ -131,7 +136,7 @@ def generate(project_root, output_dir, devices, backends): break if device_included: - found_devices.append(f"Device::Type::k{dev.capitalize()}") + found_devices.append(f"Device::Type::k{DEVICE_ENUM_MAP.get(dev, dev.capitalize())}") # Process Active Backends for bb in backends: @@ -143,7 +148,8 @@ def generate(project_root, output_dir, devices, backends): provider_root = CCL_PROVIDER_BACKENDS.get(bb) if provider_root: for dev in devices: - provider_path = f"{provider_root}/{dev}/api.h" + provider_dir = dev + provider_path = f"{provider_root}/{provider_dir}/api.h" if os.path.exists(os.path.join(src_dir, provider_path)): manifest_lines.append(f'#include "{provider_path}"') diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index bfda536..16a8bd4 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -68,21 +68,21 @@ if(WITH_NVIDIA) ) endif() -# Ali PPU -if(WITH_ALI) - list(APPEND DEVICE_LIST "ali") +# T-Head +if(WITH_THEAD) + list(APPEND DEVICE_LIST "t-head") - set(ALI_PATTERNS + set(THEAD_PATTERNS "devices/cuda/*.cc" "devices/cuda/*.cpp" "devices/cuda/*.cu" - "devices/ali/*.cc" - "devices/ali/*.cpp" - "devices/ali/*.cu" + "devices/t-head/*.cc" + "devices/t-head/*.cpp" + "devices/t-head/*.cu" ) - file(GLOB ALI_SOURCES ${ALI_PATTERNS}) + file(GLOB THEAD_SOURCES ${THEAD_PATTERNS}) - target_sources(infiniccl PRIVATE ${ALI_SOURCES}) + target_sources(infiniccl PRIVATE ${THEAD_SOURCES}) target_link_libraries(infiniccl PRIVATE CUDA::cudart CUDA::cuda_driver) set_target_properties(infiniccl PROPERTIES diff --git a/src/backend_device_map.h b/src/backend_device_map.h index 51ac898..8e94048 100644 --- a/src/backend_device_map.h +++ b/src/backend_device_map.h @@ -18,7 +18,7 @@ struct IsSupportedCombination : std::true_type {}; template <> -struct IsSupportedCombination +struct IsSupportedCombination : std::true_type {}; template <> diff --git a/src/backends/ccl/nccl/ali/api.h b/src/backends/ccl/nccl/ali/api.h deleted file mode 100644 index 92cfbbc..0000000 --- a/src/backends/ccl/nccl/ali/api.h +++ /dev/null @@ -1,24 +0,0 @@ -#ifndef INFINI_CCL_BACKENDS_CCL_NCCL_ALI_API_H_ -#define INFINI_CCL_BACKENDS_CCL_NCCL_ALI_API_H_ - -#include "backends/ccl/nccl/api.h" -#include "devices/ali/runtime_.h" - -namespace infini::ccl { - -template <> -struct NcclDataTypeTraits { -#if defined(__CUDA_BF16_TYPES_EXIST__) - static constexpr ncclDataType_t kBFloat16 = ncclBfloat16; -#else - static constexpr ncclDataType_t kBFloat16 = ncclNumTypes; -#endif -}; - -template <> -struct CclApi - : NcclApi {}; - -} // namespace infini::ccl - -#endif // INFINI_CCL_BACKENDS_CCL_NCCL_ALI_API_H_ diff --git a/src/backends/ccl/nccl/t-head/api.h b/src/backends/ccl/nccl/t-head/api.h new file mode 100644 index 0000000..e2b47ad --- /dev/null +++ b/src/backends/ccl/nccl/t-head/api.h @@ -0,0 +1,24 @@ +#ifndef INFINI_CCL_BACKENDS_CCL_NCCL_T_HEAD_API_H_ +#define INFINI_CCL_BACKENDS_CCL_NCCL_T_HEAD_API_H_ + +#include "backends/ccl/nccl/api.h" +#include "devices/t-head/runtime_.h" + +namespace infini::ccl { + +template <> +struct NcclDataTypeTraits { +#if defined(__CUDA_BF16_TYPES_EXIST__) + static constexpr ncclDataType_t kBFloat16 = ncclBfloat16; +#else + static constexpr ncclDataType_t kBFloat16 = ncclNumTypes; +#endif +}; + +template <> +struct CclApi + : NcclApi {}; + +} // namespace infini::ccl + +#endif // INFINI_CCL_BACKENDS_CCL_NCCL_T_HEAD_API_H_ diff --git a/src/device.h b/src/device.h index 4dbf0de..2d65196 100644 --- a/src/device.h +++ b/src/device.h @@ -22,7 +22,7 @@ class Device { kKunlun = 7, kHygon = 8, kQy = 9, - kAli = 10, + kTHead = 10, kCount }; @@ -68,7 +68,7 @@ class Device { {Type::kKunlun, "kunlun"}, {Type::kHygon, "hygon"}, {Type::kQy, "qy"}, - {Type::kAli, "ali"}, + {Type::kTHead, "t-head"}, }}}; static constexpr ConstexprMap; + Device::Type::kTHead, Device::Type::kQy>; // Deferred computation of active devices. The `Filter` and `FilterList` // evaluation are nested inside a class template so that `DeviceEnabled` @@ -165,7 +165,7 @@ struct DevicePriority { }; template <> -struct DevicePriority { +struct DevicePriority { static constexpr int value = 5; }; diff --git a/src/devices/ali/data_type_.h b/src/devices/ali/data_type_.h deleted file mode 100644 index b4600f5..0000000 --- a/src/devices/ali/data_type_.h +++ /dev/null @@ -1,26 +0,0 @@ -#ifndef INFINI_CCL_DEVICES_ALI_DATA_TYPE_H_ -#define INFINI_CCL_DEVICES_ALI_DATA_TYPE_H_ - -// clang-format off -#include -#include -// clang-format on - -#include "data_type_impl.h" -#include "devices/ali/device_.h" - -namespace infini::ccl { - -template <> -struct TypeMap { - using type = half; -}; - -template <> -struct TypeMap { - using type = __nv_bfloat16; -}; - -} // namespace infini::ccl - -#endif // INFINI_CCL_DEVICES_ALI_DATA_TYPE_H_ diff --git a/src/devices/ali/caster_.cuh b/src/devices/t-head/caster_.cuh similarity index 58% rename from src/devices/ali/caster_.cuh rename to src/devices/t-head/caster_.cuh index 18e8cd7..7ac5c30 100644 --- a/src/devices/ali/caster_.cuh +++ b/src/devices/t-head/caster_.cuh @@ -1,5 +1,5 @@ -#ifndef INFINI_CCL_DEVICES_ALI_CASTER_CUH_ -#define INFINI_CCL_DEVICES_ALI_CASTER_CUH_ +#ifndef INFINI_CCL_DEVICES_T_HEAD_CASTER_CUH_ +#define INFINI_CCL_DEVICES_T_HEAD_CASTER_CUH_ #include "caster.h" #include "data_type_.h" @@ -7,58 +7,58 @@ namespace infini::ccl { template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static float Apply(half x) { return __half2float(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static half Apply(float x) { return __float2half(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static float Apply(__nv_bfloat16 x) { return __bfloat162float(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static __nv_bfloat16 Apply(float x) { return __float2bfloat16(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static __nv_bfloat16 Apply(int x) { return __int2bfloat16_rn(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static half Apply(int x) { return __int2half_rn(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static __nv_bfloat16 Apply(double x) { return __double2bfloat16(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static half Apply(double x) { return __double2half(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static half Apply(__nv_bfloat16 x) { return __half(x); } }; } // namespace infini::ccl -#endif // INFINI_CCL_DEVICES_ALI_CASTER_CUH_ +#endif // INFINI_CCL_DEVICES_T_HEAD_CASTER_CUH_ diff --git a/src/devices/t-head/data_type_.h b/src/devices/t-head/data_type_.h new file mode 100644 index 0000000..6175967 --- /dev/null +++ b/src/devices/t-head/data_type_.h @@ -0,0 +1,26 @@ +#ifndef INFINI_CCL_DEVICES_T_HEAD_DATA_TYPE_H_ +#define INFINI_CCL_DEVICES_T_HEAD_DATA_TYPE_H_ + +// clang-format off +#include +#include +// clang-format on + +#include "data_type_impl.h" +#include "devices/t-head/device_.h" + +namespace infini::ccl { + +template <> +struct TypeMap { + using type = half; +}; + +template <> +struct TypeMap { + using type = __nv_bfloat16; +}; + +} // namespace infini::ccl + +#endif // INFINI_CCL_DEVICES_T_HEAD_DATA_TYPE_H_ diff --git a/src/devices/ali/device_.h b/src/devices/t-head/device_.h similarity index 62% rename from src/devices/ali/device_.h rename to src/devices/t-head/device_.h index 86a5519..c845754 100644 --- a/src/devices/ali/device_.h +++ b/src/devices/t-head/device_.h @@ -1,5 +1,5 @@ -#ifndef INFINI_CCL_DEVICES_ALI_DEVICE_H_ -#define INFINI_CCL_DEVICES_ALI_DEVICE_H_ +#ifndef INFINI_CCL_DEVICES_T_HEAD_DEVICE_H_ +#define INFINI_CCL_DEVICES_T_HEAD_DEVICE_H_ #include @@ -9,10 +9,10 @@ namespace infini::ccl { template <> -struct DeviceEnabled : std::true_type {}; +struct DeviceEnabled : std::true_type {}; template <> -MemorySpace GetMemorySpace(const void* ptr) { +MemorySpace GetMemorySpace(const void* ptr) { if (!ptr) { return MemorySpace::kHost; } @@ -29,4 +29,4 @@ MemorySpace GetMemorySpace(const void* ptr) { } // namespace infini::ccl -#endif // INFINI_CCL_DEVICES_ALI_DEVICE_H_ +#endif // INFINI_CCL_DEVICES_T_HEAD_DEVICE_H_ diff --git a/src/devices/ali/runtime_.h b/src/devices/t-head/runtime_.h similarity index 75% rename from src/devices/ali/runtime_.h rename to src/devices/t-head/runtime_.h index c12066c..fa4e31b 100644 --- a/src/devices/ali/runtime_.h +++ b/src/devices/t-head/runtime_.h @@ -1,5 +1,5 @@ -#ifndef INFINI_CCL_DEVICES_ALI_RUNTIME_H_ -#define INFINI_CCL_DEVICES_ALI_RUNTIME_H_ +#ifndef INFINI_CCL_DEVICES_T_HEAD_RUNTIME_H_ +#define INFINI_CCL_DEVICES_T_HEAD_RUNTIME_H_ #include @@ -7,7 +7,7 @@ #include // clang-format on -#include "devices/ali/device_.h" +#include "devices/t-head/device_.h" #include "devices/cuda/runtime_.h" #include "logging.h" #include "return_status_impl.h" @@ -15,10 +15,10 @@ namespace infini::ccl { template <> -struct Runtime : CudaRuntime> { +struct Runtime : CudaRuntime> { using Stream = cudaStream_t; - static constexpr Device::Type kDeviceType = Device::Type::kAli; + static constexpr Device::Type kDeviceType = Device::Type::kTHead; static constexpr auto Check = [](auto status, ReturnStatus err_code = ReturnStatus::kSystemError) { @@ -52,8 +52,8 @@ struct Runtime : CudaRuntime> { static constexpr auto StreamSynchronize = cudaStreamSynchronize; }; -static_assert(Runtime::Validate()); +static_assert(Runtime::Validate()); } // namespace infini::ccl -#endif // INFINI_CCL_DEVICES_ALI_RUNTIME_H_ +#endif // INFINI_CCL_DEVICES_T_HEAD_RUNTIME_H_ From e802903e128ed0d8d66ada4c121f2a708ae82079 Mon Sep 17 00:00:00 2001 From: Zimin Li Date: Sat, 10 Oct 2026 18:35:36 +0800 Subject: [PATCH 4/9] refactor: simplify T-Head dependency discovery Collect T-Head SDK hints lazily and reuse them across device, NCCL, and toolkit discovery. Keep NCCL search suffixes scoped to their platform-specific paths. --- CMakeLists.txt | 105 +++++++++++++++++++++++++------------------------ 1 file changed, 53 insertions(+), 52 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index c4a393f..f06f5d1 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -33,6 +33,38 @@ option(AUTO_DETECT_DEVICES "Automatically detect available devices" ON) option(AUTO_DETECT_BACKENDS "Automatically detect available backends" OFF) option(BUILD_EXAMPLES "Build internal examples" ON) +function(_collect_thead_hints out_var) + if(NOT DEFINED _THEAD_HINTS_MEMO) + set(_hints + /usr/local/PPU_SDK + /usr/local/PPU_SDK/CUDA_SDK + ) + + foreach(_env IN ITEMS PPU_SDK PPU_HOME PPU_PATH) + if(DEFINED ENV{${_env}} AND NOT "$ENV{${_env}}" STREQUAL "") + list(APPEND _hints + "$ENV{${_env}}" + "$ENV{${_env}}/CUDA_SDK" + ) + endif() + endforeach() + + foreach(_env IN ITEMS CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) + if(DEFINED ENV{${_env}} AND NOT "$ENV{${_env}}" STREQUAL "") + list(APPEND _hints + "$ENV{${_env}}" + "$ENV{${_env}}/.." + ) + endif() + endforeach() + + list(REMOVE_DUPLICATES _hints) + set(_THEAD_HINTS_MEMO "${_hints}" PARENT_SCOPE) + endif() + + set(${out_var} "${_THEAD_HINTS_MEMO}" PARENT_SCOPE) +endfunction() + # ========================================================= # --- AUTO-DETECTION: DEVICES --- # ========================================================= @@ -218,20 +250,8 @@ if(AUTO_DETECT_DEVICES) if(WITH_THEAD) set(THEAD_FOUND TRUE) else() - set(_THEAD_SDK_HINTS) - foreach(_thead_env PPU_SDK PPU_HOME PPU_PATH) - if(DEFINED ENV{${_thead_env}} AND NOT "$ENV{${_thead_env}}" STREQUAL "") - list(APPEND _THEAD_SDK_HINTS "$ENV{${_thead_env}}") - endif() - endforeach() - foreach(_thead_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) - if(DEFINED ENV{${_thead_cuda_env}} AND NOT "$ENV{${_thead_cuda_env}}" STREQUAL "") - list(APPEND _THEAD_SDK_HINTS "$ENV{${_thead_cuda_env}}" "$ENV{${_thead_cuda_env}}/..") - endif() - endforeach() - list(APPEND _THEAD_SDK_HINTS /usr/local/PPU_SDK /usr/local/PPU_SDK/CUDA_SDK) - - foreach(_thead_hint ${_THEAD_SDK_HINTS}) + _collect_thead_hints(_THEAD_HINTS) + foreach(_thead_hint IN LISTS _THEAD_HINTS) if(EXISTS "${_thead_hint}/CUDA_SDK" OR EXISTS "${_thead_hint}/bin/ppu-smi" OR EXISTS "${_thead_hint}/ppu-smi/bin/ppu-smi") set(THEAD_FOUND TRUE) endif() @@ -315,19 +335,14 @@ if(AUTO_DETECT_BACKENDS) set(_NCCL_HINTS) set(_NCCL_HEADER_NAMES nccl.h) set(_NCCL_LIBRARY_NAMES nccl) + set(_NCCL_INCLUDE_SUFFIXES include) + set(_NCCL_LIBRARY_SUFFIXES lib lib64) if(WITH_THEAD) - foreach(_thead_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) - if(DEFINED ENV{${_thead_cuda_env}} AND NOT "$ENV{${_thead_cuda_env}}" STREQUAL "") - list(APPEND _NCCL_HINTS "$ENV{${_thead_cuda_env}}") - endif() - endforeach() - foreach(_thead_env PPU_SDK PPU_HOME PPU_PATH) - if(DEFINED ENV{${_thead_env}} AND NOT "$ENV{${_thead_env}}" STREQUAL "") - list(APPEND _NCCL_HINTS "$ENV{${_thead_env}}/CUDA_SDK" "$ENV{${_thead_env}}") - endif() - endforeach() - list(APPEND _NCCL_HINTS /usr/local/PPU_SDK/CUDA_SDK /usr/local/PPU_SDK) + _collect_thead_hints(_THEAD_HINTS) + set(_NCCL_HINTS ${_THEAD_HINTS}) + list(APPEND _NCCL_INCLUDE_SUFFIXES targets/x86_64-linux/include) + list(APPEND _NCCL_LIBRARY_SUFFIXES targets/x86_64-linux/lib) list(REMOVE_DUPLICATES _NCCL_HINTS) message(STATUS "T-Head detected. Searching for NCCL in ${_NCCL_HINTS}") elseif(WITH_ILUVATAR) @@ -348,16 +363,17 @@ if(AUTO_DETECT_BACKENDS) list(APPEND _NCCL_HINTS /opt/dtk /opt/dtk/rccl) set(_NCCL_HEADER_NAMES rccl/rccl.h rccl.h) set(_NCCL_LIBRARY_NAMES rccl) + list(APPEND _NCCL_LIBRARY_SUFFIXES rccl/lib) message(STATUS "Hygon detected. Searching for NCCL-compatible RCCL in ${_NCCL_HINTS}") endif() find_path(AUTO_NCCL_INC NAMES ${_NCCL_HEADER_NAMES} HINTS ${_NCCL_HINTS} - PATH_SUFFIXES include targets/x86_64-linux/include + PATH_SUFFIXES ${_NCCL_INCLUDE_SUFFIXES} QUIET) find_library(AUTO_NCCL_LIB NAMES ${_NCCL_LIBRARY_NAMES} HINTS ${_NCCL_HINTS} - PATH_SUFFIXES lib lib64 targets/x86_64-linux/lib rccl/lib + PATH_SUFFIXES ${_NCCL_LIBRARY_SUFFIXES} QUIET) if(AUTO_NCCL_INC AND AUTO_NCCL_LIB) @@ -433,18 +449,8 @@ endif() # --- DEPENDENCY DISCOVERY --- # ========================================================= if(WITH_THEAD) - set(_THEAD_CUDA_HINTS) - foreach(_thead_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) - if(DEFINED ENV{${_thead_cuda_env}} AND NOT "$ENV{${_thead_cuda_env}}" STREQUAL "") - list(APPEND _THEAD_CUDA_HINTS "$ENV{${_thead_cuda_env}}") - endif() - endforeach() - foreach(_thead_env PPU_SDK PPU_HOME PPU_PATH) - if(DEFINED ENV{${_thead_env}} AND NOT "$ENV{${_thead_env}}" STREQUAL "") - list(APPEND _THEAD_CUDA_HINTS "$ENV{${_thead_env}}/CUDA_SDK" "$ENV{${_thead_env}}") - endif() - endforeach() - list(APPEND _THEAD_CUDA_HINTS /usr/local/PPU_SDK/CUDA_SDK /usr/local/PPU_SDK /usr/local/cuda) + _collect_thead_hints(_THEAD_CUDA_HINTS) + list(APPEND _THEAD_CUDA_HINTS /usr/local/cuda) list(REMOVE_DUPLICATES _THEAD_CUDA_HINTS) set(THEAD_CUDA_ROOT "") @@ -625,34 +631,29 @@ if(WITH_NCCL) set(_NCCL_HEADER_NAMES nccl.h) set(_NCCL_LIBRARY_NAMES nccl) set(NCCL_COMPILE_DEFINITIONS) + set(_NCCL_INCLUDE_SUFFIXES include) + set(_NCCL_LIBRARY_SUFFIXES lib lib64) if(WITH_THEAD) - foreach(_thead_cuda_env CUDA_SDK CUDA_HOME CUDA_PATH CUDA_TOOLKIT_ROOT) - if(DEFINED ENV{${_thead_cuda_env}} AND NOT "$ENV{${_thead_cuda_env}}" STREQUAL "") - list(APPEND _NCCL_HINTS "$ENV{${_thead_cuda_env}}") - endif() - endforeach() - foreach(_thead_env PPU_SDK PPU_HOME PPU_PATH) - if(DEFINED ENV{${_thead_env}} AND NOT "$ENV{${_thead_env}}" STREQUAL "") - list(APPEND _NCCL_HINTS "$ENV{${_thead_env}}/CUDA_SDK" "$ENV{${_thead_env}}") - endif() - endforeach() - list(APPEND _NCCL_HINTS /usr/local/PPU_SDK/CUDA_SDK /usr/local/PPU_SDK) + set(_NCCL_HINTS ${_THEAD_CUDA_HINTS}) + list(APPEND _NCCL_INCLUDE_SUFFIXES targets/x86_64-linux/include) + list(APPEND _NCCL_LIBRARY_SUFFIXES targets/x86_64-linux/lib) elseif(WITH_HYGON) list(APPEND _NCCL_HINTS "${HYGON_DTK_ROOT}" "${HYGON_DTK_ROOT}/rccl") set(_NCCL_HEADER_NAMES rccl/rccl.h rccl.h) set(_NCCL_LIBRARY_NAMES rccl) set(NCCL_COMPILE_DEFINITIONS INFINI_CCL_USE_RCCL) + list(APPEND _NCCL_LIBRARY_SUFFIXES rccl/lib) endif() list(REMOVE_DUPLICATES _NCCL_HINTS) find_library(NCCL_LIB NAMES ${_NCCL_LIBRARY_NAMES} HINTS ${_NCCL_HINTS} - PATH_SUFFIXES lib lib64 targets/x86_64-linux/lib rccl/lib + PATH_SUFFIXES ${_NCCL_LIBRARY_SUFFIXES} REQUIRED) find_path(NCCL_INC NAMES ${_NCCL_HEADER_NAMES} HINTS ${_NCCL_HINTS} - PATH_SUFFIXES include targets/x86_64-linux/include + PATH_SUFFIXES ${_NCCL_INCLUDE_SUFFIXES} REQUIRED) include_directories(${NCCL_INC}) From c3261e71243799cd70f4cb38aaef2394e1a86362 Mon Sep 17 00:00:00 2001 From: Zimin Li Date: Tue, 13 Oct 2026 15:36:01 +0800 Subject: [PATCH 5/9] refactor: normalize T-Head device identifier Use `thead` as the public device name and source-directory identifier, derive `kThead` through standard bridge generation, and remove the T-Head enum mapping. Also retain the memoized T-Head hints in function scope so first-use backend detection resolves SDK paths correctly. --- CMakeLists.txt | 1 + scripts/gen_bridge.py | 7 +----- src/CMakeLists.txt | 8 +++---- src/backend_device_map.h | 2 +- src/backends/ccl/nccl/t-head/api.h | 24 --------------------- src/backends/ccl/nccl/thead/api.h | 24 +++++++++++++++++++++ src/device.h | 10 ++++----- src/devices/t-head/data_type_.h | 26 ----------------------- src/devices/{t-head => thead}/caster_.cuh | 24 ++++++++++----------- src/devices/thead/data_type_.h | 26 +++++++++++++++++++++++ src/devices/{t-head => thead}/device_.h | 10 ++++----- src/devices/{t-head => thead}/runtime_.h | 14 ++++++------ 12 files changed, 86 insertions(+), 90 deletions(-) delete mode 100644 src/backends/ccl/nccl/t-head/api.h create mode 100644 src/backends/ccl/nccl/thead/api.h delete mode 100644 src/devices/t-head/data_type_.h rename src/devices/{t-head => thead}/caster_.cuh (59%) create mode 100644 src/devices/thead/data_type_.h rename src/devices/{t-head => thead}/device_.h (63%) rename src/devices/{t-head => thead}/runtime_.h (76%) diff --git a/CMakeLists.txt b/CMakeLists.txt index f06f5d1..f052c92 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -59,6 +59,7 @@ function(_collect_thead_hints out_var) endforeach() list(REMOVE_DUPLICATES _hints) + set(_THEAD_HINTS_MEMO "${_hints}") set(_THEAD_HINTS_MEMO "${_hints}" PARENT_SCOPE) endif() diff --git a/scripts/gen_bridge.py b/scripts/gen_bridge.py index 4ea917e..fc52e88 100644 --- a/scripts/gen_bridge.py +++ b/scripts/gen_bridge.py @@ -47,11 +47,6 @@ "cncl": "backends/ccl/cncl", } -# Map public device names to their C++ enum names. -DEVICE_ENUM_MAP = { - "t-head": "THead", -} - # ================================================================= # LOGIC # ================================================================= @@ -136,7 +131,7 @@ def generate(project_root, output_dir, devices, backends): break if device_included: - found_devices.append(f"Device::Type::k{DEVICE_ENUM_MAP.get(dev, dev.capitalize())}") + found_devices.append(f"Device::Type::k{dev.capitalize()}") # Process Active Backends for bb in backends: diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 16a8bd4..600629f 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -70,15 +70,15 @@ endif() # T-Head if(WITH_THEAD) - list(APPEND DEVICE_LIST "t-head") + list(APPEND DEVICE_LIST "thead") set(THEAD_PATTERNS "devices/cuda/*.cc" "devices/cuda/*.cpp" "devices/cuda/*.cu" - "devices/t-head/*.cc" - "devices/t-head/*.cpp" - "devices/t-head/*.cu" + "devices/thead/*.cc" + "devices/thead/*.cpp" + "devices/thead/*.cu" ) file(GLOB THEAD_SOURCES ${THEAD_PATTERNS}) diff --git a/src/backend_device_map.h b/src/backend_device_map.h index 8e94048..d7e365c 100644 --- a/src/backend_device_map.h +++ b/src/backend_device_map.h @@ -18,7 +18,7 @@ struct IsSupportedCombination : std::true_type {}; template <> -struct IsSupportedCombination +struct IsSupportedCombination : std::true_type {}; template <> diff --git a/src/backends/ccl/nccl/t-head/api.h b/src/backends/ccl/nccl/t-head/api.h deleted file mode 100644 index e2b47ad..0000000 --- a/src/backends/ccl/nccl/t-head/api.h +++ /dev/null @@ -1,24 +0,0 @@ -#ifndef INFINI_CCL_BACKENDS_CCL_NCCL_T_HEAD_API_H_ -#define INFINI_CCL_BACKENDS_CCL_NCCL_T_HEAD_API_H_ - -#include "backends/ccl/nccl/api.h" -#include "devices/t-head/runtime_.h" - -namespace infini::ccl { - -template <> -struct NcclDataTypeTraits { -#if defined(__CUDA_BF16_TYPES_EXIST__) - static constexpr ncclDataType_t kBFloat16 = ncclBfloat16; -#else - static constexpr ncclDataType_t kBFloat16 = ncclNumTypes; -#endif -}; - -template <> -struct CclApi - : NcclApi {}; - -} // namespace infini::ccl - -#endif // INFINI_CCL_BACKENDS_CCL_NCCL_T_HEAD_API_H_ diff --git a/src/backends/ccl/nccl/thead/api.h b/src/backends/ccl/nccl/thead/api.h new file mode 100644 index 0000000..14dbd2b --- /dev/null +++ b/src/backends/ccl/nccl/thead/api.h @@ -0,0 +1,24 @@ +#ifndef INFINI_CCL_BACKENDS_CCL_NCCL_THEAD_API_H_ +#define INFINI_CCL_BACKENDS_CCL_NCCL_THEAD_API_H_ + +#include "backends/ccl/nccl/api.h" +#include "devices/thead/runtime_.h" + +namespace infini::ccl { + +template <> +struct NcclDataTypeTraits { +#if defined(__CUDA_BF16_TYPES_EXIST__) + static constexpr ncclDataType_t kBFloat16 = ncclBfloat16; +#else + static constexpr ncclDataType_t kBFloat16 = ncclNumTypes; +#endif +}; + +template <> +struct CclApi + : NcclApi {}; + +} // namespace infini::ccl + +#endif // INFINI_CCL_BACKENDS_CCL_NCCL_THEAD_API_H_ diff --git a/src/device.h b/src/device.h index 2d65196..71bb3a9 100644 --- a/src/device.h +++ b/src/device.h @@ -22,7 +22,7 @@ class Device { kKunlun = 7, kHygon = 8, kQy = 9, - kTHead = 10, + kThead = 10, kCount }; @@ -68,7 +68,7 @@ class Device { {Type::kKunlun, "kunlun"}, {Type::kHygon, "hygon"}, {Type::kQy, "qy"}, - {Type::kTHead, "t-head"}, + {Type::kThead, "thead"}, }}}; static constexpr ConstexprMap; + Device::Type::kThead, Device::Type::kQy>; // Deferred computation of active devices. The `Filter` and `FilterList` // evaluation are nested inside a class template so that `DeviceEnabled` @@ -165,7 +165,7 @@ struct DevicePriority { }; template <> -struct DevicePriority { +struct DevicePriority { static constexpr int value = 5; }; diff --git a/src/devices/t-head/data_type_.h b/src/devices/t-head/data_type_.h deleted file mode 100644 index 6175967..0000000 --- a/src/devices/t-head/data_type_.h +++ /dev/null @@ -1,26 +0,0 @@ -#ifndef INFINI_CCL_DEVICES_T_HEAD_DATA_TYPE_H_ -#define INFINI_CCL_DEVICES_T_HEAD_DATA_TYPE_H_ - -// clang-format off -#include -#include -// clang-format on - -#include "data_type_impl.h" -#include "devices/t-head/device_.h" - -namespace infini::ccl { - -template <> -struct TypeMap { - using type = half; -}; - -template <> -struct TypeMap { - using type = __nv_bfloat16; -}; - -} // namespace infini::ccl - -#endif // INFINI_CCL_DEVICES_T_HEAD_DATA_TYPE_H_ diff --git a/src/devices/t-head/caster_.cuh b/src/devices/thead/caster_.cuh similarity index 59% rename from src/devices/t-head/caster_.cuh rename to src/devices/thead/caster_.cuh index 7ac5c30..1b262f7 100644 --- a/src/devices/t-head/caster_.cuh +++ b/src/devices/thead/caster_.cuh @@ -1,5 +1,5 @@ -#ifndef INFINI_CCL_DEVICES_T_HEAD_CASTER_CUH_ -#define INFINI_CCL_DEVICES_T_HEAD_CASTER_CUH_ +#ifndef INFINI_CCL_DEVICES_THEAD_CASTER_CUH_ +#define INFINI_CCL_DEVICES_THEAD_CASTER_CUH_ #include "caster.h" #include "data_type_.h" @@ -7,58 +7,58 @@ namespace infini::ccl { template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static float Apply(half x) { return __half2float(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static half Apply(float x) { return __float2half(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static float Apply(__nv_bfloat16 x) { return __bfloat162float(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static __nv_bfloat16 Apply(float x) { return __float2bfloat16(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static __nv_bfloat16 Apply(int x) { return __int2bfloat16_rn(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static half Apply(int x) { return __int2half_rn(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static __nv_bfloat16 Apply(double x) { return __double2bfloat16(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static half Apply(double x) { return __double2half(x); } }; template <> -struct HardwareCastImpl { +struct HardwareCastImpl { __host__ __device__ static half Apply(__nv_bfloat16 x) { return __half(x); } }; } // namespace infini::ccl -#endif // INFINI_CCL_DEVICES_T_HEAD_CASTER_CUH_ +#endif // INFINI_CCL_DEVICES_THEAD_CASTER_CUH_ diff --git a/src/devices/thead/data_type_.h b/src/devices/thead/data_type_.h new file mode 100644 index 0000000..e9b825e --- /dev/null +++ b/src/devices/thead/data_type_.h @@ -0,0 +1,26 @@ +#ifndef INFINI_CCL_DEVICES_THEAD_DATA_TYPE_H_ +#define INFINI_CCL_DEVICES_THEAD_DATA_TYPE_H_ + +// clang-format off +#include +#include +// clang-format on + +#include "data_type_impl.h" +#include "devices/thead/device_.h" + +namespace infini::ccl { + +template <> +struct TypeMap { + using type = half; +}; + +template <> +struct TypeMap { + using type = __nv_bfloat16; +}; + +} // namespace infini::ccl + +#endif // INFINI_CCL_DEVICES_THEAD_DATA_TYPE_H_ diff --git a/src/devices/t-head/device_.h b/src/devices/thead/device_.h similarity index 63% rename from src/devices/t-head/device_.h rename to src/devices/thead/device_.h index c845754..24930fd 100644 --- a/src/devices/t-head/device_.h +++ b/src/devices/thead/device_.h @@ -1,5 +1,5 @@ -#ifndef INFINI_CCL_DEVICES_T_HEAD_DEVICE_H_ -#define INFINI_CCL_DEVICES_T_HEAD_DEVICE_H_ +#ifndef INFINI_CCL_DEVICES_THEAD_DEVICE_H_ +#define INFINI_CCL_DEVICES_THEAD_DEVICE_H_ #include @@ -9,10 +9,10 @@ namespace infini::ccl { template <> -struct DeviceEnabled : std::true_type {}; +struct DeviceEnabled : std::true_type {}; template <> -MemorySpace GetMemorySpace(const void* ptr) { +MemorySpace GetMemorySpace(const void* ptr) { if (!ptr) { return MemorySpace::kHost; } @@ -29,4 +29,4 @@ MemorySpace GetMemorySpace(const void* ptr) { } // namespace infini::ccl -#endif // INFINI_CCL_DEVICES_T_HEAD_DEVICE_H_ +#endif // INFINI_CCL_DEVICES_THEAD_DEVICE_H_ diff --git a/src/devices/t-head/runtime_.h b/src/devices/thead/runtime_.h similarity index 76% rename from src/devices/t-head/runtime_.h rename to src/devices/thead/runtime_.h index fa4e31b..f228a77 100644 --- a/src/devices/t-head/runtime_.h +++ b/src/devices/thead/runtime_.h @@ -1,5 +1,5 @@ -#ifndef INFINI_CCL_DEVICES_T_HEAD_RUNTIME_H_ -#define INFINI_CCL_DEVICES_T_HEAD_RUNTIME_H_ +#ifndef INFINI_CCL_DEVICES_THEAD_RUNTIME_H_ +#define INFINI_CCL_DEVICES_THEAD_RUNTIME_H_ #include @@ -7,7 +7,7 @@ #include // clang-format on -#include "devices/t-head/device_.h" +#include "devices/thead/device_.h" #include "devices/cuda/runtime_.h" #include "logging.h" #include "return_status_impl.h" @@ -15,10 +15,10 @@ namespace infini::ccl { template <> -struct Runtime : CudaRuntime> { +struct Runtime : CudaRuntime> { using Stream = cudaStream_t; - static constexpr Device::Type kDeviceType = Device::Type::kTHead; + static constexpr Device::Type kDeviceType = Device::Type::kThead; static constexpr auto Check = [](auto status, ReturnStatus err_code = ReturnStatus::kSystemError) { @@ -52,8 +52,8 @@ struct Runtime : CudaRuntime static constexpr auto StreamSynchronize = cudaStreamSynchronize; }; -static_assert(Runtime::Validate()); +static_assert(Runtime::Validate()); } // namespace infini::ccl -#endif // INFINI_CCL_DEVICES_T_HEAD_RUNTIME_H_ +#endif // INFINI_CCL_DEVICES_THEAD_RUNTIME_H_ From ed2fb49ea7e23311408b4255e8e5413a558282be Mon Sep 17 00:00:00 2001 From: Zimin Li Date: Wed, 14 Oct 2026 10:09:17 +0000 Subject: [PATCH 6/9] docs: update documentation to include T-Head as a supported platform --- .github/pull_request_template.md | 2 ++ README.md | 6 ++++-- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/.github/pull_request_template.md b/.github/pull_request_template.md index 326ede4..bb23ef0 100644 --- a/.github/pull_request_template.md +++ b/.github/pull_request_template.md @@ -59,6 +59,7 @@ Please check all the platforms and/or backends this PR affects (i.e., code is to - [ ] Moore Threads GPU - [ ] Cambricon MLU - [ ] HYGON DCU +- [ ] T-Head PPU ### Backend @@ -110,6 +111,7 @@ See `CONTRIBUTING.md` § Pull Requests for the official testing requirements and - [ ] Moore Threads GPU - [ ] Cambricon MLU - [ ] HYGON DCU +- [ ] T-Head PPU ### Test Involved Backend diff --git a/README.md b/README.md index 0158be0..3ffeb67 100644 --- a/README.md +++ b/README.md @@ -138,6 +138,7 @@ cmake .. -DWITH_NVIDIA=ON -DWITH_OMPI=ON | `WITH_MOORE` | Enable Moore Threads GPU support | `OFF` | | `WITH_CAMBRICON` | Enable Cambricon MLU support | `OFF` | | `WITH_HYGON` | Enable HYGON DCU support | `OFF` | +| `WITH_THEAD` | Enable T-Head PPU support | `OFF` | | `WITH_CPU` | CPU support (always enabled) | `ON` (internal, not user‑settable) | | **Backend (Communication) Options** ||| | `WITH_OMPI` | Enable OpenMPI backend | `ON` if no backend specified, otherwise `OFF` | @@ -189,7 +190,7 @@ After having a successful build and a complete `cluster.yaml`, we are ready for | `nodes[].ip` | Yes | Node | Node IP address or hostname. Use `localhost` or `127.0.0.1` for the local node. | | `nodes[].user` | No | Node | SSH user for this node. Overrides `common_user`. | | `nodes[].dir` | No | Node | Node-specific project source directory. Overrides `common_dir` for this node and is useful when the project path differs across hosts. | -| `nodes[].type` | Yes | Node | Architecture/build label used in build, install, and wrapper paths. Common values include `cpu`, `nvidia`, `iluvatar`, `metax`, `moore`, `cambricon`, and `hygon`. | +| `nodes[].type` | Yes | Node | Architecture/build label used in build, install, and wrapper paths. Common values include `cpu`, `nvidia`, `iluvatar`, `metax`, `moore`, `cambricon`, `hygon`, and `t-head`. | | `nodes[].slots` | No | Node | Number of processes to launch on this node. This usually matches the number of devices assigned to the node. Defaults to `8`. | | `nodes[].cmake_flags` | No | Node | Node-specific CMake options used during `--build`, such as `-DUSE_CUDA=ON` or `-DUSE_MACA=ON`. Overrides global `cmake_flags` for this node. | | `nodes[].backend_env` | No | Node | Node-specific runtime environment variables, such as `CUDA_VISIBLE_DEVICES`, `UCX_TLS`, or `UCX_NET_DEVICES`. Overrides or prepends to global `backend_env` values for this node. | @@ -346,6 +347,7 @@ export LD_LIBRARY_PATH=${INFINI_INSTALL}/lib:$LD_LIBRARY_PATH | **Moore Threads** | Full | Requires MUSA SDK and at least one of `MACA_ROOT` (default `/usr/local/musa`), `MACA_PATH`, and `MUSA_HOME` to be set. | | **Cambricon** | Full | Requires CNToolKit and `NEUWARE_HOME` to be set. | | **HYGON** | Full | Requires HYGON DTK and HYHAL. | +| **T-Head** | Full | Requires T-Head SDK. | @@ -356,7 +358,7 @@ export LD_LIBRARY_PATH=${INFINI_INSTALL}/lib:$LD_LIBRARY_PATH |---------|---------------|----------------------|---------------| | **OpenMPI** | Full | `WITH_OMPI=ON` | The default backend. Requires the OpenMPI development package.| | **MPICH** | Full | `WITH_MPICH=ON` | Requires the MPICH development package.| -| **NCCL** | Partial | `WITH_NCCL=ON` | Requires NVIDIA or Iluvatar NCCL, or HYGON RCCL. Currently available when `WITH_NVIDIA=ON`, `WITH_ILUVATAR=ON`, or `WITH_HYGON=ON`.| +| **NCCL** | Partial | `WITH_NCCL=ON` | Requires NVIDIA, Iluvatar, or T-Head NCCL, or HYGON RCCL. Currently available when `WITH_NVIDIA=ON`, `WITH_ILUVATAR=ON`, `WITH_HYGON=ON` or `WITH_THEAD=ON`.| | **MCCL** | Partial | `WITH_MCCL=ON` | Requires MetaX or Moore MCCL. Currently available when `WITH_METAX=ON` or `WITH_MOORE=ON`.| | **CNCL** | Partial | `WITH_CNCL=ON` | Requires Cambricon CNCL. Available only when `WITH_CAMBRICON=ON`.| From f4afd56cfe2fe9d48dea7f9e11a3251af77dc4ce Mon Sep 17 00:00:00 2001 From: Zimin Li Date: Wed, 14 Oct 2026 10:51:20 +0000 Subject: [PATCH 7/9] style: clang-format `src/devices/thead/runtime_.h` --- src/devices/thead/runtime_.h | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/src/devices/thead/runtime_.h b/src/devices/thead/runtime_.h index f228a77..3026e0f 100644 --- a/src/devices/thead/runtime_.h +++ b/src/devices/thead/runtime_.h @@ -7,15 +7,16 @@ #include // clang-format on -#include "devices/thead/device_.h" #include "devices/cuda/runtime_.h" +#include "devices/thead/device_.h" #include "logging.h" #include "return_status_impl.h" namespace infini::ccl { template <> -struct Runtime : CudaRuntime> { +struct Runtime + : CudaRuntime> { using Stream = cudaStream_t; static constexpr Device::Type kDeviceType = Device::Type::kThead; From 09195b29f8f6c4bbdecfebc081b6eadbdbed18b2 Mon Sep 17 00:00:00 2001 From: Zimin Li Date: Wed, 14 Oct 2026 19:03:06 +0800 Subject: [PATCH 8/9] fix: add return statement to avoid warning --- examples/ccl/all_gather.cc | 2 ++ 1 file changed, 2 insertions(+) diff --git a/examples/ccl/all_gather.cc b/examples/ccl/all_gather.cc index c3a41b1..a7aeddb 100644 --- a/examples/ccl/all_gather.cc +++ b/examples/ccl/all_gather.cc @@ -67,6 +67,8 @@ void PrintResult(const std::vector& output, const char* mode, int size, std::cout << "[r" << source << ": " << output[offset] << "] "; } std::cout << std::endl; + + return; } void WorkerThread(ThreadArgs args) { From 1be91dd4c2e06f7259368b90d2c1081a0825cfd7 Mon Sep 17 00:00:00 2001 From: Zimin Li Date: Wed, 14 Oct 2026 11:08:37 +0000 Subject: [PATCH 9/9] style: clang-format `examples/ccl/all_gather.cc` --- examples/ccl/all_gather.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/ccl/all_gather.cc b/examples/ccl/all_gather.cc index a7aeddb..fbf51d1 100644 --- a/examples/ccl/all_gather.cc +++ b/examples/ccl/all_gather.cc @@ -67,7 +67,7 @@ void PrintResult(const std::vector& output, const char* mode, int size, std::cout << "[r" << source << ": " << output[offset] << "] "; } std::cout << std::endl; - + return; }