NCCL源码解析: P2P 连接的建立

文章目录

  • 前言
  • 概括
  • 详解
    • ncclTransportP2pSetup()

前言

NCCL 源码解析总目录

我尽量在每个函数之前介绍每个函数的作用,建议先不要投入到函数内部实现,先把函数作用搞清楚,有了整体框架,再回归到细节。

习惯: 我的笔记习惯:为了便于快速理解,函数调用关系通过缩进表示,也可能是函数展开,根据情况而定。

如下

// 调用 proxyConnInit
NCCLCHECK(proxyConnInit(peer, connectionPool, proxyState, (ncclProxyInitReq*) op->reqBuff, (ncclProxyInitResp*) op->respBuff, &op->connection));
// 对函数 proxyConnInit 进行展开,可方便看参数
static ncclResult_t proxyConnInit(struct ncclProxyLocalPeer* peer, struct ncclProxyConnectionPool* connectionPool, struct ncclProxyState* proxyState, ncclProxyInitReq* req, ncclProxyInitResp* resp, struct 

如有问题,请留言指正。

图后面再补;
有些遗漏之处,还没涉及,后面补;
闲话后面再补。

概括

recvpeer 表示本卡作为接收端的对端
sendpeer 表示本卡作为发送端的对端

对于每个 channel ,卡与卡之间要建立通信,先通过调用 selectTransport<0>() 建立接收通道,0 表示与 recvpeer 建立通信,再通过selectTransport<1>() 建立发送通道,1表示与 sendpeer 建立通信。
建立通道时会遍历 NTRANSPORTS 4种情况:P2P、共享内存、网络、collNet(collective Network, 还没看,不了解)

struct ncclTransport* ncclTransports[NTRANSPORTS] = {&p2pTransport,&shmTransport,&netTransport,&collNetTransport
};

本文重点关注 P2P。

接口如下:

struct ncclTransport p2pTransport = {"P2P",p2pCanConnect,{ p2pSendSetup, p2pSendConnect, p2pSendFree, NULL, p2pSendProxySetup, NULL, p2pSendProxyFree, NULL },{ p2pRecvSetup, p2pRecvConnect, p2pRecvFree, NULL, p2pRecvProxySetup, NULL, p2pRecvProxyFree, NULL }
};

发送建立流程为 p2pCanConnect() -> p2pSendSetup() -> p2pSendProxySetup()
接收建立流程为 p2pCanConnect() -> p2pRecvSetup() -> p2pRecvProxySetup()

先检查两个卡支不支持 P2P,主要检查两项:设备支不支持、路径支不支持,路径类型要小于 PATH_PXB,即不通过主桥的路径;
然后 p2pSendSetup() 填充一下 p2pConnectInfo, 向 proxy 线程请求 ncclProxyMsgSetup;
proxy 线程调用 p2pSendProxySetup(), 在本卡内申请显存,返回首地址以及相应的句柄devIpc, 其他进程或者线程可以通过这个句柄,获得此显存的操作地址。(我猜是让对端 GPU 卡也来操作这段内存,从而完成通信,还没看到那一步,完事来更新)
接收与发送机制一样。

详解

ncclTransportP2pSetup()

此P2P 非彼 P2P。ncclTransportP2pSetup 的 P2P 是广义上的两个设备之间的通信设置,包含 P2P、网络以及共享内存等。
建立两卡通信的入口函数。
因为要与 proxy 双线程操作,但是又是同步的,所以下文把两个线程的操作线性展开了,请注意。

ncclTransportP2pSetup(comm, &ringGraph, 0)
ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, int connIndex, int* highestTransportType/*=NULL*/)
{// 信息保存在 data[i] 中, i 为 rank// data[i] 大小为 2 * 64 个 connect, 先存放 recv, 再存放 sendrecvData[i] = data[i];// recvData[]  所有 recvChannels 的 ncclConnect 缓冲区的首地址// 首先 <0> 表示处理的是接收,处理与前一个 rank 的连接selectTransport<0>(comm, graph, recvData[i]+recvChannels++, c, recvPeer, connIndex, &type)static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclConnect* connect, int channelId, int peer, int connIndex, int* transportType){struct ncclPeerInfo* myInfo = comm->peerInfo+comm->rank;struct ncclPeerInfo* peerInfo = comm->peerInfo+peer;struct ncclConnector* connector = (type == 1) ? comm->channels[channelId].peers[peer]->send + connIndex :comm->channels[channelId].peers[peer]->recv + connIndex;NCCLCHECK(transportComm->setup(comm, graph, myInfo, peerInfo, connect, connector, channelId, connIndex)){NCCLCHECK(ncclCalloc(&resources, 1));recv->transportResources = resources;struct p2pConnectInfo* info = (struct p2pConnectInfo*)connectInfo;// 如果使用nvlink, 且两个GPU 计算能力一样,(gpu1->gpu.cudaCompCap == 80), 那么 useRead = 1// 如果通过参数 P2P_READ_ENABLE 设置该值,  P2P 使用 read 而不是 writeinfo->read = useRead; for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) if (!(info->read && p == NCCL_PROTO_SIMPLE)) recvSize += comm->buffSizes[p];// 如果同一个进程内的,且 DirectDisable 没有设置,P2P_USE_CUDA_MEMCPY 参数没有设置,并且ncclCuMemEnable 为假// 那么{resources->type = P2P_DIRECT;recv->conn.flags |= info->read ? NCCL_DIRECT_READ : NCCL_DIRECT_WRITE;}// 建立到 proxy 的连接, 连接信息在 recv->proxyConnNCCLCHECK(ncclProxyConnect(comm, TRANSPORT_P2P, 0, tpProxyRank, &recv->proxyConn));// 请求 proxy 执行 ncclProxyMsgSetup// recvSize += comm->buffSizes[p];// 发送数据 4字节 recvSize  10485760 = 4096 + NCCL_NUM_PROTOCOLS 3 类型的缓冲区大小// 接收数据缓冲区 info->p2pBuff// 要接收的大小 sizeof(struct ncclP2pBuff)// info->p2pBuff 保存 buf 信息NCCLCHECK(ncclProxyCallBlocking(comm, &recv->proxyConn, ncclProxyMsgSetup, &recvSize, sizeof(int), &info->p2pBuff, sizeof(struct ncclP2pBuff)));// 下面为 proxy 线程// proxy 线程接收数据进行处理{// op->connection : 设备与 proxy 连接的控制对象// proxyState : rank 的 ncclProxyState// op->reqBuff : proxy 本地的接收缓冲区首地址, 按照  op->reqSize 大小申请// op->reqSize : 客户端发送的发送数据的大小// op->respBuff: proxy 本地的发送缓冲区的首地址,按照 op->respSize 大小申请// p2pRecvProxySetup : 设备申请内存,首地址信息存入 respBuffNCCLCHECK(op->connection->tcomm->proxySetup(op->connection, proxyState, op->reqBuff, op->reqSize, op->respBuff, op->respSize, &done));static ncclResult_t p2pRecvProxySetup(struct ncclProxyConnection* connection, struct ncclProxyState* proxyState, void* reqBuff, int reqSize, void* respBuff, int respSize, int* done) {// 获取设备侧告知的 recvSize 的值 10485760int size = *((int*)reqBuff);struct ncclP2pBuff* p2pBuff = (struct ncclP2pBuff*)respBuff;NCCLCHECK(ncclP2pAllocateShareableBuffer(size, &p2pBuff->ipcDesc, &p2pBuff->directPtr));ncclResult_t ncclP2pAllocateShareableBuffer(size_t size, ncclIpcDesc *ipcDesc, void **ptr) {// 在设备侧申请内存,地址保存在 ptrNCCLCHECK(ncclCudaCalloc((char **)ptr, size));// cudaIpcGetMemHandle : 获取现有设备内存分配的进程间内存句柄// 获取指向使用cudaMalloc创建的现有设备内存分配的基址的指针,并将其导出以供另一个进程使用// __host__ cudaError_t cudaIpcGetMemHandle ( cudaIpcMemHandle_t* handle, void* devPtr )// 获取现有设备内存分配的进程间内存句柄。// 参数:// handle - 指向用户分配的 cudaIpcMemHandle 以返回句柄的指针。// devPtr - 指向先前分配的设备内存的基指针cudaError_t res = cudaIpcGetMemHandle(&ipcDesc->devIpc, *ptr);// cudaIpcOpenMemHandle : 打开从另一个进程导出的进程间内存句柄并返回可用于本地进程的设备指针// __host__ cudaError_t cudaIpcOpenMemHandle ( void** devPtr, cudaIpcMemHandle_t handle, unsigned int  flags )// 打开从另一个进程导出的进程间内存句柄并返回可用于本地进程的设备指针。// 参数// devPtr - 返回设备指针// handle - cudaIpcMemHandle 打开// flags - 此操作的标志。必须指定为cudaIpcMemLazyEnablePeerAccess}p2pBuff->size = size;connection->transportResources = p2pBuff->directPtr;}}// 下面不是 proxy 线程// 设备收到 proxy 返回的信息: 设备内部申请的缓冲区首地址,以及地址句柄 ipcDesc->devIpc// comm->peerInfo AllGather1 时保存的所有 rank 的信息: rank cudaDev hostHash pidHash busId// info->rank = myInfo->rank// p2pBuff : info->p2pBuff// devMem : (void**)&resources->recvDevMem 设备接收资源的接收缓冲区内存指针地址// ipcPtr : &resources->recvMemIpc 设备接收资源的接收 内存Ipc 指针地址NCCLCHECK(p2pMap(comm, myInfo, comm->peerInfo+info->rank, &info->p2pBuff, (void**)&resources->recvDevMem, &resources->recvMemIpc));static ncclResult_t p2pMap(struct ncclComm *comm, struct ncclPeerInfo* myInfo, struct ncclPeerInfo* peerInfo, struct ncclP2pBuff* p2pBuff, void** devMem, void** ipcPtr){// 如果 ncclCuMemEnable 为假,且两个 GPU 设备在同一进程中// 那么{// 如果本设备与对端设备不是同一设备if (peerInfo->cudaDev != myInfo->cudaDev) {// 如果可以从设备直接访问 peerDevice,则可以通过调用 cudaDeviceEnablePeerAccess() 来启用访问cudaError_t err = cudaDeviceEnablePeerAccess(peerInfo->cudaDev, 0);}// 把 proxy 从设备申请的内存首地址赋值给 *devMem,即 resources->recvDevMem// resources->recvDevMem = p2pBuff->directPtr;*devMem = p2pBuff->directPtr;// 同一个设备不用 ipc*ipcPtr = NULL;}else{if ((myInfo->pidHash == peerInfo->pidHash) && (peerInfo->cudaDev == myInfo->cudaDev)) {// 同一个进程,同一个设备// Same PID and GPU*devMem = p2pBuff->directPtr;*ipcPtr = NULL;} else {// 不同进程或者不同设备// Different PID or different GPUNCCLCHECK(ncclP2pImportShareableBuffer(comm, comm->topParentRanks[peerInfo->rank], p2pBuff->size, &p2pBuff->ipcDesc, devMem));ncclResult_t ncclP2pImportShareableBuffer(struct ncclComm *comm, int tpPeer, size_t size, ncclIpcDesc *ipcDesc, void **devMemPtr) {// cudaIpcOpenMemHandle : 打开从另一个进程导出的进程间内存句柄并返回可用于本地进程的设备指针// __host__ cudaError_t cudaIpcOpenMemHandle ( void** devPtr, cudaIpcMemHandle_t handle, unsigned int  flags )// 打开从另一个进程导出的进程间内存句柄并返回可用于本地进程的设备指针。// 参数// devPtr - 返回设备指针// handle - cudaIpcMemHandle 打开// flags - 此操作的标志。必须指定为cudaIpcMemLazyEnablePeerAccess// 通过 ipcDesc->devIpc 获取设备内存首地址 devMemPtrCUDACHECK(cudaIpcOpenMemHandle(devMemPtr, ipcDesc->devIpc, cudaIpcMemLazyEnablePeerAccess));}// devMem 已经赋值为设备内存首地址*ipcPtr = *devMem;}}}}}// 发送// 信息保存在 data[i] 中, i 为 rank// data[i] 大小为 2 * 64 个 connect, 先存放 recv, 再存放 send// sendData[]  所有 recvChannels 的发送 ncclConnect 缓冲区的首地址sendData[i] = recvData[i] + recvChannels;// 调用发送,处理与后一个 rank 的连接NCCLCHECKGOTO(selectTransport<1>(comm, graph, sendData[i]+sendChannels++, c, sendPeer, connIndex, &type), ret, fail);static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclConnect* connect, int channelId, int peer, int connIndex, int* transportType) {NCCLCHECK(transportComm->setup(comm, graph, myInfo, peerInfo, connect, connector, channelId, connIndex));ncclResult_t p2pSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* myInfo, struct ncclPeerInfo* peerInfo, struct ncclConnect* connectInfo, struct ncclConnector* send, int channelId, int connIndex){NCCLCHECK(ncclCalloc(&resources, 1));send->transportResources = resources;info->read = useRead;if (graph && connIndex == 1) info->read = 0;const char* useReadStr = info->read ? "/read" : "";// For P2P Read the SIMPLE buffer is tagged on the end of the ncclSendMem structureif (info->read) // 只有读的时候,使用缓冲区 NCCL_PROTO_SIMPLEsendSize += comm->buffSizes[NCCL_PROTO_SIMPLE];info->rank = myInfo->rank;resources->type = P2P_DIRECT;send->conn.flags |= info->read ? NCCL_DIRECT_READ : NCCL_DIRECT_WRITE;// 与接收一样的操作// 设备收到 proxy 返回的信息保存在 p2pBuff中: 设备内部申请的缓冲区首地址,以及地址句柄 ipcDesc->devIpcNCCLCHECK(ncclProxyCallBlocking(comm, &send->proxyConn, ncclProxyMsgSetup, &sendSize, sizeof(int), &info->p2pBuff, sizeof(struct ncclP2pBuff)));// p2pMap : 根据接收到的信息做一个发送缓冲区的首地址解析,得到 sendDevMem 或者 sendMemIpc// info->rank = myInfo->rank// p2pBuff : info->p2pBuff// devMem : (void**)&resources->recvDevMem 设备接收资源的接收缓冲区内存指针地址// ipcPtr : &resources->recvMemIpc 设备接收资源的接收 内存Ipc 指针地址NCCLCHECK(p2pMap(comm, myInfo, comm->peerInfo+info->rank, &info->p2pBuff, (void**)&resources->sendDevMem, &resources->sendMemIpc));}}if (sendPeer == recvPeer) {if (recvChannels+sendChannels) {NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, recvPeer, bootstrapTag, data[i], sizeof(struct ncclConnect)*(recvChannels+sendChannels)), ret, fail);NCCLCHECKGOTO(bootstrapRecv(comm->bootstrap, recvPeer, bootstrapTag, data[i], sizeof(struct ncclConnect)*(recvChannels+sendChannels)), ret, fail);sendData[i] = data[i];recvData[i] = data[i]+sendChannels;}} else {// 如果 sendPeer recvPeer 不是同一个// 假设 0 -> 1 -> 2, 当前 rank 为 1// sendPeer = 2, recvPeer = 0// sendPeer : 我作为发送的对端 rank// recvPeer : 我作为接收的对端 rankif (recvChannels) // 向前一个 rank 发送 recvChannels 个接收连接信息NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, recvPeer, bootstrapTag, recvData[i], sizeof(struct ncclConnect)*recvChannels), ret, fail);if (sendChannels) // 向后一个 rank 发送 recvChannels 个发送连接信息NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, sendPeer, bootstrapTag, sendData[i], sizeof(struct ncclConnect)*sendChannels), ret, fail);if (sendChannels) // 接收后一个 rank 的接收连接信息到 sendDataNCCLCHECKGOTO(bootstrapRecv(comm->bootstrap, sendPeer, bootstrapTag, sendData[i], sizeof(struct ncclConnect)*sendChannels), ret, fail);if (recvChannels) // 接收前一个 rank 的接收发送信息NCCLCHECKGOTO(bootstrapRecv(comm->bootstrap, recvPeer, bootstrapTag, recvData[i], sizeof(struct ncclConnect)*recvChannels), ret, fail);}
}

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/news/670901.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

2024年第九届信号与图像处理国际会议(ICSIP 2024)

2024第九届信号与图像处理国际会议&#xff08;ICSIP 2024&#xff09;将于2024年7月12-14日在中国南京召开。ICSIP每年召开一次&#xff0c;在过去的七年中吸引了1200多名与会者&#xff0c;是展示信号和图像处理领域最新进展的领先国际会议之一。本次将汇集来自亚太国家、北美…

LeetCode:13.罗马数字转整数

13. 罗马数字转整数 - 力扣&#xff08;LeetCode&#xff09; 目录 思路&#xff1a; 官解代码&#xff1a; 作者辣眼代码: 每日表情包&#xff1a; 思路&#xff1a; 思路已经很明了了&#xff0c;题目已经给出一般规则和特殊规则&#xff08;而且题目确保给定的是正确的…

活动图(Activity Diagram)

一、定义 动态图。显示人或对象的活动&#xff0c;其方式类似于流程图 二、构成 包含有&#xff1a; 初始节点(开始)最后一个节点(结束)活动转换判定(决策),同步条分岔或汇合泳道 1、 初始节点(开始) 实心圆表示初始节点 2、最后一个节点(结束) 圆圈内加一个实心圆来表…

判断字符串是否包含正则表达式默认的特殊字符c++

判断字符串是否包含正则表达式默认的特殊字符 业务描述&#xff1a; 上层配置的字符列表中&#xff0c;既有准确的字符串&#xff0c;又有可以进行正则匹配的字符串&#xff0c;这时候需要区分出来那些是正则匹配的字符串。 思路: 判断字符串中&#xff0c;是否存在正则表达…

SF相关1111

AndroidQ 图形系统&#xff08;11&#xff09;UI刷新&#xff0c;SurfaceFlinger&#xff0c;Vsync机制总结_android viewrootimpl surfaceflinger hw-CSDN博客

C语言之字符逆序(牛客网)

个人主页&#xff08;找往期文章包括但不限于本期文章中不懂的知识点&#xff09;&#xff1a;我要学编程(ಥ_ಥ)-CSDN博客 字符逆序__牛客网 题目&#xff1a; 思路&#xff1a;既然有空格就不能用scanf函数来接收字符了。因为scanf函数遇到空格会停止读取。我们可以用get…

解决:VSCode 连接服务器时出错:Could not establish connection to : XHR failed

对于 VSCode 1.86.0 版本&#xff0c;若出现这个错误&#xff0c;简易方法是回退到 1.85.x 版本&#xff0c;VSCode 1.85.x 官方下载地址实际上很多 VSCode 连接不上服务器的问题&#xff0c;都可以靠三个步骤暴力解决 SSH 连接服务器&#xff0c;然后删除已有的服务器端的 VS…

UE5 获得频谱让nigara随音乐律动

参考视频:UE - Niagara实现可视化音乐动态粒子效果 案例演示及教程_哔哩哔哩_bilibili 先创建一个Niagara 在Properties的Sim Target改为GPU,Calculate Bounds Mode改为Fixed模式 生成的数量改为1000 这里的BoxSize可以选择修改,具体作用是粒子初始生成的范围 Drag,阻力,用来限…

Android rom定制 修改system分区的容量大小

1、写在前面 系统ROM定制化,预置app太多,会导致系统rom很大,原生系统system分区已经不够用了,要加大系统systemui分区 2.修改system分区的容量大小的核心类 device/mediatekprojects/$project/BoardConfig.mk build/make/core/Makefile3、修改system 分区的容量大小的核…

Java上(2024尚硅谷)

day01.[环境变量,HelloWorld] 1.会常用的dos命令 2.会安装java所需要的环境(jdk) 3.会配置java的环境变量 4.知道java开发三步骤 5.会java的入门程序(HelloWorld) 6.会三种注释方式 7.知道Java入门程序所需要注意的地方 8.知道println和print的区别第一章 Java概述 1.1 JavaS…

康佳智能电视LED32IS97N 刷机升级方法,及刷机数据,务必先查物料号确认一致才能刷机

升级步骤&#xff1a; 康佳电视串号检测平台 &#xff1a; http://service.kkapp.com/KKAFTERSALE/pages/sn/sn.jsp 务必确认自己的主程序软件号&#xff08;物料号&#xff09; 强制刷机方法&#xff1a; 1、下载刷机数据包&#xff0c;先把文件 NAND_TARGET.rar 解压&#…

【自定义序列化器】⭐️通过继承JsonSerializer和实现WebMvcConfigurer类完成自定义序列化

目录 前言 解决方案 具体实现 一、自定义序列化器 二、两种方式指定作用域 1、注解 JsonSerialize() 2、实现自定义全局配置 WebMvcConfigurer 三、拓展 WebMvcConfigurer接口 章末 前言 小伙伴们大家好&#xff0c;上次做了自定义对象属性拷贝&#x…

Dijkstra算法(求最短路)

简介&#xff1a; 迪杰斯特拉算法(Dijkstra)是由荷兰计算机科学家狄克斯特拉于1959年提出的&#xff0c;因此又叫狄克斯特拉算法。是从一个顶点到其余各顶点的最短路径算法&#xff0c;解决的是有权图中最短路径问题。 特点&#xff1a; 迪杰斯特拉算法采用的是一种贪心策略&a…

【Java基础】关于Java基础的一些有趣的常识!

前言 今天看到了一篇文章&#xff0c;是关于茶余饭后的Java常识的一些有趣解答&#xff0c;我觉得写的很有趣很易懂&#xff0c;所以截取了其中我觉得比较有趣的问题分享给大家。原文&#xff1a;饭后茶余的java常识 - 知乎 (zhihu.com) 1. Java语言的特点有哪些&#xff1f; …

MySQL的DML语言

DML&#xff1a;Data Manipulation Language&#xff08;数据操作语言&#xff09; DML语言用来对数据库中表的数据记录进行增、删、改操作。 一、添加数据命令 注意&#xff1a; 插入数据时&#xff0c;指定的字段顺序需要与值的顺序是一一对应的。 字符串和日期型数据应该包…

#pragma once和条件编译

#pragma once 和 #ifndef 是 C/C 中用于防止头文件被多次包含的两种不同的预处理器指令。 一、那么为什么要防止头文件被重复包含 头文件的重复包含问题需要避免的原因主要有以下几点&#xff1a; 编译效率&#xff1a; 如果头文件被重复包含多次&#xff0c;编译器需要重复解…

【并发编程】手写线程池阻塞队列

&#x1f4dd;个人主页&#xff1a;五敷有你 &#x1f525;系列专栏&#xff1a;并发编程 ⛺️稳重求进&#xff0c;晒太阳 示意图 步骤1&#xff1a;自定义任务队列 变量定义 用Deque双端队列来承接任务用ReentrantLock 来做锁并声明两个条件变量 Condition fullWai…

网站不收录,与服务器不备案有关吗

随着互联网的快速发展&#xff0c;网站已经成为企业、个人和机构宣传和展示自己的重要平台。然而&#xff0c;许多网站在建设完成后却面临着不收录的问题&#xff0c;这给网站的管理者和拥有者带来了很大的困扰。其中&#xff0c;一些人认为&#xff0c;网站不收录的原因与服务…

代码随想录算法训练营第23天(回溯1)| 回溯算法的理论77.组合问题

回溯的理论基础 其实在讲解二叉树的时候&#xff0c;就给大家介绍过回溯&#xff0c;这次正式开启回溯算法&#xff0c;大家可以先看视频&#xff0c;对回溯算法有一个整体的了解。 题目链接/文章讲解 视频讲解 回溯的总结&#xff1a; 树的深度&#xff08;递归的层数&…

什么是Instagram Reels?用好Reels 让你的流量暴涨!

Instagram Reels是Instagram在2020年全新推出的短视频功能&#xff0c;旨在与TikTok展开竞争。作为跨境卖家的你&#xff0c;利用 Reels 这一神器&#xff0c;将为你带去更多的流量。那该如何利用好这一神器呢&#xff1f;本篇文章&#xff0c;大白将带大家深入了解 Reels 并用…