CUDA并行计算系列之常量内存与性能测量

红太狼 2022-12-11 10:29 184阅读 0赞

CUDA学习计划：

1. 最基本的GPU程序；

2. GPU硬件组成和软件概念；

3. 共享内存的使用；

4. 常量内存与性能测量；

5. GPU程序调试；

6. 提高CUDA程序运行效率的方法。

# 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.常量内存的优势和不足 #

优势：

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

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

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

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

缺点：

1. 当所有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);