feat(qy): enable InfiniCCL NCCL backend for Denglin QY (2026.11.12)

CMakeLists.txt

@@ -15,6 +15,7 @@ option(WITH_CAMBRICON   "Enable Cambricon MLU support"  OFF)
 option(WITH_ASCEND      "Enable Ascend NPU support"     OFF)
 option(WITH_HYGON       "Enable Hygon DCU support"      OFF)
 option(WITH_THEAD       "Enable T-Head PPU support"     OFF)
+option(WITH_QY          "Enable Denglin QY GPU support" OFF)
 
 set(WITH_CPU ON CACHE INTERNAL "CPU backend is always enabled")
 
@@ -45,6 +46,7 @@ include(cmake/devices/Cambricon.cmake)
 include(cmake/devices/Ascend.cmake)
 include(cmake/devices/Hygon.cmake)
 include(cmake/devices/THead.cmake)
+include(cmake/devices/Qy.cmake)
 
 # Backend modules provide backend probes and required dependency resolution.
 include(cmake/backends/MPI.cmake)
@@ -91,6 +93,8 @@ if(AUTO_DETECT_DEVICES)
     # T-Head
     infiniccl_resolve_device_detection(
         WITH_THEAD "T-Head" infiniccl_detect_thead THEAD_FOUND)
+    infiniccl_maybe_enable_device(
+        WITH_QY "QY" infiniccl_detect_qy QY_FOUND)
 
     # NVIDIA Final Resolution: resolve conflicts.
     if(THEAD_FOUND AND NVIDIA_FOUND AND NOT NVIDIA_REQUESTED)
@@ -117,7 +121,7 @@ if(AUTO_DETECT_BACKENDS)
     endif()
 
     # Detect NCCL Dependencies
-    if(WITH_NVIDIA OR WITH_THEAD OR WITH_ILUVATAR OR WITH_HYGON)
+    if(WITH_NVIDIA OR WITH_THEAD OR WITH_ILUVATAR OR WITH_HYGON OR WITH_QY)
         infiniccl_probe_nccl(NCCL_FOUND)
         if(NCCL_FOUND)
             set(WITH_NCCL ON)
@@ -214,6 +218,10 @@ if(WITH_HYGON)
     infiniccl_configure_hygon()
 endif()
 
+if(WITH_QY)
+    infiniccl_configure_qy()
+endif()
+
 if(WITH_OMPI OR WITH_MPICH)
     find_package(MPI REQUIRED)
 endif()

cmake/backends/NCCL.cmake

@@ -34,6 +34,11 @@ function(_infiniccl_find_nccl prefix required include_output library_output defi
         set(_definitions INFINI_CCL_USE_RCCL)
         list(APPEND _library_suffixes rccl/lib)
         message(STATUS "Hygon detected. Searching for NCCL-compatible RCCL in ${_hints}")
+    elseif(WITH_QY)
+        if(QY_SDK)
+            list(APPEND _hints "${QY_SDK}" "${QY_ROOT}")
+        endif()
+        message(STATUS "QY detected. Searching for NCCL in ${_hints}")
     endif()
     list(REMOVE_DUPLICATES _hints)
 
@@ -74,10 +79,10 @@ function(infiniccl_probe_nccl output)
 endfunction()
 
 function(infiniccl_configure_nccl include_output library_output definitions_output)
-    if(NOT WITH_NVIDIA AND NOT WITH_THEAD AND NOT WITH_ILUVATAR AND NOT WITH_HYGON)
+    if(NOT WITH_NVIDIA AND NOT WITH_THEAD AND NOT WITH_ILUVATAR AND NOT WITH_HYGON AND NOT WITH_QY)
         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`.")
+            "NCCL backend requires NVIDIA, T-Head, Iluvatar, Hygon, or QY GPU support. "
+            "Please enable `WITH_NVIDIA`, `WITH_THEAD`, `WITH_ILUVATAR`, `WITH_HYGON`, or `WITH_QY`.")
     endif()
 
     _infiniccl_find_nccl(_INFINICCL_NCCL TRUE _include _library _definitions)

cmake/devices/Qy.cmake

@@ -0,0 +1,46 @@
+include_guard(GLOBAL)
+
+function(infiniccl_detect_qy output)
+    set(_found FALSE)
+    foreach(_qy_env QY_ROOT QY_HOME QY_PATH)
+        if(DEFINED ENV{${_qy_env}} AND EXISTS "$ENV{${_qy_env}}/sdk/include/cuda_runtime.h")
+            set(_found TRUE)
+            break()
+        endif()
+    endforeach()
+    if(NOT _found AND EXISTS "/usr/local/dlgpu/sdk/include/cuda_runtime.h")
+        set(_found TRUE)
+    endif()
+    if(NOT _found AND EXISTS "/usr/local/denglin/sdk/include/cuda_runtime.h")
+        set(_found TRUE)
+    endif()
+    set(${output} ${_found} PARENT_SCOPE)
+endfunction()
+
+function(infiniccl_configure_qy)
+    set(QY_ROOT "")
+    foreach(_qy_env QY_ROOT QY_HOME QY_PATH)
+        if(NOT QY_ROOT AND DEFINED ENV{${_qy_env}} AND NOT "$ENV{${_qy_env}}" STREQUAL "")
+            set(QY_ROOT "$ENV{${_qy_env}}")
+        endif()
+    endforeach()
+    if(NOT QY_ROOT AND EXISTS "/usr/local/dlgpu/sdk/bin/dlcc")
+        set(QY_ROOT "/usr/local/dlgpu")
+    endif()
+    if(NOT QY_ROOT AND EXISTS "/usr/local/denglin/sdk/bin/dlcc")
+        set(QY_ROOT "/usr/local/denglin")
+    endif()
+    if(NOT QY_ROOT OR NOT EXISTS "${QY_ROOT}/sdk/include/cuda_runtime.h")
+        message(FATAL_ERROR
+            "`WITH_QY` is `ON` but Denglin QY SDK was not found. "
+            "Set `QY_ROOT` (expected `sdk/include/cuda_runtime.h`).")
+    endif()
+    set(QY_ROOT "${QY_ROOT}" CACHE PATH "Denglin QY SDK root" FORCE)
+    set(QY_SDK "${QY_ROOT}/sdk" PARENT_SCOPE)
+    set(QY_SDK "${QY_ROOT}/sdk" CACHE PATH "Denglin QY SDK path" FORCE)
+    find_library(QY_CURT_LIB NAMES curt HINTS "${QY_ROOT}/sdk/lib" "${QY_ROOT}/sdk/lib64" REQUIRED)
+    set(QY_CURT_LIB "${QY_CURT_LIB}" PARENT_SCOPE)
+    include_directories("${QY_ROOT}/sdk/include")
+    link_directories("${QY_ROOT}/sdk/lib" "${QY_ROOT}/sdk/lib64")
+    message(STATUS "QY: SDK ${QY_ROOT}/sdk, curt ${QY_CURT_LIB}")
+endfunction()

src/CMakeLists.txt

@@ -250,6 +250,23 @@ if(WITH_ASCEND)
 endif()
 
 # Hygon
+
+if(WITH_QY)
+    list(APPEND DEVICE_LIST "qy")
+
+    set(QY_PATTERNS
+        "devices/cuda/*.cc"
+        "devices/cuda/*.cpp"
+        "devices/qy/*.cc"
+        "devices/qy/*.cpp"
+    )
+    file(GLOB QY_SOURCES ${QY_PATTERNS})
+    target_sources(infiniccl PRIVATE ${QY_SOURCES})
+    target_include_directories(infiniccl PRIVATE "${QY_SDK}/include")
+    target_compile_definitions(infiniccl PRIVATE WITH_QY=1)
+    target_link_libraries(infiniccl PRIVATE ${QY_CURT_LIB})
+endif()
+
 if(WITH_HYGON)
     list(APPEND DEVICE_LIST "hygon")
 

src/backend_device_map.h

@@ -29,6 +29,10 @@ template <>
 struct IsSupportedCombination<BackendType::kNccl, Device::Type::kHygon>
     : std::true_type {};
 
+template <>
+struct IsSupportedCombination<BackendType::kNccl, Device::Type::kQy>
+    : std::true_type {};
+
 template <>
 struct IsSupportedCombination<BackendType::kMccl, Device::Type::kMetax>
     : std::true_type {};

src/backends/ccl/nccl/qy/api.h

@@ -0,0 +1,15 @@
+#ifndef INFINI_CCL_BACKENDS_CCL_NCCL_QY_API_H_
+#define INFINI_CCL_BACKENDS_CCL_NCCL_QY_API_H_
+
+#include "backends/ccl/nccl/api.h"
+#include "devices/qy/runtime_.h"
+
+namespace infini::ccl {
+
+template <>
+struct CclApi<BackendType::kNccl, Device::Type::kQy>
+    : NcclApi<Device::Type::kQy> {};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_BACKENDS_CCL_NCCL_QY_API_H_

src/device.h

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

src/devices/qy/data_type_.h

@@ -0,0 +1,26 @@
+#ifndef INFINI_CCL_DEVICES_QY_DATA_TYPE_H_
+#define INFINI_CCL_DEVICES_QY_DATA_TYPE_H_
+
+// clang-format off
+#include <cuda_bf16.h>
+#include <cuda_fp16.h>
+// clang-format on
+
+#include "data_type_impl.h"
+#include "devices/qy/device_.h"
+
+namespace infini::ccl {
+
+template <>
+struct TypeMap<Device::Type::kQy, DataType::kFloat16> {
+  using type = half;
+};
+
+template <>
+struct TypeMap<Device::Type::kQy, DataType::kBFloat16> {
+  using type = __nv_bfloat16;
+};
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_DEVICES_QY_DATA_TYPE_H_

src/devices/qy/device_.h

@@ -0,0 +1,32 @@
+#ifndef INFINI_CCL_DEVICES_QY_DEVICE_H_
+#define INFINI_CCL_DEVICES_QY_DEVICE_H_
+
+#include <cuda_runtime.h>
+
+#include "device.h"
+#include "devices/cuda/checks.h"
+
+namespace infini::ccl {
+
+template <>
+struct DeviceEnabled<Device::Type::kQy> : std::true_type {};
+
+template <>
+MemorySpace GetMemorySpace<Device::Type::kQy>(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_QY_DEVICE_H_

src/devices/qy/runtime_.h

@@ -0,0 +1,75 @@
+#ifndef INFINI_CCL_DEVICES_QY_RUNTIME_H_
+#define INFINI_CCL_DEVICES_QY_RUNTIME_H_
+
+#include <utility>
+
+// clang-format off
+#include <cuda_runtime.h>
+// clang-format on
+
+#include "devices/cuda/runtime_.h"
+#include "devices/qy/device_.h"
+#include "logging.h"
+#include "return_status_impl.h"
+
+namespace infini::ccl {
+
+template <>
+struct Runtime<Device::Type::kQy> : CudaRuntime<Runtime<Device::Type::kQy>> {
+  using Stream = cudaStream_t;
+
+  static constexpr Device::Type kDeviceType = Device::Type::kQy;
+
+  // Denglin overloads several cuda* entry points; wrap in forwarding lambdas
+  // so a single candidate is selected (same pattern as InfiniRT QY).
+  static constexpr auto Check =
+      [](auto status, ReturnStatus err_code = ReturnStatus::kSystemError) {
+        if (status != cudaSuccess) {
+          LOG(cudaGetErrorString(static_cast<cudaError_t>(status)));
+          return err_code;
+        }
+        return ReturnStatus::kSuccess;
+      };
+
+  static constexpr auto Malloc = [](auto &&...args) {
+    return cudaMalloc(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto Memcpy = [](auto &&...args) {
+    return cudaMemcpy(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto Free = [](auto &&...args) {
+    return cudaFree(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto MemcpyHostToDevice = cudaMemcpyHostToDevice;
+
+  static constexpr auto MemcpyDeviceToHost = cudaMemcpyDeviceToHost;
+
+  static constexpr auto Memset = [](auto &&...args) {
+    return cudaMemset(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto GetDevice = [](auto &&...args) {
+    return cudaGetDevice(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto SetDevice = [](auto &&...args) {
+    return cudaSetDevice(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto DeviceSynchronize = [](auto &&...args) {
+    return cudaDeviceSynchronize(std::forward<decltype(args)>(args)...);
+  };
+
+  static constexpr auto StreamSynchronize = [](auto &&...args) {
+    return cudaStreamSynchronize(std::forward<decltype(args)>(args)...);
+  };
+};
+
+static_assert(Runtime<Device::Type::kQy>::Validate());
+
+}  // namespace infini::ccl
+
+#endif  // INFINI_CCL_DEVICES_QY_RUNTIME_H_