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

最新推荐文章于 2024-08-23 11:14:15 发布

Innerpeace_yu

最新推荐文章于 2024-08-23 11:14:15 发布

阅读量1.2w

点赞数 2

cuda并行编程专栏收录该内容

15 篇文章 4 订阅

订阅专栏

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

2017-01-13 11:41:28| 分类： HPC&CUDA优化 | 标签：hpc gpu cuda |举报 |字号订阅

下载LOFTER 我的照片书 |

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 (R) Cuda command line profiler
Copyright (c) 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

$ 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

$ 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 <stdio.h>
#include <math.h>
#include <time.h>

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

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;
	int dimx, dimy;
	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;
}