【GPU】NVIDIA GPU 系列

NVIDIA GPU story

GPU诞生契机

开始GPU故事的讨论之前,我们先来看一下CPU,CPU是顺序处理器,可以相继处理一条条指令,现在CPU一般都是多核心,可以同时处理多条指令和多个线程。
CPU 擅长执行并行性有限的程序和繁重的计算,如数学计算和串行排序算法。

但是,举个简单的例子,在图像处理任务中,有 1024 * 1024 * 3 个像素需要同时处理,以便进行渲染,CPU 通过串行处理(或半串行处理,我这么说是因为 CPU 可以通过利用线程中的指令级并行性(ILP)和同时执行多个线程来加快处理速度)来降低处理速度。

ILP:Instruction-Level Parallelism,指令级并行,是指在没有改变程序执行顺序的前提下,计算机体系结构和编译器技术通过同时执行多个指令来提高性能的技术。简单来说,ILP 是指在一个时钟周期内执行多个指令的能力,目的是提高处理器的利用率和吞吐量。

因此,计算机架构师们想到了使用协处理器,它可以提供更高的并行性,但只能进行简单的处理,比如在灰度滤波器下改变每个像素的值,而且只能专门用于图形处理

早期的 GPU 是专为图形处理而设计的特定功能协处理器。 它们是功能固定的图形流水线。

早期传统 GPU 如何用于通用计算

GPU 只计算每个像素的颜色,这也被称为 “渲染”(Rendering)。 使用 DirectXOpenGL 等图形 API 可以对它们进行编程。

在传统 GPU 上执行图形处理(或映射其他计算)的转换层

在深入探讨早期 GPU 硬件结构的更多细节之前,我们需要先定义一些术语,包括着色器(shader)、顶点(vertex)和像素(pixel)。

  • 着色器(shader):执行图形相关任务的计算机程序,可以分为顶点着色器、像素着色器、几何着色器等
  • 顶点(vertex):描述特定属性的数据结构,例如二维或三维空间中一个点的位置,或一个表面上多个点的位置。在计算机图形学中,顶点通常表示一个三维空间中的点,该点具有位置坐标(x, y, z)以及其他属性,如颜色、纹理坐标、法线向量等。顶点可以用来定义几何对象的形状,如多边形网格中的点。
  • 顶点着色器(vertex shader):将虚拟空间中每个顶点的三维位置转换为其在屏幕上显示的二维坐标的程序。
  • 像素(片段)着色器(pixel or fragment shader):计算单个像素或片段的颜色、亮度、对比度和其他属性的程序。

传统 GPU(图形处理单元)的设计架构中,特别是在早期 GPU 中,处理顶点(vertex)和片段(fragment,也称像素 pixel)的计算单元是分开的。

CUDA诞生之前的GPU

下表列出了英伟达™(NVIDIA®)公司从 1998 年到 2004 年推出的不同微体系结构的 GPU,以及晶体管数量、制造工艺、支持的 OpenGL 和 DirectX 版本和使用 GPU 芯片供电的 GPU 显卡等规格。 这些 GPU 都无法使用 CUDA 进行编程。摩尔定律效应在晶体管数量逐年增加的表格中显而易见。

1998-2004 GPU

GeForce 256 GPU 显卡通常被称为 “第一代 GPU”。 它拥有一个固定功能的 32 位浮点顶点变换和光照处理器,以及一个固定功能的整数像素片段流水线,该流水线采用 OpenGL 和 Microsoft DX7 API 编程。

GeForce 3 系列推出了第一个可编程顶点处理器,执行顶点着色器,以及一个可配置的 32 位浮点片段流水线,该流水线采用 DX8 和 OpenGL 编程。

CUDA诞生之后的GPU

2006-2022 GPU

首先介绍一下CUDA,CUDA 是计算统一设备架构(Compute Unified Device Architecture)的缩写,是一种并行计算框架平台和API,允许软件使用某些类型的 NVIDIA GPU 进行通用计算。C/C++程序员可以使用 "CUDA C/C++"对 NVIDIA GPU进行编程,其中代码将通过 NVCC 编译器编译为 PTX

NVCC:NVCC 是 NVidia Cuda Compiler 的缩写,这是一款基于 LLVM 的编译器。

PTX:PTX 是并行线程执行(Parallel Thread Execution)的缩写,它是一种低级并行线程执行虚拟机和指令集架构(ISA),用于英伟达™(NVIDIA®)CUDA 环境。 使用 PTX 对 GPU 进行编程就像为 CPU 开发汇编程序一样。

下图显示了 NVCC 如何编译 CUDA C/C++ 程序:

NVCC编译CUDA程序流程

如图所示,CPU代码启动并配置GPU,然后GPU执行Kernel并对数据进行操作,最后CPU从GPU接收数据并继续执行其他任务。

为了更深入地了解 CUDA C/C++ 程序在成为可执行程序之前会发生什么,请看下图:

CUDA程序编译

  • p.cu:这代表了一个CUDA源文件,通常扩展名为.cu。这个文件包含了 CPU code 和 GPU code,其中 GPU code 被标记为__global__
  • NVCC Compiler:这是NVIDIA编译器,专门用于编译CUDA程序。
  • Translation to PTX:这表示编译过程的第一步,即将.cu源代码编译成PTX代码。PTX代码是针对特定架构的GPU指令集,它比原始的C代码更接近硬件。
  • GPU Driver translates PTX to binary code:这是与GPU硬件交互的软件接口。它负责管理GPU资源,如内存和计算能力,并且将PTX代码转换成GPU可以执行的二进制代码(SASS)。SASS代码是最终被GPU硬件执行的指令。
  • GPU:这是图形处理单元,它是执行CUDA程序的主要硬件。GPU执行SASS代码,并进行并行计算。

接下来将详细介绍从2006年到2022年不同架构的NVIDIA GPU:

200620102012201420162017201820202022
TeslaFermiKeplerMaxwellPascalVolta(Workstation)Turing(Consumer, workstation)Ampere(data center)Hopper(data center)
引入CUDA并行计算平台首个完整GPU计算架构SM改名为SMX,SM中CUDA core数量192个SM中CUDA core数量下降至128个NVLink1.0NVLink2.0、Tensor Core1.0Tensor Core2.0、RT Core1.0Tensor Core3.0、RT Core2.0、NVLink3.0、MIG1.0Tensor Core4.0、NVLink4.0、MIG2.0

Tesla 2006——英伟达™(NVIDIA®)图形处理器微架构的重大变革

2006 年,NVIDIA 推出了一种新的微架构,实现了统一着色器模型。 其目的是更好地管理硬件资源利用率以及顶点和片段(像素)阶段之间的负载平衡。顶点和片段处理器之间的工作负载不平衡,原因在于像素和顶点的数量不匹配。有了统一着色器模型,就无需为顶点和像素设置不同的阶段。

硬件的改变方式是用于所有片段处理,从而简化了 GPU 的设计。 随后,GPU 内部的内核也变得更加简单。 内核变成了顺序内核和标量内核,一次只能处理一个计算任务。 如今,这些内核被称为 CUDA 内核。

这些内核被归入 SM 名下,取代了顶点和片段单元的阶段。 只需交换内核,就能轻松应对负载平衡挑战。
每个 SM 以 32 个线程为一组接收线程,这些线程的组合被称为 warp。 warp中的所有线程在同一时间执行相同的指令,但数据不同。 这就是它被称为 SIMT 的原因。

Tesla架构图

  • System memory:系统内存,GPU 通过 PCIe 接口与之交互。
  • Vertex Work Distribution:将顶点数据分发到不同的 Streaming Processors 进行处理。
  • Pixel Work Distribution:将像素数据分发到不同的 Streaming Processors 进行处理。
  • Compute Work Distribution:将计算任务分发到不同的 Streaming Processors 进行处理。
  • Streaming Processor Array (SPA):流式处理器阵列,包含多个 TPC(Thread Processing Cluster)。
  • TPC (Thread Processing Cluster):线程处理集群,每个 TPC 包含多个 SM(Streaming Multi-Processor)。
  • ROP (Raster Operation Processor)‌:这是处理光栅化操作的核心,例如像素着色和深度缓冲等;
  • DRAM (Dynamic Random Access Memory)‌:这是GPU的显存,用于存储纹理、缓冲区和其他大量数据;

TPC细节如下:

Tesla TPC detail

SM细节如下:

Tesla SM detail

  • Round-Robin调度策略(时间片轮转调度)是指轮流让每个warp执行一小段时间,这样每个warp都有机会运行,不会有一个warp一直占用资源。
  • Greedy-Fetch策略是指从一个warp中不断地获取指令或数据,直到在缓存中找不到所需的数据(即发生缓存缺失)。这时,就需要从更慢的内存中获取数据。这种策略是在尽可能多地利用当前warp,直到它因为等待数据而停下来,然后可能会切换到另一个不需要等待数据的warp。
  • 一种是公平地让每个warp都有执行的机会(Round-Robin),另一种是在一个 warp上尽可能多地执行直到它需要等待数据(Greedy-Fetch)。
Tesla(或 CUDA-powered)之后的 GPU 执行模式

每个 grid(cuda kernel)中都包含多个线程块 thread block(它们的另一个名称:CTA,Cooperative Thread Array,协作线程数组),每个 CTA 都有特定数量的线程。当 CUDA 程序在 GPU 上启动执行时,并不能确定哪个 SM 将为哪个 CTA 服务。
下图显示了执行过程。

GPU 程序程序执行模式

warp调度

warp是具有连续线程 ID 的一组线程,设计warp的理念是为了让硬件调度更轻便、更容易。warp中的线程数由硬件结构决定。

请注意,GPU 并行程序员无需了解任何有关 warp 的知识!

下图显示了 warp 调度是如何让硬件调度变得更轻松、更简单的。

warp调度

这张图的最上方显示了一台织布机正在编织布料,这实际上是在类比 GPU 中的多线程调度。就像织布机同时处理多条线一样,GPU 的 SM 也能够同时处理多个线程组。
然而,由于硬件限制,任何给定的时间只能有一个 warp 在执行。所以,SM 会在不同的 warp 之间快速切换,给人一种同时执行的感觉。这种调度方法称为时间切片(Time-slicing),使得所有 warp 都能获得执行的机会,从而提高了整体性能。

SIMT 架构的挑战——分歧

当 warp 中的一部分线程需要执行一条指令,而其他线程必须执行另一条指令时,就会出现这种情况。 如下图所示,这种情况可能是由于代码中的分支造成的。

如图所示,这给性能带来了挑战,因为它限制了并行性,并使这些线程组的执行序列化

分支影响并行度

随着代码的执行,线程开始分叉(diverge),然后在执行 Z 之前重新汇聚(reconverge)。这种分支行为会导致 GPU 的流水线停顿,因为不是所有线程都在同一时间执行相同的指令。当遇到分支时,一些线程会暂时停止,等待其他线程完成他们的分支路径后再重新汇合。这种情况降低了 GPU 的性能,因为它无法充分利用所有可用的硬件资源。

分歧会导致序列化,从而降低性能

在避免stalls方面,GPU 与 CPU 有何不同?

CPU 通常专注于降低延迟(latency),而 GPU 则更关注吞吐量(throughput)

  • CPU 的设计目标:降低延迟
    • Cache(高速缓存):
      CPU 使用多层次的高速缓存(L1、L2、L3)来减少访问内存的时间
    • Speculative Execution(推测执行):
      CPU 通过推测执行来猜测哪些指令可能会被执行,并提前加载数据或将指令排队执行。如果预测正确,可以节省时间;如果预测错误,则会撤销推测的结果。
    • Out-of-Order Execution(乱序执行):
      CPU 会在不影响结果的情况下,重新排列指令的执行顺序,以充分利用计算单元。
    • Branch Prediction(分支预测):
      CPU 会预测程序中的分支指令结果,提前执行可能的路径,以减少分支带来的延迟。
  • GPU 的设计目标:提高吞吐量
    • 高吞吐量(High Throughput):
      GPU 通过同时处理大量线程来最大化吞吐量。即使某些线程在等待数据或其他资源,GPU 也可以切换到其他线程继续执行。
    • Saturating Memory Bus(饱和内存总线):
      GPU 设计了大量的内存带宽来支持数千个并发线程。即使每个线程的延迟较高,但总体来看,GPU 可以在单位时间内完成更多的计算任务。
    • Parallel Processing(并行处理):
      GPU 通过并行处理大量数据来提高吞吐量。例如,在图像处理中,每个像素可以被视为一个独立的任务,GPU 可以同时处理多个像素。

CPU vs. GPU 的应用场景

  • Latency-Oriented Programs(延迟导向程序):
    CPU 更适合处理延迟敏感的应用,如操作系统调度、数据库事务处理、实时控制系统等。这类应用要求响应时间尽可能短。
  • Throughput-Oriented Programs(吞吐量导向程序):
    GPU 更适合处理吞吐量敏感的应用,如图像渲染、机器学习、科学计算等。这类应用通常需要处理大量的数据,并且对单位时间内完成的工作量有较高要求。

Fermi microarchitecture (2010)

回到正题,我们继续讨论CUDA之后的GPU架构,在上文我们已讨论了Tesla的硬件架构以及CUDA提出以后,程序在GPU上的执行模式,以及warp调度、SIMT挑战、GPU与CPU避免stall的不同思想等问题。

Fermi架构设计的创新:

  • 增加SP数量:
    在 Fermi 架构中,NVIDIA 将每个 SM 内的 SP(又称为 CUDA core)数量翻倍。这种增加使得每个 SM 可以同时处理更多的线程,从而提高了并行处理能力。
  • 半线程组(Half-Warps):
    传统的线程组(Warp)通常包含 32 个线程。在 Fermi 架构中,增加了对半线程组的支持,即可以同时处理 16 个线程(半个 Warp)。这种设计使得在某些情况下,即使不是所有线程都可以完全填充一个 Warp,也能更有效地利用硬件资源。
  • 64 位浮点运算:
    Fermi 架构本身不支持单个 CUDA 核心进行 64 位浮点运算。但是,通过组合两个 CUDA core,可以实现 64 位浮点运算的能力。这种方式虽然不如专用的 64 位核心那样高效,但在一定程度上弥补了缺少 64 位支持的不足。
  • 32 位 ALU:
    Fermi 架构将 ALU 从以前的 24 位升级到了 32 位。这种改进提高了单精度浮点运算的性能。
  • 增强的 C++ 特性:
    Fermi 架构增加了更多的 C++ 特性支持,使得 CUDA 编程更加灵活和强大。
  • TPC 被 GPC 替换:
    Fermi 架构中,原先的 Thread Processing Clusters (TPCs) 被替换为 Graphics Processing Clusters (GPCs)。GPC 是更高层次的组织结构,每个 GPC 包含多个 SMs,进一步优化了并行处理能力。

Fermi 架构图

  • Giga Thread Engine是负责将线程块(CTA)调度到 SM 的单元

下图显示了Fermi SM 的结构:
dual warp scheduler 可以同时选择两个 warp,并从每个 warp 向一组 16 个core、16 个 LD/ST 单元或 4 个 SFU 发出一条指令。

Fermi SM detail

如上图所示,shared memory/L1 cache 可用于为单个线程缓存数据(用于寄存器溢出/L1 cache)和/或在线程块的线程间共享数据。

Kepler microarchitecture (2012)

亮点:

  • 节能,降低时钟频率,并将核心时钟与 GPU 卡频率统一起来。
  • 引入Hyper-Q/ MPS(多进程服务)技术,使单个 GPU 上可以同时启动多个 CPU 线程内核。其目的是提高 GPU 的利用率。
  • 取消了硬件调度程序,改用软件调度程序。此外,更多的 SM 被打包,每个 SM 都需要更多的资源。
  • SM的改进:Kepler架构中集成了更多的流式多处理器(SM),每个SM的核心数从32个到192个,每个SM都拥有更多的资源,并且为这种 SM 创造了一个新的缩写:“SMX”,代表下一代 SM。
  • 动态并行性(Dynamic Parallelism):它是对 CUDA 的扩展,使 CUDA kernel 能够通过启动新的 kernel 来创建新的线程网格(thread grid)。在早期的 GPU 中,kernel只能从主机(CPU)代码中启动。动态并行的优势在于减少主机负担和主机与设备之间的通信。

下图显示了具有动态并行性的 GPU 与没有动态并行性的旧 GPU 之间的差异。

动态并行有无的区别

SMX相比于SM做出的改进;

  • Warp 调度器的数量:
    Kepler 的 SMX 配备了四个 Warp Schedulers,每个 Warp 调度器能够在单个时钟周期内处理一个完整的 Warp(32 个线程);相比之下,Fermi 架构中的每个 SM 只有一个 Warp 调度器,并且每次只能处理半个 Warp(16 个线程)。
  • 双派发单元(Double Dispatch Units):
    每个 Warp 调度器都配备了两个 dispatch unit,这意味着当一个 Warp 中的两条指令相互独立时,可以同时执行这两条指令。这种设计使得 Kepler 的 SMX 能够在一个时钟周期内处理更多的指令,从而提高了计算效率。
  • CUDA core共享:
    在 Kepler 的 SMX 中,每列 32 个 CUDA 核心是由两个调度器共享的。这意味着每个调度器可以访问这些 CUDA 核心的一部分,从而实现更高的并行度。
    这种共享设计使得即使在处理依赖性较弱的指令时,仍然可以充分利用计算资源。

Kepler SMX detail

Kepler架构图

Maxwell microarchitecture (2014)

  • Maxwell 也注重能效,将 CUDA core数量从 Kepler SMX 的 192 个减少到 SMM 的 128 个,使内核数量与warp尺寸保持一致。这改进了芯片分区,节省了面积和功耗。
  • 与 Kepler 相比,Maxwell 采用了更简单的调度逻辑。
  • Maxwell 将 L2 cache 的大小从 256KB 增加到 2MB。

下图显示了 Maxwell 的整体架构和 SMM 的架构。

Maxwell架构图

Maxwell SMM detail

Pascal Microarchitecture (2016)

  • 在 Pascal 中,SM 与上一代相比没有变化。
  • 16 纳米工艺允许在相同的芯片面积上封装更多的晶体管。 这些晶体管用于增加寄存器文件(RF)和共享内存(或Scratchpad Memory)的大小。
  • 引入统一内存,在 "Page Migration Machine(PMM)"的帮助下,CPU 和 GPU 可以访问相同的内存地址空间。
  • NVLink 也与 Pascal 一起推出。它是一种基于线缆的高带宽通信协议,比 PCIe 更快,代表 NVidia Link。每个 GPU 可以拥有多个 NVLink。此外,GPU 还可以使用网状网络进行通信,而不是连接到中心枢纽进行数据交换。
  • 指令和线程级抢占(Instruction-Level Preemption 和 Thread-Level Preemption)在这种微架构中成为可能。这种能力使得在执行过程中可以中断当前的指令或线程,并恢复到中断之前的状态,这对于需要精确控制执行流程的应用非常有用。
  • 此外,HBM 2 内存(三维堆叠内存)被用作 GPU 的主内存。

Pascal架构图

Volta (workstation) microarchitecture 2017

亮点:

  • NVLink 2.0
  • 第一代张量核
    • 张量核将两个 4*4 的 16 位矩阵相乘,然后通过融合乘加(FMA)操作将第三个 16 位或 32 位矩阵添加到结果中,得到 32 位浮点结果,该结果可优化降级为 16 位浮点结果。
    • systolic array是由处理元件(PE)组成的阵列,在将数据输出到内存之前对数据流进行协调,从而实现计算、内存和 I/O 带宽之间的平衡。 将张量核集成到 GPU 的 SM 中,可以加速深度学习应用的训练和推理过程。

下图显示了张量核的功能:
Tensor core

Turing (consumer) microarchitecture 2018

Turing微体系结构发生了巨大变化,趋向于
(1)人工智能(AI),在 SM 中增加了张量核,
(2)图像处理,在 SM 中增加了光线跟踪核。
Turing 与 Tesla 之前的架构非常相似,因为它采用了分层架构,如下图所示:

Turing架构图

变化如下:

  1. SP 或 CUDA 内核成为超标量处理器(superscalar processors),能够像我们在 CPU 中看到的那样并行执行整数和浮点运算,例如英特尔的奔腾(Pentium)
  2. 配备 16 个内存控制器的 GDDR6X 内存子系统可提供极大的带宽
  3. 由于每个线程都有自己的指令指针 (IP),因此 warp 中的线程不会共享它们的指令指针 (IP),因此它们可以独立调度,这意味着细粒度调度(线程级调度与 warp 级调度相比)需要更多的硬件开销。

Ampere microarchitecture (2020)

随着Ampere的推出,英伟达还推出了第三代张量核。新的张量内核(TC)对稀疏性进行了优化,这意味着它们能以优化的方式更快地处理稀疏计算。32 位张量浮点运算和 64 位浮点运算(FP64)为高性能计算提供了更高的精度,为人工智能应用提供了高达 20 倍的速度提升。
下图显示了张量内核如何提高计算性能。

支持稀疏计算

借助 Ampere 微体系结构,英伟达™(NVIDIA®)推出了一项名为 "多实例 GPU(MIG)"的新功能。 该功能可将每个 GPU 划分为多个较小的 GPU 实例。 这些较小的 GPU 实例拥有各自的高带宽内存、高速缓存和计算内核,在硬件层面上相互安全隔离。 下图显示了一个 GPU 可以划分为多个较小的 GPU,其计算和内存资源完全隔离。

MIG示意图

其他亮点:

  • 推出了 NVLink 版本 3,与 PCIe 相比,速度提高了 10 倍。
  • 专门用于图形处理的光线追踪(RT)内核有了更多改进
  • GPU 的内存带宽更大,二级缓存也大大增加(比上一代产品增加了 7 倍)。

Hopper microarchitecture (2022) — data centers

与Ampere微架构相比,该架构改进了不同的组件,

  • 内存系统升级为HBM2
  • 第二代MIG技术
  • 新的机密计算支持
  • 第四代英伟达NVLink
  • 第三代NVSwitch
  • 新的NVLink交换系统和PCIe Gen 5。

新的 SM 通过引入第四代张量内核、新的 DPX(动态编程)指令、更快的 IEEE 64 位和 32 位浮点运算、新的线程块集群、分布式共享内存和新的异步执行功能(包括张量内存加速器(TMA))提高了性能。

参考文章

在C++中调用英伟达GPU进行计通常使用CUDA(Compute Unified Device Architecture)框架。CUDA是英伟达提供的一种并行计平台和编程模型,允许开发者使用C++等高级编程语言来编写在GPU上运行的程序。以下是一个简单的示例,展示了如何在C++中使用CUDA调用GPU进行计: 1. 首先,确保你已经安装了CUDA工具包,并且你的系统上有一个兼容的英伟达GPU。 2. 编写CUDA内核函数: ```cpp // kernel.cu extern "C" __global__ void add(int *a, int *b, int *c, int n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) { c[idx] = a[idx] + b[idx]; } } ``` 3. 在C++代码中调用这个内核函数: ```cpp // main.cpp #include <iostream> #include <cuda_runtime.h> extern void add(int *a, int *b, int *c, int n); int main() { int n = 1024; int *a, *b, *c; size_t size = n * sizeof(int); // 分配主机内存 a = (int *)malloc(size); b = (int *)malloc(size); c = (int *)malloc(size); // 初始化输入数据 for (int i = 0; i < n; i++) { a[i] = i; b[i] = i * 2; } // 分配设备内存 int *d_a, *d_b, *d_c; cudaMalloc((void **)&d_a, size); cudaMalloc((void **)&d_b, size); cudaMalloc((void **)&d_c, size); // 将数据从主机复制到设备 cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice); // 调用内核函数 add<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n); // 将结果从设备复制回主机 cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost); // 打印结果 for (int i = 0; i < n; i++) { std::cout << a[i] << " + " << b[i] << " = " << c[i] << std::endl; } // 释放内存 free(a); free(b); free(c); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); return 0; } ``` 4. 编译和运行: ```sh nvcc -c kernel.cu -o kernel.o g++ main.cpp kernel.o -o main -lcudart ./main ``` 这个示例展示了如何在C++中使用CUDA调用英伟达GPU进行简单的加法计。你可以根据需要扩展这个示例,编写更复杂的内核函数和处理更复杂的数据结构。
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

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

余额充值