当前位置: 首页 > news >正文

【分布式】小白看Ring算法 - 03

相关系列

【分布式】NCCL部署与测试 - 01
【分布式】入门级NCCL多机并行实践 - 02
【分布式】小白看Ring算法 - 03
【分布式】大模型分布式训练入门与实践 - 04

概述

NCCL(NVIDIA Collective Communications Library)是由NVIDIA开发的一种用于多GPU间通信的库。NCCL的RING算法是NCCL库中的一种通信算法,用于在多个GPU之间进行环形通信。

RING算法的基本思想是将多个GPU连接成一个环形结构,每个GPU与相邻的两个GPU进行通信。数据沿着环形结构传递,直到到达发送方的位置。这样的环形结构可以有效地利用GPU之间的带宽,提高通信的效率。

RING算法的步骤如下:

数据拷贝
数据沿着环形路径传递
传输完成
进行下一轮通信/结束通信过程
初始化
通信缓冲区
等待
接收方

Scatter-Reduce

以Scatter-Reduce为例,假设有4张GPU,RANK_NUM=4。
则需要根据RANK_NUM把每张CPU划分为4个chunk。
为什么要这么划分?

在 NCCL 中,划分 chunk 的数量与 GPU 的数量相关联,这是因为 chunk 的目的是将大的消息划分为多个小的数据块,以便并行处理和降低通信的延迟。这种划分通常会基于 GPU 的数量,以确保每个 GPU 可以处理到一部分数据块,从而提高整体的通信效率。

  1. 并行性: 划分 chunk 可以增加通信的并行性。每个 GPU 处理自己的数据块,不同的 GPU 可以并行地执行通信操作,从而提高整体的吞吐量。
  2. 减少延迟: 较小的数据块通常可以更快地传输,因此通过划分 chunk,可以减少每个通信操作的延迟。这对于一些对通信延迟敏感的应用程序是至关重要的。
  3. 资源分配: NCCL 可能会根据 GPU 的数量来分配适当的资源,例如内存等。通过划分 chunk,可以更好地管理这些资源。
  4. Load Balancing: 均衡负载是分布式系统中的一个关键问题。通过根据 GPU 的数量划分 chunk,可以更容易地实现负载均衡,确保每个 GPU 处理的工作量相对均匀。

划分了chunk以后,我们一次RING的通路将会走通4块GPU,每次只传输一块chunk的数据。这样需要走很多次通路才能把所有数据传输完。
假如 ringIx=0,第一次循环到第三次循环时:
在这里插入图片描述

我们将绿色视为这次循环需要传输的数据。
数据ABCD在不同的GPU中流通。
最终达到以下情况,scatter-reduce就完成了:
在这里插入图片描述
将图中蓝色部分输出,就完成了一次ring算法下的Scatter-Reduce。

当然,如果要做All-Reduce,此时不需要继续按照原来的规则计算类,理论上只需要再算一次All-Gather,就能把蓝色的块分发给其他几块卡。All-Reduce的相关讲解网络上很多。此处就不讲了。

NCCL代码流程

1
1
1
1
2
2
2
2
4
4
4
4
5
5
5
5
6
6
6
6
7
7
7
7
8
8
8
8
9
9
9
9
10
10
10
10
11
11
11
11
12
12
12
12
13
13
13
13
rank0:fillInfo
bootstrap AllGather
rank1:fillInfo
rank2:fillInfo
rank3:fillInfo
rank0:getSystem
rank1:getSystem
rank2:getSystem
rank3:getSystem
rank0:computePath
rank1:computePath
rank2:computePath
rank3:computePath
rank0:search channel
rank1:search channel
rank2:search channel
rank3:search channel
bootstrap AllGather
rank0:connect
rank1:connect
rank2:connect
rank3:connect
rank0:setupChannel
rank1:setupChannel
rank2:setupChannel
rank3:setupChannel
rank0:p2pSetup
rank1:p2pSetup
rank2:p2pSetup
rank3:p2pSetup
rank0:tuneModel
rank1:tuneModel
rank2:tuneModel
rank3:tuneModel
rank0:p2pChannel
rank1:p2pChannel
rank2:p2pChannel
rank3:p2pChannel
bootstrap IntraNodeBarrier
rank0:NetProxy
rank1:NetProxy
rank2:NetProxy
rank3:NetProxy

fillInfo:
这段代码在init.cc中

static ncclResult_t fillInfo(struct ncclComm* comm, struct ncclPeerInfo* info, uint64_t commHash) {info->rank = comm->rank;CUDACHECK(cudaGetDevice(&info->cudaDev));info->hostHash=getHostHash()+commHash;info->pidHash=getPidHash()+commHash;// Get the device MAJOR:MINOR of /dev/shm so we can use that// information to decide whether we can use SHM for inter-process// communication in a container environmentstruct stat statbuf;SYSCHECK(stat("/dev/shm", &statbuf), "stat");info->shmDev = statbuf.st_dev;info->busId = comm->busId;NCCLCHECK(ncclGpuGdrSupport(&info->gdrSupport));return ncclSuccess;
}

这段代码的目的是为了获取和存储与通信相关的信息,以便在NCCL通信中使用。其中包括设备标识、主机哈希、进程ID哈希、共享内存设备标识、总线ID以及对GDR的支持情况等。

在initTransportsRank中,搜索完信息并作第一次AllGather, 收集所有通信节点的信息。
然后再为通信组分配额外的内存,以存储每个通信节点的信息(包括一个额外的用于表示CollNet root的位置)。
遍历节点和复制信息时,需要检查是否存在相同主机哈希和总线ID的重复GPU。如果是,发出警告并返回ncclInvalidUsage错误。

后面的一系列过程就是计算路径,然后这里涉及一些搜索算法,通常会将BFS搜索到的路径都存在一个位置,选择更优的路径。
搜索时也会根据实际情况判断选择ring算法或者tree算法。
搜索内容可能是无穷的,因此NCCL设置了一个超时时间,超过该时间则终端搜索。
完成路径的计算后,再做一次AllGather。

来到scatter-reduce的实现部分:

		size_t realChunkSize;if (Proto::Id == NCCL_PROTO_SIMPLE) {realChunkSize = min(chunkSize, divUp(size-gridOffset, nChannels));realChunkSize = roundUp(realChunkSize, (nthreads-WARP_SIZE)*sizeof(uint64_t)/sizeof(T));}else if (Proto::Id == NCCL_PROTO_LL)realChunkSize = size-gridOffset < loopSize ? args->coll.lastChunkSize : chunkSize;else if (Proto::Id == NCCL_PROTO_LL128)realChunkSize = min(divUp(size-gridOffset, nChannels*minChunkSizeLL128)*minChunkSizeLL128, chunkSize);realChunkSize = int(realChunkSize);ssize_t chunkOffset = gridOffset + bid*int(realChunkSize);

这里涉及了NCCL协议的通信模式:
一共有三种,分别是NCCL_PROTO_SIMPLE、NCCL_PROTO_LL和NCCL_PROTO_LL128。

NCCL_PROTO_SIMPLE:

描述: 使用简单的通信协议。
差异点: 计算realChunkSize时,采用了一些特殊的逻辑,其中min(chunkSize, divUp(size-gridOffset, nChannels))用于确定实际的块大小,并通过roundUp调整为合适的大小。这可能涉及到性能和资源的考虑,以及对通信模式的调整。

NCCL_PROTO_LL:

描述: 使用连续链表(Linked List,LL)的通信协议。
差异点: 在计算realChunkSize时,首先检查size-gridOffset < loopSize条件,如果为真,则使用args->coll.lastChunkSize,否则使用默认的chunkSize。这可能与LL协议的特性有关,具体考虑了循环的情况。
NCCL_PROTO_LL128:

描述: 使用连续链表的通信协议,每次传输128字节。
差异点: 计算realChunkSize时,采用了min(divUp(size-gridOffset, nChannels*minChunkSizeLL128)*minChunkSizeLL128, chunkSize)的逻辑。这考虑了128字节的限制,以及对通信块大小的一些限制。
总体来说,这三种协议模式的区别主要体现在计算realChunkSize的逻辑上,这可能受到性能、资源利用、通信模式等方面的不同考虑。具体选择哪种协议模式通常取决于系统的特性和应用场景的需求。

Protocol ModeDescriptionCalculation of realChunkSize
NCCL_PROTO_SIMPLEUses a simple communication protocol.realChunkSize = roundUp(min(chunkSize, divUp(size-gridOffset, nChannels)), (nthreads-WARP_SIZE)*sizeof(uint64_t)/sizeof(T))
NCCL_PROTO_LLUses a linked list (LL) communication protocol.realChunkSize = size-gridOffset < loopSize ? args->coll.lastChunkSize : chunkSize
NCCL_PROTO_LL128Uses a linked list (LL) communication protocol, with each transfer involving 128 bytes.realChunkSize = min(divUp(size-gridOffset, nChannels*minChunkSizeLL128)*minChunkSizeLL128, chunkSize)

最后是正式计算部分:

 /////////////// begin ReduceScatter steps ///////////////ssize_t offset;int nelem = min(realChunkSize, size-chunkOffset);int rankDest;// step 0: push data to next GPUrankDest = ringRanks[nranks-1];offset = chunkOffset + rankDest * size;prims.send(offset, nelem);// k-2 steps: reduce and copy to next GPUfor (int j=2; j<nranks; ++j) {rankDest = ringRanks[nranks-j];offset = chunkOffset + rankDest * size;prims.recvReduceSend(offset, nelem);}// step k-1: reduce this buffer and data, which will produce the final resultrankDest = ringRanks[0];offset = chunkOffset + rankDest * size;prims.recvReduceCopy(offset, chunkOffset, nelem, /*postOp=*/true);

ssize_t offset; int nelem = min(realChunkSize, size-chunkOffset); int rankDest;:

offset 是一个偏移量变量,用于指定数据在通信缓冲区中的位置。
nelem 表示每次操作的元素个数,取 realChunkSize 和 size-chunkOffset 的较小值。
rankDest 是目标GPU的排名。

第一步:将数据推送到下一个GPU。
计算目标GPU的排名 rankDest 和在通信缓冲区中的偏移量 offset。
调用 prims.send 函数,将数据从当前GPU发送到目标GPU。
// k-2 steps: reduce and copy to next GPU:

第2到第k-1步:
将数据在环形路径上经过各个GPU节点,依次进行Reduce操作,并将结果复制到下一个GPU。
通过循环,依次计算目标GPU的排名 rankDest 和在通信缓冲区中的偏移量 offset。
调用 prims.recvReduceSend 函数,接收数据并执行Reduce操作,然后将结果发送到下一个GPU。

第k-1步:
将最后一个GPU的数据进行Reduce操作,得到最终的结果。
计算目标GPU的排名 rankDest 和在通信缓冲区中的偏移量 offset。
调用 prims.recvReduceCopy 函数,接收数据并执行Reduce操作,然后将结果复制到指定的位置,最终产生最终的ReduceScatter结果。

在实际运行中,我们在host端的代码只是规定计算流,当这些定义好的原子操作加入到stream中去以后,就由固定的流来分配实际运行的情况了。

加入Barria,在本地(intra-node)执行一个屏障操作,确保同一节点内的所有GPU都达到了同步点。

 // Compute time models for algorithm and protocol combinationsNCCLCHECK(ncclTopoTuneModel(comm, minCompCap, maxCompCap, &treeGraph, &ringGraph, &collNetGraph));// Compute nChannels per peer for p2pNCCLCHECK(ncclTopoComputeP2pChannels(comm));if (ncclParamNvbPreconnect()) {// Connect p2p when using NVB pathint nvbNpeers;int* nvbPeers;NCCLCHECK(ncclTopoGetNvbGpus(comm->topo, comm->rank, &nvbNpeers, &nvbPeers));for (int r=0; r<nvbNpeers; r++) {int peer = nvbPeers[r];int delta = (comm->nRanks + (comm->rank-peer)) % comm->nRanks;for (int c=0; c<comm->p2pnChannelsPerPeer; c++) {int channelId = (delta+comm->p2pChannels[c]) % comm->p2pnChannels;if (comm->channels[channelId].peers[peer].recv[0].connected == 0) { // P2P uses only 1 connectorcomm->connectRecv[peer] |= (1<<channelId);}}delta = (comm->nRanks - (comm->rank-peer)) % comm->nRanks;for (int c=0; c<comm->p2pnChannelsPerPeer; c++) {int channelId = (delta+comm->p2pChannels[c]) % comm->p2pnChannels;if (comm->channels[channelId].peers[peer].send[0].connected == 0) { // P2P uses only 1 connectorcomm->connectSend[peer] |= (1<<channelId);}}}NCCLCHECK(ncclTransportP2pSetup(comm, NULL, 0));free(nvbPeers);}NCCLCHECK(ncclCommSetIntraProc(comm, intraProcRank, intraProcRanks, intraProcRank0Comm));/* Local intra-node barrier */NCCLCHECK(bootstrapBarrier(comm->bootstrap, comm->intraNodeGlobalRanks, intraNodeRank, intraNodeRanks, (int)intraNodeRank0pidHash));if (comm->nNodes) NCCLCHECK(ncclProxyCreate(comm));

以上就是整个scatter-reduce的流程。

相关系列

【分布式】NCCL部署与测试 - 01
【分布式】入门级NCCL多机并行实践 - 02
【分布式】小白看Ring算法 - 03
【分布式】大模型分布式训练入门与实践 - 04

相关文章:

【分布式】小白看Ring算法 - 03

相关系列 【分布式】NCCL部署与测试 - 01 【分布式】入门级NCCL多机并行实践 - 02 【分布式】小白看Ring算法 - 03 【分布式】大模型分布式训练入门与实践 - 04 概述 NCCL&#xff08;NVIDIA Collective Communications Library&#xff09;是由NVIDIA开发的一种用于多GPU间…...

使用Git bash切换Gitee、GitHub多个Git账号

Git是分布式代码管理工具&#xff0c;使用命令行的方式提交commit、revert回滚代码。这里介绍使用Git bash软件来切换Gitee、GitHub账号。     假设在gitee.com上的邮箱是alicefoxmail.com 、用户名为alice&#xff1b;在github上的邮箱是bobfoxmail.com、用户名为bob。 账号…...

【RtpRtcp】1: webrtc m79:audio的ChannelReceive 创建并使用

m79中,RtpRtcp::Create 的调用很少 不知道谁负责创建ChannelReceiveclass ChannelReceive : public ChannelReceiveInterface,public MediaTransportAudioSinkInterface {接收编码后的音频帧:接收rtcp包:...

Ubuntu系统安装docker

1.检查是否安装老版本 检查卸载老版本docker ubuntu下自带了docker的库&#xff0c;不需要添加新的源。 但是ubuntu自带的docker版本太低&#xff0c;需要先卸载旧的再安装新的。 apt-get remove docker docker-engine docker.io containerd runc 如果不能正常卸载&#x…...

如何访问linux上的web服务

1.获取服务运行端口 例如8080 2.如果时vmware 需要先配置转发端口和主机ip 主机ip需要未使用的 例如&#xff1a; 3.查看虚拟机防火墙设置 centos8 为例 &#xff1a; firewall-cmd --zonepublic --list-ports 查看放通端口 如果没有放通 firewall-cmd --zonepublic --add-p…...

Vatee万腾的数字化掌舵:Vatee科技解决方案的全面引领

随着数字化时代的到来&#xff0c;Vatee万腾凭借其卓越的科技实力和全面的解决方案&#xff0c;成功地在数字化探索的航程中掌舵引领。 首先&#xff0c;Vatee万腾以其强大的数字化科技实力成为行业的引领者。vatee万腾不仅在人工智能、大数据分析、云计算等前沿领域取得了显著…...

YOLOv5 第Y6周 模型改进

&#x1f368; 本文为[&#x1f517;365天深度学习训练营学习记录博客 &#x1f366; 参考文章&#xff1a;365天深度学习训练营 &#x1f356; 原作者&#xff1a;[K同学啊] &#x1f680; 文章来源&#xff1a;[K同学的学习圈子](https://www.yuque.com/mingtian-fkmxf/zxwb4…...

Unity Android FireBase bugly报错查询

报错如下图&#xff0c;注意&#xff0c;标红的三处 使用的il2cpp和架构是arm64-v8a 那我们就可以根据这些去找对应的符号表&#xff0c;在unity安装目录下 Unity2020.3.33f1\Editor\Data\PlaybackEngines\AndroidPlayer\Variations\il2cpp\Release\Symbols\arm64-v8a 找到l…...

React中如何解决点击<Tree>节点前面三角区域不触发onClick事件

React中如何解决点击节点前面三角区域不触发onClick事件&#xff0c;如何区别‘左边’和‘右边’区域点击逻辑呢&#xff1f;&#xff08;Tree引用开源组件TDesign&#xff09; 只需要在onClick里面加限制一下就行&#xff1a; <TreeexpandMutexactivabletransitiondata{t…...

如何利用4G路由器构建茶饮连锁店物联网

随着年轻消费群体的增长&#xff0c;加上移动互联网营销的助推&#xff0c;各类新式奶茶消费风靡大街小巷&#xff0c;也促进了品牌奶茶连锁店的快速扩张。 在店铺快速扩张的局势下&#xff0c;品牌总部对于各间连锁店的零售统计、营销规划、物流调配、卫生监测、安全管理等事务…...

【2024系统架构设计】 系统架构设计师第二版-大数据架构理论设计与实践

目录 1 传统数据库的数据过载问题 2 大数据处理系统 3 Lambda架构 4 Kappa架构...

正整数分解

题目编号&#xff1a;Exp08-Basic01&#xff0c;GJBook3-12-05 题目名称&#xff1a;正整数分解 题目描述&#xff1a;正整数n&#xff0c;按第一项递减的顺序依次输出其和等于n的所有不增的正整数和式。 输入&#xff1a;一个正整数n&#xff08;0<n≤15&#xff09;。 …...

基于51单片机电子钟闹钟LCD1602显示proteus仿真设计

基于51单片机的LCD1602电子钟闹钟proteus仿真设计 基于51单片机的LCD1602电子钟闹钟proteus仿真设计功能介绍&#xff1a;仿真图&#xff1a;原理图&#xff1a;设计报告&#xff1a;程序&#xff1a;器件清单&#xff1a;资料清单&&下载链接&#xff1a; 基于51单片机…...

第三节-Android10.0 Binder通信原理(三)-ServiceManager篇

1、概述 在Android中&#xff0c;系统提供的服务被包装成一个个系统级service&#xff0c;这些service往往会在设备启动之时添加进Android系统&#xff0c;当某个应用想要调用系统某个服务的功能时&#xff0c;往往是向系统发出请求&#xff0c;调用该服务的外部接口。在上一节…...

使用XHProf查找PHP性能瓶颈

使用XHProf查找PHP性能瓶颈 XHProf是facebook 开发的一个测试php性能的扩展&#xff0c;本文记录了在PHP应用中使用XHProf对PHP进行性能优化&#xff0c;查找性能瓶颈的方法。 下载 网上很多是编译安装xhprof-0.9.4版本&#xff0c;应该是用php5&#xff0c;在php8.0下编译x…...

矩阵论(Matrix)

​ 大纲 矩阵微积分&#xff1a;多元微积分的一种特殊表达&#xff0c;尤其是在矩阵空间上进行讨论的时候逆矩阵(inverse matrix)矩阵分解&#xff1a;特征分解&#xff08;Eigendecomposition&#xff09;&#xff0c;又称谱分解&#xff08;Spectral decomposition&#xf…...

解决Emmy Lua插件在IDEA或 Reder 没有代码提示的问题(设置文件关联 增加对.lua.txt文件的支持)

目录 Reder版本2019.x Reder版本2021.1.5x Reder版本2019.x 解决Emmy Lua插件在IDEA或 Reder 没有代码提示的问题(设置文件关联 增加对.lua.txt文件的支持) Reder版本2021.1.5x 解决Emmy Lua插件在IDEA或 Reder 没有代码提示的问题(设置文件关联 增加对.lua.txt文件的支持)…...

macos端文件夹快速访问工具 Default Folder X 最新for mac

Default Folder X 是一款实用的工具&#xff0c;提供了许多增强功能和快捷方式&#xff0c;使用户能够更高效地浏览和管理文件。它的快速导航、增强的文件对话框、自定义设置和快捷键等功能&#xff0c;可以大大提升用户的工作效率和文件管理体验。 快速导航和访问&#xff1a;…...

树形 DP:树的直径

leetCode 104.二叉树的最大深度104. 二叉树的最大深度 - 力扣&#xff08;LeetCode&#xff09; class Solution { public:int maxDepth(TreeNode* root) {if(root nullptr) return 0;int lDepth maxDepth(root->left);int rDepth maxDepth(root->right);return max(l…...

【Python百宝箱】第三维度的魔法:探索Python游戏世界

Python在游戏开发中的魔力 前言 游戏开发一直是计算机科学中最引人入胜和具有挑战性的领域之一。随着技术的不断进步&#xff0c;开发者们寻找着更快、更灵活的工具来实现他们的创意。在这个探索的过程中&#xff0c;Python以其简洁、易学和强大的特性成为了游戏开发的热门选…...

网络编程(Modbus进阶)

思维导图 Modbus RTU&#xff08;先学一点理论&#xff09; 概念 Modbus RTU 是工业自动化领域 最广泛应用的串行通信协议&#xff0c;由 Modicon 公司&#xff08;现施耐德电气&#xff09;于 1979 年推出。它以 高效率、强健性、易实现的特点成为工业控制系统的通信标准。 包…...

React Native 导航系统实战(React Navigation)

导航系统实战&#xff08;React Navigation&#xff09; React Navigation 是 React Native 应用中最常用的导航库之一&#xff0c;它提供了多种导航模式&#xff0c;如堆栈导航&#xff08;Stack Navigator&#xff09;、标签导航&#xff08;Tab Navigator&#xff09;和抽屉…...

QMC5883L的驱动

简介 本篇文章的代码已经上传到了github上面&#xff0c;开源代码 作为一个电子罗盘模块&#xff0c;我们可以通过I2C从中获取偏航角yaw&#xff0c;相对于六轴陀螺仪的yaw&#xff0c;qmc5883l几乎不会零飘并且成本较低。 参考资料 QMC5883L磁场传感器驱动 QMC5883L磁力计…...

Ascend NPU上适配Step-Audio模型

1 概述 1.1 简述 Step-Audio 是业界首个集语音理解与生成控制一体化的产品级开源实时语音对话系统&#xff0c;支持多语言对话&#xff08;如 中文&#xff0c;英文&#xff0c;日语&#xff09;&#xff0c;语音情感&#xff08;如 开心&#xff0c;悲伤&#xff09;&#x…...

深度学习习题2

1.如果增加神经网络的宽度&#xff0c;精确度会增加到一个特定阈值后&#xff0c;便开始降低。造成这一现象的可能原因是什么&#xff1f; A、即使增加卷积核的数量&#xff0c;只有少部分的核会被用作预测 B、当卷积核数量增加时&#xff0c;神经网络的预测能力会降低 C、当卷…...

Chromium 136 编译指南 Windows篇:depot_tools 配置与源码获取(二)

引言 工欲善其事&#xff0c;必先利其器。在完成了 Visual Studio 2022 和 Windows SDK 的安装后&#xff0c;我们即将接触到 Chromium 开发生态中最核心的工具——depot_tools。这个由 Google 精心打造的工具集&#xff0c;就像是连接开发者与 Chromium 庞大代码库的智能桥梁…...

LOOI机器人的技术实现解析:从手势识别到边缘检测

LOOI机器人作为一款创新的AI硬件产品&#xff0c;通过将智能手机转变为具有情感交互能力的桌面机器人&#xff0c;展示了前沿AI技术与传统硬件设计的完美结合。作为AI与玩具领域的专家&#xff0c;我将全面解析LOOI的技术实现架构&#xff0c;特别是其手势识别、物体识别和环境…...

rknn toolkit2搭建和推理

安装Miniconda Miniconda - Anaconda Miniconda 选择一个 新的 版本 &#xff0c;不用和RKNN的python版本保持一致 使用 ./xxx.sh进行安装 下面配置一下载源 # 清华大学源&#xff08;最常用&#xff09; conda config --add channels https://mirrors.tuna.tsinghua.edu.cn…...

深入解析光敏传感技术:嵌入式仿真平台如何重塑电子工程教学

一、光敏传感技术的物理本质与系统级实现挑战 光敏电阻作为经典的光电传感器件&#xff0c;其工作原理根植于半导体材料的光电导效应。当入射光子能量超过材料带隙宽度时&#xff0c;价带电子受激发跃迁至导带&#xff0c;形成电子-空穴对&#xff0c;导致材料电导率显著提升。…...

深度解析云存储:概念、架构与应用实践

在数据爆炸式增长的时代&#xff0c;传统本地存储因容量限制、管理复杂等问题&#xff0c;已难以满足企业和个人的需求。云存储凭借灵活扩展、便捷访问等特性&#xff0c;成为数据存储领域的主流解决方案。从个人照片备份到企业核心数据管理&#xff0c;云存储正重塑数据存储与…...