Adapt seven operators to Hygon machines.

Author: zhuyue

README.md

@@ -19,7 +19,7 @@ InfiniCore 是一个跨平台统一编程工具集，为不同芯片平台的功
   - 摩尔线程 GPU；
   - 天数智芯 GPU；
   - 沐曦 GPU；
-  - 曙光 DCU；
+  - 海光 DCU；
 - 华为昇腾 NPU；
 - 寒武纪 MLU；
 - 昆仑芯 XPU；
@@ -50,7 +50,7 @@ python scripts/install.py [XMAKE_CONFIG_FLAGS]
 | `--metax-gpu=[y\|n]`     | 是否编译沐曦 GPU 接口实现         | n
 | `--moore-gpu=[y\|n]`     | 是否编译摩尔线程 GPU 接口实现     | n
 | `--iluvatar-gpu=[y\|n]`  | 是否编译沐曦 GPU 接口实现         | n
-| `--sugon-dcu=[y\|n]`     | 是否编译曙光 DCU 接口实现         | n
+| `--hygon-dcu=[y\|n]`     | 是否编译海光 DCU 接口实现         | n
 | `--kunlun-xpu=[y\|n]`    | 是否编译昆仑 XPU 接口实现         | n
 | `--ninetoothed=[y\|n]`   | 是否编译九齿实现                 | n
 | `--ccl=[y\|n]`           | 是否编译 InfiniCCL 通信库接口实现 | n

include/infinicore.h

@@ -45,7 +45,7 @@ typedef enum {
     INFINI_DEVICE_MOORE = 5,
     INFINI_DEVICE_ILUVATAR = 6,
     INFINI_DEVICE_KUNLUN = 7,
-    INFINI_DEVICE_SUGON = 8,
+    INFINI_DEVICE_HYGON = 8,
     INFINI_DEVICE_TYPE_COUNT
 } infiniDevice_t;
 

src/infiniccl-test/main.cpp

@@ -12,7 +12,7 @@ void printUsage() {
     std::cout << "infiniccl-test --<device>" << std::endl
               << std::endl;
     std::cout << "  --<device>" << std::endl;
-    std::cout << "    Specify the device type --(nvidia|cambricon|ascend|metax|moore|iluvatar|kunlun|sugon)." << std::endl
+    std::cout << "    Specify the device type --(nvidia|cambricon|ascend|metax|moore|iluvatar|kunlun|hygon)." << std::endl
               << std::endl;
     std::cout << "The program will run tests on all visible devices of the specified device type."
               << " Use Environmental Variables such as CUDA_VSIBLE_DEVICES to limit visible device IDs.";
@@ -44,7 +44,7 @@ ParsedArgs parseArgs(int argc, char *argv[]) {
         else PARSE_DEVICE("--moore", INFINI_DEVICE_MOORE)
         else PARSE_DEVICE("--iluvatar", INFINI_DEVICE_ILUVATAR)
         else PARSE_DEVICE("--kunlun", INFINI_DEVICE_KUNLUN)
-        else PARSE_DEVICE("--sugon", INFINI_DEVICE_SUGON)
+        else PARSE_DEVICE("--hygon", INFINI_DEVICE_HYGON)
         else {
             printUsage();
         }

src/infiniccl/cuda/infiniccl_cuda.h

@@ -4,7 +4,7 @@
 #include "../infiniccl_impl.h"
 
 // Windows does not support CUDA
-#if (defined(ENABLE_NVIDIA_API) || defined(ENABLE_ILUVATAR_API)) && defined(ENABLE_CCL) && !defined(_WIN32)
+#if (defined(ENABLE_NVIDIA_API) || defined(ENABLE_ILUVATAR_API) || defined(ENABLE_HYGON_API)) && defined(ENABLE_CCL) && !defined(_WIN32)
 INFINICCL_DEVICE_API_IMPL(cuda)
 #else
 INFINICCL_DEVICE_API_NOOP(cuda)

src/infiniccl/infiniccl.cc

@@ -20,6 +20,7 @@ __C infiniStatus_t infinicclCommInitAll(
     switch (device_type) {
         COMM_INIT_ALL(INFINI_DEVICE_NVIDIA, cuda);
         COMM_INIT_ALL(INFINI_DEVICE_ILUVATAR, cuda);
+        COMM_INIT_ALL(INFINI_DEVICE_HYGON, cuda);
         COMM_INIT_ALL(INFINI_DEVICE_ASCEND, ascend);
         COMM_INIT_ALL(INFINI_DEVICE_CAMBRICON, cambricon);
         COMM_INIT_ALL(INFINI_DEVICE_METAX, metax);
@@ -44,6 +45,7 @@ __C infiniStatus_t infinicclCommDestroy(infinicclComm_t comm) {
     switch (comm->device_type) {
         COMM_DESTROY(INFINI_DEVICE_NVIDIA, cuda);
         COMM_DESTROY(INFINI_DEVICE_ILUVATAR, cuda);
+        COMM_DESTROY(INFINI_DEVICE_HYGON, cuda);
         COMM_DESTROY(INFINI_DEVICE_ASCEND, ascend);
         COMM_DESTROY(INFINI_DEVICE_CAMBRICON, cambricon);
         COMM_DESTROY(INFINI_DEVICE_METAX, metax);
@@ -75,6 +77,7 @@ __C infiniStatus_t infinicclAllReduce(
     switch (comm->device_type) {
         ALL_REDUCE(INFINI_DEVICE_NVIDIA, cuda);
         ALL_REDUCE(INFINI_DEVICE_ILUVATAR, cuda);
+        ALL_REDUCE(INFINI_DEVICE_HYGON, cuda);
         ALL_REDUCE(INFINI_DEVICE_ASCEND, ascend);
         ALL_REDUCE(INFINI_DEVICE_CAMBRICON, cambricon);
         ALL_REDUCE(INFINI_DEVICE_METAX, metax);

src/infiniop-test/src/main.cpp

@@ -22,7 +22,7 @@ void printUsage() {
     std::cout << "    Path to the test gguf file" << std::endl
               << std::endl;
     std::cout << "  --<device>[:id]" << std::endl;
-    std::cout << "    (Optional) Specify the device type --(cpu|nvidia|cambricon|ascend|metax|moore|iluvatar|kunlun|sugon) and device ID (optional). CPU by default." << std::endl
+    std::cout << "    (Optional) Specify the device type --(cpu|nvidia|cambricon|ascend|metax|moore|iluvatar|kunlun|hygon) and device ID (optional). CPU by default." << std::endl
               << std::endl;
     std::cout << "  --warmup <warmups>" << std::endl;
     std::cout << "    (Optional) Number of warmups to perform before timing. Default to 0." << std::endl
@@ -78,7 +78,7 @@ ParsedArgs parseArgs(int argc, char *argv[]) {
             PARSE_DEVICE("--moore", INFINI_DEVICE_MOORE)
             PARSE_DEVICE("--iluvatar", INFINI_DEVICE_ILUVATAR)
             PARSE_DEVICE("--kunlun", INFINI_DEVICE_KUNLUN)
-            PARSE_DEVICE("--sugon", INFINI_DEVICE_SUGON)
+            PARSE_DEVICE("--hygon", INFINI_DEVICE_HYGON)
             else if (arg == "--warmup" && i + 1 < argc) {
                 args.warmups = std::stoi(argv[++i]);
             }

src/infiniop/devices/handle.cc

@@ -5,7 +5,7 @@
 #ifdef ENABLE_CPU_API
 #include "cpu/cpu_handle.h"
 #endif
-#if defined(ENABLE_NVIDIA_API) || defined(ENABLE_ILUVATAR_API)
+#if defined(ENABLE_NVIDIA_API) || defined(ENABLE_ILUVATAR_API) || defined(ENABLE_HYGON_API)
 #include "nvidia/nvidia_handle.h"
 #endif
 #ifdef ENABLE_CAMBRICON_API
@@ -62,6 +62,9 @@ __C infiniStatus_t infiniopCreateHandle(infiniopHandle_t *handle_ptr) {
 #ifdef ENABLE_METAX_API
         CREATE(INFINI_DEVICE_METAX, metax);
 #endif
+#ifdef ENABLE_HYGON_API
+        CREATE(INFINI_DEVICE_HYGON, hygon);
+#endif
 
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -101,6 +104,9 @@ __C infiniStatus_t infiniopDestroyHandle(infiniopHandle_t handle) {
 #endif
 #ifdef ENABLE_METAX_API
         DELETE(INFINI_DEVICE_METAX, metax);
+#endif
+#ifdef ENABLE_HYGON_API
+        DELETE(INFINI_DEVICE_HYGON, hygon);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;

src/infiniop/devices/nvidia/nvidia_common.cu

@@ -104,4 +104,16 @@ infiniStatus_t Handle::create(InfiniopHandle **handle_ptr, int device_id) {
 
 } // namespace iluvatar
 
+namespace hygon {
+
+Handle::Handle(int device_id)
+    : nvidia::Handle(INFINI_DEVICE_HYGON, device_id) {}
+
+infiniStatus_t Handle::create(InfiniopHandle **handle_ptr, int device_id) {
+    *handle_ptr = new Handle(device_id);
+    return INFINI_STATUS_SUCCESS;
+}
+
+} // namespace hygon
+
 } // namespace device

src/infiniop/devices/nvidia/nvidia_handle.h

@@ -35,6 +35,17 @@ struct Handle : public nvidia::Handle {
 
 } // namespace iluvatar
 
+namespace hygon {
+
+struct Handle : public nvidia::Handle {
+    Handle(int device_id);
+
+public:
+    static infiniStatus_t create(InfiniopHandle **handle_ptr, int device_id);
+};
+
+} // namespace hygon
+
 } // namespace device
 
 #endif // __INFINIOP_CUDA_HANDLE_H__

src/infiniop/devices/nvidia/nvidia_kernel_common.cuh

@@ -1,5 +1,8 @@
-#ifdef ENABLE_SUGON_CUDA_API
-#define INFINIOP_CUDA_KERNEL __launch_bounds__(512) __global__ void
+#ifndef __INFINIOP_CUDA_KERNEL_COMMON_CUH__
+#define __INFINIOP_CUDA_KERNEL_COMMON_CUH__
+
+#if defined(ENABLE_HYGON_API)
+#define INFINIOP_CUDA_KERNEL __launch_bounds__(1024) __global__ void
 #else
 #define INFINIOP_CUDA_KERNEL __global__ void
 #endif
@@ -15,8 +18,14 @@
 
 #define CHECK_CUDA(API) CHECK_INTERNAL(API, cudaSuccess)
 
+#ifdef ENABLE_HYGON_API
+// Hygon DCU uses different bfloat16 type definitions
+using cuda_bfloat16 = __nv_bfloat16;
+using cuda_bfloat162 = __nv_bfloat162;
+#else
 using cuda_bfloat16 = nv_bfloat16;
 using cuda_bfloat162 = nv_bfloat162;
+#endif
 
 namespace device::nvidia {
 
@@ -41,7 +50,7 @@ exp_(const float val) {
     return expf(val);
 }
 
-#ifndef ENABLE_ILUVATAR_API
+#if !defined(ENABLE_ILUVATAR_API) && !defined(ENABLE_HYGON_API)
 __forceinline__ __device__ long double
 exp_(const long double val) {
     return expl(val);
@@ -62,3 +71,5 @@ __forceinline__ __device__ __nv_bfloat16
 exp_(const __nv_bfloat16 x) {
     return hexp(x);
 }
+
+#endif // __INFINIOP_CUDA_KERNEL_COMMON_CUH__