CUDA学习四：常量内存与事件

作者：JackGao24 博客园
作者：JackGao16 CSDN
文章链接： http://blog.csdn.net/u013108511/article/details/77579244
邮箱：[e-mail]

这篇博客主要将问题聚焦在一下的两点：

１、通过GPU上特殊的内存区域（常量内存）来加速应用程序的执行；
２、通过时间来测量CUDA应用程序的性能。

1、常量内存

  由于GPU上众多的计算资源，因此性能的瓶颈可能并不在芯片的数学计算的吞吐量上，而在于芯片上的内存带宽，因此有必要研究一些手段开减少计算问题时的内存通讯量。
  继以前的全局内存和共享内存，本篇博客重点介绍的是常量内存。常量内存用于保存在核函数执行期间不会发生变化的数据，NVIDIA硬件提供了64K的常量内存并且采用了不同于标准全局内存大的处理模式，在某些情况下常量内存取代全局内存能有效的减少内存的带宽。

1.1、光线跟踪的简介

  光线跟踪是从三维的对象场景中生成二维图像的一种方式。
  原理：在场景中选择一个位置放上一台假想的相机，这台数字相机包含一个传感器来生成图像，因此我们需要判断那些光将接触到这个传感器，图像中的像素与命中传感器的光线有着相同的颜色和强度。反过来想可以想象成像素发出的光进入到场景中的过程。
  光纤跟踪中的大部分计算都是光线与场景中的物体的相交计算。实现基本的光线跟踪器是容易的，如果需要机器视觉：作为人工智能的一个方面，像人一样，视觉的经验的获得可能会伴随着听觉触觉等的配合吗？大脑的容量是多少？不断在大脑中更新的模型是一个怎样复杂的结构？语音、图像、文字的理解的模型的共享，共享学习是否是机器学习的未来的一个研究方向呢？我为什么会想到这些？的话也可以在光线跟踪器中构建更为发杂的模型以生成更为真实的图像。

1.2、在GPU上实现光线的跟踪

  我们只构造最简单的光线追踪器，这样的我们可以将重点放在常量内存的使用上来。只支持一组包含球状物的场景，相机被固定在Z轴，也就是xoy的一个平面，也不计算照明的效果。

1.2.1 以往的实现方式

下面的代码是未使用常量内存的实现:

#include "cuda.h"#include "../common/book.h"#include "../common/cpu_bitmap.h"#define DIM 1024#define INF 2e10f#define SPHERES 20#define rnd(x) (x * rand()/ RAND_MAX)struct Sphere {    float r,g,b;    float radius;    float x,y,z;    __device__ 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;    }};__global__ void kernel( Sphere *s, unsigned char *ptr ) {    // map from threadIdx/BlockIdx to pixel position    int x = threadIdx.x + blockIdx.x * blockDim.x;    int y = threadIdx.y + blockIdx.y * blockDim.y;    int offset = x + y * blockDim.x * gridDim.x;    float   ox = (x - DIM/2);    float   oy = (y - DIM/2);    float   r=0, g=0, b=0;    float   maxz = -INF;    for(int i=0; i<SPHERES; i++) {        float   n;        float   t = s[i].hit( ox, oy, &n );        if (t > maxz) {            float fscale = n;            r = s[i].r * fscale;            g = s[i].g * fscale;            b = s[i].b * fscale;            maxz = t;        }    }     ptr[offset*4 + 0] = (int)(r * 255);    ptr[offset*4 + 1] = (int)(g * 255);    ptr[offset*4 + 2] = (int)(b * 255);    ptr[offset*4 + 3] = 255;}struct DataBlock {    unsigned *dev_bitmap;    Sphere *s;};int main(  void ) {    //记录起止时间    cudaEvent_t start, stop;    HANDLE_ERROR(cudaEventCreate(&start));    HANDLE_ERROR(cudaEventCreate(&stop));    HANDLE_ERROR(cudaEventRecord(start,0));    CPUBitmap bitmap( DIM, DIM);    unsigned char *dev_bitmap;    Sphere *s;    //在GPU 上分配内存以计算输出位图    HANDLE_ERROR(cudaMalloc((void**)&dev_bitmap, bitmap.image_size()));    //在Sphere数据集分配内存    HANDLE_ERROR(cudaMalloc((void**)&s, sizeof(Sphere)*SPHERES));    //分配临时内存并对其进行初始化,并复制到GPU的内存上,然后释放临时的内存    Sphere *temp_s = (Sphere *) malloc (sizeof(Sphere) * SPHERES);    for(int i = 0; i < SPHERES; i++) {        temp_s[i].r = rnd(1.0f);        temp_s[i].g = rnd(1.0f);        temp_s[i].b = rnd(1.0f);        temp_s[i].x = rnd(1000.0f)-500;        temp_s[i].y = rnd(1000.0f)-500;        temp_s[i].z = rnd(1000.0f)-500;        temp_s[i].radius = rnd(100.0f) + 20;    }    //通过cudaMemcpy()将这个球面的数组复制到GPU,然后释放临时缓冲区    HANDLE_ERROR(cudaMemcpy(s,temp_s,sizeof(Sphere)*SPHERES,cudaMemcpyHostToDevice));    //释放CPU临时的内存空间    free(temp_s);    //在球面的数据中生成一张位图    dim3 grids(DIM/16, DIM/16);    dim3 threads(16,16);    kernel<<<grids, threads>>>(s,dev_bitmap);    //把输出图像从GPU中复制出来,并显示它,释放所有已经分配但还未释放的内存    HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(), dev_bitmap,bitmap.image_size(),cudaMemcpyDeviceToHost));    //获取结束时间并打印时间代价    HANDLE_ERROR(cudaEventRecord(stop,0));    HANDLE_ERROR(cudaEventSynchronize(stop));    float elapsedTime;    HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,start,stop));    printf( "Time to generate:  %3.1f ms\n", elapsedTime );    HANDLE_ERROR( cudaEventDestroy( start ) );    HANDLE_ERROR( cudaEventDestroy( stop ) );    bitmap.display_and_exit();    //释放内存    cudaFree(dev_bitmap);    cudaFree(s); }

程序流程解释：
  CUDA C编程有一套固定的流程，记住该流程可以更好的把握整个程序的脉络，程序的流程是：
1、在GPU上为输出位图分配内存，用于保存输出的结果；
2、在GPU上为Sphere数据集分配内存，用于存储输入数据；
3、在CPU上为输入的数据分配内存，进行随机的初始化操作之后，将其复制到GPU分配的内存中；
4、GPU上的计算部分
5、将GPU上计算结果的位图从GPU复制到CPU并进行显示。

打印出的图像效果:

处理时间:(这里只给出一个时间,多次运行,时间的区间保持在均值5.5ms,这个数值用来和后续的常量内存实现方式进行比较)

1.2.2 常量内存的实现方式

#include "cuda.h"#include "../common/book.h"#include "../common/cpu_bitmap.h"#define DIM 1024#define INF 2e10f#define SPHERES 20#define rnd(x) (x * rand()/ RAND_MAX)struct Sphere {    float r,g,b;    float radius;    float x,y,z;    __device__ 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[SPHERES];__global__ void kernel(unsigned char *ptr ) {    // map from threadIdx/BlockIdx to pixel position    int x = threadIdx.x + blockIdx.x * blockDim.x;    int y = threadIdx.y + blockIdx.y * blockDim.y;    int offset = x + y * blockDim.x * gridDim.x;    float   ox = (x - DIM/2);    float   oy = (y - DIM/2);    float   r=0, g=0, b=0;    float   maxz = -INF;    for(int i=0; i<SPHERES; i++) {        float   n;        float   t = s[i].hit( ox, oy, &n );        if (t > maxz) {            float fscale = n;            r = s[i].r * fscale;            g = s[i].g * fscale;            b = s[i].b * fscale;            maxz = t;        }    }     ptr[offset*4 + 0] = (int)(r * 255);    ptr[offset*4 + 1] = (int)(g * 255);    ptr[offset*4 + 2] = (int)(b * 255);    ptr[offset*4 + 3] = 255;}struct DataBlock {    unsigned *dev_bitmap;    //Sphere *s;};//__constant__ Sphere s[SPHERES];int main(  void ) {    //记录起止时间    cudaEvent_t start, stop;    HANDLE_ERROR(cudaEventCreate(&start));    HANDLE_ERROR(cudaEventCreate(&stop));    HANDLE_ERROR(cudaEventRecord(start,0));    CPUBitmap bitmap( DIM, DIM);    unsigned char *dev_bitmap;    //__constant__ Sphere s[SPHERES];    //在GPU 上分配内存以计算输出位图    HANDLE_ERROR(cudaMalloc((void**)&dev_bitmap, bitmap.image_size()));    //在Sphere数据集分配内存    HANDLE_ERROR(cudaMalloc((void**)&s, sizeof(Sphere)*SPHERES));    //分配临时内存并对其进行初始化,并复制到GPU的内存上,然后释放临时的内存    Sphere *temp_s = (Sphere *) malloc (sizeof(Sphere) * SPHERES);    for(int i = 0; i < SPHERES; i++) {        temp_s[i].r = rnd(1.0f);        temp_s[i].g = rnd(1.0f);        temp_s[i].b = rnd(1.0f);        temp_s[i].x = rnd(1000.0f)-500;        temp_s[i].y = rnd(1000.0f)-500;        temp_s[i].z = rnd(1000.0f)-500;        temp_s[i].radius = rnd(100.0f) + 20;    }    HANDLE_ERROR(cudaMemcpyToSymbol(s,temp_s,sizeof(Sphere)*SPHERES));    //通过cudaMemcpy()将这个球面的数组复制到GPU,然后释放临时缓冲区    //HANDLE_ERROR(cudaMemcpy(s,temp_s,sizeof(Sphere)*SPHERES,cudaMemcpyHostToDevice));    //释放CPU临时的内存空间    free(temp_s);    //在球面的数据中生成一张位图    dim3 grids(DIM/16, DIM/16);    dim3 threads(16,16);    kernel<<<grids, threads>>>(dev_bitmap);    //把输出图像从GPU中复制出来,并显示它,释放所有已经分配但还未释放的内存    HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(), dev_bitmap,bitmap.image_size(),cudaMemcpyDeviceToHost));    //获取结束时间并打印时间代价    HANDLE_ERROR(cudaEventRecord(stop,0));    HANDLE_ERROR(cudaEventSynchronize(stop));    float elapsedTime;    HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,start,stop));    printf( "Time to generate:  %3.1f ms\n", elapsedTime );    HANDLE_ERROR( cudaEventDestroy( start ) );    HANDLE_ERROR( cudaEventDestroy( stop ) );    bitmap.display_and_exit();    //释放内存    cudaFree(dev_bitmap);    cudaFree(s); }

程序解释：
常量内存的设定：
1、常量内存的声明和定义的时候，在变量的前面加上一个_ _ constant _ _ 修饰符；
2、在GPU上内存的申请也由原来的利用cudaMalloc()指针的动态分配修改为在常量内存中的静态地址分配空间：

  HANDLE_ERROR(cudaMemcpyToSymbol(s,temp_s,sizeof(Sphere)*SPHERES));

  上面的这行代码同样也是将主机内存复制到GPU中，但与参数为cudaMemcpyHostToDevice()的cudaMemcpy()之间的唯一差别在于上面的这行代码将内存内数据复制到常量内存中，而cudaMemcpy()会复制到全局的内存。
使用事件来测试性能：
1、在GPU上获取一个时间戳：首先创建一个事件，然后记录一个事件

//记录起止时间cudaEvent_t start, stop;HANDLE_ERROR(cudaEventCreate(&start));HANDLE_ERROR(cudaEventCreate(&stop));HANDLE_ERROR(cudaEventRecord(start,0));

2、记录当前时间，但是当GPU和CPU异构计算同时执行的时候时间的记录便会变得相当的复杂，这个时候我们引入一个同步的语句，在stop时间之前的GPU工作已经完成了才能安全的读取stop中保存的时间戳。

//获取结束时间并打印时间代价HANDLE_ERROR(cudaEventRecord(stop,0));HANDLE_ERROR(cudaEventSynchronize(stop));

程序运行效果图:

运行时间:(多次运行的时间开销保持在均值4.3ms的一个范围内)

2、常量内存带来的性能提升原因

_ _ constant _ _把变量的访问限制为只读，这种限制带来的汇报就是从常量数据中获取相同的数据更节约内存带宽，主要原因是：
1、对常量内存的单次操作可以广播到其他的“临近”的线程，节约15次的读取操作（本实验中每次的广播范围是半个线程块16个线程的半线程束）
线程束：一个包含32个线程的集合，这个线程的集合被编织在一起并且步调一致的形式执行，在程序中的每一行，线程束的每个线程将在不同的数据上执行相同的指令。
2、常量内存的数据将缓存起来，后续的线程束访问时也将会命中，大大减少内存流量。

使用常量数据也可能带来不好的影响：当线程束中的16个线程取址是不同的地址时，对于常量内存这些操作会被串行化，需要16倍的时间来发出请求，而全局内存的内存访问，这些请求可以同时发出，这个时候全局内存访问可能更快。