CUDA优化实例（一）占用率

CUDA优化实例（一）占用率

• 前言
• 实验
• 分析与结论

---

前言

占用率是指活跃的线程占总线程的比率，占用率越高，kernel效率越高。有较多的线程同时参与运行就会有较多的活跃的，就会有较高的占用率。所以怎么才能同时具有较多的同时参与运行的线程呢，这跟kernel所需要的资源量和设备的计算能力以及块的组织（能影响块在SM的分布）有关。我们且保证kernel不需要任何资源，我的电脑计算能力大于3即一个SM上最多可放下16个块，此时，占用率就只和块的大小有关了。下面我就用实例来看块是怎么影响占用率的。

实验

实验只需一个文件blocksize.cu：

/*
*       This program aims to speed up the kernel by adjusting the blocksize 
*   which could change the occupancy of the SM
*
*
*
*   time:3,5,2018
*       author:jieeeeeeee
*/



#include <iostream>
#include <stdlib.h>
#include <stdio.h>
#define CHECK(res) { if(res != cudaSuccess){printf("Error ：%s:%d , ", __FILE__,__LINE__);   \
printf("code : %d , reason : %s \n", res,cudaGetErrorString(res));exit(-1);}}




__global__ void kernel(int *d_in,int *d_out)
{
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;   
    d_out[idx] = d_in[idx];     
}




int main(int argc, char **argv)
{
    if(argc == 2)
        std::cout<<"the 1 dimension blocksize is :  "<<argv[1]<<std::endl;
    else{
        std::cout<<"argc error !"<<std::endl;
        return 1;
    }

    unsigned int nSize = atoi(argv[1]);


    int N = 1<<14;//16KB

    int *a_in,*a_out;

    a_in = new int[N];
    a_out = new int[N];

    int *d_in,*d_out;

    cudaMalloc((int**)&d_in, N*sizeof(int));
    cudaMalloc((int**)&d_out, N*sizeof(int));

    //initialize the a_in
    for(int i = 0;i<N;i++){
        //generate a random integer in (0,100)
        int a = rand()%100;
        a_in[i] = a;
    }

    cudaMemcpy(d_in,a_in,N*sizeof(int),cudaMemcpyHostToDevice);

    unsigned int blockSize = nSize; 
    unsigned int gridSize = N/nSize;


    kernel<<<gridSize,blockSize>>>(d_in,d_out);
    CHECK(cudaDeviceSynchronize());

    cudaMemcpy(a_out,d_out,N*sizeof(int),cudaMemcpyDeviceToHost);

  // Check for errors 
  for (int i = 0; i < N; i++)
    if(a_out[i]!=a_in[i]) std::cout << "Checked error" << std::endl;



  cudaFree(a_in);
  cudaFree(a_out);
    return 0; 

}

编译：nvcc --arch=sm_50 blocksize.cu
查看性能：nvprof --metrics achieved_occupancy./a.out xx xx是块的大小
实验结果：

块大小	消耗时间/us	占用率/%
1	92	27
2	47	27
4	25	28
8	13	29
16	8	31
32	5	35
64	3.4	67
128	3.0	79
256	3.0	78
512	3.1	81
1024	3.1	84

分析与结论

• 1.块的大小从1-32，占用率相差不多，耗时成线性下降的原因是：因为kernel不消耗资源，所以一个SM有最大的16个块同时运行，但由于块大小小于32每个块只有一个线程束，也就是说每个SM只有16个线程束在运行，一个块最多的线程束64个，所以占有率只有接近25%，一个线程数的运行时间相同但运行的线程成线性增加，所以耗时成线性降低。
• 2.大于32小于128占有率成线性上升是因为，一个SM最多占16个块，当一个块有128个线程才能达到SM最大的线程数2048，所以在128之前，SM中的线程都是不满的，更达不到80%的实际占满时的占用率。
• 3.大于128后，占用率成不在怎么变化，耗时也不再怎么变化，因为资源用尽，所以每个SM中的线程都被充分启动，所以都接近最大占用率。
• 4.占用率最大一般是80%，因为程束调度器调度活跃的线程束，而调度器本身也是有延迟的（有速度消耗的）所以，占用率在计算时会把这部分延迟算进去，所以占用率一般不会达到90%。