GPU 编程入门到精通（五）之 GPU 程序优化进阶

最新推荐文章于 2024-07-19 18:54:02 发布
最新推荐文章于 2024-07-19 18:54:02 发布
阅读量106 收藏
点赞数
文章标签： python 人工智能


博主因为工作其中的须要，開始学习 GPU 上面的编程，主要涉及到的是基于 GPU 的深度学习方面的知识。鉴于之前没有接触过 GPU 编程。因此在这里特地学习一下 GPU 上面的编程。

有志同道合的小伙伴，欢迎一起交流和学习。我的邮箱: caijinping220@gmail.com 。使用的是自己的老古董笔记本上面的 Geforce 103m 显卡，尽管显卡相对于如今主流的系列已经很的弱，可是对于学习来说。还是能够用的。本系列博文也遵从由简单到复杂，记录自己学习的过程。

0. 文件夹

1. 数组平方和并行化进阶

GPU 编程入门到精通（四）之 GPU 程序优化 这篇博文中提到了 grid、block、thread 三者之间的关系。知道了他们之间是逐渐包括的关系。我们在上面的程序中通过使用 512 个线程达到了 493 倍左右的性能提升，那么是不是能够继续得到提升呢？？？

答案是肯定的，这就要进一步考虑 GPU 的并行化处理了。前面的程序仅仅是使用了单个 block 下的 512 个线程，那么。我们可不能够使用多个 block 来实现？？？

对。就是利用这个思想。达到进一步的并行化。

这里使用 8 个 block * 64 threads = 512 threads 实现。

  • 首先，改动主函数宏定义。定义块数量：

      // ======== define area ========
  #define DATA_SIZE 1048576    // 1M
  #define BLOCK_NUM 8        // block num
  #define THREAD_NUM 64        // thread num

      通过在程序中加入 block 和 threads 的宏定义，这两个定义是我们在后面会用到的。

    他们决定了计算平方和使用的 CUDA 核心数。

  • 接下来，改动内核函数：

      _global__ static void squaresSum(int *data, int *sum, clock_t *time)
  {
      const int tid = threadIdx.x;
      const int bid = blockIdx.x;

      for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
          tmp_sum += data[i] * data[i];
      }

      sum[bid * THREAD_NUM + tid] = tmp_sum;
  }

      注意：这里的内存遍历方式和前面讲的是一致的。理解一下。

    同一时候记录的时间是一个块的開始和结束时间。由于这里我们最后须要计算的是最早開始和最晚结束的两个时间差，即求出最糟糕的时间。

  • 然后。就是主函数里面的详细实现了：

      // malloc space for datas in GPU
  cudaMalloc((void**) &sum, sizeof(int) * THREAD_NUM * BLOCK_NUM);

  // calculate the squares's sum
  squaresSum<<<BLOCK_NUM, THREAD_NUM, 0>>>(gpuData, sum, time);

      这里边。sum 数组的长度计算方式变化了，可是大小没有变化。另在在调用 GPU 内核函数的时候，參数发生了变化。须要告诉 GPU block 数 和 thread 数。只是这边共享内存没有使用。

  • 最后，在 CPU 中计算部分和

      // print result
  int tmp_result = 0;
  for (int i = 0; i < THREAD_NUM * BLOCK_NUM; ++i) {
      tmp_result += result[i];
  }

编译执行以后。得到例如以下结果：

advance

性能与直接使用 512 个线程基本一致。由于受到 GPU 内存带宽的限制，GPU 编程入门到精通（四）之 GPU 程序优化 中的优化。已经接近极限，所以通过 block 方式，效果不明显。

2. 线程同步和共享内存

前面的程序。计算求和的工作在 CPU 中完毕。总共须要在 CPU 中做 512 次加法运算。那么有没有办法降低 CPU 中运行加法的次数呢？？？

能够通过同步和共享内存技术，实如今 GPU 上的 block 块内求取部分和。这样最后仅仅须要在 CPU 计算 16 个和就能够了。

详细实现方法例如以下：

  • 首先，在改动内核函数，定义一块共享内存，用 __shared__ 指示：

      __global__ static void squaresSum(int *data, int *sum, clock_t *time)
  {
      // define of shared memory
      __shared__ int shared[BLOCK_NUM];

      const int tid = threadIdx.x;
      const int bid = blockIdx.x;

      if (tid == 0) time[bid] = clock();

      shared[tid] = 0;
      // 把部分和结果放入共享内存中
      for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
          shared[tid] += data[i] * data[i];
      }
      // 同步操作。必须等之前的线程都执行结束，才干继续后面的程序
      __syncthreads();
      // 同步完毕之后。将部分和加到 shared[0] 上面。这里全都在一个线程内完毕
      if (tid == 0) {
          for (int i = 1; i < THREAD_NUM; i++) {
              shared[0] += shared[i];
          }
          sum[bid] = shared[0];
      }

      if (tid == 0) time[bid + BLOCK_NUM] = clock();
  }

      利用 __shared__ 声明的变量是 shared memory。每一个 block 中。各个 thread 之间对于共享内存是共享的。利用的是 GPU 上的内存，所以速度非常快。不必操心 latency 的问题。

  __syncthreads() 函数是 CUDA 的内部函数，表示全部 threads 都必须同步到这个点。才会运行接下来的代码。我们要做的就是等待每一个 thread 计算结束以后。再来计算部分和，所以同步是不可缺少的环节。把每一个 block 的部分和计算到 shared[0] 里面。

  • 接下来，改动 main 函数:

      // calculate the squares's sum
  squaresSum<<<BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int)>>>(gpuData, sum, time);

    编译执行后结果例如以下：

    shared_memory

    事实上和前一版程序相比，时间上没有什么优势，原因在于，我们须要在 GPU 中额外执行求和的这部分代码。导致了执行周期的变长，只是对应的，在 CPU 中的执行时间会降低。

3. 加法树

我们在这个程序中，仅仅当每一个 block 的 thread0 的时候，计算求和的工作，这样做影响了运行的效率，事实上求和能够并行化处理的，也就是通过加法树来实现并行化。举个样例，要计算 8 个数的和。我们不是必需用一个 for 循环。逐个相加。而是能够通过第一级流水线实现两两相加。变成 4 个数，第二级流水实现两两相加，变成 2 个数。第三级流水实现两两相加，求得最后的和。

以下通过加法树的方法，实现最后的求和，改动内核函数例如以下：

__global__ static void squaresSum(int *data, int *sum, clock_t *time)
{
    __shared__ int shared[BLOCK_NUM];
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;

    int offset = THREAD_NUM / 2;

    if (tid == 0) time[bid] = clock();

    shared[tid] = 0;

    for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
        shared[tid] += data[i] * data[i];
    }

    __syncthreads();
    while (offset > 0) {
        if (tid < offset) {
            shared[tid] += shared[tid + offset];
        }
        offset >>= 1;
        __syncthreads();
    }

    if (tid == 0) {
        sum[bid] = shared[0];
        time[bid + BLOCK_NUM] = clock();
    }
}

此程序实现的就是上诉描写叙述的加法树的结构。注意这里第二个 __syncthreads() 的使用，也就是说，要进行下一级流水线的计算。必须建立在前一级必须已经计算完成的情况下。

主函数部分不许要改动，最后编译执行结果例如以下:

add_tree

性能有一部分的改善。

通过使用 GPU 的并行化编程。确实对性能会有非常大程度上的提升。

因为受限于 Geforce 103m 的内存带宽，程序仅仅能优化到这一步，关于是否还有其它的方式优化，有待进一步学习。

4. 总结

通过这几篇博文的讨论，数组平方和的代码优化到这一阶段。

从但线程到多线程，再到共享内存，通过使用这几种 GPU 上面的结构，做到了程序的优化。例如以下给出数组平方和的完整代码：

/* *******************************************************************
##### File Name: squareSum.cu
##### File Func: calculate the sum of inputs's square
##### Author: Caijinping
##### E-mail: caijinping220@gmail.com
##### Create Time: 2014-5-7
* ********************************************************************/

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>

// ======== define area ========
#define DATA_SIZE 1048576    // 1M
#define BLOCK_NUM 8            // block num
#define THREAD_NUM 64        // thread num

// ======== global area ========
int data[DATA_SIZE];

void printDeviceProp(const cudaDeviceProp &prop);
bool InitCUDA();
void generateData(int *data, int size);
__global__ static void squaresSum(int *data, int *sum, clock_t *time);

int main(int argc, char const *argv[])
{
    // init CUDA device
    if (!InitCUDA()) {
        return 0;
    }
    printf("CUDA initialized.\n");

    // generate rand datas
    generateData(data, DATA_SIZE);

    // malloc space for datas in GPU
    int *gpuData, *sum;
    clock_t *time;
    cudaMalloc((void**) &gpuData, sizeof(int) * DATA_SIZE);
    cudaMalloc((void**) &sum, sizeof(int) * BLOCK_NUM);
    cudaMalloc((void**) &time, sizeof(clock_t) * BLOCK_NUM * 2);
    cudaMemcpy(gpuData, data, sizeof(int) * DATA_SIZE, cudaMemcpyHostToDevice);

    // calculate the squares's sum
    squaresSum<<<BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int)>>>(gpuData, sum, time);

    // copy the result from GPU to HOST
    int result[BLOCK_NUM];
    clock_t time_used[BLOCK_NUM * 2];
    cudaMemcpy(&result, sum, sizeof(int) * BLOCK_NUM, cudaMemcpyDeviceToHost);
    cudaMemcpy(&time_used, time, sizeof(clock_t) * BLOCK_NUM * 2, cudaMemcpyDeviceToHost);

    // free GPU spaces
    cudaFree(gpuData);
    cudaFree(sum);
    cudaFree(time);

    // print result
    int tmp_result = 0;
    for (int i = 0; i < BLOCK_NUM; ++i) {
        tmp_result += result[i];
    }

    clock_t min_start, max_end;
    min_start = time_used[0];
    max_end = time_used[BLOCK_NUM];
    for (int i = 1; i < BLOCK_NUM; ++i)    {
        if (min_start > time_used[i]) min_start = time_used[i];
        if (max_end < time_used[i + BLOCK_NUM]) max_end = time_used[i + BLOCK_NUM];
    }
    printf("(GPU) sum:%d time:%ld\n", tmp_result, max_end - min_start);

    // CPU calculate
    tmp_result = 0;

    for (int i = 0; i < DATA_SIZE; ++i)    {
        tmp_result += data[i] * data[i];
    }

    printf("(CPU) sum:%d\n", tmp_result);

    return 0;
}

__global__ static void squaresSum(int *data, int *sum, clock_t *time)
{
    __shared__ int shared[BLOCK_NUM];
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;

    int offset = THREAD_NUM / 2;

    if (tid == 0) time[bid] = clock();

    shared[tid] = 0;

    for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
        shared[tid] += data[i] * data[i];
    }

    __syncthreads();
    while (offset > 0) {
        if (tid < offset) {
            shared[tid] += shared[tid + offset];
        }
        offset >>= 1;
        __syncthreads();
    }

    if (tid == 0) {
        sum[bid] = shared[0];
        time[bid + BLOCK_NUM] = clock();
    }
}

// ======== used to generate rand datas ========
void generateData(int *data, int size)
{
    for (int i = 0; i < size; ++i) {
        data[i] = rand() % 10;
    }
}


void printDeviceProp(const cudaDeviceProp &prop)
{
    printf("Device Name : %s.\n", prop.name);
    printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
    printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
    printf("regsPerBlock : %d.\n", prop.regsPerBlock);
    printf("warpSize : %d.\n", prop.warpSize);
    printf("memPitch : %d.\n", prop.memPitch);
    printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
    printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
    printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
    printf("totalConstMem : %d.\n", prop.totalConstMem);
    printf("major.minor : %d.%d.\n", prop.major, prop.minor);
    printf("clockRate : %d.\n", prop.clockRate);
    printf("textureAlignment : %d.\n", prop.textureAlignment);
    printf("deviceOverlap : %d.\n", prop.deviceOverlap);
    printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
}

bool InitCUDA()
{
    //used to count the device numbers
    int count;    

    // get the cuda device count
    cudaGetDeviceCount(&count);
    if (count == 0) {
        fprintf(stderr, "There is no device.\n");
        return false;
    }

    // find the device >= 1.X
    int i;
    for (i = 0; i < count; ++i) {
        cudaDeviceProp prop;
        if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
            if (prop.major >= 1) {
                //printDeviceProp(prop);
                break;
            }
        }
    }

    // if can't find the device
    if (i == count) {
        fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
        return false;
    }

    // set cuda device 
    cudaSetDevice(i);

    return true;
}


欢迎大家和我一起讨论和学习 GPU 编程。
caijinping220@gmail.com
http://blog.csdn.net/xsc_c
weixin_34007886
关注
​使用PotPlayer播放器查看软解和硬解4K高清视频时的CPU及GPU占用情况​
dvlinker的技术专栏
03-17 2万+
本文主要演示了如何使用PotPlayer视频播放器去观测视频软解和硬解时的CPU与GPU的占用情况，以供参考。
GPU编程与优化
02-25
GPU编程与优化--大众高性能计算。扫描版，完整，433页。
gpu编程优化
SophieZhou的专栏
03-22 2217
近几年比较火的一个东西就是大规模、大数据等概念了。用到的工具除了hadoop等分布式平台外，还有一种是gpu，前者更多是用来做真正大规模的数据吧，比如大规模的网络数据的，但是如果数据是图像数据的话，可以用gpu来做。 gpu卡是由很多个处理单位组成，每个处理单元就像cpu cores一样，这个单元叫SM，不同的gpu卡有不同的SM版本。每个SM均像一个多核的cpu一样，有存储和寄存器以及
GPU编程之进击的优化-李修宇.docx
12-01
第一章 设备微架构 1.0 CUDA设备 1.0.0 核心微架构 1.0.1 指令编码格式 1.0.2 寄存器文件结构 1.0.3 指令流水线 1.0.4 Pascal架构（待续） 1.1 GCN设备 1.1.0 核心微架构 1.1.1 指令编码格式 1.1.2 寄存器文件结构 1.1.3 指令流水线 1.2 GPU设备上的条件分支 第二章 GPU矩阵乘法的高效实现 2.0 前言 2.1 指令级并行和数据预取 2.2 双缓冲区 2.3 宽数据内存事务 2.4 二级数据预取 2.5 细节调优 第三章 基于GPU的大规模稀疏矩阵直接求解器 3.0 简介 3.1 基于quotient graph的符号分析 3.1.1 顶点重排序 3.1.2 构建消去树 3.1.3 寻找超结点 3.1.4 符号分解 3.2 多波前法 3.3 超节点方法 3.4 多波前+超节点方法的并行分解算法 小结 参考资料 第四章 CNN中的卷积计算 前言 1 基于矩阵乘法的卷积算法 2 类矩阵乘法的卷积算法 小结 第五章 基于GPU的LU分解（待续） 5.0 一般实现 5.1 分块实现 5.2 使用动态并行 5.3 多GPU版本 第六章 GPU上的光线追踪（待续） 6.1 kd-tree算法介绍及内核实现 6.2 less-分支版本的kd-tree算法及内核实现 6.3 ropes-kdtree算法及内核实现 第七章 LBM流体计算（待续） 尾章 GPU编程优化技术总结 4.1.0 CUDA设备上的优化技术 4.1.1 访存优化 4.1.2 指令优化 4.1.3 内核调用优化 4.2.0 GCN设备上的优化技术 4.2.1 访存优化 4.2.2 指令优化 4.2.3 内核调用优化 4.3 构建性能可移植的程序
GPU高级调试与优化
格友
11-02 3150
GPU的历史很短，只有十几年。但它发展迅猛，凭借强大的并行计算能力和高效率的固定硬件单元，在人工智能、虚拟和增强现实（VR/AR）、3D游戏、视频编解码等领域大显身手。而且这种趋势还在延续，基于GPU的应用和创新势头正猛。但是从系统架构来看，针对GPU的架构转型还在进行过程中，目前GPU依然还处于外设的地位，还没有摆脱从属身份。因为这个根本特征，对GPU编程并不像对CPU编程那样直接，而调试和优化
GPU编程与CG语言：入门指南
《GPU编程与CG语言之阳春白雪下里巴人》是...《GPU编程与CG语言之阳春白雪下里巴人》是一本结合了学术深度和实践指导的实用教材，为读者提供了丰富的学习资源和人生启示，无论是对GPU编程入门者还是进阶者都极具价值。
从入门到进阶，史上最全Python精华
m0_72444380的博客
09-23 631
去年我在使用Medium平台时，只用它来读别人写的文章，也就是在消费内容，从而读了一大堆关于Python的文章。随后，文章深入分析了6种数据可视化类型，并给出了使用PythonMatplotlib的例子，实现了以下内容两种散点图，线图，柱状图，条形图和箱形图。“迭代器指的是可在其上迭代（循环）的对象，其作用是为数据抽象出一个容器，使其表现出与迭代对象类似的功能。”该教程使用Python，带领你体验完全交互的例子，即Bokeh交互可视化库，并附有公共数据集。......
OpenGL从入门到精通课程讲义.pdf
07-18
### OpenGL从入门到精通课程讲义知识点概览 #### 一、OpenGL简介及核心概念 - **OpenGL**：一种跨语言、跨平台的应用程序接口（API），用于渲染2D、3D矢量图形。 - **模型与网格(Mesh)**：模型通常由多个网格组成...
matlab 入门到精通.zip
10-29
《MATLAB入门到精通》教程全面解析 MATLAB（矩阵实验室）是一款强大的数学计算软件，广泛应用于科研、工程和教育领域。本教程“matlab 入门到精通.zip”是一套完整的MATLAB学习资源，包括视频教学和配套的学习资料...
GPU编程之进击的优化系列四--GPU编程优化技术总结
weixin_34185560的博客
03-06 399
第四章 GPU编程优化技术总结 4.1.0 CUDA设备上的优化技术 4.1.1 访存优化 4.1.2 指令优化 4.1.3 内核调用优化 4.2.0 GCN设备上的优化技术 4.2.1 访存优化 4.2.2 指令优化 4.2.3 内核调用优化 4.3 构建性能可移植的程序 第四章 GPU编程优化技术总结 我们在两个章节分别讲述针对CUDA和GCN这两大...
GPU编程优化笔记
已迁移至 https://seekstar.github.io/
12-31 268
GPU结构概览 参考：https://blog.csdn.net/fengtian12345/article/details/80546799 硬件上，GPU由多个SM(Steaming Multiprocessor)构成，SM有多个warp，warp有多个SP(streaming processor)，一个SP对应一个线程。 一个warp中的SP执行相同的指令。 软件上，一个kernel(就是核函数的kernel)对应一个grid，一个grid有多个block，一个block有多个thread。 bloc
GPU 编程入门到精通（四）之 GPU 程序优化
weixin_34234829的博客
08-05 173
博主因为工作其中的须要，開始学习 GPU 上面的编程，主要涉及到的是基于 GPU 的深度学习方面的知识，鉴于之前没有接触过 GPU 编程。因此在这里特地学习一下 GPU 上面的编程。有志同道合的小伙伴。欢迎一起交流和学习，我的邮箱: caijinping220@gmail.com 。使用的是自己的老古董笔记本上面的 Geforce 103m 显卡，尽管显卡相对于如今主流的系列...
(10)GPU编程与优化——NVIDIA Turing计算平台
AI天才研究院
09-03 1349
作者：禅与计算机程序设计艺术 1.简介 近年来，随着研究者们对计算机硬件的要求越来越高、数据规模越来越大、AI任务越来越复杂，图形处理单元（Graphics Processing Unit，GPU）已成为深度学习等高性能计算领域的一个重要组件。NVIDIA公司推出了基于CUDA的通用计算平台CUDA-X
GPU编程和优化-最佳实践分享
最新发布
SunGlassGreat
07-19 550
为了加速向量数据库的查询性能，可以采取多种策略。此外，NVIDIA的BMF技术还可以加速GPU上的视频处理流水线，提高视频处理的实时性和质量。用户体验优化则包括推荐结果的个性化和可读性的提升，通过优化推荐结果的展示方式和提高推荐结果的相关性来实现。最新的实践还包括利用深度学习、机器学习等高级技术，通过自动学习用户的喜好和动态变化的兴趣，实现更加智能化的推荐。NVIDIA NeMo是一个开源的语音和语言模型开发套件，提供了丰富的可重用组件和预训练模型，支持在NVIDIA GPU上进行高效的全周期开发。
GPU并行运算与CUDA编程 – 优化篇
qihongminghik的博客
02-23 407
CPU传输到GPU中的数据受限会被存储到DRAM，GPU在访问数据时会从上往下访问，如果遇到cache miss，就会向下一级访问。对于立即数“1.0”在默认情况下是双精度的浮点数，加上后缀“f”后，“1.0f”就会被编译为单精度浮点数，计算时更加快速。一些固定大小的SM组成一个bank，当不同的线程访问同一个bank的不同word时，就会出现bank conflict（如上图）。SM可以让同一个block中的所有线程自由访问，如果多个线程需要访问同一组数据，可以考虑把数据存放在SM中再进行计算。
GPU编程 CUDA C++ 使用统一内存编程之【优化使用统一内存】
行子的博客
04-06 446
为了在使用统一内存时获得较高性能，需要避免缺页异常、保持数据的布局性（让相关的数据尽量靠近对应的处理器），但避免内存抖动（即频繁地在不同的处理器之间传输数据）。可以使用CUDA运行时API函数cudaMemPrefetchAsync()手动的给编译器一些提示。cudaMemPrefetchAsync()函数的作用是在CUDA流stream中将统一内存。设备dstDevice（主机的设备GPU号用cudaCpuDeviceId表示）中的内存区域，从而。防止（或减少）缺页异常，提高数据的局部性。
GPU编程优化篇-姜雪伟-专题视频课程
海洋个人博客
08-01 521
本课程从可编程流水线到GPU内部运行原理一一给读者介绍，本课程涉及到的Shader编程技术都是实战中经常使用的，结合着案例讲解让读者更容易接受，另外对GPU编程的优化也做了进一步的讲解。...
GPU加速02:超详细Python Cuda零基础入门教程，没有显卡也能学！
WANGJUNAIJIAO的博客
12-16 7878
Python Numba库可以调用CUDA进行GPU编程，CPU端被称为主机，GPU端被称为设备，运行在GPU上的函数被称为核函数，调用核函数时需要有执行配置，以告知CUDA以多大的并行粒度来计算。使用GPU编程时要合理地将数据在主机和设备间互相拷贝。GPU程序执行流程初始化，并将必要的数据拷贝到GPU设备的显存上。使用某个执行配置，以一定的并行粒度调用CUDA核函数。CPU和GPU异步计算。将GPU计算结果拷贝回主机。
NVIDIA GPU编程高级指南：从入门到进阶
"GPU编程的高级指南，适用于初学者和进阶学习者，主要涵盖NVIDIA GPU的编程技术。" 本文档，即《GPU编程的高级指南》，是NVIDIA公司发布的一份详细的技术参考资料，主要面向那些希望深入理解和利用GPU进行图形处理...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值