1. 常量内存的概念

常量内存用于保存在核函数执行期间不会发生变化的数据。

由于GPU上包含数百个计算单元，因此主要的计算瓶颈已经不在芯片的数学计算吞吐量，而在于芯片的内存带宽（即数据从存储器到计算核心的传输延迟）。

特性：

   （1）常量内存的数据不能修改；
   （2）使用\_\_constant\_\_修饰符；

2.定义和使用常量内存

使用__constant__修饰符来定义常量内存变量:

//定义一个结构体
#define INF 2e10f
struct Sphere{
    float r,g,b;
    float radius;
    float x,y,z;
    __deice__ float hit(float ox,float oy,float *n)
    {
        float dx=ox-x;
        float dy=oy-y;
        if(dx*dx+dy*dy<radius*radius)
        {
            float dz=sqrtf(radius*radius-dx*dx-dy*dy);
            *n=dz/sqrtf(radius*radius);
            return dz+z;
        }
        return -INF;
    }
}
//核函数中定义常量内存变量
__constant__ Sphere s[20];

当从主机内存复制到GPU上的常量内存时，使用cudaMemcpyToSymbol()函数，而非cudaMemcpy()。

//CPU中分配内存
Sphere *temp_s=(Sphere*)malloc(sizeof(Sphere)*20);
//初始化操作（该处省略，编程时需补上）
//在GPU中为Sphere数据分配内存
HANDLE_ERROR(cudaMalloc((void**)&s,sizeof(Sphere)*20));
//数据传输到GPU中的常量内存中
HANDLE_ERROR(cudaMemcpyToSymbol(s,temp_s,sizeof(Sphere)*20));
//传输完毕释放临时内存
free(temp_s);

3.常量内存的优势和不足

优势：

对常量内存的单次读操作可以广播到其他邻近线程，这将节约15此读取操作；（同一半线程束之间）

    所谓的邻近线程设计线程束Warp的概念，32个线程是一个线程束，NVIDIA硬件将把单词内存读取操作广播到每个半线程束（Half-Warp，即16个线程），所以会节约15次读取操作。

常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生额外的内存通信量；（不同半线程束之间）

    硬件会主动把这个常量数据缓存在GPU中，在第一次从常量内存读取某个地址数据后，当其他半线程束请求同一个地址时，将命中缓存。

缺点：

当所有16个线程读相同的地址时，会极大地提升性能；但当分别读取不同的地址时，实际反而会降低性能；

    如果半线程束的所有16个线程需要访问常量内存中不同的数据，那么这16次不同的读取操作会串行化，从而需要16倍的时间发出请求；而从全局内存中读取，那么请求会同时发出。

4.测量运行性能

//定义事件变量
hipEvent_t start, stop;
float dt_ms = 0.0f;
//创建事件
hipEventCreate(&start);
hipEventCreate(&stop);
//开始记录事件
hipEventRecord(start,NULL);
//运行内容在此
//获取事件运行时间
hipEventRecord(stop,NULL);
hipEventSynchronize(start);
hipEventSynchronize(stop);
hipEventElapsedTime(&dt_ms, start, stop);
printf("Transfer CPU -> GPU, time: %lf ms\n", dt_ms);