【Cuda并行编程之二】Cuda Memory Hierarchy_Cuda内存层次结构

要想编写高效的程序，那么一定要对内存结构有比较深刻的认识，就像C/C++里面的堆内存，栈内存，全局存储区，静态存储区，常量区等。Cuda是并行计算框架，而GPU的内存有限，那么如果想编写高效的Cuda程序，首先要对其内存结构有一个简单的认识。

首先我们先上一张图，然后通过解释一些名词和代码来进行解释。

各种存储器比较：

存储器	位置	拥有缓存	访问权限	变量生存周期
register	GPU片内	N/A	device可读/写	与thread相同
local memory	板载显存	无	device可读/写	与thread相同
shared memory	GPU片内	N/A	device可读/写	与block相同
constant memory	板载显存	有	device可读，host可读写	可在程序中保持
texture memory	板载显存	有	device可读，host可读写	可在程序中保持
global memory	板载显存	无	device可读写，host可读写	可在程序中保持
host memory	主机内存	无	host可读写	可在程序中保持
pinned memory	主机内存	无	host可读写	可在程序中保持

registers：寄存器。它是GPU片上告诉缓存器，执行单元可以以极低的延迟访问寄存器。寄存器的基本单元是寄存器文件（register file），每个寄存器文件大小为32bit。寄存器文件数量虽然客观，但是平均分给并行执行的线程，每个线程拥有的数量就非常有限了。编程时，不要为每个线程分配过多的私有变量。下面程序中，aBegin,aEnd,aStep,a等变量都是寄存器变量，每个线程都会维护这些变量。

__global__ void registerDemo(float *B,float *A ,int wA)
{
    int aBegin = wA*BLOCK_SIZE * blockIdx.y;
    int aEnd = aBegin + wA - 1;
    int aStep = BLOCK_SIZE;

    for(int a=aBegin;a<=aEnd;a+=aStep)
    {
        //...
    }
}

local memory：局部存储器。对于每个线程，局部存储器也是私有的。如果寄存器被消耗完，数据将被存储在局部存储器中。如果每个线程用了过多的寄存器，或声明了大型结构体或数组，或者编译期无法确定数组的大小，线程的私有数据就有可能被分配到local memory中。一个线程的输入和中间变量将被保存在寄存器或者局部存储器中。局部存储器中的数据将被保存在显存中，而不是片上的寄存器或者缓存中，因此对local memory的访问速度比较慢。

shared memory：共享存储器。共享存储器也是GPU片内的告诉存储器。它是一个块可以被同一block中的所有线程访问的可读写存储器。访问共享存储器的速度几乎和访问寄存器一样快。是实现线程间通信的延迟最小的方法。共享存储器可用于实现多种功能，如用于保存共用的计数器（例如计算循环迭代次数）或者block的公共结果（例如规约的结果）。

static variable使用shared memory：

#include<iostream>
#include<stdio.h>

#if 1
__global__ void example(float *u)
{
    int i = threadIdx.x;
    <strong>__shared__ int tmp[4];</strong>
    tmp[i] = u[i];
    u[i] = tmp[i] * tmp[i] + tmp[3-i] ;
}
#endif

#if 1

int main()
{
    float host_u[4] = {1,2,3,4};
    float * dev_u ;
    size_t size = 4*sizeof(float);

    cudaMalloc(&dev_u , size);
    cudaMemcpy(dev_u,host_u,size,cudaMemcpyHostToDevice);

    example<<<1,4>>> (dev_u);

    cudaMemcpy(host_u , dev_u , size , cudaMemcpyDeviceToHost);

    cudaFree(dev_u);

    for(int i=0;i<4;i++)
        printf("%f
",host_u[i]);
    return 0;
}

#endif

dynamic variable使用shared memory：

#include<iostream>
#include<stdio.h>

<strong>extern __shared__ int tmp[];</strong>

__global__ void example(float *u)
{
    int i = threadIdx.x;
    tmp[i] = u[i];
    u[i] = tmp[i] * tmp[i] + tmp[3-i];
}

int main()
{
    float host_u[4] = {1,2,3,4};
    float * dev_u;
    <strong>size_t size = 4*sizeof(float);</strong>

    cudaMalloc(&dev_u,size);
    cudaMemcpy(dev_u , host_u ,size , cudaMemcpyHostToDevice);
    example<<<1,4,<strong>size</strong>>>>(dev_u);

    cudaMemcpy(host_u, dev_u,size,cudaMemcpyDeviceToHost);
    cudaFree(dev_u);
    for(int i=0;i<4;i++)
        printf("%f ",host_u[i]);
    return 0;
}

global memory：全局存储器位于显存（占据了显存的绝大部分），CPU、GPU都可以进行读写访问。整个网格中的任意线程都能读写全局存储器的任意位置由于全局存储器是可写的。全局存储器能够提供很高的带宽，但同时也具有较高的访存延迟。显存中的全局存储器也称为线性内存。线性内存通常使用cudaMalloc()函数分配，cudaFree()函数释放，并由cudaMemcpy()进行主机端与设备端的数据传输。

此外，也可以使用__device__关键字定义的变量分配全局存储器，这个变量应该在所有函数外定义，必须对使用这个变量的host端和device端函数都可见才能成功编译。在定义__device__变量的同时可以对其赋值。

static variable使用global memory：

global_mem_static.cu:

#include<stdio.h>
#include<iostream>

__device__ float devU[4];
__device__ float devV[4];

//__global__ function
__global__ void addUV()
{
    int i = threadIdx.x;
    devU[i] += devV[i];
}

int main()
{
    float hostU[4] = {1,2,3,4};
    float hostV[4] = {5,6,7,8};

    int size = 4* sizeof(float);

    //cudaMemcpyToSymbol:将数据复制到__constant__或者__device__变量中
    //cudaMemcpyFromSymbol:同上相反
    //cudaMalloc:在设备端分配内存
    //cudaMemcpy:数据拷贝
    //cudaFree():内存释放
    //cudaMemset():内存初始化
    cudaMemcpyToSymbol(devU,hostU,size,0,cudaMemcpyHostToDevice);
    cudaMemcpyToSymbol(devV,hostV,size,0,cudaMemcpyHostToDevice);

    addUV<<<1,4>>>();

    cudaMemcpyFromSymbol( hostU,devU,size,0,cudaMemcpyDeviceToHost );

    for(int i=0;i<4;i++)
        printf("hostU[%d] = %f
",i,hostU[i]);
    return 0;
}

结果：

dynamic variable使用global memory：

global_mem_dynamic.cu:

#include<iostream>
#include<stdio.h>

__global__ void add4f(float *u , float *v)
{
    int i = threadIdx.x;
    u[i] += v[i];
}

void print(float * U ,int size)
{
    for(int i=0;i<4;i++)
    {
        printf("U[%d] = %f
",i,U[i]);
    }
}

int main()
{
    float hostU[4] = {1,2,3,4};
    float hostV[4] = {5,6,7,8};

    float * devU ;
    float * devV ;
    int size = sizeof(float) * 4;

    //在设备内存上分配空间
    cudaMalloc( &devU,size );
    cudaMalloc( &devV,size );
    //数据拷贝
    cudaMemcpy( devU ,hostU ,size ,cudaMemcpyHostToDevice );
    cudaMemcpy( devV ,hostV ,size ,cudaMemcpyHostToDevice );

    add4f<<<1,4>>> (devU,devV);
    //数据返回
    cudaMemcpy(hostU,devU,size,cudaMemcpyDeviceToHost);

    print(hostU,size);
    //释放空间
    cudaFree(devV);
    cudaFree(devU);
    return 0;
}

结果：

host memory : 主机端内存，即CPU对应的我们普通意义上的内存。主机端内存分为两种：可分页内存（pageable memory）和页锁定（page-locked 或pinned）内存。可分页内存即为通过操作系统API（malloc()，new()）分配的存储器空间：而页锁定内存始终不会被分配到低俗的虚拟内存中，能够保证存在于屋里内存中，并且能够通过DMA加速与设备端的通信。

constant memory：常数存储器。它是只读的地址空间。常熟存储器中的数据位于显存，但拥有缓存加速。常数存储器的空间较小，在Cuda程序中用于存储需要频繁访问的只读参数。当来自同一half-warp的线程访问常数存储器中的同一数据时，如果发生缓存命中，那么只需要一个周期就可以获得数据。常数存储器有缓存机制，用以节约带宽，加快访问速度。每个SM拥有8KB的常数存储器缓存。常数存储器是只读的，因此不存在缓存一致性问题。

constant memory的使用：

#include<iostream>

using namespace std;

__constant__ int devVar = 100;

__global__ void xminus(int *a)
{
    int i = threadIdx.x;
    a[i] = devVar+i;
}

int main()
{
    int *h_a = (int*)malloc(4*10) ;
    int *d_a ;
    cudaMalloc(&d_a, 4*10) ;
    cudaMemset(d_a, 0, 40) ;

    xminus<<<1,4>>>(d_a);

    cudaMemcpy(h_a, d_a, 4*10, cudaMemcpyDeviceToHost) ;

    for(int i = 0; i < 4 ; i++)
        cout << h_a[i] << " " ;
    cout << endl ;
}

结果：

texture memory：纹理存储器。在此不做过多介绍。

再结合一个图解释一下GPU里面的一些概念作为结尾：

Thread：线程。即一个GPU核心处理的单个单元。

Block：块。一个块里面有多个线程组成。block是软件级别的概念。

Grid：格。即二维的格，由多个block组成。

SM：Streaming multiprocessor。一个GPU里面有多个SM，一个SM里面有多个SP（streaming processor），是硬件级别的概念。

Warp：Warp是SM调度和执行的基本单位。warp是32个并列的线程，软件级别的概念。CPU上执行一条指令时候都是一个线程的，但是GPU则是以warp为单位。SM执行一条指令，那么这条指令使得32个线程同时执行，而每个线程都会操作自己的内存处理自己相应的数据，因此就达到了执行一条指令操作多个数据，也就是SIMD（single instruction and multiple data）。

注明出处：http://blog.csdn.net/lavorange/article/details/20465869