From def6af815851018e3e67f589d7a9fb17a1c85b6a Mon Sep 17 00:00:00 2001 From: baihuawei Date: Fri, 30 Jul 2021 16:22:48 +0800 Subject: [PATCH] ascend support mpi --- mindspore/ccsrc/runtime/device/CMakeLists.txt | 22 +++ .../ascend/distribute/ascend_collective.cc | 94 +++++++++++++ .../ascend/distribute/ascend_collective.h | 74 ++++++++++ .../distribute/collective_group_wrapper.cc | 28 ++++ .../distribute/collective_group_wrapper.h | 36 +++++ .../ascend/distribute/mpi_collective_group.cc | 132 ++++++++++++++++++ .../ascend/distribute/mpi_collective_group.h | 74 ++++++++++ .../device/ascend/distribute/mpi_pycc.cc | 41 ++++++ .../device/ascend/distribute/mpi_pycc.h | 44 ++++++ .../runtime/hccl_adapter/hccl_adapter.cc | 46 ++---- .../runtime/hccl_adapter/plugin/hccl_plugin.h | 17 +-- mindspore/ccsrc/utils/dlopen_macro.h | 55 ++++++++ 12 files changed, 616 insertions(+), 47 deletions(-) create mode 100644 mindspore/ccsrc/runtime/device/ascend/distribute/ascend_collective.cc create mode 100644 mindspore/ccsrc/runtime/device/ascend/distribute/ascend_collective.h create mode 100644 mindspore/ccsrc/runtime/device/ascend/distribute/collective_group_wrapper.cc create mode 100644 mindspore/ccsrc/runtime/device/ascend/distribute/collective_group_wrapper.h create mode 100644 mindspore/ccsrc/runtime/device/ascend/distribute/mpi_collective_group.cc create mode 100644 mindspore/ccsrc/runtime/device/ascend/distribute/mpi_collective_group.h create mode 100644 mindspore/ccsrc/runtime/device/ascend/distribute/mpi_pycc.cc create mode 100644 mindspore/ccsrc/runtime/device/ascend/distribute/mpi_pycc.h create mode 100644 mindspore/ccsrc/utils/dlopen_macro.h diff --git a/mindspore/ccsrc/runtime/device/CMakeLists.txt b/mindspore/ccsrc/runtime/device/CMakeLists.txt index 68534970005..c7e99adbbea 100644 --- a/mindspore/ccsrc/runtime/device/CMakeLists.txt +++ b/mindspore/ccsrc/runtime/device/CMakeLists.txt @@ -17,6 +17,8 @@ endif() if(ENABLE_D) file(GLOB_RECURSE D_SRC_LIST RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "ascend/*.cc" "kernel_adjust.cc") + list(REMOVE_ITEM D_SRC_LIST "ascend/distribute/mpi_collective_group.cc" + "ascend/distribute/collective_group_wrapper.cc" "ascend/distribute/mpi_pycc.cc") endif() if(ENABLE_TDTQUE) file(GLOB_RECURSE TDT_SRC_LIST RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} @@ -43,6 +45,13 @@ if(ENABLE_MPI) pybind11_add_module(_ms_mpi "gpu/mpi/mpi_initializer.cc") target_link_libraries(_ms_mpi PRIVATE mindspore::pybind11_module mindspore::ompi) endif() + + if(ENABLE_D) + set_property(SOURCE "ascend/distribute/mpi_pycc.cc" + PROPERTY COMPILE_DEFINITIONS SUBMODULE_ID=mindspore::SubModuleId::SM_DEVICE) + pybind11_add_module(_ascend_mpi "ascend/distribute/mpi_pycc.cc") + target_link_libraries(_ascend_mpi PRIVATE mindspore::pybind11_module mindspore::ompi) + endif() endif() # gpu @@ -84,4 +93,17 @@ if(ENABLE_D) set_property(SOURCE ${GE_SRC_LIST} PROPERTY COMPILE_DEFINITIONS SUBMODULE_ID=mindspore::SubModuleId::SM_GE) target_include_directories(_mindspore_runtime_device_obj PRIVATE ${CMAKE_BINARY_DIR}/proto/ge) add_dependencies(_mindspore_runtime_device_obj graph) + if(ENABLE_MPI) + set(ASCEND_PATH /usr/local/Ascend) + set(ASCEND_TOOLKIT_RUNTIME_PATH ${ASCEND_PATH}/ascend-toolkit/latest/fwkacllib/lib64) + set(ASCEND_RUNTIME_PATH ${ASCEND_PATH}/fwkacllib/lib64) + find_library(HCCL hccl ${ASCEND_RUNTIME_PATH} ${ASCEND_TOOLKIT_RUNTIME_PATH}) + file(GLOB_RECURSE ASCEND_COLLECTIVE_LIST "ascend/distribute/mpi_collective_group.cc" + "ascend/distribute/collective_group_wrapper.cc") + set_property(SOURCE ${ASCEND_COLLECTIVE_LIST} + PROPERTY COMPILE_DEFINITIONS SUBMODULE_ID=mindspore::SubModuleId::SM_DEVICE) + add_library(ascend_collective SHARED ${ASCEND_COLLECTIVE_LIST}) + target_link_libraries(ascend_collective PRIVATE ${HCCL} mindspore::ompi) + target_link_libraries(_ascend_mpi PRIVATE ascend_collective) + endif() endif() diff --git a/mindspore/ccsrc/runtime/device/ascend/distribute/ascend_collective.cc b/mindspore/ccsrc/runtime/device/ascend/distribute/ascend_collective.cc new file mode 100644 index 00000000000..fa28e1a2e17 --- /dev/null +++ b/mindspore/ccsrc/runtime/device/ascend/distribute/ascend_collective.cc @@ -0,0 +1,94 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "runtime/device/ascend/distribute/ascend_collective.h" +#include "utils/log_adapter.h" + +static constexpr const char *kAscendCollectiveFileName = "libascend_collective.so"; +namespace mindspore { +namespace device { +namespace ascend { +namespace collective { +HcclCollectiveGroup &HcclCollectiveGroup::instance() { + static HcclCollectiveGroup instance; + return instance; +} + +void HcclCollectiveGroup::FinalizeCollective() { + MS_LOG(INFO) << "Finalize Collective"; + if (collective_handle_ != nullptr) { + MS_EXCEPTION_IF_NULL(finalize_mpi_); + finalize_mpi_(); + if (dlclose(collective_handle_) != 0) { + MS_LOG(EXCEPTION) << "Closing libascend_collective.so handle failed."; + } + } +} + +bool HcclCollectiveGroup::InitCollective() { + MS_LOG(INFO) << "InitCollective"; + if (inited_) { + return true; + } + collective_handle_ = dlopen(kAscendCollectiveFileName, RTLD_NOW); + if (collective_handle_ == nullptr) { + MS_LOG(EXCEPTION) + << "Loading libascend_collective.so failed. Many reasons could cause this:\n1.libascend_collective.so is not " + "installed.\n2.hccl is not " + "installed or found.\n3.mpi is not installed or found"; + } + init_mpi_ = DlsymFuncObj(InitMPI, collective_handle_); + finalize_mpi_ = DlsymFuncObj(FinalizeMPI, collective_handle_); + get_group_comm_ = DlsymFuncObj(GetGroupComm, collective_handle_); + get_group_size_ = DlsymFuncObj(GetGroupSize, collective_handle_); + get_rank_id_by_group_ = DlsymFuncObj(GetRankIdByGroup, collective_handle_); + get_device_id_ = DlsymFuncObj(GetDeviceId, collective_handle_); + create_comm_for_group_ = DlsymFuncObj(CreateCommForGroup, collective_handle_); + destroy_hccl_comm_ = DlsymFuncObj(DestroyHcclComm, collective_handle_); + MS_EXCEPTION_IF_NULL(init_mpi_); + init_mpi_(); + inited_ = true; + MS_LOG(INFO) << "InitCollective success"; + return true; +} +HcclComm HcclCollectiveGroup::GetGroupComm(const std::string &name) { + MS_EXCEPTION_IF_NULL(get_group_comm_); + return get_group_comm_(name); +} +int HcclCollectiveGroup::GetRankSize(const std::string &name) const { + MS_EXCEPTION_IF_NULL(get_group_size_); + return get_group_size_(name); +} +int HcclCollectiveGroup::GetRankId(const std::string &name) const { + MS_EXCEPTION_IF_NULL(get_rank_id_by_group_); + return get_rank_id_by_group_(name); +} +int HcclCollectiveGroup::GetDeviceId() const { + MS_EXCEPTION_IF_NULL(get_device_id_); + return get_device_id_(); +} +void HcclCollectiveGroup::CreateCommGroup(const std::string &name, const std::vector &ranks) { + MS_EXCEPTION_IF_NULL(create_comm_for_group_); + create_comm_for_group_(name, ranks); +} +void HcclCollectiveGroup::DestroyCommGroup() { + MS_EXCEPTION_IF_NULL(destroy_hccl_comm_); + destroy_hccl_comm_(); +} +} // namespace collective +} // namespace ascend +} // namespace device +} // namespace mindspore diff --git a/mindspore/ccsrc/runtime/device/ascend/distribute/ascend_collective.h b/mindspore/ccsrc/runtime/device/ascend/distribute/ascend_collective.h new file mode 100644 index 00000000000..d8e17afaf8b --- /dev/null +++ b/mindspore/ccsrc/runtime/device/ascend/distribute/ascend_collective.h @@ -0,0 +1,74 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_ASCEND_COLLECTIVE_H +#define MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_ASCEND_COLLECTIVE_H + +#include +#include +#include +#include +#include "hccl/hccl_types.h" +#include "utils/utils.h" +#include "utils/dlopen_macro.h" +namespace mindspore { +namespace device { +namespace ascend { +namespace collective { + +ORIGIN_METHOD(InitMPI, void); +ORIGIN_METHOD(FinalizeMPI, void); +ORIGIN_METHOD(GetGroupComm, HcclComm, const std::string &); +ORIGIN_METHOD(GetGroupSize, int, const std::string &); +ORIGIN_METHOD(GetRankIdByGroup, int, const std::string &); +ORIGIN_METHOD(GetDeviceId, int); +ORIGIN_METHOD(CreateCommForGroup, bool, const std::string &, const std::vector &); +ORIGIN_METHOD(DestroyHcclComm, void); + +class HcclCollectiveGroup { + public: + HcclCollectiveGroup(HcclCollectiveGroup const &) = delete; + HcclCollectiveGroup &operator=(const HcclCollectiveGroup &) = delete; + static HcclCollectiveGroup &instance(); + bool InitCollective(); + void FinalizeCollective(); + HcclComm GetGroupComm(const std::string &name); + int GetDeviceId() const; + int GetRankId(const std::string &name = kHcclWorldGroup) const; + int GetRankSize(const std::string &name = kHcclWorldGroup) const; + void CreateCommGroup(const std::string &name, const std::vector &ranks); + void DestroyCommGroup(); + const void *collective_handle() const { return collective_handle_; } + + private: + HcclCollectiveGroup() = default; + ~HcclCollectiveGroup() = default; + bool inited_ = false; + void *collective_handle_ = nullptr; + InitMPIFunObj init_mpi_ = nullptr; + FinalizeMPIFunObj finalize_mpi_ = nullptr; + GetGroupCommFunObj get_group_comm_ = nullptr; + GetGroupSizeFunObj get_group_size_ = nullptr; + GetRankIdByGroupFunObj get_rank_id_by_group_ = nullptr; + GetDeviceIdFunObj get_device_id_ = nullptr; + CreateCommForGroupFunObj create_comm_for_group_ = nullptr; + DestroyHcclCommFunObj destroy_hccl_comm_ = nullptr; +}; +} // namespace collective +} // namespace ascend +} // namespace device +} // namespace mindspore +#endif // MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_ASCEND_COLLECTIVE_H diff --git a/mindspore/ccsrc/runtime/device/ascend/distribute/collective_group_wrapper.cc b/mindspore/ccsrc/runtime/device/ascend/distribute/collective_group_wrapper.cc new file mode 100644 index 00000000000..a7beab312c5 --- /dev/null +++ b/mindspore/ccsrc/runtime/device/ascend/distribute/collective_group_wrapper.cc @@ -0,0 +1,28 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "runtime/device/ascend/distribute/collective_group_wrapper.h" + +void InitMPI() { MPICollective::instance().Init(); } +void FinalizeMPI() { MPICollective::instance().FinalizeMPI(); } +int GetRankIdByGroup(const std::string &name) { return MPICollective::instance().GetRankIdByGroup(name); } +int GetGroupSize(const std::string &name) { return MPICollective::instance().GetGroupSize(name); } +int GetDeviceId() { return MPICollective::instance().GetDeviceId(); } +HcclComm GetGroupComm(const std::string &name) { return MPICollective::instance().GetGroupComm(name); } +bool CreateCommForGroup(const std::string &name, const std::vector &ranks) { + return MPICollective::instance().CreateCommGroup(name, ranks); +} +void DestroyHcclComm() { MPICollective::instance().DestroyHcclComm(); } diff --git a/mindspore/ccsrc/runtime/device/ascend/distribute/collective_group_wrapper.h b/mindspore/ccsrc/runtime/device/ascend/distribute/collective_group_wrapper.h new file mode 100644 index 00000000000..419794d2ed3 --- /dev/null +++ b/mindspore/ccsrc/runtime/device/ascend/distribute/collective_group_wrapper.h @@ -0,0 +1,36 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_COLLECTIVE_GROUP_WRAPPER_H +#define MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_COLLECTIVE_GROUP_WRAPPER_H + +#include +#include +#include "runtime/device/ascend/distribute/mpi_collective_group.h" +#ifndef EXPORT_WRAPPER +#define EXPORT_WRAPPER __attribute__((visibility("default"))) +#endif +using MPICollective = mindspore::device::ascend::collective::MPICollective; + +extern "C" EXPORT_WRAPPER void InitMPI(); +extern "C" EXPORT_WRAPPER void FinalizeMPI(); +extern "C" EXPORT_WRAPPER int GetRankIdByGroup(const std::string &name); +extern "C" EXPORT_WRAPPER int GetGroupSize(const std::string &name); +extern "C" EXPORT_WRAPPER int GetDeviceId(); +extern "C" EXPORT_WRAPPER HcclComm GetGroupComm(const std::string &name); +extern "C" EXPORT_WRAPPER bool CreateCommForGroup(const std::string &name, const std::vector &ranks); +extern "C" EXPORT_WRAPPER void DestroyHcclComm(); +#endif // MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_COLLECTIVE_GROUP_WRAPPER_H diff --git a/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_collective_group.cc b/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_collective_group.cc new file mode 100644 index 00000000000..c8c578a5917 --- /dev/null +++ b/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_collective_group.cc @@ -0,0 +1,132 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "hccl/hccl.h" +#include "runtime/rt.h" +#include "runtime/device/ascend/distribute/mpi_collective_group.h" +namespace mindspore { +namespace device { +namespace ascend { +namespace collective { +MPICollective::MPICollective() : mpi_inited_(false), rank_id_(0), local_rank_id_(0), rank_size_(0) {} +void MPICollective::FinalizeMPI() { + group_info_.clear(); + group_comm_.clear(); + int finalized; + MPI_Finalized(&finalized); + if (finalized == 0) { + MPI_Finalize(); + } +} +void MPICollective::DestroyHcclComm() { + for (auto &it : group_comm_) { + CHECK_RET(HcclCommDestroy(it.second), HCCL_SUCCESS, "HcclCommDestroy failed"); + } +} +MPICollective &MPICollective::instance() { + static MPICollective instance; + return instance; +} +int MPICollective::GetRankIdByGroup(const std::string &name) { + CHECK_RET(group_info_.count(name), 1, "Failed to get MPI group rank by group name " + name); + return group_info_[name].first; +} +int MPICollective::GetGroupSize(const std::string &name) { + CHECK_RET(group_info_.count(name), 1, "Failed to get MPI group size by group name " + name); + return group_info_[name].second; +} +HcclComm MPICollective::GetGroupComm(const std::string &name) { + CHECK_RET(group_comm_.count(name), 1, "Failed to get MPI group comm by group name " + name); + return group_comm_[name]; +} +int MPICollective::GetDeviceId() { return local_rank_id_; } +bool MPICollective::Init() { + int init_flag = 0; + CHECK_RET(MPI_Initialized(&init_flag), MPI_SUCCESS, "Check mpi initialized fail!"); + if (init_flag == 0) { + CHECK_RET(MPI_Init(nullptr, nullptr), MPI_SUCCESS, "Failed to init mpi!"); + } + + CHECK_RET(MPI_Comm_group(MPI_COMM_WORLD, &comm_group_world_), MPI_SUCCESS, "comm_group_world_ init fail!"); + + CHECK_RET(MPI_Comm_rank(MPI_COMM_WORLD, &rank_id_), MPI_SUCCESS, "Failed to init mpi rank id!"); + + CHECK_RET(MPI_Comm_size(MPI_COMM_WORLD, &rank_size_), MPI_SUCCESS, "Failed to init mpi rank size!"); + AssignLocalRankID(); + group_info_["hccl_world_group"] = {rank_id_, rank_size_}; + mpi_inited_ = true; + return true; +} + +bool MPICollective::CreateCommGroup(const std::string &name, const std::vector &ranks) { + CHECK_RET(mpi_inited_, true, "HcclCollectiveGroup has not been inited."); + CHECK_RET(ranks.empty(), false, "Ranks is empty."); + std::vector group_ranks(ranks.begin(), ranks.end()); + CHECK_RET(group_comm_.count(name), 0, "Group comm has already been created."); + CHECK_RET(rtSetDevice(local_rank_id_), RT_ERROR_NONE, "Call rtSetDevice error."); + HcclRootInfo rootInfo; + if (static_cast(rank_id_) == ranks[0]) { + CHECK_RET(HcclGetRootInfo(&rootInfo), HCCL_SUCCESS, "HcclGetRootInfo failed."); + } + MPI_Group mpi_group = MPI_GROUP_NULL; + CHECK_RET(MPI_Group_incl(comm_group_world_, group_ranks.size(), group_ranks.data(), &mpi_group), MPI_SUCCESS, + "Create mpi group failed!"); + MPI_Comm mpi_group_comm; + + CHECK_RET(MPI_Comm_create(MPI_COMM_WORLD, mpi_group, &mpi_group_comm), MPI_SUCCESS, "Create mpi comm fail!"); + + CHECK_RET(MPI_Bcast(&rootInfo, sizeof(rootInfo), MPI_BYTE, 0, mpi_group_comm), MPI_SUCCESS, + "Mpi reduce_scatter failed!"); + + HcclComm group_hcomm = nullptr; + int group_rank[1]; + int global_rank[1] = {rank_id_}; + CHECK_RET(MPI_Group_translate_ranks(comm_group_world_, 1, global_rank, mpi_group, group_rank), MPI_SUCCESS, + "Failed to translate global rank to group rank."); + if (group_rank[0] == MPI_UNDEFINED) { + return false; + } + + CHECK_RET(HcclCommInitRootInfo(ranks.size(), &rootInfo, static_cast(group_rank[0]), &group_hcomm), + HCCL_SUCCESS, "HcclCommInitRootInfo failed."); + group_comm_[name] = group_hcomm; + group_info_[name] = {group_rank[0], static_cast(ranks.size())}; + return true; +} +void MPICollective::AssignLocalRankID() { + char host_name[MAX_HOSTNAME_LEN] = {0}; + CHECK_RET(gethostname(host_name, MAX_HOSTNAME_LEN), MPI_SUCCESS, "Getting host name failed!"); + size_t host_hash = std::hash()(host_name); + + const int kRankSize = rank_size_; + size_t all_host_hashs[kRankSize]; + all_host_hashs[rank_id_] = host_hash; + CHECK_RET(MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, all_host_hashs, sizeof(size_t), MPI_BYTE, MPI_COMM_WORLD), + MPI_SUCCESS, "MPI_Allgather host hash failed."); + for (int global_rank = 0; global_rank < kRankSize; global_rank++) { + if (global_rank == rank_id_) { + break; + } + if (all_host_hashs[global_rank] == all_host_hashs[rank_id_]) { + local_rank_id_++; + } + } + return; +} +} // namespace collective +} // namespace ascend +} // namespace device +} // namespace mindspore diff --git a/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_collective_group.h b/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_collective_group.h new file mode 100644 index 00000000000..3bdcb3a32c2 --- /dev/null +++ b/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_collective_group.h @@ -0,0 +1,74 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_MPI_COLLECTIVE_INIT_H +#define MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_MPI_COLLECTIVE_INIT_H + +#include +#include +#include +#include +#include +#include +#include +#include "hccl/hccl_types.h" +#include "pybind11/pybind11.h" +namespace mindspore { +namespace device { +namespace ascend { +namespace collective { +constexpr int MAX_HOSTNAME_LEN = 1024; +class MPICollective { + public: + MPICollective(MPICollective const &) = delete; + MPICollective &operator=(const MPICollective &) = delete; + static MPICollective &instance(); + void AssignLocalRankID(); + bool Init(); + void FinalizeMPI(); + int GetRankIdByGroup(const std::string &name); + int GetGroupSize(const std::string &name); + HcclComm GetGroupComm(const std::string &name); + int GetDeviceId(); + bool CreateCommGroup(const std::string &name, const std::vector &ranks); + void DestroyHcclComm(); + + private: + MPICollective(); + ~MPICollective() = default; + bool mpi_inited_; + int rank_id_; + int local_rank_id_; + int rank_size_; + MPI_Group comm_group_world_; + std::map> group_info_; + std::map group_comm_; +}; +#define CHECK_RET(expression, result, message) \ + { \ + auto ret = (expression); \ + if (ret != result) { \ + std::ostringstream oss; \ + oss << "Error in file " << __FILE__ << " | Error on line " << __LINE__ \ + << " | Ascend collective Error: " << message << " | Error Number " << ret; \ + pybind11::pybind11_fail(oss.str()); \ + } \ + } +} // namespace collective +} // namespace ascend +} // namespace device +} // namespace mindspore +#endif // MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_COLLECTIVE_INIT_H diff --git a/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_pycc.cc b/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_pycc.cc new file mode 100644 index 00000000000..b9719aab1e6 --- /dev/null +++ b/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_pycc.cc @@ -0,0 +1,41 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "runtime/device/ascend/distribute/mpi_pycc.h" +#include +namespace mindspore { +namespace device { +namespace ascend { +namespace collective { +MpiPycc &MpiPycc::instance() { + static MpiPycc instance; + return instance; +} + +int MpiPycc::GetDeviceID() { return GetDeviceId(); } +int MpiPycc::GetRankId(const std::string &group) { return GetRankIdByGroup(group); } +int MpiPycc::GetRankSize(const std::string &group) { return GetGroupSize(group); } + +// cppcheck-suppress syntaxError +PYBIND11_MODULE(_ascend_mpi, mpi_initializer) { + mpi_initializer.def("get_device_id", &MpiPycc::GetDeviceID, "get device id"); + mpi_initializer.def("get_rank_id", &MpiPycc::GetRankId, "get rank id"); + mpi_initializer.def("get_rank_size", &MpiPycc::GetRankSize, "get rank size"); +} +} // namespace collective +} // namespace ascend +} // namespace device +} // namespace mindspore diff --git a/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_pycc.h b/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_pycc.h new file mode 100644 index 00000000000..5282a3bbc23 --- /dev/null +++ b/mindspore/ccsrc/runtime/device/ascend/distribute/mpi_pycc.h @@ -0,0 +1,44 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_MPI_PYCC_H +#define MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_MPI_PYCC_H + +#include +#include "runtime/device/ascend/distribute/collective_group_wrapper.h" + +namespace mindspore { +namespace device { +namespace ascend { +namespace collective { +class MpiPycc { + public: + MpiPycc(MpiPycc const &) = delete; + MpiPycc &operator=(const MpiPycc &) = delete; + static MpiPycc &instance(); + static int GetDeviceID(); + static int GetRankId(const std::string &group); + static int GetRankSize(const std::string &group); + + private: + MpiPycc() = default; + ~MpiPycc() = default; +}; +} // namespace collective +} // namespace ascend +} // namespace device +} // namespace mindspore +#endif // MINDSPORE_MINDSPORE_CCSRC_RUNTIME_DEVICE_ASCEND_DISTRIBUTE_MPI_PYCC_H diff --git a/mindspore/ccsrc/runtime/hccl_adapter/hccl_adapter.cc b/mindspore/ccsrc/runtime/hccl_adapter/hccl_adapter.cc index cbb5dd43fc5..2fe05293e58 100644 --- a/mindspore/ccsrc/runtime/hccl_adapter/hccl_adapter.cc +++ b/mindspore/ccsrc/runtime/hccl_adapter/hccl_adapter.cc @@ -32,22 +32,6 @@ static constexpr const char *kHcclPluginFileName = "libhccl_plugin.so"; static constexpr const char *kHcclDeployModeEnv = "DEPLOY_MODE"; static constexpr const char *kHcclAlgoEnv = "HCCL_ALGO"; -inline static std::string GetDlErrorMsg() { - const char *result = dlerror(); - return (result == nullptr) ? "Unknown" : result; -} - -template -static T DlsymWithCast(void *handle, const char *symbol_name) { - T symbol = reinterpret_cast(dlsym(handle, symbol_name)); - if (symbol == nullptr) { - MS_LOG(EXCEPTION) << "Dlsym symbol " << symbol_name << " failed, result = " << GetDlErrorMsg(); - } - return symbol; -} - -#define DlsymFuncObj(func_name) DlsymWithCast(plugin_handle_, k##func_name##Name); - static std::map GenHcclOptions(uint32_t device_id, std::string_view rank_id, std::string_view rank_file) { auto env_deploy_mode = mindspore::common::GetEnv(kHcclDeployModeEnv); @@ -92,21 +76,21 @@ void HcclAdapter::InitPlugin() { MS_LOG(EXCEPTION) << "Dlopen " << kHcclPluginFileName << " failed, result = " << GetDlErrorMsg(); } - init_hcom_graph_adapter_ = DlsymFuncObj(InitHcomGraphAdapter); - finalize_hcom_graph_adapter_ = DlsymFuncObj(FinalizeHcomGraphAdapter); - get_hccl_kernel_info_store_ = DlsymFuncObj(GetHcclKernelInfoStore); - get_all_kernel_builder_ = DlsymFuncObj(GetAllKernelBuilder); - init_hccl_comm_ = DlsymFuncObj(HcclCommInitClusterInfo); - finalize_hccl_comm_ = DlsymFuncObj(HcclCommDestroy); - launch_hccl_broadcast_ = DlsymFuncObj(HcclBroadcast); - launch_hccl_all_reduce_ = DlsymFuncObj(HcclAllReduce); - hccl_create_group_ = DlsymFuncObj(HcomCreateGroup); - hccl_destroy_group_ = DlsymFuncObj(HcomDestroyGroup); - hccl_get_rank_id_ = DlsymFuncObj(HcomGetRankId); - hccl_get_rank_size_ = DlsymFuncObj(HcomGetRankSize); - hccl_exec_initialize_ = DlsymFuncObj(HcomExecInitialize); - hccl_exec_finalize_ = DlsymFuncObj(HcomExecFinalize); - hccl_exec_enqueue_op_ = DlsymFuncObj(HcomExecEnqueueOperation); + init_hcom_graph_adapter_ = DlsymFuncObj(InitHcomGraphAdapter, plugin_handle_); + finalize_hcom_graph_adapter_ = DlsymFuncObj(FinalizeHcomGraphAdapter, plugin_handle_); + get_hccl_kernel_info_store_ = DlsymFuncObj(GetHcclKernelInfoStore, plugin_handle_); + get_all_kernel_builder_ = DlsymFuncObj(GetAllKernelBuilder, plugin_handle_); + init_hccl_comm_ = DlsymFuncObj(HcclCommInitClusterInfo, plugin_handle_); + finalize_hccl_comm_ = DlsymFuncObj(HcclCommDestroy, plugin_handle_); + launch_hccl_broadcast_ = DlsymFuncObj(HcclBroadcast, plugin_handle_); + launch_hccl_all_reduce_ = DlsymFuncObj(HcclAllReduce, plugin_handle_); + hccl_create_group_ = DlsymFuncObj(HcomCreateGroup, plugin_handle_); + hccl_destroy_group_ = DlsymFuncObj(HcomDestroyGroup, plugin_handle_); + hccl_get_rank_id_ = DlsymFuncObj(HcomGetRankId, plugin_handle_); + hccl_get_rank_size_ = DlsymFuncObj(HcomGetRankSize, plugin_handle_); + hccl_exec_initialize_ = DlsymFuncObj(HcomExecInitialize, plugin_handle_); + hccl_exec_finalize_ = DlsymFuncObj(HcomExecFinalize, plugin_handle_); + hccl_exec_enqueue_op_ = DlsymFuncObj(HcomExecEnqueueOperation, plugin_handle_); } void HcclAdapter::FinalizePlugin() { diff --git a/mindspore/ccsrc/runtime/hccl_adapter/plugin/hccl_plugin.h b/mindspore/ccsrc/runtime/hccl_adapter/plugin/hccl_plugin.h index 7ba10c0b8bd..f3f521cf078 100644 --- a/mindspore/ccsrc/runtime/hccl_adapter/plugin/hccl_plugin.h +++ b/mindspore/ccsrc/runtime/hccl_adapter/plugin/hccl_plugin.h @@ -22,6 +22,7 @@ #include #include "external/ge/ge_api_types.h" #include "hccl/hccl.h" +#include "utils/dlopen_macro.h" constexpr const char *kHcclOpsKernelInfoStore = "ops_kernel_info_hccl"; @@ -38,22 +39,6 @@ using OptionsType = std::map; using OpsKernelBuilderMap = std::map>; using HExecCallBack = std::function; -#define PLUGIN_METHOD(name, return_type, params...) \ - extern "C" { \ - __attribute__((visibility("default"))) return_type Plugin##name(params); \ - } \ - constexpr const char *k##name##Name = "Plugin" #name; \ - using name##FunObj = std::function; \ - using name##FunPtr = return_type (*)(params); - -#define ORIGIN_METHOD(name, return_type, params...) \ - extern "C" { \ - return_type name(params); \ - } \ - constexpr const char *k##name##Name = #name; \ - using name##FunObj = std::function; \ - using name##FunPtr = return_type (*)(params); - PLUGIN_METHOD(InitHcomGraphAdapter, ge::Status, const OptionsType &); PLUGIN_METHOD(FinalizeHcomGraphAdapter, ge::Status); PLUGIN_METHOD(GetHcclKernelInfoStore, void, std::shared_ptr *); diff --git a/mindspore/ccsrc/utils/dlopen_macro.h b/mindspore/ccsrc/utils/dlopen_macro.h new file mode 100644 index 00000000000..8f18ad5ac42 --- /dev/null +++ b/mindspore/ccsrc/utils/dlopen_macro.h @@ -0,0 +1,55 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef MINDSPORE_CCSRC_UTILS_DLOPEN_MACRO_H +#define MINDSPORE_CCSRC_UTILS_DLOPEN_MACRO_H +#include +#include +#include +#include "utils/log_adapter.h" + +#define PLUGIN_METHOD(name, return_type, params...) \ + extern "C" { \ + __attribute__((visibility("default"))) return_type Plugin##name(params); \ + } \ + constexpr const char *k##name##Name = "Plugin" #name; \ + using name##FunObj = std::function; \ + using name##FunPtr = return_type (*)(params); + +#define ORIGIN_METHOD(name, return_type, params...) \ + extern "C" { \ + return_type name(params); \ + } \ + constexpr const char *k##name##Name = #name; \ + using name##FunObj = std::function; \ + using name##FunPtr = return_type (*)(params); + +inline static std::string GetDlErrorMsg() { + const char *result = dlerror(); + return (result == nullptr) ? "Unknown" : result; +} + +template +static T DlsymWithCast(void *handle, const char *symbol_name) { + T symbol = reinterpret_cast(dlsym(handle, symbol_name)); + if (symbol == nullptr) { + MS_LOG(EXCEPTION) << "Dlsym symbol " << symbol_name << " failed, result = " << GetDlErrorMsg(); + } + return symbol; +} + +#define DlsymFuncObj(func_name, plugin_handle) DlsymWithCast(plugin_handle, k##func_name##Name); +#endif // MINDSPORE_CCSRC_UTILS_DLOPEN_MACRO_H