模型部署——cuda编程入门

最新推荐文章于 2025-05-05 14:12:43 发布
最新推荐文章于 2025-05-05 14:12:43 发布
阅读量790 收藏 18
点赞数 24
文章标签： cuda nsight systems c++
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/m0_52835296/article/details/147433112
版权

CUDA中的线程与线程束

在这里插入图片描述

  • kernel是在device上线程中并行执行的函数，核函数用__global__符号声明，在调用时需要用<<<grid_size, block_size>>>来指定kernel要执行的线程数量。在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号thread ID，这个ID值可以通过核函数的内置变量threadIdx来获得。
  • synchronize是同步的意思，有几种synchronize
    cudaDeviceSynchronize: CPU与GPU端完成同步，CPU不执行之后的语句，直到这个语句以前的所有cuda操作结束
    cudaStreamSynchronize: 跟cudaDeviceSynchronize很像，但是这个是针对某一个stream的。只同步指定的stream中的cpu/gpu操作，其他的不管
    cudaThreadSynchronize: 现在已经不被推荐使用的方法
    __syncthreads: 线程块内同步
  • 核函数编写和调用举例
#include <cuda_runtime.h>
#include <stdio.h>

// 核函数
__global__ void print_idx_kernel(){
    printf("block idx: (%3d, %3d, %3d), thread idx: (%3d, %3d, %3d)\n",
         blockIdx.z, blockIdx.y, blockIdx.x,
         threadIdx.z, threadIdx.y, threadIdx.x);
}

void print_one_dim(){
    int inputSize = 8;
    int blockDim = 4;
    int gridDim = inputSize / blockDim;

    dim3 block(blockDim);
    dim3 grid(gridDim);
    
    // 核函数调用
    print_idx_kernel<<<grid, block>>>();
    cudaDeviceSynchronize();
}

.cu与.cpp的相互引用及Makefile

编译器：gcc g++ nvcc
举个例子：

nvcc print_index.cu -o app -I /usr/local/cuda/include

获取编译器选项：

g++ --help
nvcc --help

Makefile编写（是否可以使用CMakeLists.txt?）
.cpp中不能直接调用核函数，需要在.cu中提供调用接口

使用CUDA进行MATMUL计算

host端与device端数据传输

在这里插入图片描述
host端与device端数据传输代码实现：

void MatmulOnDevice(float *M_host, float *N_host, float* P_host, int width, int blockSize){
    /* 设置矩阵大小 */
    int size = width * width * sizeof(float);

    /* 分配M, N在GPU上的空间*/
    float *M_device;
    float *N_device;

    cudaMalloc(&M_device, size);
    cudaMalloc(&N_device, size);

    /* 分配M, N拷贝到GPU上*/
    cudaMemcpy(M_device, M_host, size, cudaMemcpyHostToDevice);
    cudaMemcpy(N_device, N_host, size, cudaMemcpyHostToDevice);

    /* 分配P在GPU上的空间*/
    float *P_device;
    cudaMalloc(&P_device, size);

    /* 调用kernel来进行matmul计算, 在这个例子中我们用的方案是：将一个矩阵切分成多个blockSize * blockSize的大小 */
    dim3 dimBlock(blockSize, blockSize);
    dim3 dimGrid(width / blockSize, width / blockSize);
    MatmulKernel <<<dimGrid, dimBlock>>> (M_device, N_device, P_device, width);

    /* 将结果从device拷贝回host*/
    cudaMemcpy(P_host, P_device, size, cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();

    /* Free */
    // free与malloc的顺序是反着的
    cudaFree(P_device);
    cudaFree(N_device);
    cudaFree(M_device);
}

cuda core矩阵乘法核函数编写

/* matmul的函数实现*/
__global__ void MatmulKernel(float *M_device, float *N_device, float *P_device, int width){
    /* 
        我们设定每一个thread负责P中的一个坐标的matmul
        所以一共有width * width个thread并行处理P的计算
    */
    // 确定负责计算的结果元素的索引
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    float P_element = 0;

    /* 对于每一个P的元素，我们只需要循环遍历width次M和N中的元素就可以了*/
    for (int k = 0; k < width; k ++){
        float M_element = M_device[y * width + k];
        float N_element = N_device[k * width + x];
        P_element += M_element * N_element;
    }

    P_device[y * width + x] = P_element;
}

cuda core 每个线程执行核函数计算一个结果元素
GPU刚开始执行核函数的时候，会存在一个warmup阶段，耗时会比较长
CPU与GPU的浮点运算会存在误差，误差控制在e-4以内是ok的
CUDA中规定，一个block中可以分配的thread的数量最大是1024个线程，如果大于1024会显示配置错误
为什么block size = 1的时候比等于16的时候慢很多？

cuda中的error handler

在这里插入图片描述

获取GPU的硬件信息

利用cuda runtime api打印硬件信息 & LOG
#include <stdio.h>
#include <cuda_runtime.h>
#include <string>

#include "utils.hpp"

int main(){
    int count;
    int index = 0;
    cudaGetDeviceCount(&count);
    while (index < count) {
        cudaSetDevice(index);
        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, index);
        LOG("%-40s",             "*********************Architecture related**********************");
        LOG("%-40s%d%s",         "Device id: ",                   index, "");
        LOG("%-40s%s%s",         "Device name: ",                 prop.name, "");
        LOG("%-40s%.1f%s",       "Device compute capability: ",   prop.major + (float)prop.minor / 10, "");
        LOG("%-40s%.2f%s",       "GPU global meory size: ",       (float)prop.totalGlobalMem / (1<<30), "GB");
        LOG("%-40s%.2f%s",       "L2 cache size: ",               (float)prop.l2CacheSize / (1<<20), "MB");
        LOG("%-40s%.2f%s",       "Shared memory per block: ",     (float)prop.sharedMemPerBlock / (1<<10), "KB");
        LOG("%-40s%.2f%s",       "Shared memory per SM: ",        (float)prop.sharedMemPerMultiprocessor / (1<<10), "KB");
        LOG("%-40s%.2f%s",       "Device clock rate: ",           prop.clockRate*1E-6, "GHz");
        LOG("%-40s%.2f%s",       "Device memory clock rate: ",    prop.memoryClockRate*1E-6, "Ghz");
        LOG("%-40s%d%s",         "Number of SM: ",                prop.multiProcessorCount, "");
        LOG("%-40s%d%s",         "Warp size: ",                   prop.warpSize, "");

        LOG("%-40s",             "*********************Parameter related************************");
        LOG("%-40s%d%s",         "Max block numbers: ",           prop.maxBlocksPerMultiProcessor, "");
        LOG("%-40s%d%s",         "Max threads per block: ",       prop.maxThreadsPerBlock, "");
        LOG("%-40s%d:%d:%d%s",   "Max block dimension size:",     prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2], "");
        LOG("%-40s%d:%d:%d%s",   "Max grid dimension size: ",     prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2], "");
        index ++;
        printf("\n");
    }
    return 0;
}
Roofline model（待补充）

Nsight system and Nsight compute

谷歌搜索下载:官网链接

Nsight system

参考链接

安装目录:
ls /usr/local/bin |grep nsys
nsys
nsys-ui

启动GUI界面
sudo ./nsys-ui(不加sudo会存在权限问题）

举个例子：
配置可执行文件以及感兴趣内容：
在这里插入图片描述可视化分析：
在这里插入图片描述详细使用手册：官网文档

Nsight compute
查看可安装版本：
sudo apt policy nsight-compute-2022.2.1
安装：
sudo apt install nsight-compute-2022.2.1
查看安装位置：
dpkg -L nsight-compute-2022.2.1
路径：/opt/nvidia/nsight-compute/2022.2.1/
文件：ncu ncu-ui等

启动：
sudo ./ncu-ui

举个例子：
基本配置：replay mode: application
在这里插入图片描述选择感兴趣内容：
在这里插入图片描述launch即可，第一次运行会比较慢，会重复运行很多次。
结果：
在这里插入图片描述不知道为什么roofline model没有正常显示出来，需要查一查？

扩展知识

共享内存以及BANK CONFLICT

shared memory

硬件结构

SM（Streaming Multiprocessor）
在CUDA编程模型中，线程被组织成线程块（block），多个线程块组成一个网格（grid）。每个线程块被分配到一个SM中执行，而SM内部的warp调度器会将线程块中的线程分成多个warp进行执行。
当一个warp中的线程需要等待某些操作（例如内存访问）完成时，SM可以切换到另一个warp继续执行，从而提高计算效率。
在这里插入图片描述

核函数编写
#include "cuda_runtime_api.h"
#include "utils.hpp"

#define BLOCKSIZE 16

/* 
    使用shared memory把计算一个tile所需要的数据分块存储到访问速度快的memory中
*/
__global__ void MatmulSharedStaticKernel(float *M_device, float *N_device, float *P_device, int width){
    __shared__ float M_deviceShared[BLOCKSIZE][BLOCKSIZE];
    __shared__ float N_deviceShared[BLOCKSIZE][BLOCKSIZE];
    /* 
        对于x和y, 根据blockID, tile大小和threadID进行索引
    */
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    float P_element = 0.0;

    int ty = threadIdx.y;
    int tx = threadIdx.x;
    /* 对于每一个P的元素，我们只需要循环遍历width / tile_width 次就okay了，这里有点绕，画图理解一下*/
    for (int m = 0; m < width / BLOCKSIZE; m ++) {
        M_deviceShared[ty][tx] = M_device[y * width + (m * BLOCKSIZE + tx)];
        N_deviceShared[ty][tx] = N_device[(m * BLOCKSIZE + ty)* width + x];
        __syncthreads(); // 上述两句所有thread都会执行，等待所有thread执行完成

        for (int k = 0; k < BLOCKSIZE; k ++) {
            P_element += M_deviceShared[ty][k] * N_deviceShared[k][tx];
        }
        __syncthreads();
    }

    P_device[y * width + x] = P_element;
}

__global__ void MatmulSharedDynamicKernel(float *M_device, float *N_device, float *P_device, int width, int blockSize){
    /* 
        声明动态共享变量的时候需要加extern，同时需要是一维的 
        注意这里有个坑, 不能够像这样定义： 
            __shared__ float M_deviceShared[];
            __shared__ float N_deviceShared[];
        因为在cuda中定义动态共享变量的话，无论定义多少个他们的地址都是一样的。
        所以如果想要像上面这样使用的话，需要用两个指针分别指向shared memory的不同位置才行
    */

    extern __shared__ float deviceShared[];
    int stride = blockSize * blockSize;
    /* 
        对于x和y, 根据blockID, tile大小和threadID进行索引
    */
    int x = blockIdx.x * blockSize + threadIdx.x;
    int y = blockIdx.y * blockSize + threadIdx.y;

    float P_element = 0.0;

    int ty = threadIdx.y;
    int tx = threadIdx.x;
    /* 对于每一个P的元素，我们只需要循环遍历width / tile_width 次就okay了 */
    for (int m = 0; m < width / blockSize; m ++) {
        deviceShared[ty * blockSize + tx] = M_device[y * width + (m * blockSize + tx)];
        deviceShared[stride + (ty * blockSize + tx)] = N_device[(m * blockSize + ty)* width + x];
        __syncthreads();

        for (int k = 0; k < blockSize; k ++) {
            P_element += deviceShared[ty * blockSize + k] * deviceShared[stride + (k * blockSize + tx)];
        }
        __syncthreads();
    }

    if (y < width && x < width) {
        P_device[y * width + x] = P_element;
    }
}

动态共享内存比静态共享内存速度慢，没有特殊情况下，使用静态共享内存。

cuda event进行时间测算

在这里插入图片描述

BANK CONFLICT（存储体冲突）

在shared memory中什么是bank?

在这里插入图片描述

什么时候会发生bank conflict

在这里插入图片描述

按行存储，按列访问的时候，会发生bank conflict:
在这里插入图片描述

如何减缓bank conflict

在这里插入图片描述

代码实现
#include "cuda_runtime_api.h"
#include "utils.hpp"

#define BLOCKSIZE 16

/* 
    使用shared memory把计算一个tile所需要的数据分块存储到访问速度快的memory中
*/
__global__ void MatmulSharedStaticConflictPadKernel(float *M_device, float *N_device, float *P_device, int width){
    /* 添加一个padding，可以防止bank conflict发生，结合图理解一下*/
    __shared__ float M_deviceShared[BLOCKSIZE][BLOCKSIZE + 1];
    __shared__ float N_deviceShared[BLOCKSIZE][BLOCKSIZE + 1];
    /* 
        对于x和y, 根据blockID, tile大小和threadID进行索引
    */
    int x = blockIdx.x * BLOCKSIZE + threadIdx.x;
    int y = blockIdx.y * BLOCKSIZE + threadIdx.y;

    float P_element = 0.0;

    int ty = threadIdx.y;
    int tx = threadIdx.x;
    /* 对于每一个P的元素，我们只需要循环遍历width / tile_width 次就okay了，这里有点绕，画图理解一下*/
    for (int m = 0; m < width / BLOCKSIZE; m ++) {
        /* 这里为了实现bank conflict, 把tx与tx的顺序颠倒，同时索引也改变了*/
        M_deviceShared[tx][ty] = M_device[x * width + (m * BLOCKSIZE + ty)];
        N_deviceShared[tx][ty] = M_device[(m * BLOCKSIZE + tx)* width + y];

        __syncthreads();

        for (int k = 0; k < BLOCKSIZE; k ++) {
            P_element += M_deviceShared[tx][k] * N_deviceShared[k][ty];
        }
        __syncthreads();
    }

    /* 列优先 */
    P_device[x * width + y] = P_element;
}

__global__ void MatmulSharedDynamicConflictPadKernel(float *M_device, float *N_device, float *P_device, int width, int blockSize){
    /* 
        声明动态共享变量的时候需要加extern，同时需要是一维的 
        注意这里有个坑, 不能够像这样定义： 
            __shared__ float M_deviceShared[];
            __shared__ float N_deviceShared[];
        因为在cuda中定义动态共享变量的话，无论定义多少个他们的地址都是一样的。
        所以如果想要像上面这样使用的话，需要用两个指针分别指向shared memory的不同位置才行
    */

    extern __shared__ float deviceShared[];
    int stride = (blockSize + 1) * blockSize;
    /* 
        对于x和y, 根据blockID, tile大小和threadID进行索引
    */
    int x = blockIdx.x * blockSize + threadIdx.x;
    int y = blockIdx.y * blockSize + threadIdx.y;

    float P_element = 0.0;

    int ty = threadIdx.y;
    int tx = threadIdx.x;
    /* 对于每一个P的元素，我们只需要循环遍历width / tile_width 次就okay了 */
    for (int m = 0; m < width / blockSize; m ++) {
        /* 这里为了实现bank conflict, 把tx与tx的顺序颠倒，同时索引也改变了*/
        deviceShared[tx * (blockSize + 1) + ty]             = M_device[x * width + (m * blockSize + ty)];
        deviceShared[stride + (tx * (blockSize + 1) + ty)]  = N_device[(m * blockSize + tx) * width + y];

        __syncthreads();

        for (int k = 0; k < blockSize; k ++) {
            P_element += deviceShared[tx * (blockSize + 1) + k] * deviceShared[stride + (k * (blockSize + 1 ) + ty)];
        }
        __syncthreads();
    }

    /* 列优先 */
    P_device[x * width + y] = P_element;
}

STREAM和EVENT

什么是stream

在这里插入图片描述

参考下述链接，理解cuda编程的一些基础概念：
理解CUDA中的thread,block,grid和warp
cuda stream的使用

多流编程实现

单流：
在这里插入图片描述
多流：
在这里插入图片描述
利用nsight systems进行分析：
在这里插入图片描述

如何利用多流进行隐藏访存和核函数执行延迟的调度

举一个栗子：
在这里插入图片描述

使用CUDA进行预处理/后处理

双线性插值

在这里插入图片描述

双线性插值的cuda实现

在这里插入图片描述

查看图片大小：
identity xx.png
可视化图片：
feh xx.png
实战千问2大模型第一天——Qwen2-7B（知识问答）的部署和fastapi封装
qq_34717531的博客
09-10 1543
Qwen2-7B是由阿里云通义千问团队开发的一系列大型语言模型中的一个。这个模型是基于Transformer架构的，包含自注意力和因果掩码机制，专注于处理和生成自然语言文本。Qwen2-7B是在超大规模的预训练数据上训练得到的，这些数据不仅包括中英文，还包含其他25种语言的高质量数据。这个模型不仅在语言理解和生成方面表现出色，还在编程、数学解题等多个领域显示了其先进的性能。
NVIDIA DOCA产品体验——DPU及DOCA轻松入门与实战
weixin_42553583的博客
01-17 977
在人工智能之前的时代，CPU是计算机的核心，承载了所有计算和调度，并与内存直接交互。到了人工智能时代，算力、算法和数据是三大核心要素。单纯依靠CPU无法承载AI所需的巨大算力，于是有了GPU。虽然解决了算力问题，但是面对海量的训练数据需要处理，无论CPU还是GPU,都显得不够专业，所以还需要一个专门数据处理器，于是就有了DPU。DPU中文名称数据处理器，其目的就是应对人工智能时代爆发式的数据增长。所以，在未来，整个数据中心不再只是基于 CPU，还会引入用于加速计算的 GPU 和用于加速数据处理的 DPU。
cuda入门——记录
Downsiren的博客
06-02 268
cuda学习过程重要的东西记录
Cuda异步计算并行编程原理和存储管理
tugouxp的专栏
12-17 2247
优化时遵守的一些原则：1.Grid一定要给足block.2.Block内一定要给足thread,目的是提高并发WARP的数目，隐藏延迟.3.Block内线程的数目一定是warpsize的整数倍.每个warp的执行上下文（execution context，如程序计数器 和 寄存器等）在warp的整个生命周期内都被保存在片上内存（on-chip memory）。因此从一个执行上下文切换到另一个执行上下文是无开销的。
AI模型推理（１）——入门篇
chenxy02的博客
04-30 3656
本文主要介绍AI模型推理的相关基础概念，为后续云原生模型推理服务的学习做准备。
CUDA和TensorRT入门
weixin_40620310的博客
04-17 2219
1、对于深度神经网络的推理，TRT可以充分发挥GPU的算力，以及节省GPU的存储空间；2、要多参考官方源码的sample案例，尝试替换现有模型，再深入了解API进行网络的搭建；3、如果要使用自定义组件，至少先了解CUDA基本架构以及常用属性；4、推荐使用FP16（定义很少变量，明显能提高速度，精度影响不大）和INT8（更大的潜力，可能导致精度下降）这两种量化模式；5、在不同架构的GPU或者不同的软件版本的设备上，引擎不能通用，要重新生成一个；
CUDA~入门
whaosoft143ai的博客
04-16 319
Hello WorldGPU编程涉及到多个设备（CPU，GPU，内存，显存），因此首先明确概念Host：CPU + 内存Device：GPU + 显存A "hello world" example 和普通C程序的区别函数 返回值之前的 前缀，告诉编译器，这个函数到时候会跑在GPU上函数调用时的 ，在CPU上调用GPU实现的函数， 是因为可能会调用若干个GPU”核心“ 同时跑这个函数"hello world" 例子里，GPU实现的函数并没有接收任何参数，也没有具体的返回值。实际有意义的函数，通常会有
大模型部署：使用Ollama本地部署Llama 3.1大模型
2401_82469710的博客
01-13 1664
我们今天就在百度智能云 GPU 服务器上来部署体验最新 Llama3.1 8B 模型。
深度学习部署神器——triton inference server入门教程指北
老潘的博客
10-19 5846
triton作为一个NVIDIA开源的商用级别的服务框架，个人认为很好用而且很稳定**，API接口的变化也不大，我从2020年的20.06切换到2022年的22.06，两个大版本切换，一些涉及到代码的工程变动很少，稍微修改修改就可以直接复用，很方便。 > 本系列讲解的版本也是基于22.06。
一文详解AI模型部署及工业落地方式
3D视觉工坊
06-29 3510
点击上方“3D视觉工坊”，选择“星标”干货第一时间送达Hello大家好，我是老潘，好久不见各位~最近在复盘今年上半年做的一些事情，不管是训练模型、部署模型搭建服务，还是写一些组件代码等，零...
服务器模型部署全攻略：从零到运行
2401_85390073的博客
07-03 1325
预训练属于迁移学习的范畴。现有的神经网络在进行训练时，一般基于后向传播（BackPropagation，BP）算法，先对网络中的参数进行随机初始化，再利用[随机梯度下降（Stochastic GradientDescent，SGD）]等优化算法不断优化模型参数。而预训练的思想是，模型参数不再是随机初始化的，而是通过一些任务进行预先训练，得到一套模型参数，然后用这套参数对模型进行初始化，再进行训练。在正式探讨自然语言处理的预训练之前，回顾一下计算机视觉领域的预训练过程。
CUDA编程入门指南：新手直达高性能计算
整个教程以清晰的结构和易于理解的语言展开，旨在使读者不仅能够快速入门，还能逐步深入理解CUDA编程的复杂性，并学会如何将CUDA技术融入到实际项目中，提升计算效率。无论是对GPU编程有兴趣的开发者还是希望利用GPU...
本地部署Qwen2大模型之六：几种AI大模型部署模式的比较
F8_11的博客
12-28 1200
在动手实践本地部署Qwen2大模型的过程中，我遇到了很多的实际问题，花了很多时间和心思来解决。也正是解决这些问题的经历，让我逐渐对AI大模型的部署有了比较清晰的了解，也形成了几种不同部署模式的概念。本文尝试对几种典型的AI大模型部署模式进行对比分析，为后期围绕AI开发应用提供参考。
本地（家庭电脑）使用大模型与微调（简单入门版）
qwwewq111的博客
01-28 1663
大模型微调训练 利用LLaMA-Factory-main对模型进行本地训练 包括工具安装过程
＜C++＞冒泡排序、归并排序详解 时间复杂度 与应用
weixin_51710308的博客
05-05 442
简单但效率低，适合教学或小数据。
深入理解C语言中的整形提升与算术转换
qinyuz的博客
05-01 668
在C语言中，整形提升（Integer Promotion）是一个重要但容易被忽视的概念。它指的是在表达式中，任何小于。是无符号char，赋值-1会转换为255（0xFF）长度为标准操作数长度，使用标准长度可以提高运算效率。都是有符号char，值为-1（内存表示：0xFF）：符号扩展为0xFFFFFFFF（即-1）：零扩展为0x000000FF（即255））操作数在使用前都会被自动转换为普通整型。：确保不同大小的整型在运算时有一致的表现。有符号类型：高位补符号位。无符号类型：高位补0。是2字节，所以前两个。
VUE3报错 ReferenceError: Cannot access ‘openInit‘ before initialization
最新发布
dy1717的博客
05-05 267
vue3 ReferenceError: Cannot access xxx before initialization问题报错处理
STL之list容器
Laydya的博客
05-03 1402
1.list的底层是双向链表结构，双向链表中的每个元素在互不相关的独立结点中，在结点中通过指针指向前一个元素和后一个元素2.list是可以在常数范围内在任意位置的插入和删除的序列式容器，并且该容器可以前后双向迭代3.vector的优点：支持随机访问，间接的就支持了排序/二分查找/堆排序等缺点：1.头部和中部的插入效率低，需要挪动数据2.插入数据空间不够需要增容，增容需要开新空间，拷贝数据，释放旧空间list的优点：没有增容代价，且任意位置插入都是0（1）；
C++日志系统实现(二)
A_zyx0808的博客
05-04 790
使用缓冲池、线程池实现异步日志系统。
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

打赏作者

兔子的倔强

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

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

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

打赏作者

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

抵扣说明：

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

余额充值