feat: add HYGON platform support across runtime, device, build, and launcher stack

.github/pull_request_template.md

@@ -58,6 +58,7 @@ Please check all the platforms and/or backends this PR affects (i.e., code is to
 - [ ] MetaX GPU
 - [ ] Moore Threads GPU
 - [ ] Cambricon MLU
+- [ ] HYGON DCU
 
 ### Backend
 
@@ -107,6 +108,7 @@ See `CONTRIBUTING.md` § Pull Requests for the official testing requirements and
 - [ ] MetaX GPU
 - [ ] Moore Threads GPU
 - [ ] Cambricon MLU
+- [ ] HYGON DCU
 
 ### Test Involved Backend
 

CMakeLists.txt

@@ -12,6 +12,7 @@ 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)
 
 set(WITH_CPU ON CACHE INTERNAL "CPU backend is always enabled")
 
@@ -176,6 +177,44 @@ if(AUTO_DETECT_DEVICES)
         set(WITH_CAMBRICON OFF)
         message(STATUS "Cambricon environment not detected.")
     endif()
+
+    # Hygon DCU
+    if(NOT WITH_HYGON)
+        set(HYGON_FOUND FALSE)
+
+        if((DEFINED ENV{DTKROOT} OR DEFINED ENV{DTK_ROOT} OR DEFINED ENV{ROCM_PATH})
+           AND EXISTS "/dev/kfd")
+            set(HYGON_FOUND TRUE)
+        elseif(EXISTS "/opt/dtk" AND EXISTS "/dev/kfd")
+            set(HYGON_FOUND TRUE)
+        else()
+            find_program(HYGON_SMI_PATH
+                NAMES hy-smi rocm-smi
+                HINTS /opt/hyhal/bin /opt/dtk/bin
+            )
+            if(HYGON_SMI_PATH)
+                execute_process(
+                    COMMAND ${HYGON_SMI_PATH}
+                    RESULT_VARIABLE SMI_RESULT
+                    OUTPUT_QUIET
+                    ERROR_QUIET
+                )
+                if(SMI_RESULT EQUAL 0)
+                    set(HYGON_FOUND TRUE)
+                endif()
+            endif()
+        endif()
+
+        if(HYGON_FOUND)
+            set(WITH_HYGON ON)
+        endif()
+    endif()
+
+    if(WITH_HYGON)
+        message(STATUS "Hygon environment detected.")
+    else()
+        message(STATUS "Hygon environment not detected.")
+    endif()
 endif()
 
 # =========================================================
@@ -374,6 +413,51 @@ if(WITH_CAMBRICON)
     find_library(CAMBRICON_PAPI_LIB       NAMES cnpapi     HINTS "${NEUWARE_HOME}/lib64" REQUIRED)
 endif()
 
+if(WITH_HYGON)
+    set(HYGON_DTK_ROOT "")
+    foreach(_hygon_env DTKROOT DTK_ROOT ROCM_PATH)
+        if(NOT HYGON_DTK_ROOT AND DEFINED ENV{${_hygon_env}} AND NOT "$ENV{${_hygon_env}}" STREQUAL "")
+            set(HYGON_DTK_ROOT "$ENV{${_hygon_env}}")
+        endif()
+    endforeach()
+
+    if(NOT HYGON_DTK_ROOT AND DEFINED ENV{HIP_PATH} AND NOT "$ENV{HIP_PATH}" STREQUAL "")
+        get_filename_component(HYGON_DTK_ROOT "$ENV{HIP_PATH}/.." ABSOLUTE)
+    endif()
+
+    if(NOT HYGON_DTK_ROOT AND EXISTS "/opt/dtk")
+        set(HYGON_DTK_ROOT "/opt/dtk")
+    endif()
+
+    if(NOT HYGON_DTK_ROOT)
+        message(FATAL_ERROR "`WITH_HYGON` is `ON` but DTK was not found. Set `DTKROOT`, `DTK_ROOT`, `ROCM_PATH`, or `HIP_PATH`.")
+    endif()
+
+    set(HYGON_HIP_HINTS "${HYGON_DTK_ROOT}" "${HYGON_DTK_ROOT}/hip")
+    find_path(HYGON_HIP_INCLUDE_DIR
+        NAMES hip/hip_runtime.h
+        HINTS ${HYGON_HIP_HINTS}
+        PATH_SUFFIXES include hip/include
+        REQUIRED
+    )
+    find_library(HYGON_HIP_RUNTIME_LIB
+        NAMES amdhip64
+        HINTS ${HYGON_HIP_HINTS}
+        PATH_SUFFIXES lib lib64 hip/lib
+        REQUIRED
+    )
+    find_program(HYGON_HIPCC_COMPILER
+        NAMES hipcc
+        HINTS ${HYGON_HIP_HINTS}
+        PATH_SUFFIXES bin hip/bin
+        REQUIRED
+    )
+
+    set(CMAKE_CXX_COMPILER "${HYGON_HIPCC_COMPILER}" CACHE FILEPATH "Hygon DTK hipcc compiler" FORCE)
+    set(HYGON_HIP_DEFINITIONS __HIP_PLATFORM_AMD__)
+    include_directories(${HYGON_HIP_INCLUDE_DIR})
+endif()
+
 if(WITH_OMPI OR WITH_MPICH)
     find_package(MPI REQUIRED)
 endif()

README.md

@@ -137,6 +137,7 @@ cmake .. -DWITH_NVIDIA=ON -DWITH_OMPI=ON
 | `WITH_METAX`      | Enable MetaX GPU support  | `OFF` |
 | `WITH_MOORE`      | Enable Moore Threads GPU support  | `OFF` |
 | `WITH_CAMBRICON`  | Enable Cambricon MLU support  | `OFF` |
+| `WITH_HYGON`      | Enable HYGON DCU 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` |
@@ -187,7 +188,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`, and `cambricon`. |
+| `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[].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. |
@@ -343,6 +344,7 @@ export LD_LIBRARY_PATH=${INFINI_INSTALL}/lib:$LD_LIBRARY_PATH
 | **MetaX**      | Full | Requires MACA SDK and `MACA_PATH` (default `/opt/maca`) to be set. |
 | **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. |
 
 </details>
 

examples/CMakeLists.txt

@@ -49,6 +49,12 @@ foreach(source_file ${EXAMPLE_SOURCES})
         target_link_libraries(${target_name} PRIVATE ${CAMBRICON_RUNTIME_LIB})
     endif()
 
+    if(WITH_HYGON)
+        target_include_directories(${target_name} PRIVATE ${HYGON_HIP_INCLUDE_DIR})
+        target_compile_definitions(${target_name} PRIVATE ${HYGON_HIP_DEFINITIONS})
+        target_link_libraries(${target_name} PRIVATE ${HYGON_HIP_RUNTIME_LIB})
+    endif()
+
     if(WITH_OMPI OR WITH_MPICH)
         target_link_libraries(${target_name} PRIVATE MPI::MPI_CXX)
     endif()

examples/utils.h

@@ -55,6 +55,7 @@ struct Metrics {
               << bus_bw << " GB/s (Bus BW)" << std::endl;
     std::cout << "Alg Bandwidth:  " << std::fixed << std::setprecision(2)
               << alg_bw << " GB/s" << std::endl;
+    return;
   }
 };
 

src/CMakeLists.txt

@@ -199,6 +199,24 @@ if(WITH_CAMBRICON)
     )
 endif()
 
+# Hygon
+if(WITH_HYGON)
+    list(APPEND DEVICE_LIST "hygon")
+
+    set(HYGON_PATTERNS
+        "devices/hygon/*.cc"
+        "devices/hygon/*.cpp"
+    )
+    file(GLOB_RECURSE HYGON_SRCS ${HYGON_PATTERNS})
+
+    set_source_files_properties(${HYGON_SRCS} PROPERTIES LANGUAGE CXX)
+
+    target_sources(infiniccl PRIVATE ${HYGON_SRCS})
+    target_include_directories(infiniccl PRIVATE ${HYGON_HIP_INCLUDE_DIR})
+    target_compile_definitions(infiniccl PRIVATE ${HYGON_HIP_DEFINITIONS})
+    target_link_libraries(infiniccl PRIVATE ${HYGON_HIP_RUNTIME_LIB})
+endif()
+
 # =========================================================
 # --- BACKEND SECTIONS (Communication Protocols) ---
 # =========================================================

src/backends/mpi/ompi/comm_instance.h

@@ -19,6 +19,7 @@ struct OmpiInstance : public BackendCommInstance {
       INFINI_CHECK_MPI(MPI_Comm_free(&handle));
       handle = MPI_COMM_NULL;
     }
+    return;
   }
 };
 

src/communicator.h

@@ -24,6 +24,7 @@ class Communicator {
   void set_world_info(int rank, int size) {
     global_rank_ = rank;
     global_size_ = size;
+    return;
   }
 
   auto intra_comm() const { return intra_comm_.get(); }
@@ -32,10 +33,12 @@ class Communicator {
 
   void set_intra_comm(std::unique_ptr<BackendCommInstance> inst) {
     intra_comm_ = std::move(inst);
+    return;
   }
 
   void set_inter_comm(std::unique_ptr<BackendCommInstance> inst) {
     inter_comm_ = std::move(inst);
+    return;
   }
 
   BackendType intra_comm_backend() const {
@@ -52,7 +55,10 @@ class Communicator {
 
   int device_id() const { return device_id_; }
 
-  void set_device_id(int id) { device_id_ = id; }
+  void set_device_id(int id) {
+    device_id_ = id;
+    return;
+  }
 
   Device::Type device_type() const { return device_type_; }
 

src/device.h

@@ -156,6 +156,11 @@ struct DevicePriority<Device::Type::kCambricon> {
   static constexpr int value = 5;
 };
 
+template <>
+struct DevicePriority<Device::Type::kHygon> {
+  static constexpr int value = 5;
+};
+
 enum class MemorySpace : std::uint8_t { kHost = 0, kDevice = 1, kUnknown };
 
 template <Device::Type kDev>

src/devices/hygon/caster_.h

@@ -0,0 +1,80 @@
+#ifndef INFINI_CCL_DEVICES_HYGON_CASTER_H_
+#define INFINI_CCL_DEVICES_HYGON_CASTER_H_
+
+#include "caster.h"
+#include "data_type_.h"
+#include "traits.h"
+
+namespace infini::ccl {
+
+// DTK exposes half/bfloat16 compound operators to `hipcc`, but they are not
+// usable on host staging buffers in the MPI backend. Force these types through
+// the bridge.
+template <typename S, typename Op>
+struct SupportsOp<half, S, Op, void> : std::false_type {};
+
+template <typename S, typename Op>
+struct SupportsOp<hip_bfloat16, S, Op, void> : std::false_type {};
+
+template <>
+struct HardwareCastImpl<Device::Type::kHygon, float, half> {
+  __host__ __device__ static float Apply(half x) { return __half2float(x); }
+};
+
+template <>
+struct HardwareCastImpl<Device::Type::kHygon, half, float> {
+  __host__ __device__ static half Apply(float x) { return __float2half(x); }
+};
+
+template <>
+struct HardwareCastImpl<Device::Type::kHygon, float, hip_bfloat16> {
+  __host__ __device__ static float Apply(hip_bfloat16 x) {
+    return static_cast<float>(x);
+  }
+};
+
+template <>
+struct HardwareCastImpl<Device::Type::kHygon, hip_bfloat16, float> {
+  __host__ __device__ static hip_bfloat16 Apply(float x) {
+    return hip_bfloat16(x);
+  }
+};
+
+template <>
+struct HardwareCastImpl<Device::Type::kHygon, hip_bfloat16, int> {
+  __host__ __device__ static hip_bfloat16 Apply(int x) {
+    return hip_bfloat16(static_cast<float>(x));
+  }
+};
+
+template <>
+struct HardwareCastImpl<Device::Type::kHygon, half, int> {
+  __host__ __device__ static half Apply(int x) {
+    return __float2half(static_cast<float>(x));
+  }
+};
+
+template <>
+struct HardwareCastImpl<Device::Type::kHygon, hip_bfloat16, double> {
+  __host__ __device__ static hip_bfloat16 Apply(double x) {
+    return hip_bfloat16(static_cast<float>(x));
+  }
+};
+
+template <>
+struct HardwareCastImpl<Device::Type::kHygon, half, double> {
+  __host__ __device__ static half Apply(double x) {
+    return __float2half(static_cast<float>(x));
+  }
+};
+
+template <>
+struct HardwareCastImpl<Device::Type::kHygon, half, hip_bfloat16> {
+  __host__ __device__ static half Apply(hip_bfloat16 x) {
+    return __float2half(static_cast<float>(x));
+  }
+};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_DEVICES_HYGON_CASTER_H_

src/devices/hygon/data_type_.h

@@ -0,0 +1,26 @@
+#ifndef INFINI_CCL_DEVICES_HYGON_DATA_TYPE_H_
+#define INFINI_CCL_DEVICES_HYGON_DATA_TYPE_H_
+
+// clang-format off
+#include <hip/hip_bfloat16.h>
+#include <hip/hip_fp16.h>
+// clang-format on
+
+#include "data_type_impl.h"
+#include "devices/hygon/device_.h"
+
+namespace infini::ccl {
+
+template <>
+struct TypeMap<Device::Type::kHygon, DataType::kFloat16> {
+  using type = half;
+};
+
+template <>
+struct TypeMap<Device::Type::kHygon, DataType::kBFloat16> {
+  using type = hip_bfloat16;
+};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_DEVICES_HYGON_DATA_TYPE_H_

src/devices/hygon/device_.h

@@ -0,0 +1,36 @@
+#ifndef INFINI_CCL_DEVICES_HYGON_DEVICE_H_
+#define INFINI_CCL_DEVICES_HYGON_DEVICE_H_
+
+// clang-format off
+#include <hip/hip_runtime.h>
+// clang-format on
+
+#include "device.h"
+
+namespace infini::ccl {
+
+template <>
+struct DeviceEnabled<Device::Type::kHygon> : std::true_type {};
+
+template <>
+MemorySpace GetMemorySpace<Device::Type::kHygon>(const void *ptr) {
+  if (!ptr) {
+    return MemorySpace::kHost;
+  }
+
+  hipPointerAttribute_t attr;
+  hipError_t status = hipPointerGetAttributes(&attr, ptr);
+  if (status != hipSuccess) {
+    (void)hipGetLastError();
+    return MemorySpace::kHost;
+  }
+
+  return attr.type == hipMemoryTypeDevice || attr.type == hipMemoryTypeArray ||
+                 attr.type == hipMemoryTypeManaged
+             ? MemorySpace::kDevice
+             : MemorySpace::kHost;
+}
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_DEVICES_HYGON_DEVICE_H_