nvshmem介绍及其在deepep中的使用

已于 2025-05-17 11:27:12 修改
阅读量1.2k 收藏 9
点赞数 30
分类专栏： NCCL 文章标签： nvshmem nccl nvlink rdma
于 2025-05-05 18:17:35 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/fengcai_ke/article/details/147644895
版权

nvshmem简介

介绍

nvshmem是nvidia基于openshmem标准实现的高性能通信库，特点如下：

  • 支持对称内存，所有GPU的虚拟地址一致，可通过指针直接访问
  • 支持host-size和device-size api发起通信操作

集群中参与nvshmem操作的每个进程认为是一个PE(processing element)，每个PE分配唯一ID。PE之间可通过对称内存进行通信和数据共享，对称内存通过nvshmem_malloc进行分配，使用其他api分配的内存都是私有的。
在这里插入图片描述
官方文档：https://docs.nvidia.com/nvshmem/api/introduction.html

transport

nvshmem支持p2p, ibrc, ibgda等transport，可在编译和运行时通过环境变量控制使能哪个transport。
p2p只能用于node内PE通信，ibrc/ibgda可用于node内和node间PE通信。
编译时
p2p没有编译选项控制，其他transport通过如下环境变量控制

NVSHMEM_IBRC_SUPPORT //默认使能 
NVSHMEM_UCX_SUPPORT
NVSHMEM_LIBFABRIC_SUPPORT 
NVSHMEM_IBDEVX_SUPPORT 
NVSHMEM_IBGDA_SUPPORT

运行时
p2p通过NVSHMEM_DISABLE_P2P控制，ibgda通过NVSHMEM_IB_ENABLE_IBGDA控制，其他transport默认使能

NVSHMEM_DISABLE_P2P
NVSHMEM_IB_ENABLE_IBGDA

ibrc和ibgda的区别可参考：https://developer.nvidia.com/zh-cn/blog/improving-network-performance-of-hpc-systems-using-nvidia-magnum-io-nvshmem-and-gpudirect-async/

API

初始化 api
用来初始化相关资源，其中flag指定bootstrap模式，attributes指定额外信息，比如myrank和nranks

void nvshmem_init(void)
int nvshmemx_init_attr(unsigned int flags, nvshmemx_init_attr_t *attributes)

内存管理 api
用来分配对称内存，需所有PE同时调用

void *nvshmem_malloc(size_t size)
void nvshmem_free(void *ptr)
void *nvshmem_align(size_t alignment, size_t size)

单边操作 api
提供了host-size和device-size api，device-size api还区分thread/warp/block等使用多个的cuda core进行操作

//TYPENAME和TYPE为nvshem支持的数据类型，比如int，long等

//操作多个元素
void nvshmem_TYPENAME_put(TYPE *dest, const TYPE *source, size_t nelems, int pe)
void nvshmemx_TYPENAME_put_on_stream(TYPE *dest, const TYPE *source, size_t nelems, int pe, cudaStream_t stream)
__device__ void nvshmem_TYPENAME_put(TYPE *dest, const TYPE *source, size_t nelems, int pe)
__device__ void nvshmemx_TYPENAME_put_block(TYPE *dest, const TYPE *source, size_t nelems, int pe)
__device__ void nvshmemx_TYPENAME_put_warp(TYPE *dest, const TYPE *source, size_t nelems, int pe)

//操作单个元素
void nvshmem_TYPENAME_p(TYPE *dest, TYPE value, int pe)
__device__ void nvshmem_TYPENAME_p(TYPE *dest, TYPE value, int pe)

集合通信 api
nvshmem提供了reducescatter, allreduce, allgather, broadcat和all2all等集合通信api。
提供了host-size和device-size api，device-size api还区分thread/warp/block等使用多个的cuda core进行操作

int nvshmem_TYPENAME_alltoall(nvshmem_team_t team, TYPE *dest, const TYPE *source, size_t nelems)
int nvshmemx_TYPENAME_alltoall_on_stream(nvshmem_team_t team, TYPE *dest, const TYPE *source, size_t nelems, cudaStream_t stream)
__device__ int nvshmem_TYPENAME_alltoall(nvshmem_team_t team, TYPE *dest, const TYPE *source, size_t nelems)
__device__ int nvshmemx_TYPENAME_alltoall_block(nvshmem_team_t team, TYPE *dest, const TYPE *source, size_t nelems)
__device__ int nvshmemx_TYPENAME_alltoall_warp(nvshmem_team_t team, TYPE *dest, const TYPE *source, size_t nelems)

host-size api集合通信也支持调用nccl api，比如all2all调用nccl的send/recv实现

template <typename TYPE>
int nvshmemi_alltoall_on_stream(nvshmem_team_t team, TYPE *dest, const TYPE *source, size_t nelems,
                                cudaStream_t stream) {
#ifdef NVSHMEM_USE_NCCL
    nvshmemi_team_t *teami = nvshmemi_team_pool[team];
    int team_n_pes = nvshmem_team_n_pes(team);
    if (nvshmemi_use_nccl && nvshmemi_get_nccl_dt<TYPE>() != ncclNumTypes &&
        ((nccl_version >= 2700 && team_n_pes <= 4096 /* NCCL limit for Group API */) ||
         (nccl_version >= 2800 && team_n_pes <= 32768 /* NCCL limit for Group API */))) {
        size_t rank_offset = nelems * sizeof(TYPE);
        NCCL_CHECK(nccl_ftable.GroupStart());
        for (int pe = 0; pe < team_n_pes; pe++) {
            NCCL_CHECK(nccl_ftable.Send(((char *)source) + pe * rank_offset, nelems,
                                        nvshmemi_get_nccl_dt<TYPE>(), pe,
                                        (ncclComm_t)teami->nccl_comm, stream));
            NCCL_CHECK(nccl_ftable.Recv(((char *)dest) + pe * rank_offset, nelems,
                                        nvshmemi_get_nccl_dt<TYPE>(), pe,
                                        (ncclComm_t)teami->nccl_comm, stream));
        }
        NCCL_CHECK(nccl_ftable.GroupEnd());
    } else
#endif /* NVSHMEM_USE_NCCL */
    {
        nvshmemi_call_alltoall_on_stream_kernel<TYPE>(team, dest, source, nelems, stream);
    }
    return 0;
}

nvshmem例子
通过此例子熟悉nvshmem用法，功能为每个PE将自己的id发送到相邻的下一个PE
在这里插入图片描述

#include <stdio.h>
#include <cuda.h>
#include <nvshmem.h>
#include <nvshmemx.h>

__global__ void simple_shift(int *destination) {
    int mype = nvshmem_my_pe();
    int npes = nvshmem_n_pes();
    int peer = (mype + 1) % npes;

    nvshmem_int_p(destination, mype, peer);
}

int main(void) {
    int mype_node, msg;
    cudaStream_t stream;

    nvshmem_init();
    mype_node = nvshmem_team_my_pe(NVSHMEMX_TEAM_NODE);
    cudaSetDevice(mype_node);
    cudaStreamCreate(&stream);

    int *destination = (int *) nvshmem_malloc(sizeof(int));

    simple_shift<<<1, 1, 0, stream>>>(destination);
    nvshmemx_barrier_all_on_stream(stream);
    cudaMemcpyAsync(&msg, destination, sizeof(int), cudaMemcpyDeviceToHost, stream);

    cudaStreamSynchronize(stream);
    printf("%d: received message %d\n", nvshmem_my_pe(), msg);

    nvshmem_free(destination);
    nvshmem_finalize();
    return 0;
}

//执行结果
mpirun --mca btl_tcp_if_include eth0 --allow-run-as-root --prefix /root/host/openmpi-4.1.3/build/ --hostfile /root/hostfile -x NVSHMEM_BOOTSTRAP=MPI -x NVSHMEM_IB_GID_INDEX=3 -x LD_LIBRARY_PATH=/root/deepep/nvshmem_src/install/lib:/root/host/openmpi-4.1.3/build/lib:/root/deepep/gdrcopy-2.4.4/build/lib -x NVSHMEM_ENABLE_NIC_PE_MAPPING=1 -x NVSHMEM_HCA_LIST=^mlx5_3:1,mlx5_4 /root/deepep/nvshmem_src/install/bin/examples/ring-bcast
1: received message 0
5: received message 4
4: received message 3
2: received message 1
7: received message 6
0: received message 7
3: received message 2
6: received message 5

nvshmem处理流程

nvshmem初始化
a. bootstrap初始化
nvshmem支持MPI, PMI, SHMEM和UID等多种bootstrap模式，可通过nvshmemx_init_attr的参数flag或者环境变量NVSHMEM_BOOTSTRAP指定，bootstrap提供了allgather，alltoall等接口在集群中传递信息

b. memory初始化
使用cuda vmm机制实现内存管理，初始化只是申请虚拟地址，后续申请内存时再分配物理内存，并映射到对应的虚拟地址，可参考：https://developer.nvidia.com/blog/introducing-low-level-gpu-virtual-memory-management/

每个PE申请虚拟地址起始值为heap_base_，大小为heap_size_*p2p_npes，即包含node内所有PE，每个PE对应的虚拟地址大小为heap_size_(128G)。
不管同node还是跨node，每个PE的虚拟地址都是相同的。
nvshmem_malloc申请内存时将不同PE的物理内存映射到PE对应的虚拟地址，实现对其他PE内存的访问
在这里插入图片描述
c. transport初始化
以ibrc和ibgda为例，transport初始化主要是rdma qp的建立。
对ibrc来说，一个网卡建立的qp个数为: n_pes * qp_count，其中n_pes为集群PE个数，qp_count为网卡和每个peer建立的qp个数，代码中为固定值2，一个qp给proxy使用，一个qp给host-size api使用

对ibgda来说，一个网卡建立的qp个数为: n_pes * num_rc_eps_per_pe，其中num_rc_eps_per_pe可通过环境变量NVSHMEM_IBGDA_NUM_RC_PER_PE指定

nvshmem_malloc
分配内存主要流程如下：

  • 分配PINNED物理内存，将物理内存映射到本PE对应的虚拟地址
  • 获取物理内存句柄，通过bootstrap传播给其他PE，其他PE导入此句柄，实现同node对本PE内存访问
  • 通过ibv_reg_mr将虚拟地址注册到网卡，获取对应的lkey和rkey，将虚拟地址和key通过bootstrap传播给其他PE，实现跨node其他PE对本PE内存访问

比如调用nvshmem_malloc分配100M内存
在这里插入图片描述

PE0访问PE7内存的地址为：0x10020000000+128G*7=0x1e020000000
PE0访问PE8内存的地址为：0x10020000000
PE0访问PE15内存的地址为: 0x10020000000

nvshmem_int_p
a. 主机侧发起

  • 如果peer在本node，则调用cudaMemcpyAsync实现
  • 如果peer在其他node，则由主机cpu调用ibv_post_send通知网卡进行数据发送(使用第一个qp)

b. 设备侧发起

  • 如果peer在本node，则调用cude kernel实现内存拷贝
  • 如果peer不在本node，还要分两种情况 使能igbda，由cuda core调用ibgda_post_send通知网卡进行数据发送未使能igbda，由cuda core通知proxy线程，proxy线程调用ibv_post_send通知网卡进行数据发送(使用第二个qp)

nvshmem在deepep中的使用

deepep提供了高效的all2all通信，但并未直接使用nvshmem中的集合通信all2all api，而是使用nvshmem的ibrc和ibgda transport实现node间rank的通信，使用cuda ipc实现node内rank的通信。

deepep使用nvshmemx_init_attr进行nvshmem的初始化，通过flag指定bootstrap模式为UID，通过attr指定root_unique_id，rank和num_ranks等信息。
root_unique_id为root rank的ip和端口组合，通过pytorch的dist.all_gather_object传播给其他rank，其他rank通过root_unique_id和root rank交互，从而获取集群所有rank的连接方式。

nvshmemx_uniqueid_t root_unique_id;
nvshmemx_init_attr_t attr;
std::memcpy(&root_unique_id, root_unique_id_val.data(), sizeof(nvshmemx_uniqueid_t));
nvshmemx_set_attr_uniqueid_args(rank, num_ranks, &root_unique_id, &attr);
nvshmemx_init_attr(NVSHMEMX_INIT_WITH_UNIQUEID, &attr);

编译环境变量

CUDA_HOME=/path/to/cuda \
GDRCOPY_HOME=/path/to/gdrcopy \
NVSHMEM_SHMEM_SUPPORT=0 \
NVSHMEM_UCX_SUPPORT=0 \
NVSHMEM_USE_NCCL=0 \
NVSHMEM_MPI_SUPPORT=0 \
NVSHMEM_IBGDA_SUPPORT=1 \
NVSHMEM_PMIX_SUPPORT=0 \
NVSHMEM_TIMEOUT_DEVICE_POLLING=0 \
NVSHMEM_USE_GDRCOPY=1 \
cmake -S . -B build/ -DCMAKE_INSTALL_PREFIX=/path/to/your/dir/to/install

高带宽kernel
node间同号rank通过nvshmem ibrc通信，node内rank使用cuda ipc通信。
node0上8个rank都需要获取unique_id，并传播给同号rank。
在这里插入图片描述
通信时node内通过同号rank发送，到了同号rank后再通过nvlink转发到目的rank。

低延时kernel
node间和node内rank通信均使用nvshmem ibgda，即使node内rank通信也不用p2p transport，减少rdma到nvlink报文的转发带来的延迟

os.environ['NVSHMEM_DISABLE_P2P'] = '1'
os.environ['NVSHMEM_IB_ENABLE_IBGDA'] = '1'
os.environ['NVSHMEM_IBGDA_NIC_HANDLER'] = 'gpu'
os.environ['NVSHMEM_IBGDA_NUM_RC_PER_PE'] = f'{num_qps_per_rank}'
# Make sure QP depth is always larger than the number of on-flight WRs, so that we can skip WQ slot check
os.environ['NVSHMEM_QP_DEPTH'] = '1024'
# NOTES: NVSHMEM initialization requires at least 256 MiB
os.environ['NVSHMEM_CUMEM_GRANULARITY'] = f'{2 ** 29}'

只有rank 0获取unique_id，并传播给所有rank。
在这里插入图片描述

每个网卡创建的qp数量为: num_rc_eps_per_pe * n_pes，num_rc_eps_per_pe为本rank专家个数。
在两机16卡的情况下，假设专家个数288(256路由专家+32冗余专家)，则每个卡上专家个数为18，即num_rc_eps_per_pe为18，网卡创建的qp数量为288。

通信时通过同号qp发给其他所有rank。

Support NVLink protocol for intranode low-latency kernels (conflict with hook-based overlapping)
低延时模式同主机gpu通信为什么不用nvlink？nvlink会占用SM，而使用rdma可由网卡自动完成数据传输，不占用sm，可做到计算和通信overlap

Enable IBGDA for the low latency mode, which refers to “no package forwarding between NVLink and RDMA”
低延时模式跨主机gpu通信为什么不用nvlink？需要先转到同号卡，同号卡再转给目的卡，增加了延迟

CUDA专题11—告别内存等待：深度解析CUDA同步域的原理与实战
探索人工智能革命，深入算法原理与创新应用，未来科技无限可能。
04-02 45
部分CUDA应用程序可能因内存栅栏/刷新操作需等待超出CUDA内存一致性模型要求范围的内存事务，从而导致性能下降。以观察上述示例。根据CUDA内存一致性模型的保证，断言条件必然成立，这意味着线程1对变量x的写入必须在线程2对b的写入之前，对线程3可见。通过原子变量a的release-acquire操作所提供的内存序仅能确保x对线程2可见（因其作用域限于设备内部），而无法保证对线程3的可见性。
多节点多 GPU：大规模使用 NVIDIA cuFFTMp FFT
qq_27815483的博客
07-11 1165
本文主要介绍了多节点多 GPU：大规模使用 NVIDIA cuFFTMp FFT
DeepSeek-DeepEP通信（一）
TH_NUM的博客
03-11 427
抽象层次高：隐藏底层通信细节（如 CUDA IPC 或 InfiniBand），简化分布式 GPU 编程。进程级隔离：需显式传递句柄并协调同步，适合紧密协作的多进程应用（如多进程推理服务）。定位：多 GPU/多节点的高性能通信库，基于 PGAS（分区全局地址空间） 模型。适用场景：大规模并行计算（如 AI 训练、科学模拟），需多 GPU 协同作业。适用范围：单机多 GPU，例如共享数据的多个独立进程（如 MPI 进程）。定位：同一节点（单机）内多进程间的 GPU 内存共享。减少数据移动，提升吞吐量。
【GPU】Nvidia CUDA 编程高级教程——利用蒙特卡罗法求解近似值（NVSHMEM）
从善若水的博客
12-17 1660
`NVSHMEM`是一个并行编程模型，用于在多个 NVIDIA GPU 之间进行高效和可扩展的通信。**NVSHMEM 依托于 OpenSHMEM 构建而成**，可为横跨多个 GPU 内存的数据提供全局地址空间，并可通过细粒度的由 GPU 发起的操作、由 CPU 发起的操作和 CUDA 流操作访问该空间。NVSHMEM 为许多应用提供了令人信服的多 GPU 编程模型，对于具有高密度 GPU 和复杂互连的现代 GPU 服务器（例如NVIDIA DGX A100 服务器 的 NVIDIA NVSwitch）来说
NVSHMEM介绍、InfiniBand GPUDirect、和NVshmem使用案例说明
爱串门的小马驹博客
03-11 1765
NVIDIA NVSHMEM 是基于 NVIDIA 的“共享内存”库，它提供了一个易于使用的 CPU 端接口，用于分配对称分布在 NVIDIA GPU 集群中的固定内存。。是NVIDIA提供的一种并行编程接口，专为NVIDIA GPU集群设计，以实现高效且可扩展的通信。NVSHMEM基于OpenSHMEM规范，为跨越多个GPU内存的数据创建了全局地址空间，支持GPU间的数据共享和并行计算。
DeepEP - 针对混合专家（MoE）和专家并行（EP）的通信库
技术文档翻译分享
03-02 854
一、关于 DeepEP 二、性能 1、带有 NVLink 和 RDMA 转发的常规内核 2、低延迟内核，纯RDMA 三、快速开始 1、需要 2、下载并安装 NVSHMEM 依赖 3、开发 4、安装 四、网络配置 1、交通隔离 2、自适应路由 3、拥塞控制 五、接口和示例 1、在模型训练或推理预填充中的示例用途 2、在推理解码中的示例使用 六、通知 未定义行为的 PTX 使用 自动调优您的集群
【GPU】Nvidia CUDA 编程高级教程——NVSHMEM 内存模型
从善若水的博客
01-15 4143
NVSHMEM 的内存分配 API nvshmem_malloc()，其工作方式有点类似于标准的cudaMalloc()，但cudaMalloc()会返回一个本地 GPU 的私有地址[^1]。**使用nvshmem_malloc()分配的对象称为对称数据对象**。**每个对称数据对象在所有 PE 上都有一个名称、类型和大小相同的对应数据对象。由nvshmem_malloc()返回的指针对应的虚拟地址称为对称地址**。在 NVSHMEM 通信例程中使用对称地址对其他 PE 进行远程访问是合法的（对称地址也可以
NVSHMEM安装及使用
jkjgj的博客
10-14 2851
NVSHMEM 是 NVIDIA 提供的一个并行编程库，专门用于在 GPU 集群或多 GPU 系统中实现高效的共享内存访问。它通过使用共享内存来实现 GPU 间的数据通信，适合大规模并行计算应用，尤其是那些需要在多个 GPU 之间快速交换数据的场景。
【nvidia CUDA 高级编程】NVSHMEM 直方图——分布式方法
从善若水的博客
01-15 1206
另一种方法就是**分割直方图**并把各个部分分配到不同的GPU上。**当输入数据中的一个条目不属于该 GPU 上的直方图位置时，我们将自动递增远程 PE 中的相关直方图条目**。然后我们必须在最后把直方图的各个部分拼接起来。我们将这种方法称为“分布式”方法。
【nvidia CUDA 高级编程】NVSHMEM 直方图——复制式方法
从善若水的博客
01-15 1081
在多个 GPU 上分配工作量的一种方法与我们在 𝜋 估算器上使用的方法相同：给定 𝑁 个整数，我们即可把它们均匀地分配到所有 GPU 上，然后可以对所有 PE 进行归约。我们将此称为**“复制式”方法**，因为在所有 GPU 上都存在完整的直方图副本。我们将第一步，即增加每个直方图桶内的计数值，命名为“列表”步骤；将合并所有 PE 上的结果的第二步命名为“结合”步骤，并分别计算时间（以便与下一个方法进行比较）。
通过全站创新推动高性能计算
专注于人工智能领域的小何尚
06-09 492
高性能计算 (HPC) 已成为科学发现的重要工具。无论是发现新的救命药物、应对气候变化，还是对我们的世界进行精确模拟，这些解决方案都需要巨大且快速增长的处理能力。 它们越来越超出传统计算方法的范围。这就是为什么业界接受 NVIDIA GPU 加速计算的原因。 与人工智能相结合，它为科学进步带来了数百万倍的性能飞跃。 如今，有 2,700 个应用程序可以从 NVIDIA GPU 加速中受益，而且这个数字还在不断增加，得到了不断壮大的 300 万开发者社区的支持。在整个 HPC 应用程序范围内实现多倍加速需要在
DeepSeek DeepEP学习（二）normal notify dispatch
KIDGIN7439的专栏
03-06 1437
dispatch的过程分为两步，第一步通过notify dispatch计算一些meta信息，比如每台机器将会收到多少token，然后再执行实际的dispatch，本节主要介绍notify dispatch。
深度解读DeepSeek：开源周（Open Source Week）技术解读
造梦先森Kai的专栏
03-22 1679
深度解读DeepSeek：开源周（Open Source Week）技术解读。
DeepEP通信库
m0_49448331的博客
02-25 978
DeepEP 是一个高效的专家并行通信库，专为 MoE 模型的训练和推理优化。它通过提供高吞吐量和低延迟的 GPU 内核（也称为 MoE 分发和组合操作），显著降低了 MoE 模型在分布式训练和推理中的通信瓶颈。根据 DeepSeek 的描述，DeepEP 旨在支持大规模 AI 工作负载，减少训练成本，提高性能，并使先进 AI 研究更加可访问。
CUDA之GPU生态系统
silencez_w的博客
12-14 4862
易于编程和性能的巨大飞跃是 CUDA 平台被广泛采用的关键原因之一。CUDA 平台成功的第二大原因是拥有广泛而丰富的生态系统。 与任何新平台一样，CUDA 的成功取决于可用于 CUDA 生态系统的工具、库、应用程序和合作伙伴。任何新的计算平台都需要开发人员将应用程序移植到新平台。为此，开发人员需要最先进的工具和开发环境。 应用程序开始扩展后，在数据中心级别需要更多工具。NVIDIA 致力于为开发者和企业提供最先进的工具和生态系统服务。 图 1：CUDA 生态系统：使 CUDA 平台成为最佳开发者选择的构建
C/C++数据结构-二叉树的存储和遍历
最新发布
05-22
代码实现了二叉树的存储和遍历功能，采用链式存储结构（Node结构体）表示二叉树节点，包含数据域和左右子节点指针。核心功能包括： 创建二叉树：支持递归法通过前序序列（如ABD##E##C#FG###，#表示空节点）动态构建二叉树； 遍历算法： 递归遍历：实现前序、中序、后序遍历（根左右、左根右、左右根）； 非递归遍历：基于栈模拟递归过程，完成三种深度优先遍历； 层次遍历：利用队列实现广度优先遍历（BFS）。 内存管理：使用malloc分配节点内存（可优化为C++的new）。 代码通过分离递归与迭代逻辑，全面展示了二叉树的核心操作，但需注意输入处理（如换行符）以避免栈溢出异常。
cmd-bat-批处理-脚本-jscript-jspopup.zip
05-22
cmd-bat-批处理-脚本-jscript-jspopup.zip
cmd-bat-批处理-脚本-02.substring.zip
05-22
cmd-bat-批处理-脚本-02.substring.zip
markdown语法排版
09-04
Markdown是一种轻量级的标记语言，用于简单而直观地排版文本。下面是一些常用的Markdown语法排版示例： 1. 标题：通过在文本前面添加井号(#)来表示不同级别的标题，例如： `# 一级标题`、`## 二级标题`等。 2. 强调：使用星号(*)或下划线(_)将文本包围起来，表示斜体或粗体，例如：`*斜体*`、`**粗体**`。 3. 列表：使用短横线(-)、加号(+)或星号(*)加空格来表示无序列表，使用数字加点加空格表示有序列表。 4. 链接：使用方括号([])包裹链接文字，紧跟着小括号(())包裹链接地址，例如：`[链接文字](链接地址)`。 5. 图片：与链接类似，只需在链接语法前面加一个感叹号(!)，例如：`![图片描述](图片链接)`。 6. 引用：在文本前面添加大于号(>)来表示引用，可以嵌套使用多个大于号表示多层引用。 7. 代码块：使用反引号(\`)将代码包裹起来，单行代码使用一个反引号，多行代码使用三个反引号。 8. 表格：使用竖线(|)分隔单元格，使用破折号(-)表示表头和表格内容的分隔线。 这只是Markdown语法的一小部分，更多详细的语法和用法可以参考Markdown的官方文档或其他相关教程。
评论 9
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

分享放大价值

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值