【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);

  //Malloc
  float* 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);

  //cudaMalloc
  float *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 grid
  dim3 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}$ 的方矩阵）。

用 nvprof 检测活跃的线程束

完整内容https://face2ai.com/CUDA-F-3-3-并行性表现/