CUDA 学习笔记 —— （八）内存模型

已于 2022-04-14 17:59:54 修改
阅读量1.3k 收藏 3
点赞数
分类专栏： GPU 文章标签： CUDA 内存
于 2022-03-31 09:47:26 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-NC-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/chongbin007/article/details/123838980
版权

文章目录

GPU device内存

https://face2ai.com/CUDA-F-4-1-内存模型概述/

内存种类

一个SM中有一个shared memory，这个时一个block中所有线程都可访问的。
每个线程有自己的local memory和一组寄存器。
所有的block都可以访问一个叫做global memory和一个constant memory(只读），和一个纹理寄存器（只读）。
在这里插入图片描述
在这里插入图片描述
速度从快到慢

  1. Register file
  2. Shared Memory
  3. Constant Memory
  4. Texture Memory
  5. Local Memory and Global Memory

不同内存使用方式
在这里插入图片描述
总结：

  1. 一般我们操作可编程内存常用的时shared_memory和global_memory。声明变量的时候默认是在register，而声明数组默认是在local memory。
  2. GPU也有缓存，一级二级，这些是不可编程的，硬件决定。texture memory应该也是不可编程的。

Global memory

https://developer.nvidia.com/blog/how-access-global-memory-efficiently-cuda-c-kernels/
声明可以使用__device__标识符
使用cudaMalloc , cudaFree and cudaMemcpy中host to device.
这些操作的都是global meomory。他是常驻DRAM的。

比如下面我们声明后动态赋值

__device__ int globalArray[256];
 
void foo()
{
    ...
    int *myDeviceMemory = 0;
    cudaError_t result = cudaMalloc(&myDeviceMemory, 256 * sizeof(int));
    ...
}
连续访问与跨越访问global memory

https://blog.csdn.net/Bruce_0712/article/details/65444997
因为对 Global memory 访问没有缓存，因此显存的性能对GPU至关重要。为了能够高效的访问显存，读取和存储必须对齐，宽度为4Byte。如果没有正确的对齐，读写将被编译器拆分为多次操作，极大的影响效率。

每个 thread 一次读取的内存数据量，可以是 32 bits、64 bits、或 128 bits。不过，32 bits 的效率是最好的。64 bits 的效率会稍差，而一次读取 128 bits 的效率则比一次读取 32 bits 要显著来得低（但仍比 non-coalesced 的存取要好）。

每组16 Threads 同时访问连续且对齐的64/128 Byte称为Coalesced访问模式，这是达到带宽的理路峰值的必要条件。

看个例子

template <typename T>
__global__ void offset(T* a, int s)
{
  int i = blockDim.x * blockIdx.x + threadIdx.x + s;
  a[i] = a[i] + 1;
}

template <typename T>
__global__ void stride(T* a, int s)
{
  int i = (blockDim.x * blockIdx.x + threadIdx.x) * s;
  a[i] = a[i] + 1;
}
 for (int i = 0; i <= 32; i++) {
   checkCuda( cudaMemset(d_a, 0, n * sizeof(T)) );

   checkCuda( cudaEventRecord(startEvent,0) );
   offset<<<n/blockSize, blockSize>>>(d_a, i);
   stride<<<n/blockSize, blockSize>>>(d_a, i);
   checkCuda( cudaEventRecord(stopEvent,0) );
   checkCuda( cudaEventSynchronize(stopEvent) );

   checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
   printf("%d, %f\n", i, 2*nMB/ms);
 }

我们看到使用这两种方式调用kernel函数唯一区别就是内存的访问方式int i = (blockDim.x * blockIdx.x + threadIdx.x) * s;这里一个是使用内存下标连续访问，而第二种是跳跃访问，运行结果发现

Offset, Bandwidth (GB/s):
0~ 32 均为45.955883左右

Stride, Bandwidth (GB/s):
1, 45.687134
2, 23.113905
3, 15.409269
4, 11.591246
5, 9.256516
6, 7.697044
7, 6.603973
8, 5.765683
9, 5.769941
10, 5.761431
11, 5.723443
12, 5.628603
13, 5.564459
14, 5.410318
15, 5.321867
16, 5.343708
17, 5.170417
18, 5.079649
19, 4.963469
20, 4.985641
21, 4.734848
22, 4.686563
23, 4.526443
24, 4.653068
25, 4.278477
26, 4.273797
27, 4.098898
28, 4.105360
29, 3.809118
30, 3.794317
31, 3.660965
32, 4.018775

可以看到跨越访问在global memeory会损失性能。
但是当我们使用2D数组的时候，就不可避免的会跨越访问，这时，我们可以借助shared Memor处理转置关系来提高性能。

Shared Memory

on chip所以速度非常快，大概是local和globa lmemory的100倍。

shared memory会被分配给每个线程block上，因为block也是在逻辑上分配的，所以shared也是逻辑上分配给每个block。

因为每个block上的不同线程访问同一个block，所以有竞争问题。同步问题可以使用__syncthreads关键字来避免。
举例：

静态共享内存

静态共享内存在创建时候指明大小。

动态共享内存

而动态内存可以不指明大小。
但是在启动kernel函数的时候来指定
dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n);
看下面代码例子，对比使用静态共享内存和动态共享内存


#include <stdio.h>
__global__ void staticReverse(int *d, int n)
{
  __shared__ int s[1000];
  int t = threadIdx.x;
  int tr = n-t-1;
  //从global memory拷贝写入shared memory
  s[t] = d[t];
  //因为数组s是所有线程共享的，如果不做同步执行下面语句则可能出现数据竞争问题
	// Will not conttinue until all threads completed
	//调用同步函数，只有当前block中所有线程都完成之后，再往下走
  __syncthreads();
  //从shared memory读，然后写回到global memory
  d[t] = s[tr];
}

__global__ void dynamicReverse(int *d, int n)
{
  extern __shared__ int s[];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}
//目的：将一个数组中的数据前后交换，实现倒序
int main(void)
{
  const int n = 1000;
  int a[n], r[n], d[n];
  
  for (int i = 0; i < n; i++) {
    a[i] = i;
    r[i] = n-i-1;
    d[i] = 0;
  }

  int *d_d;
  cudaMalloc(&d_d, n * sizeof(int)); 
  
  // run version with static shared memory
  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
  float time_gpu;
  cudaEvent_t start_GPU,stop_GPU;
  cudaEventCreate(&start_GPU);
  cudaEventCreate(&stop_GPU);
  cudaEventRecord(start_GPU,0);
  staticReverse<<<1,n>>>(d_d, n);
  //globalReverse<<<1,n>>>(d_d, n);
  cudaEventRecord(stop_GPU,0);
  cudaEventSynchronize(start_GPU);
  cudaEventSynchronize(stop_GPU);
  cudaEventElapsedTime(&time_gpu, start_GPU,stop_GPU);
  printf("\nThe time from GPU:\t%f(ms)\n", time_gpu);
  cudaDeviceSynchronize();
  cudaEventDestroy(start_GPU);
  cudaEventDestroy(stop_GPU);
  
  cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost);
  //check
  for (int i = 0; i < n; i++) {
    if (d[i] != r[i]) 
      printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
  }
    
  
  // run dynamic shared memory version
  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);

  cudaEventCreate(&start_GPU);
  cudaEventCreate(&stop_GPU);
  cudaEventRecord(start_GPU,0);
  dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n);
  cudaEventRecord(stop_GPU,0);
  cudaEventSynchronize(start_GPU);
  cudaEventSynchronize(stop_GPU);
  cudaEventElapsedTime(&time_gpu, start_GPU,stop_GPU);
  printf("\nThe time from GPU:\t%f(ms)\n", time_gpu);
  cudaDeviceSynchronize();
  cudaEventDestroy(start_GPU);
  cudaEventDestroy(stop_GPU);
  cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);
  for (int i = 0; i < n; i++) 
    if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
}

__syncthreads() 是轻量级的，并且是以block 级别做同步。

CPU Host内存

对CUDA架构而言，主机端的内存被分为两种，一种是可分页内存（pageable memroy）和页锁定内存（page-locked或 pinned）。

可分页内存 Pageable

可分页内存是使用I malloc()或者new在主机上分配

页锁定内存 Pinned(Page-locked)

页锁定内存是使用CUDA函数cudaMallocHost 或者cudaHostAlloc()在主机内存上分配.
cudaFreeHost()来释放
注意cudaMalloc()是在GPU上分配内存

页锁定内存的重要属性是：主机的操作系统将不会对这块内存进行分页和交换操作，确保该内存始终驻留在物理内存中。

由于每个页锁定内存都需要分配物理内存，并且这些内存不能交换到磁盘上，所以页锁定内存比使用标准malloc()分配的可分页内存更消耗内存空间。

主机端-设备端的数据传输带宽高；某些设备上可以通过zero-copy功能映射到设备地址空间，从GPU直接访问，省掉主存与显存间进行数据拷贝的工作；

pinned memory 不可以分配过多：导致操作系统用于分页的物理内存变少， 导致系统整体性能下降；通常由哪个cpu线程分配，就只有这个线程才有访问权限；

测试

下面我们做个测试：

  1. host分配可分页内存，然后拷贝到device再拷贝回host，重复执行10次
  2. host分配页锁定内存，然后拷贝到device再拷贝回host，重复执行10次
  3. 对比执行时间
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "iostream"
#include <stdio.h>
 
using namespace std;

#define COPY_COUNTS 10
#define MEM_SIZE 25*1024*1024
 
float cuda_host_alloc_test(int size, bool up)
{
	//耗时统计
	cudaEvent_t start, stop;
	float elapsedTime;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
 
	int *a, *dev_a;
	//在主机上分配页锁定内存
    cudaError_t cudaStatus = cudaMallocHost((void **)&a, size * sizeof(*a));
	//在设备上分配内存空间
	cudaStatus = cudaMalloc((void **)&dev_a, size * sizeof(*dev_a));
	//计时开始
	cudaEventRecord(start, 0);
 
	for (int i = 0; i < COPY_COUNTS; i++)
	{
		//从主机到设备复制数据
		cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice);
		//从设备到主机复制数据
		cudaStatus = cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost);

	}
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime, start, stop);
 
	cudaFreeHost(a);
	cudaFree(dev_a);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);
 
	return (float)elapsedTime / 1000;
 
}
 
float cuda_host_Malloc_test(int size, bool up)
{
	//耗时统计
	cudaEvent_t start, stop;
	float elapsedTime;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	int *a, *dev_a;
 
	//在主机上分配可分页内存
	a = (int*)malloc(size * sizeof(*a));
 
	//在设备上分配内存空间
	cudaError_t	cudaStatus = cudaMalloc((void **)&dev_a, size * sizeof(*dev_a));
 
	//计时开始
	cudaEventRecord(start, 0);
 
	//执行从copy host to device 然后再 device to host执行100次，记录时间
	for (int i = 0; i < COPY_COUNTS; i++) {
		//从主机到设备复制数据
		cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice);
		//从设备到主机复制数据
		cudaStatus = cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost);

	}
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime, start, stop);
 
	free(a);
	cudaFree(dev_a);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);
 
	return (float)elapsedTime / 1000;
}
 
int main()
{
	float allocTime = cuda_host_alloc_test(MEM_SIZE, true);
	cout << "页锁定内存: " << allocTime << " s" << endl;
	float mallocTime = cuda_host_Malloc_test(MEM_SIZE, true);
	cout << "可分页内存: " << mallocTime << " s" << endl;
	return 0;
}

chongbin@DESKTOP-SFAMAG4:~/c_project/6memory$ ./a.out 
页锁定内存: 0.685681 s
可分页内存: 1.43228 s

可以看到：页锁定内存的访问时间约为可分页内存的访问时间的一半。

总结与建议

  1. Pinned内存使用零拷贝，减少内存拷贝到时间，但是不要过分的分配pinned memory。因为这会削减整体系统的性能，他是直接减少物理内存的。所以建议衡量使用pinned meomory和对这个系统的性能，合理使用pinned memory.

  2. 建议在写代码时先使用pageble memory后面优化的时候再用pinned memory.

  3. 建议在内存拷贝的时候，将内存分配更多的小份拷贝。比如可以使用cudaMemcpy2D(dest, dest_pitch, src, src_pitch, w, h, cudaMemcpyHostToDevice)来做二维数组的拷贝。

参考链接：
https://blog.csdn.net/dcrmg/article/details/54975432
https://developer.nvidia.com/blog/how-optimize-data-transfers-cuda-cc/

Charles Ray
关注
  • 0
    点赞
  • 踩
  • 3
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 0
    评论
专栏目录
我的CUDA学习笔记
08-29
这是我学习CUDA过程中记录的一些笔记，大部分是翻译自CUDA programming guide，加入了一些自己的理解
计算机视觉大型攻略 —— CUDA(2)执行模型
01-06
Professional CUDA C Programming[1]是一本不错的入门书籍，虽说命名为”Professional”，但实际上非常适合入门阅读。他几乎涵盖了所有理论部分和编程技巧，更重要的是每一章都有完整的实例程序。只不过对于入门来讲，这本书有点太厚了，行文有些啰嗦。准备写几篇文章提取一下关键章节的关键部分。 上一篇写了如何写一个简单的CUDA程序。为了进一步优化程序性能，我们经常会调整Block数量和Thread数量，不断的寻找最优的组合。这一篇探讨的是最优组合背后的故事，为什么有些组合可以达到更好的性能。理解CUDA的执行模型，有助于进一步并行化线程，提高程序性能。  这一篇算
CUDA学习笔记，GPU编程，C++编程
06-08
CUDA学习笔记
【yolov4】学习笔记一
12-21
想多熟悉下pytorch,因而选择了yolov4-pytorch版本来熟悉整个流程。 代码下载地址: https://github.com/Tianxiaomo/pytorch-YOLOv4 这位大神厉害啊，这么快就复现了,好好学习膜拜中！ yolov4.weights权重百度网盘下载地址: https://pan.baidu.com/s/1dAGEW8cm-dqK14TbhhVetA 提取码:dm5b 下载后用pycharm打开工程。 个人更喜欢把参数写在代码中，所以将demo.py中main部分改为 if __name__ == '__main__': cfgfile = 'cfg
计算机视觉大型攻略 —— CUDA(3)内存模型（一）CUDA内存
plateros的专栏
02-21 1727
上篇写了执行模型，这篇继续总结[1]的第四章: 内存模型。 首先介绍CUDA内存空间以及物理内存与逻辑内存的对应关系，后续文章写如何通过改善访存策略来提升CUDA Kernel函数的性能。
CUDA进阶：内存模型
Shilong的博客
09-19 1107
CUDA内存模型
CUDA 内存模型
A-Egoist的博客
12-28 901
CUDA 内存模型 CUDA 可编程内存类型： 寄存器(Registers) 共享内存(Shared Memory) 本地内存(Local Memory) 常量内存(Constant Memory) 纹理内存(Texture Memory) 全局内存(Global Memory) 上述各种内存都有自己的作用域，生命周期和缓存行为。一个核函数中的每个线程都有自己私有的本地内存。一个线程块有自己的共享内存，对同一线程块内所有线程可见，其内容持续线程块的整个生命周期。所有线程都可以访问全局内存。所有线程都能访
【CUDA编程】CUDA内存模型
非晚非晚的博客
11-04 1799
thread：每个线程都用自己的registers（寄存器）和local memory（局部内存） block：每个线程块(block)内都有自己的shared memory（共享内存），所有线程块内的所有线程共享这段内存资源 grid：每个grid都有自己的global memory（全局内存），constant memory（常量内存）和texture memory（纹理内存），不同线程块的线程都可使用。其中常量内存和纹理内存为只读内存空间。
cuda内存模型，操作数据复制
Rolandxxx的博客
09-19 948
而Pageable memory是普通房间，在酒店房间不够时，选择性得把你的房间腾出来给其他人交换用，你就到硬盘去待着，这就可以容纳更多人了。性能要比Pageable memory好，因为Pageable memory会降低你程序的优先级，把内存交换给别人用。SharedMemory离GPU最近，其实pageable memory是最远的，但因为GPU不能直接访问pageable memory，所以就比较起来也没有意义。只需要知道，谁距离计算芯片近，谁速度就越快，空间越小，价格越贵，
CUDA编程错误检测checkCudaErrors
weixin_42269062的博客
12-14 4378
checkCudaErrors可对CUDA API进行错误检测，及时定位错误所在位置。 需要头文件：#include "helper_cuda.h" 在/home/dell/NVIDIA_CUDA-11.2_Samples/common/inc文件夹中。 （CUDA Runtime API的一个特性：之前的kernel或者CUDA函数挂掉了，会导致后续持续的返回错误） 查看核函数是否正确执行，在核函数后加上 cudaError_t cudaStatus = cudaGetLastError();
cuda立体匹配——中值滤波
06-20
立体匹配是立体视觉从图像生成三维点云的常规手段。立体匹配算法主要是通过建立一个能量代价函数，通过此能量代价函数最小化来估计像素点视差值。立体匹配算法的实质就是一个最优化求解问题，通过建立合理的能量函数...
cuda runtime-GPU内存模型和数据拷贝
zwhdldz的博客
07-30 249
CUDA Driver 是与 GPU 沟通的驱动级别底层 API，对 Driver API 的理解，有利于理解后续的 Runtime API。CUDA Driver 随显卡驱动发布，需要与 cudatoolkit 分开看，CUDA Driver 对应于 cuda.h 和 libcuda.so。Driver API 主要知识点是 Context 的管理机制 以及 CUDA 系列接口的开发习惯（错误检查方法），还有内存模型。
【CUDA编程（3）】：CUDA内存模型
Hali_Botebie的博客
09-27 62
在一个内核函数中，每个线程都有自己的本地内存，每个线程块有自己的共享内存并对块内的所有线程可见，一个线程网格中的所有线程都可以访问全局内存、常量内存和纹理内存，其中常量内存和纹理内存为只读内存空间。中寄存器的数量比较有限，一旦内核函数使用了超过硬件限制的寄存器数量，则会使用本地内存来代替多占用的寄存器，这种寄存器溢出的情况会带来性能上的不利影响，实际编程过程中我们应该避免这种情况。纹理内存是一种通过指定的只读缓存访问的全局内存，是对二维空间局部性的优化，所以使用纹理内存访问二维数据的线程可以达到最优性能。
虚拟内存、锁页内存、内存分页、分段、段页式内存管理
Rolandxxx的博客
09-17 835
当主存没有足够的空间并且有更多的数据写入内存时，一些页面将被移动到硬盘驱动器（hard drive）。使用分页内存，每个程序都有自己的逻辑内存，即内存可以分成连续的页面。所以一个程序的内存可能在它自己的逻辑内存上是连续的，但在物理内存上是分段的。逻辑地址是程序运行时由CPU生成的地址，它是一个虚拟地址，之所以说是虚拟地址，是因为它不是物理上真实存在的地址。，是一种内存管理方案，计算机通过该方案存储和检索二级存储器中的数据，二级存储器通常是指硬盘驱动器上的交换swap区域，用于主存储器。
CUDA全局内存(CUDA C编程 权威指南)
最新发布
qq_43616651的博客
11-30 172
访问共享内存必须同步。常量内存：使用__constant__来修饰常量变量，必须在全局空间内和核函数之外进行声明，大小只有64KB，常量内存时静态声明的，并对统一编译单元中所有的核函数可见。这种类型的内存分配对于主机（CPU）和设备（GPU）之间的异步内存传输非常有利，因为GPU可以直接访问它，无需将数据复制到可分页内存中的缓冲区。统一内存寻址提供了一个“单指针到数据”模型，在概念上类似于零拷贝内存，但零拷贝内存是在主机内存中进行分配，因此受到在PCIe总线上访问零拷贝内存的一个像，核函数性能具有高延迟。
linux内存锁页,Linux内存子系统——Locking Pages(内存锁定)
weixin_31098573的博客
04-30 902
你可以让系统将特定的虚拟内存页与实际页帧相"关联"，并保持这样的状态(称为锁定)。该部分内存不会被swap机制交换出来，也不会产生pagefault(因为已经分配了实际的物理内存)。为什么需要锁定内存一个背景知识pagefault用户在分配出一部分虚拟内存时，其背后可能并没有真正的物理内存与之对应，只有在用户真正需要访问内存时，系统才会为这段虚拟内存分配实际的物理内存，这个过程叫做pagefaul...
CUDA 编程简介（下）
weixin_44885334的博客
12-13 997
另外，如果 warp 里的线程同时读取同一个数据，使用 broadcast 机制，也可以在。如果需要大量数据，这种一次次的数据调度会特别慢。在使用异步预取进行了一系列重构之后，您应该看到内存传输次数减少了，但是每次传输的量增加了，并且内核执行时间大大减少了。不过，在 CUDA 的云平台上，不同的非默认 stream 上，它们还是串行执行的。不过，在 CUDA 的云平台上，不同的非默认 stream 上，它们还是串行执行的。现在的 GPU 的一次性调度的内存块大小也变大了，足够支持一个 warp 内的。
CUDA编程——zero copy
Zhang Junior 的博客
02-04 8047
零复制　　zero copy（零复制）是一种特殊形式的内存映射，它允许你将host内存直接映射到设备内存空间上。其实就是设备可以通过直接内存访问（direct memory access，DMA）方式来访问主机的锁页内存。 　　锁页主机内存　　现代操作系统都支持虚拟内存，操作系统实现虚拟内存的主要方法就是通过分页机制。操作系统将内存中暂时不使用的内容换出到外存（硬盘等大容量存储）上，从而腾出空间存
CUDA C编程5 - CUDA内存模型
DU_YULIN的博客
03-06 593
系列文章目录 文章目录系列文章目录前言一、CUDA内存模型1. 寄存器2. 本地内存3. 共享内存4. 常量内存5. 纹理内存6. 全局内存7. GPU缓存二、静态全局内存示例总结参考资料附录 前言 最近在温习CUDA C 全局内存的知识，这里对关键知识点进行总结并分享给大家。 较差的全局内存访问方式是造成内存负载效率大幅下降的原因之一。 这里主要介绍以下几点内容： 剖析核函数与全局内存的联系及其对性能的影响； 介绍全局内存访问模式 介绍如何通过核函数高效的利用全局内存 一、CUDA内存
CUDA学习笔记（四）——CUDA性能
05-19
CUDA是一个可以大大提高GPU计算性能的编程平台，但是为了实现最高的性能，需要遵循一些最佳实践和优化技巧。以下是一些常用的CUDA性能优化方法： 1. 矢量化：使用SIMD指令和数据结构（如向量和矩阵）来并行计算。 ...

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

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

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

打赏作者

Charles Ray

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

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

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

打赏作者

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

抵扣说明：

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

余额充值