CUDA性能优化—-kernel调优(nvprof工具的使用)

CUDA性能优化—-kernel调优(nvprof工具的使用)1、引言本文主要介绍并行分析,涉及掌握nvprof的几个metrics参数,所用的例子是CUDA性能优化—-线程配置一文中所提到的sumMatrix2D.cu例子。接下来本文会做一些列的试验,测试环境:TeslaM2070一块,CUDA6.0,操作系统:RedHat4.1.2-50,gccversion4.1.220080704首先回顾一下sumMatrix2D的kern…

大家好,又见面了,我是你们的朋友全栈君。

转自博客——http://subblogwujiaxing009m26.lofter.com/

1、引言

本文主要介绍并行分析,涉及掌握nvprof的几个metrics参数,所用的例子是CUDA性能优化—-线程配置一文中所提到的sumMatrix2D.cu例子。
接下来本文会做一些列的试验,测试环境:Tesla M2070一块,CUDA 6.0,
操作系统:Red Hat 4.1.2-50,gcc version 4.1.2 20080704

首先回顾一下sumMatrix2D的kernel函数:

__global__ void sumMatrix2DKernel(float *d_MatA,float *d_MatB,float *d_MatC,int nx,int ny) 
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x; 
    int idy = threadIdx.y + blockDim.y * blockIdx.y; 
    int tid = nx*idy + idx; 

    if(idx < nx && idy < ny) 
       d_MatC[tid] = d_MatA[tid] + d_MatB[tid]; 
}

输入数据矩阵的维度是nx=16384, ny=16384:

int nx = 1<<14;
int ny = 1<<14;

下面的代码用来配置main函数的参数,也就是block的维度配置:

if (argc > 2) 
{ 
    dimx = atoi(argv[1]);
    dimy = atoi(argv[2]); 
} 
dim3 block(dimx, dimy); 
dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

编译运行:

$ nvcc -O3 -arch=sm_20 sumMatrix2D.cu -o sumMatrix2D
$ ./sumMatrix2D 32 32

其中第二句命令中的”32 32″是指定block维度的参数。

查询nvprof工具版本的命令:

$ nvprof –version

输出:

nvprof: NVIDIA ® Cuda command line profiler Copyright © 2013 – 2014 NVIDIA Corporation Release version 6.0 (20)

2、Checking Active Warps with nvprof

在做各项数据比较的时候需要有个基准,这里使用四个block维度配置的时间消耗作为基准观察,分别为(32,32)(32,16)(16,32)和(16,16),其中第一个参数是x维度,第二个参数是y维度。
下面是几种配置的时间消耗输出结果:

$ ./sumMatrix2D 32 32
./sumMatrix2D Program Starting…
–sumMatrix2DOnHost() elapsed 360.000000 ms… –sumMatrix2DOnGPU<<<(512,512),(32,32)>>> elapsed 70.000000 ms…

$ ./sumMatrix2D 32 16 ./sumMatrix2D Program Starting…
–sumMatrix2DOnHost() elapsed 360.000000 ms…
–sumMatrix2DOnGPU<<<(512,1024),(32,16)>>> elapsed 40.000000 ms…

$ ./sumMatrix2D 16 32 ./sumMatrix2D Program Starting…
–sumMatrix2DOnHost() elapsed 360.000000 ms…
–sumMatrix2DOnGPU<<<(1024,512),(16,32)>>> elapsed 60.000000 ms…

$ ./sumMatrix2D 16 16 ./sumMatrix2D Program Starting…
–sumMatrix2DOnHost() elapsed 360.000000 ms…
–sumMatrix2DOnGPU<<<(1024,1024),(16,16)>>> elapsed 50.000000 ms…

比较这几个结果,不难发现,性能最差的是第一个(32,32),性能最好的是第二个(32,16),这里可以猜测到是:拥有更多的block数目并行性更好。这个猜测可以使用nvprof 的achieved_occupancy这个metric参数来验证。该参数的定义公式在CUDA性能优化—-warp深度解析有介绍,实际上就是指每个SM在每个cycle能够达到的最大active warp数目占总warp的比例。下面是使用该参数后得到的结果(注意由于输出项多,做了简化处理):
命令:

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 32 32

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 32 32
==27432== NVPROF is profiling process 27432, command: ./sumMatrix2D 32 32
–sumMatrix2DOnGPU<<<(512,512),(32,32)>>> achieved_occupancy 0.506396

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 32 16
==27454== NVPROF is profiling process 27454, command: ./sumMatrix2D 32 16
–sumMatrix2DOnGPU<<<(512,1024),(32,16)>>> achieved_occupancy 0.731333

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 16 32
==27493== NVPROF is profiling process 27493, command: ./sumMatrix2D 16 32
–sumMatrix2DOnGPU<<<(1024,512),(16,32)>>> achieved_occupancy 0.826147

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 16 16
==27545== NVPROF is profiling process 27545, command: ./sumMatrix2D 16 16
–sumMatrix2DOnGPU<<<(1024,1024),(16,16)>>> achieved_occupancy 0.819718

从上面的输出对比可以得知两点认识:
(1)由于第二个配置比第一个有更多的block数量,device就会达到更多active warp(跟鸡蛋放在多个篮子的道理差不多)。也就是第二个性能优于第一个的原因。
(2)第四个的achieved Occupancy比较高,但是却不是最快的,因此,较高的achieved Occupancy并不一定就意味着更好的性能,也就是说还有更多的因素影响着GPU的性能。

3、checking memory operations with nvprof

对于d_MatC[tid] = d_MatA[tid] + d_MatB[tid]来说共有三个memory操作:两个memory load和一个memory store。要查看这些操作的效率可以使用nvprof的两个metric参数,如果想要查看memory的throughput,则可使用gld_throughput参数,实验结果如下(注意由于输出项多,做了简化处理):
命令:

$ nvprof –metrics gld_throughput ./sumMatrix2D 32 32
–sumMatrix2DOnGPU<<<(512,512),(32,32)>>> elapsed 1090.000000 ms…
–Global Load Throughput:35.557GB/s

$ nvprof –metrics gld_throughput ./sumMatrix2D 32 16
–sumMatrix2DOnGPU<<<(512,1024),(32,16)>>> elapsed 1440.000000 ms…
–Global Load Throughput:56.396GB/s

$ nvprof –metrics gld_throughput ./sumMatrix2D 16 32
–sumMatrix2DOnGPU<<<(1024,512),(16,32)>>> elapsed 1070.000000 ms…
–Global Load Throughput:81.023GB/s

$ nvprof –metrics gld_throughput ./sumMatrix2D 16 16
–sumMatrix2DOnGPU<<<(1024,1024),(16,16)>>> elapsed 1060.000000 ms…
–Global Load Throughput:93.694GB/s

不难看到,第四个拥有最高的load throughput,但是却比第二个慢(第二个也就是第四个的一半多点),所以,较高的load throughput也不一定就有较高的性能。之后讲到memory transaction时会具体分析这种现象的原因,简单说,就是高load throughput有可能是一种假象,如果需要的数据在memory中存储格式未对齐、不连续,会导致许多额外的不必要的load操作,所以本文中的efficiency会这么低。

然后,我们可以使用nvprof的gld_efficiency来度量load efficiency,该metric参数是指我们确切需要的global load throughput与实际得到global load memory的比值。这个metric参数可以让我们知道,应用程序的load操作利用device memory bandwidth的程度,实验结果如下(注意由于输出项多,做了简化处理):
命令:

$ nvprof –metrics gld_efficiency ./sumMatrix2D 32 32 –sumMatrix2DOnGPU<<<(512,512),(32,32)>>> elapsed 1610.000000 ms…
–Global Memory Load Efficiency:100.01%

$ nvprof –metrics gld_efficiency ./sumMatrix2D 32 16
–sumMatrix2DOnGPU<<<(512,1024),(32,16)>>> elapsed 1610.000000 ms…
–Global Memory Load Efficiency:99.95%

$ nvprof –metrics gld_efficiency ./sumMatrix2D 16 32
–sumMatrix2DOnGPU<<<(1024,512),(16,32)>>> elapsed 1610.000000 ms…
–Global Memory Load Efficiency:49.89%

$ nvprof –metrics gld_efficiency ./sumMatrix2D 16 16
–sumMatrix2DOnGPU<<<(1024,1024),(16,16)>>> elapsed 1610.000000 ms…
–Global Memory Load Efficiency:49.99%

从上述结果可知,最后两个测试的load efficiency只是前两个的一半。这也可以解释,为什么较高的throughput和较高的Occupancy却没有产生较好的性能。尽管最后两个测试的load操作数目要多不少(因为二者throughput较高),但是他们的load effecitiveness却低不少(由于efficiency较低)。
观察最后两个可以发现,他们block的x维配置是warp的一半,前文曾提到,该维度应该保持为warp大小的整数倍。关于其具体原因将在后续博文详细解释。

4、Exposing More Parallelism

我们现在可以得出一个结论:blockDim.x应该是warp大小的整数倍。这样做是很容易就提升了load efficiency。现在,我们可能还有其他疑惑,比如:
(1)继续调整blockDim.x是否会继续增加load throughput?
(2)还有其他方法能增大并行性吗?
现在,我们重新整一个基准数据出来,这两个问题可以从这个基准分析个大概(此处改用了cuda的计时函数):

$ ./sumMatrix2D 64 2
–sumMatrix2DOnGPU<<<(256,8192),(64,2)>>> elapsed 33.527294 ms…

$ ./sumMatrix2D 64 4
–sumMatrix2DOnGPU<<<(256,4096),(64,4)>>> elapsed 34.802238 ms…

$ ./sumMatrix2D 64 8
–sumMatrix2DOnGPU<<<(256,2048),(64,8)>>> elapsed 36.614143 ms…

$ ./sumMatrix2D 128 2
–sumMatrix2DOnGPU<<<(128,8192),(128,2)>>> elapsed 32.602848 ms…

$ ./sumMatrix2D 128 4
–sumMatrix2DOnGPU<<<(128,4096),(128,4)>>> elapsed 34.658592 ms…

$ ./sumMatrix2D 128 8
–sumMatrix2DOnGPU<<<(128,2048),(128,8)>>> elapsed 46.740578 ms…

$ ./sumMatrix2D 256 2
–sumMatrix2DOnGPU<<<(64,8192),(256,2)>>> elapsed 32.661919 ms…

$ ./sumMatrix2D 256 4
–sumMatrix2DOnGPU<<<(64,4096),(256,4)>>> elapsed 38.260609 ms…

$ ./sumMatrix2D 256 8
–sumMatrix2DOnGPU<<<(64,2048),(256,8)>>> elapsed 0.013440 ms…
Result verification failed at elemnt 0

从上面测试数据,我们可以分析得到下面几条认识:
(1)最后一个配置(256,8)不可行,block中总共的thread数目超过了1024,这是GPU的硬件限制。
(2)最好的结果是第四个block配置(128,2)。
(3)第一个启动了最多的block,但不是最快的。
(4)因为第二个与第四个在一个block中拥有相同数目的thread,本应猜测二者有相同的表现,但是实际却是第二个略逊色,所以blockDim.x的大小 是很关键的。
(5)剩下的相对第四个都有较少的block数目,所以并行规模也是影响性能的关键因素。
现在,我们又有疑惑了,拥有block最少的应该会有一个最低的achieved Occupancy吧?而拥有最多block的应该会达到最高的achieved Occupancy吧?为了验证这些想法,我们再看一组测试数据:

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 64 2
–sumMatrix2DOnGPU<<<(256,8192),(64,2)>>> elapsed 37.495487 ms…
–Achieved Occupancy: 0.555373

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 64 4
–sumMatrix2DOnGPU<<<(256,4096),(64,4)>>> elapsed 38.886177 ms…
–Achieved Occupancy: 0.795769

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 64 8
–sumMatrix2DOnGPU<<<(256,2048),(64,8)>>> elapsed 40.603359 ms…
–Achieved Occupancy: 0.757109

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 128 2
–sumMatrix2DOnGPU<<<(128,8192),(128,2)>>> elapsed 36.666466 ms…
–Achieved Occupancy: 0.803921

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 128 4
–sumMatrix2DOnGPU<<<(128,4096),(128,4)>>> elapsed 38.689377 ms…
–Achieved Occupancy: 0.746745

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 128 8 –sumMatrix2DOnGPU<<<(128,2048),(128,8)>>> elapsed 50.706112 ms…
–Achieved Occupancy: 0.561505

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 256 2 –sumMatrix2DOnGPU<<<(64,8192),(256,2)>>> elapsed 36.828159 ms…
–Achieved Occupancy: 0.762112

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 256 4 –sumMatrix2DOnGPU<<<(64,4096),(256,4)>>> elapsed 42.040642 ms…
–Achieved Occupancy: 0.589849

$ nvprof
–metrics achieved_occupancy ./sumMatrix2D 256 8
–sumMatrix2DOnGPU<<<(64,2048),(256,8)>>> elapsed 0.015296 ms…
Result verification failed at elemnt 0 No events/metrics were profiled.
======== Error: Application returned non-zero code 1

通过上面测试数据对比分析:
(1)第一个(64,2)的achieved Occupancy竟然是最低的,尽管他有最多的block,它达到了硬件对block数量的限制。
(2)第四个(128,2)和第七个(256,2)拥有拥有不错的achieved Occupancy。
如果我们对这两个再做一个试验,再次增大,将blockDim.y设置为1,这也减少了block的大小。

$ ./sumMatrix2D 128 1
–sumMatrix2DOnGPU<<<(128,16384),(128,1)>>> elapsed 32.535934 ms…

$ ./sumMatrix2D 256 1
–sumMatrix2DOnGPU<<<(64,16384),(256,1)>>> elapsed 30.843328 ms…

这次测试有了更高的性能提升,并且(256,1)配置比(128,1)配置更好,再次查询(256,1)block配置的achieved Occupancy,load throughput和load efficiency等参数:

$ nvprof –metrics achieved_occupancy ./sumMatrix2D 256 1
–sumMatrix2DOnGPU<<<(64,16384),(256,1)>>>
Achieved Occupancy: 0.807456

$ nvprof –metrics gld_throughput ./sumMatrix2D 256 1
–sumMatrix2DOnGPU<<<(64,16384),(256,1)>>>
Global Load Throughput: 69.512GB/s

$ nvprof –metrics gld_efficiency ./sumMatrix2D 256 1
–sumMatrix2DOnGPU<<<(64,16384),(256,1)>>>
Global Memory Load Efficiency:100.21%

现在可以看出,最佳的block配置既不是拥有最高achieved Occupancy也不是最高load throughput的。所以不存在唯一metric参数来优化计算性能,我们需要从众多metric中寻求一个平衡。

5、总结

在大多数情形下,并不存在唯一的metric可以精确的优化性能。
(1)哪个metric或者event对性能的影响大小是由kernel具体的代码决定的;
(2)根据需要根据实际情况在众多相关的metric参数和event中寻求一个平衡;
(3)Grid/blcok heuristics(启发) 为调节性能提供了不错的切入点。

本文测试代码

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h> 
#include <math.h> 
#include <time.h> 
#include <memory>

#define PRECISION 1e-5 
#define HANDLE_ERROR(err) (HandleError( err, __FILE__, __LINE__ )) 

static void HandleError(cudaError_t err, const char *file, int line)
{
	if (err != cudaSuccess)
	{
		printf("%s in %s at line %d\n", cudaGetErrorString(err),
			file, line);
		exit(EXIT_FAILURE);
	}
}

void sumMatrix2DOnHost(float *h_A, float *h_B, float *hostRef, int nx, int ny)
{
	for (int i = 0; i< nx*ny; i++)
		hostRef[i] = h_A[i] + h_B[i];
}


__global__ void sumMatrix2DKernel(float *d_MatA, float *d_MatB, float *d_MatC, int nx, int ny)
{
	int idx = threadIdx.x + blockDim.x * blockIdx.x;
	int idy = threadIdx.y + blockDim.y * blockIdx.y;
	int tid = nx*idy + idx;

	if (idx < nx && idy < ny)
		d_MatC[tid] = d_MatA[tid] + d_MatB[tid];
}

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

	// set up device 
	int devID = 0; cudaDeviceProp deviceProp; 
	HANDLE_ERROR(cudaGetDeviceProperties(&deviceProp, devID)); 

	//
	printf("Using Device %d: %s\n", devID, deviceProp.name); 
	HANDLE_ERROR(cudaSetDevice(devID)); 

	// set up date size of matrix 
	int nx = 1<<14; 
	int ny = 1<<14; 
	int nxy = nx*ny;
	int nBytes = nxy * sizeof(float); 
	//
	printf("Matrix size: nx= %d, ny= %d\n",nx, ny); 

	// malloc host memory 
	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 
	for(int i=0;i<nxy;i++) 
	{ 
		h_A[i] = rand()/(float)RAND_MAX; 
		h_B[i] = rand()/(float)RAND_MAX; 
	} 
	memset(hostRef, 0, nBytes); memset(gpuRef, 0, nBytes); 
	
	// add matrix at host side for result checks 
	float iElaps; 
	clock_t iStart,iEnd; 
	iStart = clock();
	
	// time counter
	sumMatrix2DOnHost(h_A, h_B, hostRef, nx,ny); 
	iEnd = clock(); 
	//
	iElaps = (double)(iEnd-iStart)/CLOCKS_PER_SEC; 
	
	// second 
	iElaps = (double)(iEnd-iStart)/1000; 
	
	// ms 
	printf("--sumMatrix2DOnHost() elapsed %f ms..\n", iElaps); 
	
	// malloc device global memory 
	float *d_MatA, *d_MatB, *d_MatC; 
	cudaMalloc((void **)&d_MatA, nBytes); 
	cudaMalloc((void **)&d_MatB, nBytes); 
	cudaMalloc((void **)&d_MatC, nBytes);
	
	// transfer data from host to device 
	cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice); 
	cudaMemcpy(d_MatB, h_B, nBytes, cudaMemcpyHostToDevice); 
	
	 
	// invoke kernel at host side 
	int dimx = 32; 
	//int dimx = 16; 
	int dimy = 32; 
	//int dimy = 16;  

	if (argc > 2) 
	//配置block的维度 
	{ 
		dimx = atoi(argv[1]);
		dimy = atoi(argv[2]); 
	} 

	dim3 block(dimx, dimy); 
	dim3 grid((nx+block.x-1)/block.x, (ny+block.y-1)/block.y); 
	// calculate run time on GPU 

	float elapsedTime; 
	cudaEvent_t start, stop;
	HANDLE_ERROR(cudaEventCreate(&start)); 
	HANDLE_ERROR(cudaEventCreate(&stop)); 
	HANDLE_ERROR(cudaEventRecord(start, 0)); 

	sumMatrix2DKernel <<< grid, block >>>(d_MatA, d_MatB, d_MatC, nx, ny); 

	cudaDeviceSynchronize(); 
	HANDLE_ERROR(cudaEventRecord(stop, 0));
	HANDLE_ERROR(cudaEventSynchronize(stop));
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop)); 

	printf("--sumMatrix2DOnGPU<<<(%d,%d),(%d,%d)>>> elapsed %f ms..\n", grid.x, grid.y, block.x, block.y, elapsedTime); 
	// 
	// copy kernel result back to host side 
	cudaMemcpy(gpuRef, d_MatC, nBytes, cudaMemcpyDeviceToHost); 
	// check device results 
	for(int i=0; i< nxy; i++) 
	{ 
		if(fabs(gpuRef[i]-hostRef[i]) > PRECISION) 
		{ 
			fprintf(stderr,"Result verification failed at elemnt %d\n", i);
			exit(EXIT_FAILURE); 
		} 
	} 

	// free device global memory 
	cudaFree(d_MatA); 
	cudaFree(d_MatB); 
	cudaFree(d_MatC); 

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

	// reset device 
	cudaDeviceReset(); 
	//
	printf("Test Passed..\n"); return 0;

}

其他值得借鉴的文章
TX2入门(8)——优化/性能查看工具nvprof(持续补充……)
cuda nvprof 输出结果的理解和优化空间
CUDA Program Analysis

版权声明:本文内容由互联网用户自发贡献,该文观点仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 举报,一经查实,本站将立刻删除。

发布者:全栈程序员-用户IM,转载请注明出处:https://javaforall.cn/132182.html原文链接:https://javaforall.cn

【正版授权,激活自己账号】: Jetbrains全家桶Ide使用,1年售后保障,每天仅需1毛

【官方授权 正版激活】: 官方授权 正版激活 支持Jetbrains家族下所有IDE 使用个人JB账号...

(0)


相关推荐

  • mysql varchar列转成integer然后获取最大值。[通俗易懂]

    mysql varchar列转成integer然后获取最大值。[通俗易懂]https://blog.csdn.net/c_henjinxing521/article/details/51788963上面的大神写的办法可以。selectMAX(CAST(userNoasSIGNEDINTEGER))fromuserInfo;或者selectMAX(CAST(userNoasUNSIGNEDINTEGER))fromus…

  • mysql基本sql语句大全(基础用语篇)_mysql常用查询语句

    mysql基本sql语句大全(基础用语篇)_mysql常用查询语句MySQL常用语句大全一、连接MySQL格式:mysql-h主机地址-u用户名-p用户密码1、例1:连接到本机上的MYSQL。首先在打开DOS窗口,然后进入目录mysqlbin,再键入命令mysql-uroot-p,回车后提示你输密码,如果刚安装好MYSQL,超级用户root是没有密码的,故直接回车即可进入到MYSQL中了,MYSQL的提示符是:m…

  • 等价类划分法设计用例(超详细)「建议收藏」

    等价类划分法设计用例(超详细)「建议收藏」等价类划分法等价类:1、解决了不能穷举测试的问题、控制成本、控制测试用例数量2、数据值要明确,对文字敏感3、依据需求将输入划分为若干个等价类,划分等价类(需求、数据特征)等价类设计用例的难点:如何根据时间成本划分等价类等价类分为:           1、有效等价类           2、无效等价类如上图可以划分为:                 有效等价类1:[-99,99]                 无效等价类2:<-99                 无效等

    2022年10月18日
  • mysql查询记录总数_MySQL记录总条数实现查询优化「建议收藏」

    mysql查询记录总数_MySQL记录总条数实现查询优化「建议收藏」MySQL记录总条数实现查询优化发布时间:2020-05-0811:04:00来源:亿速云阅读:268作者:三月本文主要给大家介绍MySQL记录总条数实现查询优化,文章内容都是笔者用心摘选和编辑的,具有一定的针对性,对大家的参考意义还是比较大的,下面跟笔者一起了解下MySQL记录总条数实现查询优化吧。1、COUNT(*)和COUNT(COL)COUNT(*)通常是对主键进行索引扫描,而COUNT…

  • MySql必知必会实战练习(四)主键、外键、sql约束、联结表

    本博将对主键、外键、MySql数据库约束和联结表的相关特性进行总结和实战1.主键表中的每一行都应该具有可以唯一标识自己的一列(或一组列),而这个承担标识作用的列称为主键如果没有主键,数据的管理

    2021年12月29日

发表回复

您的电子邮箱地址不会被公开。

关注全栈程序员社区公众号