CUDA中各种内存及线程布局的性能实验分析（1）

cuda中提供了各种各样的的内存形式，不同的内存对程序的性能影响也不同。在编程中，根据我们的需要和数据特点，利用适当的内存，将会对程序的性能产生很大的影响。本文中，通过一个最简单的图像卷积操作进行编程分析cuda中的内存对性能的影响。
首先介绍一下图像卷积：
简单的来说，用一个小矩阵在大矩阵上进行由左向右，由上到下滑动，每滑动一次，小矩阵和大矩阵上对应的窗口进行对应位置的乘加操作得到一个数，作为结果矩阵上的一个对应位置上的值，一般来说滑动的步长为1，如下所示(图片来自ufldl)：

不同内存性能对比：

我们默认你已经安装好cuda环境。在写cuda核之前，首先可以查看一下本机的gpu基本情况，可以使用如下代码来查看：

#include <stdio.h>#include <cuda_runtime.h>#define CUDA_CHECK(condition,info) \if(condition!=cudaSuccess) \{ \    printf(info); \}int main(){    int count;    CUDA_CHECK(cudaGetDeviceCount(&count),        "cuda:get count failed!");    printf("gpu count in host:%d\n",count);    for(int i=0;i<count;++i)    {        cudaDeviceProp device_prop;        CUDA_CHECK(cudaGetDeviceProperties(&device_prop,i),            "cuda:get property failed");        printf("----------\ndevice %d:\n",i);        printf("Name:%s\n",device_prop.name);        printf("be_integrated:%d\n",            device_prop.integrated);        printf("Compute_Capability:%d.%d\n",            device_prop.major,device_prop.minor);        printf("Global_Memory:%ld MB\n",            device_prop.totalGlobalMem/1024/1024);        printf("Shared_Memory(per_block):%ld KB\n",            device_prop.sharedMemPerBlock/1024);        printf("Const_Memory:%ld KB\n",            device_prop.totalConstMem/2014);        printf("Reg_Count(per_block):%d\n",            device_prop.regsPerBlock);        printf("Warp_Count(per_block):%d\n",            device_prop.warpSize);        printf("Thread_Count(per_block):%d\n",            device_prop.maxThreadsPerBlock);        printf("can_overlap:%d\n",            device_prop.deviceOverlap);        printf("can_maphostmem:%d\n",            device_prop.canMapHostMemory);        printf("Process_Count:%d\n",            device_prop.multiProcessorCount);    }}

用nvcc编译并执行，我的执行结果为：

device 0:
Name:GeForce GTX TITAN X
be_integrated:0
Compute_Capability:5.2
Global_Memory:12204 MB
Shared_Memory(per_block):48 KB
Const_Memory:32 KB
Reg_Count(per_block):65536
Warp_Count(per_block):32
Thread_Count(per_block):1024
can_overlap:1
can_maphostmem:1
Process_Count:24

上面的数字基本都是我们需要关注的，尤其是带memory后缀的。
下面开始上程序：

__global__ void Conv(float *dev_in_data,float *dev_k_data,float *dev_out_data,int out_w,int out_h,int kw,int kh,int in_w){    int tx=threadIdx.x+blockIdx.x*blockDim.x;    int ty=threadIdx.y+blockIdx.y*blockDim.y;    float tmp_sum=0;    while(ty<out_h)    {        while(tx<out_w)        {            for(int j=ty;j<ty+kh;++j)            {                for(int i=tx;i<tx+kw;++i)                {                    tmp_sum+=dev_k_data[(j-ty)*kw+(i-tx)]*                        dev_in_data[j*in_w+i];                }            }            dev_out_data[ty*out_w+tx]=tmp_sum;            tx+=blockDim.x*gridDim.x;        }        ty+=blockDim.y*gridDim.y;    }}

上面这个是我们的核函数,main函数是这样的：

int main(){    int in_w=5000,in_h=5000;//**txg::change here**//    int kw=5,kh=5;//**txg::change here**//    int out_w,out_h;    float *in_data,*k_data,*out_data,*mapped_in_data;    //init data    if(in_data=(float*)malloc(sizeof(float)*in_w*in_h))    {        memset(in_data,0,sizeof(float)*in_w*in_h);        for(int i=0;i<in_h;++i)        {            for(int j=0;j<in_w;++j)            {                in_data[i*in_w+j]=rand()%RAND_MAX;            }        }    }    else    {        printf("alloc in_data failed!!");    }    if(k_data=(float*)malloc(sizeof(float)*kw*kh))    {        memset(k_data,0,sizeof(float)*kw*kh);        for(int i=0;i<kh;++i)        {            for(int j=0;j<kw;++j)            {                k_data[i*kw+j]=rand()%RAND_MAX;            }        }    }    else    {        printf("alloc k_data failed!!");    }    out_w=in_w-kw+1;    out_h=in_h-kh+1;    if(!(out_data=(float*)malloc(sizeof(float)*out_w*out_h)))    {            printf("alloc out_data failed!!");    }    //init_end    //gpu_info    /*    device 0:    Name:GeForce GTX TITAN X    be_integrated:0    Compute_Capability:5.2    Global_Memory:12204 MB    Shared_Memory(per_block):48 KB    Const_Memory:32 KB    Reg_Count(per_block):65536    Warp_Count(per_block):32    Thread_Count(per_block):1024    can_overlap:1    can_maphostmem:1    Process_Count:24    */    float *dev_in_data,*dev_out_data,*dev_k_data;    CUDA_CHECK(cudaMalloc((void**)&dev_in_data,sizeof(float)*in_w*in_h),"alloc in_data falied!");    CUDA_CHECK(cudaMalloc((void**)&dev_out_data,sizeof(float)*out_w*out_h),"alloc out_data falied!");    CUDA_CHECK(cudaMalloc((void**)&dev_k_data,sizeof(float)*kw*kh),"alloc k_data falied!");    CUDA_CHECK(cudaMemcpy(dev_in_data,in_data,sizeof(float)*in_w*in_h,cudaMemcpyHostToDevice),"copy in_data to device failed!!");    CUDA_CHECK(cudaMemcpy(dev_k_data,k_data,sizeof(float)*kw*kh,cudaMemcpyHostToDevice),"copy k_data to device failed!!");    dim3 tg(16,16);//**txg::change here**//    dim3 bg(32,32);//**txg::change here**//    //event    cudaEvent_t start,stop;    CUDA_CHECK(cudaEventCreate(&start),"start create failed!!");    CUDA_CHECK(cudaEventCreate(&stop),"stop create failed!!");    CUDA_CHECK(cudaEventRecord(start,0),"start record failed!!");    //call kernels    Conv<<<bg,tg>>>(dev_in_data,dev_k_data,dev_out_data,        out_w,out_h,kw,kh,in_w);    CUDA_CHECK(cudaEventRecord(stop,0),"stop record failed!!");    CUDA_CHECK(cudaEventSynchronize(stop),"stop_syn failed!!");    float elapsed_time;    CUDA_CHECK(cudaEventElapsedTime(&elapsed_time,start,stop),"elapsed time calculate failed!!");    printf("elapsed time: %0.3f s\n",elapsed_time);    //return    CUDA_CHECK(cudaMemcpy(out_data,dev_out_data,sizeof(float)*out_w*out_h,cudaMemcpyDeviceToHost),"copy out_data to host failed!!");    CUDA_CHECK(cudaFree(dev_in_data),"free dev_in_data failed!!");    CUDA_CHECK(cudaFree(dev_out_data),"free dev_out_data failed!!");    CUDA_CHECK(cudaFree(dev_k_data),"free dev_k_data failed!!");    free(in_data);    free(out_data);    free(k_data);       }

里面基本步骤是先分配cpu和gpu上的内存，初始化，调用核函数，及计时操作，都是一般的代码。
你只需要加上刚才查信息的代码的头文件及宏，就可以跑啦。
跑一跑，我们会得到程序的运行时间，在这里我们通过改变这一行，可以改变输入数据的大小：

int in_w=5000,in_h=5000;//**txg::change here**//

我这里分别改成了50，500，5000三种，对应的输入数据大小就分别是2.5KB,250KB,25M,在这里为了看上去更直观，并没有把float的字节数计入内，大小关系是一样的。
分别运行，在我的机器上的结果为：

2.5KB 26ms
250KB 294ms
25MB 2932ms

在上面的程序中我们可以看到随着数据量的变化程序运行时间在上涨，而且基本上是每100倍数据量时间上涨10倍左右。

其实在我们上面的程序中，我们在各个数据前面都没有加任何的修饰符，此时其实在gpu上我们用的是全局内存（global memory），也是gpu中容量最大的内存，但也是读取速度最慢的内存，在gpu上共存在着以下几种内存：

全局内存：容量最大，可达几g,但读取速度最慢
常量内存：全局内存的一种缓存映射，能够令线程束以广播的方式读取
共享内存：距gpu核最近的内存，读取速度很快，能达到全局内存的10倍左右，但容量最小，按每个block算的
纹理内存：全局内存的另一种缓存映射，在某些方面比较有用。
寄存器：读取速度最快的存储单元，但是按数量算的，资源最有限。

各种内存分别对应信息的如下：

Global_Memory:12204 MB
Const_Memory:32 KB
Shared_Memory(per_block):48 KB
Reg_Count(per_block):65536

可以看到，我的共享内存大小只有48K，因此，我们的输入数据是不能以此种形式存放的，但是我们的卷积核大小只有25*4=100B,因此，我们可以以共享内存的方式存放，稍改一下核函数代码：

#define ksize 1000__global__ void Conv_shared(float *dev_in_data,float *dev_k_data,float *dev_out_data,int out_w,int out_h,int kw,int kh,int in_w){    __shared__ float k_cache[ksize];    int tx=threadIdx.x+blockIdx.x*blockDim.x;    int ty=threadIdx.y+blockIdx.y*blockDim.y;    if(tx==0&&ty==0)    {        for(int i=0;i<kw*kh;++i)        {            k_cache[i]=dev_k_data[i];        }    }    __syncthreads();    float tmp_sum=0;    while(ty<out_h)    {        while(tx<out_w)        {            for(int j=ty;j<ty+kh;++j)            {                for(int i=tx;i<tx+kw;++i)                {                    tmp_sum+=k_cache[(j-ty)*kw+(i-tx)]*dev_in_data[j*in_w+i];                }            }            dev_out_data[ty*out_w+tx]=tmp_sum;            tx+=blockDim.x*gridDim.x;        }        ty+=blockDim.y*gridDim.y;    }}

主要改动如下：

__shared__ float k_cache[ksize];int tx=threadIdx.x+blockIdx.x*blockDim.x;int ty=threadIdx.y+blockIdx.y*blockDim.y;if(tx==0&&ty==0){    for(int i=0;i<kw*kh;++i)    {        k_cache[i]=dev_k_data[i];    }}

此处我用k_cache来保存来自主机的卷积核，并代替之前的全局内存的卷积核，同理跑一下程序看一下运行时间，如下：

2.5KB 27ms
250KB 86ms
25MB 729ms

和全局内存比较一下：

运行条件：kernel:5*5,tg:16*16,bg:32*32
数据量 global_memory kernel_shared
2.5KB 26ms 27ms
250KB 294ms 86ms
25MB 2932ms 729ms

可以看到在小数据量的时候，并没有提升，但随着数据量增长，计算量越来越大，运行时间提升很是时显，最多可以提升4倍左右的时间。
为了再次比较，我们将核的大小增加5倍，再次运行，看效果：

运行条件：kernel:25*25,tg:16*16,bg:32*32
数据量 global_memory kernel_shared
2.5KB 77ms 115ms
250KB 3239ms 1536ms
25MB 30369ms 19966ms

对比全局内存，共享内存在时间上依然有提升。

分析：由于在卷积过程中卷积核会多次调用，而当卷积核在全局内存上时，每次调用是很花时间的，但是当我们在开始时将卷积核拷到我们的共享内存上，那么在以后的调用中就不用去全局内存上找了，只需从共享内存上去读取，此时速度会有很大提升。
注：由于常量内存的大小也很小，而且在高版本的计算功能集上其速度和全局内存差别有限，因此在这里不做比较。纹理内存因其特殊用途以及现在已不被人常用，也不做比较。

线程布局性能分析：

为了对比每个block中线程数及布局的影响，我对线程布局进行了一些改动，并做了一些实验，要改变线程布局，我们只需改动以下几行即可：

dim3 tg(16,16);//**txg::change here**//

我将上述格局分别改成12*16，32*8，进行实验：
结果如下：

运行条件：kernel:5*5,tg:12*16,bg:32*32
数据量 global_memory kernel_shared k*5_global k*5_shared
2.5KB 26ms 27ms 77ms 116ms
250KB 324ms 101ms 3800ms 1864ms
25MB 3101ms 944ms 35095ms 24445ms

运行条件：kernel:5*5,tg:32*8,bg:32*32
数据量 global_memory kernel_shared k*5_global k*5_shared
2.5KB 26ms 28ms 77ms 115ms
250KB 161ms 54ms 1916ms 766ms
25MB 1404ms 367ms 14594ms 8969ms

为了明显对比，我们对kernel_shared的数据进行分析：

可以看到在数据量小时依然不明显，但当数据量增大时，线程数在256时效果要比192好，当线程数同为256时，32*8的格局比16*16的格局性能提升了一倍左右。
这种提升其实来自于访存合并策略带来的收益，为了进一步弄清线程格局的影响，我对线程格局进行了进一步分析实验。

为了在尽可能多的格局上进行实验，且限于最大线程数和影响（可看上面的gpu信息），我将实验稍作了改变：

kernel:2*2,tg:–,bg:16*64

在此基础上，我将tg设置成16*32，32*16，64*8，128*4，256*2并进行实验，在kernel_shared下结果如下：

运行条件：kernel:2*2,tg:–,bg:32*32
数据量 16*32 32*16 64*8 128*4 256*2
2.5KB 27ms 27ms 27ms 27ms 27ms
250KB 41ms 43ms 40ms 43ms 55ms
25MB 328ms 203ms 170ms 187ms 326ms

并作图展示：

依然关注数据量比较大的情况，可以看到运行时间呈先减少后增加的趋势，并且在64*8上运行时间最少。

分析：其实在cuda上线程去读取内存时，是以线程束的方式去读取的，一般来说一个线程束有32个线程，如果这32个线程读取的数据具有连续性，则cuda会合并其访问，即只发起一次内存访问，在计算功能集比较高的平台上，由于有缓存的存在，这样也会使得缓存命中率提高，提高访存效率。但访存合并仅限于同一个block中，所以当我们改变block中的线程格局时，当达到32*16和64*8等时会比16*32访存效率高，从而提高程序性能，但由于缓存大小的影响，访存合并的最大字节有限，所以当线程格局太多时反而会影响程序性能。

既然线程布局对程序的性能有影响，那线程块的布局有没有影响呢，我也做了一个简单的实验，为了改变线程块的布局，我们只需改变下面的这行代码即可：

dim3 bg(32,32);//**txg::change here**//

我将其改为32*16，64*8，128*4，256*2并在kernel_shared下运行，并只在25M的数据下进行实验，结果如下

运行条件：kernel:2*2,tg:256*2,bg:–
数据量 32*16 64*8 128*4 256*2
25MB 325ms 206ms 171ms 205ms

由此看到，block的格局也会对程序性能产生影响，但总体来说影响不是很大。

总结一下：我们在编写cuda程序的时候要尽可能的合理充分的利用共享内存，并且在线程布局上要尽可能让多个线程能访问连续的数据，这样能提高访存的效率。

在我们上面的实验中，我只实验了卷积核在共享内存上存放，对于输入数据和输出数据并没有借助共享内存，由于输入输出数据可能比较大，我们的共享内存放不下，但也不是没有办法利用，在后面我们将利用并行归约的思想将输入输出借助共享内存实现，从而进一步提高程序的性能。