【CUDA 基础】3.3 并行性表现

本文主要是介绍【CUDA 基础】3.3 并行性表现，希望对大家解决编程问题提供一定的参考价值，需要的开发者们随着小编来一起学习吧！

title: 【CUDA 基础】3.3 并行性表现
categories:
- CUDA
- Freshman
tags:
- nvprof
toc: true
date: 2018-04-15 21:17:52

Abstract: 本文主要通过nvprof工具来分析核函数的执行效率（资源利用率）
Keywords: nvprof

开篇废话

继续更新CUDA，前面为了加速概率论的学习停了一段CUDA，从今天开始继续CUDA和数学分析的更新，每一篇都写一点废话就相当于自己的日记了，之前很佩服那些写日记的人，因为根本不知道日记可以写些什么，但是现在看看，如果写一些文字记录自己，首先可以反思当下，其次是过一段时间以后可以看看自己到底有没有进步，这些都是有用的，所以大家可以略过我的废话，直接看正文。

本文的主要内容就是进一步理解线程束在硬件上执行的本质过程，结合上几篇关于执行模型的学习，本文相对简单，通过修改核函数的配置，来观察核函数的执行速度，以及分析硬件利用数据，分析性能，调整核函数配置是CUDA开发人员必须掌握的技能，本篇只研究对核函数的配置是如何影响效率的（也就是通过网格，块的配置来获得不同的执行效率。）
本文全文只用到下面的核函数

__global__ void sumMatrix(float * MatA,float * MatB,float * MatC,int nx,int ny)
{int ix=threadIdx.x+blockDim.x*blockIdx.x;int iy=threadIdx.y+blockDim.y*blockIdx.y;int idx=ix+iy*ny;if (ix<nx && iy<ny){MatC[idx]=MatA[idx]+MatB[idx];}
}

没有任何优化的最简单的二维矩阵加法。
全部代码：

int main(int argc,char** argv)
{//printf("strating...\n");//initDevice(0);int nx=1<<13;int ny=1<<13;int nxy=nx*ny;int nBytes=nxy*sizeof(float);//Mallocfloat* A_host=(float*)malloc(nBytes);float* B_host=(float*)malloc(nBytes);float* C_host=(float*)malloc(nBytes);float* C_from_gpu=(float*)malloc(nBytes);initialData(A_host,nxy);initialData(B_host,nxy);//cudaMallocfloat *A_dev=NULL;float *B_dev=NULL;float *C_dev=NULL;CHECK(cudaMalloc((void**)&A_dev,nBytes));CHECK(cudaMalloc((void**)&B_dev,nBytes));CHECK(cudaMalloc((void**)&C_dev,nBytes));CHECK(cudaMemcpy(A_dev,A_host,nBytes,cudaMemcpyHostToDevice));CHECK(cudaMemcpy(B_dev,B_host,nBytes,cudaMemcpyHostToDevice));int dimx=argc>2?atoi(argv[1]):32;int dimy=argc>2?atoi(argv[2]):32;double iStart,iElaps;// 2d block and 2d griddim3 block(dimx,dimy);dim3 grid((nx-1)/block.x+1,(ny-1)/block.y+1);iStart=cpuSecond();sumMatrix<<<grid,block>>>(A_dev,B_dev,C_dev,nx,ny);CHECK(cudaDeviceSynchronize());iElaps=cpuSecond()-iStart;printf("GPU Execution configuration<<<(%d,%d),(%d,%d)|%f sec\n",grid.x,grid.y,block.x,block.y,iElaps);CHECK(cudaMemcpy(C_from_gpu,C_dev,nBytes,cudaMemcpyDeviceToHost));cudaFree(A_dev);cudaFree(B_dev);cudaFree(C_dev);free(A_host);free(B_host);free(C_host);free(C_from_gpu);cudaDeviceReset();return 0;
}

可见我们用两个 $8192\times 8192$ 的矩阵相加来测试我们效率。
注意一下这里的GPU内存，一个矩阵是 $2^{13}\times 2^{13}\times 2^2=2^{28}$ 字节也就是 256M，三个矩阵就是 768M 因为我们的GPU内存就是 2G 的，所以我们没办法进行更大的矩阵计算了（无法使用原文使用的是 $2^{14}$ 的方矩阵）。