diff --git a/.github/pull_request_template.md b/.github/pull_request_template.md index 81fc28e..8efbdf7 100644 --- a/.github/pull_request_template.md +++ b/.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 diff --git a/CMakeLists.txt b/CMakeLists.txt index cec2262..ebaca2c 100644 --- a/CMakeLists.txt +++ b/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() diff --git a/README.md b/README.md index 661d5cc..071c38d 100644 --- a/README.md +++ b/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. | diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 24495ea..0065474 100644 --- a/examples/CMakeLists.txt +++ b/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() diff --git a/examples/utils.h b/examples/utils.h index 0c1e7df..3866ab3 100644 --- a/examples/utils.h +++ b/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; } }; diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 06df487..1e1e726 100644 --- a/src/CMakeLists.txt +++ b/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) --- # ========================================================= diff --git a/src/backends/mpi/ompi/comm_instance.h b/src/backends/mpi/ompi/comm_instance.h index c8d3f5d..d0f3a9a 100644 --- a/src/backends/mpi/ompi/comm_instance.h +++ b/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; } }; diff --git a/src/communicator.h b/src/communicator.h index 84e102b..fe70c87 100644 --- a/src/communicator.h +++ b/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 inst) { intra_comm_ = std::move(inst); + return; } void set_inter_comm(std::unique_ptr 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_; } diff --git a/src/device.h b/src/device.h index 64eb7c9..4809420 100644 --- a/src/device.h +++ b/src/device.h @@ -156,6 +156,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/hygon/caster_.h b/src/devices/hygon/caster_.h new file mode 100644 index 0000000..3a8a72b --- /dev/null +++ b/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 +struct SupportsOp : std::false_type {}; + +template +struct SupportsOp : std::false_type {}; + +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(hip_bfloat16 x) { + return static_cast(x); + } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static hip_bfloat16 Apply(float x) { + return hip_bfloat16(x); + } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static hip_bfloat16 Apply(int x) { + return hip_bfloat16(static_cast(x)); + } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static half Apply(int x) { + return __float2half(static_cast(x)); + } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static hip_bfloat16 Apply(double x) { + return hip_bfloat16(static_cast(x)); + } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static half Apply(double x) { + return __float2half(static_cast(x)); + } +}; + +template <> +struct HardwareCastImpl { + __host__ __device__ static half Apply(hip_bfloat16 x) { + return __float2half(static_cast(x)); + } +}; + +} // namespace infini::ccl + +#endif // INFINI_CCL_DEVICES_HYGON_CASTER_H_ diff --git a/src/devices/hygon/data_type_.h b/src/devices/hygon/data_type_.h new file mode 100644 index 0000000..1e79e6a --- /dev/null +++ b/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 +#include +// clang-format on + +#include "data_type_impl.h" +#include "devices/hygon/device_.h" + +namespace infini::ccl { + +template <> +struct TypeMap { + using type = half; +}; + +template <> +struct TypeMap { + using type = hip_bfloat16; +}; + +} // namespace infini::ccl + +#endif // INFINI_CCL_DEVICES_HYGON_DATA_TYPE_H_ diff --git a/src/devices/hygon/device_.h b/src/devices/hygon/device_.h new file mode 100644 index 0000000..9441eb8 --- /dev/null +++ b/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 +// clang-format on + +#include "device.h" + +namespace infini::ccl { + +template <> +struct DeviceEnabled : std::true_type {}; + +template <> +MemorySpace GetMemorySpace(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_ diff --git a/src/devices/hygon/runtime_.h b/src/devices/hygon/runtime_.h new file mode 100644 index 0000000..8ad05ef --- /dev/null +++ b/src/devices/hygon/runtime_.h @@ -0,0 +1,60 @@ +#ifndef INFINI_CCL_DEVICES_HYGON_RUNTIME_H_ +#define INFINI_CCL_DEVICES_HYGON_RUNTIME_H_ + +#include + +// clang-format off +#include +// clang-format on + +#include "devices/cuda/runtime_.h" +#include "devices/hygon/device_.h" +#include "logging.h" +#include "return_status_impl.h" + +namespace infini::ccl { + +template <> +struct Runtime + : CudaRuntime> { + using Stream = hipStream_t; + + static constexpr Device::Type kDeviceType = Device::Type::kHygon; + + static constexpr auto Check = + [](auto status, ReturnStatus err_code = ReturnStatus::kSystemError) { + if (status != hipSuccess) { + LOG(hipGetErrorString(static_cast(status))); + return err_code; + } + return ReturnStatus::kSuccess; + }; + + static constexpr auto Malloc = [](auto &&...args) { + return hipMalloc(std::forward(args)...); + }; + + static constexpr auto Memcpy = hipMemcpy; + + static constexpr auto Free = hipFree; + + static constexpr auto MemcpyHostToDevice = hipMemcpyHostToDevice; + + static constexpr auto MemcpyDeviceToHost = hipMemcpyDeviceToHost; + + static constexpr auto Memset = hipMemset; + + static constexpr auto GetDevice = hipGetDevice; + + static constexpr auto SetDevice = hipSetDevice; + + static constexpr auto DeviceSynchronize = hipDeviceSynchronize; + + static constexpr auto StreamSynchronize = hipStreamSynchronize; +}; + +static_assert(Runtime::Validate()); + +} // namespace infini::ccl + +#endif // INFINI_CCL_DEVICES_HYGON_RUNTIME_H_