CUDA: 优化实践,nvcc编译器选项利用

优化实践

利用nvcc编译器提供的选项，可以查看生成的PTX汇编代码，可以查看生成代码所用资源信息等，对进一步优化kernel提供一定帮助指引

---

源码

简单的向量相加源码:$c=a+b$

	//file gpu_add.cu
	
	if 0
	#define RESTRICT __restrict__
	#else
	#define RESTRICT
	#endif
	
	using Int=int;
	//using Int=uint;
	
	__global__ void add(Int n,
	                    const float * RESTRICT a,
	                    const float * RESTRICT b,
	                    float * RESTRICT c)
	{
	        Int tid=threadIdx.x+blockIdx.x*blockDim.x;
	        Int nthread=blockDim.x*gridDim.x;
	#pragma unroll 4
	        for(Int x=tid; x<n; x+=nthread)
	        {
	            c[x]=a[x]+b[x];
	        }
	        return;
	}
	
	
	void run(Int n, float *d_a, float *d_b, float *d_c)
	{
	   dim3 grid(128);
	   dim3 block(128);
	
	   add<<<grid,block>>>(n,d_a,d_b,d_c);
	}

---

细节

• 编译生成PTX

  nvcc -O3 -c -arch=sm_86 -ptx gpu_add.cu
  

  或生成更为详细信息

  nvcc -O3 -c -arch=sm_86  -lineinfo --source-in-ptx -ptx gpu_add.cu
  

• 生成目标文件,报告所用资源情况

  nvcc -O3 -c -arch=sm_86 -Xptxas -v gpu_add.cu
  

  对于Int=int和Int=uint, 所用寄存器数目不同，int是22个,uint是16个

  >>>>>>>>>>>>>>>>using Int=int
  
  ptxas info    : 0 bytes gmem
  ptxas info    : Compiling entry function '_Z3addiPKfS0_Pf' for 'sm_86'
  ptxas info    : Function properties for _Z3addiPKfS0_Pf
      0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
  ptxas info    : Used 22 registers, 384 bytes cmem[0]
  
  >>>>>>>>>>>>>>>>using Int=uint;
  
  ptxas info    : 0 bytes gmem
  ptxas info    : Compiling entry function '_Z3addjPKfS0_Pf' for 'sm_86'
  ptxas info    : Function properties for _Z3addjPKfS0_Pf
      0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
  ptxas info    : Used 16 registers, 384 bytes cmem[0]
  

• __restrict__关键字会提示nvcc加载数据不经过L2&L1(看PTX汇编可观察到)。在能加上const和__restrict__的地方应该使用这些限定字，以提示给编译器更多的优化可能

• 对Int=int, pragma unroll语句有无，循环都会被展开;对Int=uint,必须指定pragma unroll语句,循环才会显示展开. 建议用int类型做下标索引类型，以便给编译器更多的优化可能

---

文档资料

• NVCC Doc