CUDA编程：从hello world开始

已于 2024-01-09 18:00:50 修改
阅读量782 收藏 8
点赞数 9
分类专栏： cuda编程 文章标签： 深度学习 人工智能 c语言
于 2024-01-03 22:00:06 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/qq_43570025/article/details/135374018
版权

目录

第一个程序：用GPU输出Hello World

写一个CUDA C程序步骤

hello.cu

编译

运行

common.h

cuda程序执行步骤

两个数组相加

CPU上的两个数据相加

GPU上的两个数据相加

计时

nvprof

第一个程序：用GPU输出Hello World

写一个CUDA C程序步骤

  1. 用专用扩展名.cu来创建一个源文件

  2. 使用CUDA nvcc编译器来编译程序

  3. 从命令行运行可执行文件，这个文件有可在GPU上运行的内核代码

hello.cu

#include "../common/common.h"
#include <stdio.h>

/*
 * A simple introduction to programming in CUDA. This program prints "Hello
 * World from GPU! from 10 CUDA threads running on the GPU.
 */

__global__ void helloFromGPU()
{
    printf("Hello World from GPU!\n");
}

int main(int argc, char **argv)
{
    printf("Hello World from CPU!\n");

    helloFromGPU<<<1, 10>>>();
    CHECK(cudaDeviceReset());
    return 0;
}

编译

nvcc hello.cu -o hello

运行

./hello

cudaDeviceReset()函数是用来显示的释放和清空当前进程中于当前设备有关的所有资源当重置函数移除后 编译运行则只输出hello world from cpu,printf在gpu上被调用，cudaDeviceReset()函数使这些来自gpu的输出发送到主机，然后在控制台输出。没有调用cudaDeviceReset()函数就不能保证这些可以被显示

关于cudaDeviceReset(); 这句话如果没有，则不能正常的运行，因为这句话包含了隐式同步，GPU和CPU执行程序是异步的，核函数调用后成立刻会到主机线程继续，而不管GPU端核函数是否执行完毕，所以上面的程序就是GPU刚开始执行，CPU已经退出程序了，所以我们要等GPU执行完了，再退出主机线程。

common.h

#include <sys/time.h>
#include <cuda_runtime.h>
#include <stdio.h>

#ifndef _COMMON_H
#define _COMMON_H

cudaError_t ErrorCheck(cudaError_t status,const char* filename,int lineNumber)
{
    if(status != cudaSuccess)
    {
        printf("CUDA API error: \r\ncode=%d, name=%s, description=%s\r\nfile=%s, line=%d\r\n",
                status,cudaGetErrorName(status),cudaGetErrorString(status),filename,lineNumber
        );
    }
    return status;
}

#define CHECK(call)                                                            \
{                                                                              \
    const cudaError_t error = call;                                            \
    if (error != cudaSuccess)                                                  \
    {                                                                          \
        fprintf(stderr, "Error: %s:%d, ", __FILE__, __LINE__);                 \
        fprintf(stderr, "code: %d, reason: %s\n", error,                       \
                cudaGetErrorString(error));                                    \
        exit(1);                                                               \
    }                                                                          \
}


#define CHECK_CUBLAS(call)                                                     \
{                                                                              \
    cublasStatus_t err;                                                        \
    if ((err = (call)) != CUBLAS_STATUS_SUCCESS)                               \
    {                                                                          \
        fprintf(stderr, "Got CUBLAS error %d at %s:%d\n", err, __FILE__,       \
                __LINE__);                                                     \
        exit(1);                                                               \
    }                                                                          \
}

#define CHECK_CURAND(call)                                                     \
{                                                                              \
    curandStatus_t err;                                                        \
    if ((err = (call)) != CURAND_STATUS_SUCCESS)                               \
    {                                                                          \
        fprintf(stderr, "Got CURAND error %d at %s:%d\n", err, __FILE__,       \
                __LINE__);                                                     \
        exit(1);                                                               \
    }                                                                          \
}

#define CHECK_CUFFT(call)                                                      \
{                                                                              \
    cufftResult err;                                                           \
    if ( (err = (call)) != CUFFT_SUCCESS)                                      \
    {                                                                          \
        fprintf(stderr, "Got CUFFT error %d at %s:%d\n", err, __FILE__,        \
                __LINE__);                                                     \
        exit(1);                                                               \
    }                                                                          \
}

#define CHECK_CUSPARSE(call)                                                   \
{                                                                              \
    cusparseStatus_t err;                                                      \
    if ((err = (call)) != CUSPARSE_STATUS_SUCCESS)                             \
    {                                                                          \
        fprintf(stderr, "Got error %d at %s:%d\n", err, __FILE__, __LINE__);   \
        cudaError_t cuda_err = cudaGetLastError();                             \
        if (cuda_err != cudaSuccess)                                           \
        {                                                                      \
            fprintf(stderr, "  CUDA error \"%s\" also detected\n",             \
                    cudaGetErrorString(cuda_err));                             \
        }                                                                      \
        exit(1);                                                               \
    }                                                                          \
}

inline double seconds()
{
    struct timeval tp;
    struct timezone tzp;
    int i = gettimeofday(&tp, &tzp);
    return ((double)tp.tv_sec + (double)tp.tv_usec * 1.e-6);
}

#define CHECK_EXIT(sts, str) \
    if (sts) \
    {       \
        std::cerr << str << " " << __FILE__ << " "  << __LINE__ << std::endl; \
        exit(0);\
    }
#endif // _COMMON_H

cuda程序执行步骤

  1. 设置参与运算的GPU cudaSetDevice

  2. 申请显存 cudaMalloc

  3. 复制数据 cudaMemcpy（参数cudaMemcpyHostToDevice）

  4. 调用kernel函数做运算 xxxKernel

  5. 等待运算结束后同步 cudaDeviceSynchronize（kernel函数的调用是异步的）

  6. 复制运算结果到到内存 cudaMemcpy （参数cudaMemcpyDeviceToHost）

  7. 释放显存 cudaFree

两个数组相加

CPU上的两个数据相加

sumArraysOnHost.c

#include <stdlib.h>
#include <time.h>

/*
 * This example demonstrates a simple vector sum on the host. sumArraysOnHost
 * sequentially iterates through vector elements on the host.
 */

void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    for (int idx = 0; idx < N; idx++)
    {
        C[idx] = A[idx] + B[idx];
    }

}

void initialData(float *ip, int size)
{
    // generate different seed for random number
    time_t t;
    srand((unsigned) time(&t));

    for (int i = 0; i < size; i++)
    {
        ip[i] = (float)(rand() & 0xFF) / 10.0f;
    }

    return;
}

int main(int argc, char **argv)
{
    int nElem = 1024;
    size_t nBytes = nElem * sizeof(float);

    float *h_A, *h_B, *h_C;
    h_A = (float *)malloc(nBytes);
    h_B = (float *)malloc(nBytes);
    h_C = (float *)malloc(nBytes);

    initialData(h_A, nElem);
    initialData(h_B, nElem);

    sumArraysOnHost(h_A, h_B, h_C, nElem);

    free(h_A);
    free(h_B);
    free(h_C);

    return(0);
}

gcc -O2 -std=c99 -o sumArraysOnHost sumArraysOnHost.c

GPU上的两个数据相加

sumArraysOnGPU-small-case.cu

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>

/*
 * This example demonstrates a simple vector sum on the GPU and on the host.
 * sumArraysOnGPU splits the work of the vector sum across CUDA threads on the
 * GPU. Only a single thread block is used in this small case, for simplicity.
 * sumArraysOnHost sequentially iterates through vector elements on the host.
 */

void checkResult(float *hostRef, float *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;

    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef[i] - gpuRef[i]) > epsilon)
        {
            match = 0;
            printf("Arrays do not match!\n");
            printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i],
                   gpuRef[i], i);
            break;
        }
    }

    if (match) printf("Arrays match.\n\n");

    return;
}


void initialData(float *ip, int size)
{
    // generate different seed for random number
    time_t t;
    srand((unsigned) time(&t));

    for (int i = 0; i < size; i++)
    {
        ip[i] = (float)(rand() & 0xFF) / 10.0f;
    }

    return;
}


void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    for (int idx = 0; idx < N; idx++)
        C[idx] = A[idx] + B[idx];
}

__global__ void sumArraysOnGPU(float *A, float *B, float *C, const int N)
{
    int i = threadIdx.x;

    if (i < N) C[i] = A[i] + B[i];
}


int main(int argc, char **argv)
{
    printf("%s Starting...\n", argv[0]);

    // set up device
    int dev = 0;
    CHECK(cudaSetDevice(dev));

    // set up data size of vectors
    int nElem = 1 << 5;
    printf("Vector size %d\n", nElem);

    // malloc host memory
    size_t nBytes = nElem * sizeof(float);

    float *h_A, *h_B, *hostRef, *gpuRef;
    h_A     = (float *)malloc(nBytes);
    h_B     = (float *)malloc(nBytes);
    hostRef = (float *)malloc(nBytes);
    gpuRef  = (float *)malloc(nBytes);

    // initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);

    memset(hostRef, 0, nBytes);
    memset(gpuRef,  0, nBytes);

    // malloc device global memory
    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((float**)&d_A, nBytes));
    CHECK(cudaMalloc((float**)&d_B, nBytes));
    CHECK(cudaMalloc((float**)&d_C, nBytes));

    // transfer data from host to device
    CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_C, gpuRef, nBytes, cudaMemcpyHostToDevice));

    // invoke kernel at host side
    dim3 block (nElem);
    dim3 grid  (1);

    sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C, nElem);
    printf("Execution configure <<<%d, %d>>>\n", grid.x, block.x);

    // copy kernel result back to host side
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));

    // add vector at host side for result checks
    sumArraysOnHost(h_A, h_B, hostRef, nElem);

    // check device results
    checkResult(hostRef, gpuRef, nElem);

    // free device global memory
    CHECK(cudaFree(d_A));
    CHECK(cudaFree(d_B));
    CHECK(cudaFree(d_C));

    // free host memory
    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);

    CHECK(cudaDeviceReset());
    return(0);
}

计时

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>

/*
 * This example demonstrates a simple vector sum on the GPU and on the host.
 * sumArraysOnGPU splits the work of the vector sum across CUDA threads on the
 * GPU. Only a single thread block is used in this small case, for simplicity.
 * sumArraysOnHost sequentially iterates through vector elements on the host.
 * This version of sumArrays adds host timers to measure GPU and CPU
 * performance.
 */

void checkResult(float *hostRef, float *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;

    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef[i] - gpuRef[i]) > epsilon)
        {
            match = 0;
            printf("Arrays do not match!\n");
            printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i],
                   gpuRef[i], i);
            break;
        }
    }

    if (match) printf("Arrays match.\n\n");

    return;
}

void initialData(float *ip, int size)
{
    // generate different seed for random number
    time_t t;
    srand((unsigned) time(&t));

    for (int i = 0; i < size; i++)
    {
        ip[i] = (float)( rand() & 0xFF ) / 10.0f;
    }

    return;
}

void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    for (int idx = 0; idx < N; idx++)
    {
        C[idx] = A[idx] + B[idx];
    }
}
__global__ void sumArraysOnGPU(float *A, float *B, float *C, const int N)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < N) C[i] = A[i] + B[i];
}

int main(int argc, char **argv)
{
    printf("%s Starting...\n", argv[0]);

    // set up device
    int dev = 0;
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("Using Device %d: %s\n", dev, deviceProp.name);
    CHECK(cudaSetDevice(dev));

    // set up data size of vectors
    int nElem = 1 << 24;
    printf("Vector size %d\n", nElem);

    // malloc host memory
    size_t nBytes = nElem * sizeof(float);

    float *h_A, *h_B, *hostRef, *gpuRef;
    h_A     = (float *)malloc(nBytes);
    h_B     = (float *)malloc(nBytes);
    hostRef = (float *)malloc(nBytes);
    gpuRef  = (float *)malloc(nBytes);

    double iStart, iElaps;

    // initialize data at host side
    iStart = seconds();
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    iElaps = seconds() - iStart;
    printf("initialData Time elapsed %f sec\n", iElaps);
    memset(hostRef, 0, nBytes);
    memset(gpuRef,  0, nBytes);

    // add vector at host side for result checks
    iStart = seconds();
    sumArraysOnHost(h_A, h_B, hostRef, nElem);
    iElaps = seconds() - iStart;
    printf("sumArraysOnHost Time elapsed %f sec\n", iElaps);

    // malloc device global memory
    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((float**)&d_A, nBytes));
    CHECK(cudaMalloc((float**)&d_B, nBytes));
    CHECK(cudaMalloc((float**)&d_C, nBytes));

    // transfer data from host to device
    CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_C, gpuRef, nBytes, cudaMemcpyHostToDevice));

    // invoke kernel at host side
    int iLen = 512;
    dim3 block (iLen);
    dim3 grid  ((nElem + block.x - 1) / block.x);

    iStart = seconds();
    sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C, nElem);
    CHECK(cudaDeviceSynchronize());
    iElaps = seconds() - iStart;
    printf("sumArraysOnGPU <<<  %d, %d  >>>  Time elapsed %f sec\n", grid.x,
           block.x, iElaps);

    // check kernel error
    CHECK(cudaGetLastError()) ;

    // copy kernel result back to host side
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));

    // check device results
    checkResult(hostRef, gpuRef, nElem);

    // free device global memory
    CHECK(cudaFree(d_A));
    CHECK(cudaFree(d_B));
    CHECK(cudaFree(d_C));

    // free host memory
    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);

    return(0);
}

nvprof

nvprof ./sumArraysOnGPU-timer

兔子牙丫丫
关注
  • 9
    点赞
  • 踩
  • 8
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
专栏目录
CUDA 编程：基础与实践_樊哲勇1
08-03
CUDA 编程：基础与实践 CUDA（Compute Unified Device Architecture）是一种由NVIDIA公司推出的编程模型，用于利用GPU（Graphics Processing Unit）的强大计算能力来执行并行计算任务。CUDA编程是实现高性能计算的...
code : c++高性能多进程 cuda编程:GPU结构和通信速度+tiling
08-04
code for c++高性能多进程 cuda编程:GPU结构和通信速度+tiling
CUDA从入门到放弃（三）：Hello World 之第一个CUDA程序
mingo_敏
05-12 2649
vs2015中打开新建一个CUDA程序 拿个CUDA中的 hello world 示例程序过来跑一下吧！ 示例来源于：CUDA Samples 计算两个数组间的加法 #include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> cudaError_t addWithCuda(i...
[CUDA] 快速入门CUDA(1)-基本了解和HelloWorld
weixin_52157994的博客
04-10 8568
CUDA基础 文章目录CUDA基础1 CUDA简介2 GPU和CPU架构的不同之处3 查看GPU硬件信息4 需要建立的基本概念5 总结 1 CUDA简介 CUDA的全程是Computer Unified Device Architecture，是由显卡头子NVIDIA发明的。有的人对于显卡的印象在于它可以玩游戏，效果十分逼真，但从背后而言，正是因为显卡强大的图形计算能力，才使得计算机可以运行这些大型的3D游戏，并且拥有较高的画质和帧数。 2 GPU和CPU架构的不同之处 CPU具有以下特点： 对单线程有优
Cuda实战-01 入门程序HelloWorld
星辰火之梦
03-01 310
不管是学习什么编程语言，一般入门的实战程序一般都是从Hello World开始的。这一节，我们就用Hello World程序进行Cuda实战，程序的内容主要是让CPU和GPU分别打印Hello World，在CPU中打印内容为，在GPU中打印内容为，具体实现请参考程序案例。
cuda编程入门-HelloWorld
@way
11-11 928
为了提高并行运算速度，开始接触cuda编程，cuda 在安装的时候提供了很多示例，一般位于Home目录下，可以作为学习的参考，同时官网提供的入门教程，就是很好的学习资料。 cuda编程首先要转换与cpu 下编程的思想，在cpu 下，我们一般都是按照顺序进行执行，或者开几个线程处理几个任务，而在gpu 下，是分block 和 thread的，gpu由多个block（x,y ,z ）组成，每个bloc...
《cuda c编程权威指南》01- 用gpu输出hello world
一千零一夜的博客
07-30 408
cuda 输出hello world
参加【Nvidia CUDA线上训练营2023春】笔记——尝试写自己的第一个CUDA程序
扶风的博客
02-06 248
CUDA编程入门相对来说不是很难，都遵循一套流程，很容易就可以进行更改，新手入门最重要的就是转变代码编写思维，由原来的串行思路转为并行思路，充分发挥GPU加速计算的优点。
CUDA.jl：Julia中的CUDA编程
02-03
CUDA.jl：Julia中的CUDA编程
CUDA编程入门
01-07
目录1 [GPU硬件架构及运行机制]...CUDA编程需要CPU和GPU协同工作，程序中既包含host程序，又包含device程序，他们分别在CPU和GPU上运行，二者间可以进行通信。 host – CPU及其内存
CUDA Hello World小例子
03-31
用GPU和CPU初始化同样一个矩阵，并将每个元素求平方，各自计时，最后计算加速比。 GPU主要是用来做浮点数运算的，打印Hello World是毫无意义的……标题只是用来骗人的。
CUDA编程指南5.0版.pdf
09-06
CUDA编程指南 中文版5.0 CUDA编程指南 中文版
【CUDA】新手的第一个CUDA C程序：“Hello World from GPU !“
TangPlusHPC的博客
06-11 1251
CUDA是一种通用的并行计算平台和编程模型，它利用NVIDIA GPU中的并行计算引擎能更有效地解决复杂的计算问题。CUDA C是标准ANSI C的一个扩展，它带有的少数语言扩展功能使异构编程成为可能，同时也能用过API来管理设备、内存和其他任务。一个CUDA程序包含了以下两部分的混合。NVIDIA 的CUDA nvcc编译器在编译过程中将设备代码从主机代码中分离出来。主机代码是标准的C代码，使用C编译器进行编译。设备代码，也就是核函数，是用扩展的带有标记数据并行函数关键字的CUDA C语言编写的。设备代码
笔记-cuda编程:基础与实践
qq_43647367的博客
06-24 293
cuda程序框架一般为主机数据拷贝到设备，执行核函数，设备数据拷贝到主机，由于cpu是不需要数据拷贝的，所以这整体时间在于执行任务函数，而cuda正是可以加速这一点；但是cpu与gpu内存间的PCIe带宽远小于gpu的显存带宽(cpu与gpu数据传输很慢)，所以显而易见的加速gpu(相同cpu运行时间下，怎么让gpu运行时间更短)的方法有减少数据传输的时间，增加核函数的算术强度，增加并行程度；而上述都是从"宏观"或者软件层面上说的，而限于gpu硬件实现，也有一些有关速度的值得注意的地方。
cuda编程基础
weixin_41977938的博客
05-19 3735
1.向量相加的cuda实现 (按对应位置相加) #include <iostream> #include <cuda_runtime.h> #define length 10000 //__global__ int foo(int a){}表示一个内核函数， //是一组由GPU执行的并行计算任务，以foo<<>>(a)的形式或者driver API的形式调用。 //目前__global__函数必须由CPU调用，并将并行计算任务发射到GPU的任.
3. CUDA编程手册中文版---编程接口
专注于人工智能领域的小何尚
04-12 3172
第三章编程接口 CUDA C++ 为熟悉 C++ 编程语言的用户提供了一种简单的途径，可以轻松编写由设备执行的程序。 它由c++语言的最小扩展集和运行时库组成。 编程模型中引入了核心语言扩展。它们允许程序员将内核定义为 C++ 函数，并在每次调用函数时使用一些新语法来指定网格和块的维度。所有扩展的完整描述可以在 C++ 语言扩展中找到。任何包含这些扩展名的源文件都必须使用 nvcc 进行编译，如使用NVCC编译中所述。 运行时在 CUDA Runtime 中引入。它提供了在主机上执行的 C 和 C++ 函数
CUDA C 1
qq_39586331的博客
12-08 154
初识cuda c例如：使用cuda c 编写程序调用GPU输出Hello World。 使用CUDA nvcc编译器编译程序 执行可执行文件 四 编写内核函数，输出字符 1 创建源文件 2 使用CUDA nvcc编译器编译程序 3 执行可执行文件 三重尖号意味着从主线程到端代码的调用。1和10分别表示有1个块区域和10个线程 表示重置当前线程所关联过的当前设备的所
【CUDA编程】一：从简单的实例入门
qq_43257640的博客
05-31 5379
前提 拥有NVIDIA的显卡- 配置好cuda 有C/C++基础 代码1 //helloworld.cu #include <iostream> #include <stdio.h> /* * __global__ 用于修饰在gpu上运行且由cpu调用的函数 */ __global__ void k_add (int a, int b, int* c ){ *c = a+b; } /* * __host__用于修饰在cpu上运行的函数，一般缺省 */ int main(v
cuda::barrier
最新发布
12-13
CUDA的`__syncthreads()`函数在一个线程块中同步所有的线程，以确保它们都执行到这一点。但是，在某些情况下，我们可能需要跨线程块同步。 CUDA的`cuda::barrier`提供了这样的功能。它将整个GPU上的所有线程块同步起来，以确保所有线程块在之前的操作都完成之后再继续执行。 使用`cuda::barrier`可以实现全局同步，保证所有线程块执行到当前代码之前的代码都已经完成。这对于某些需要多个线程块之间通信或协作的算法或应用程序非常有用。 在使用`cuda::barrier`时，所有的线程块都需要调用`__syncthreads()`函数来执行本地同步。然后，只有当所有线程块都执行到`__syncthreads()`才能通过`cuda::barrier`进行全局同步。 `cuda::barrier`使用方法如下： ```c++ cuda::barrier global_sync_barrier; // ... __syncthreads(); // 本地同步 global_sync_barrier.sync(); // 全局同步 // ... ``` 需要注意的是，`cuda::barrier`仅适用于使用CUDA并行编程模型的GPU计算。在使用`cuda::barrier`时，应确保算法或应用程序的实现仍然正确且无死锁的。 总之，`cuda::barrier`是一种可以跨线程块同步的机制，它可以确保所有线程块在执行到当前代码之前都已经完成之前的操作。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 1
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值