[CUDA学习]5. 常量内存与事件

常量内存与事件

本文将介绍CUDA C的一些更高级的功能。具体来说，就是通过GPU上特殊的内存区域来加速应用程序的执行。这里将讨论其中的一种内存区域：常量内存(Constant Memory)。此外还将介绍一种增强CUDA C应用程序性能的方法，在这个过程中，你将学习如何通过事件来测量CUDA应用程序的性能的方法，在这个过程中你将学习如何通过事件来测量CUDA应用程序的性能。通过这些测量方法，你可以定量地分析对应用程序的修改是否会带来性能的提升(或者性能下降)。

目标

• 了解如何在CUDA C中使用常量内存
• 了解常量内存的性能特性
• 学习如何使用CUDA事件来测量应用程序的性能

常量内存

之前我们已经介绍了GPU中包含的强大的数学处理能力。实际上，正是这种强大的计算优势激发了人们开始研究如何在图形处理器上执行通用计算。由于在GPU上包含有数百个数学计算单元，因此性能瓶颈通常并不在于芯片上的数学计算的吞吐量，而是在于芯片的内存带宽。由于在图形处理器中包含了非常多的数学逻辑单元(ALU)，因此有时输入数据的速率甚至无法维持如此高速的计算速率。因此，有必要研究一些手段来减少计算问题时内存通讯量。
到目前为止，我们已经看到了CUDA C程序中可以使用全局内存和共享内存。但是，CUDA C还支持另外一种内存，即常量内存。从常量内存的名字就可以看出，常量内存用于保存在核函数执行期间不会发生变化的数据。NVIDIA硬件提供了64KB的常量内存，并且对常量内存采取了不同于标准全局内存的处理方式。在某些情况中，用常量内存来代替全局内存能有效的减少内存带宽。
下面将通过一个光线跟踪的例子来对这个问题做个简单的介绍：

光线跟踪简介

简单的说光线跟踪是从三维场景生成二维图像的一种方式。现在在GPU中主流的方法是采用一种叫光栅化(Rasterization)的技术。
那么光线跟踪如何从三维场景中生成一张二维图像的了？原理很简单：在场景中选择一个位置放置一台假象的摄像机，这台数字相机包含一个光学的传感器来生成图像。因此我们需要判断哪些光将接触这个传感器。图像中每个像素与命中传感器的光线有着相同的颜色与强度。
由于在传感器中命中的光线可能来自场景中的任意位置，因此事实也证明了采用逆向计算或许是更好的办法。也就是说，不是找出那些光线将命中某个像素，而是想象该像素发出一道射线进入场景中。按照这种思想，每个像素的行为都像是一只“观察”场景的眼睛。
我们将跟踪像素中投射出来的光线穿过场景，直到光线命中某个物体，然后计算这个像素的颜色。光线跟踪中的大部分计算都是光线与场景中物体的相交运算。

在GPU上实现光线跟踪

由于openGL和Direct X等API都不是专门为了实现光线跟踪而设计的，因此我们必须使用CUDA C来实现基本的光线跟踪器，这里只是为每个球面分配一个颜色值，然后如果它们是可见的，则通过某个预先计算的值对其着色。
那么光线跟踪器将实现哪些功能？它将从每个像素发射一道光线，并跟踪这些光线会命中那些球面。此外，它还将跟踪每道命中光线的深度。当一道光线穿过多个球面时，只有接近相机的球面才会被看到。我们的“光线跟踪器”会把相机看不到的球面隐藏起来。
这里通过一个数据对球面建模，在数据结构中包含了球面的中心坐标(x,y,z)，半径radius，以及颜色值(r,g,b)。

#define INF 2e10f

struct Sphere{
  float r,b,g;
  float radius;
  float x,y,z;
  __device__ float hit(float ox,float oy,float *n){
    float dx = ox - x;
    float dy = ox - y;
    if((dx*dx+dy*dy)<radius*radius){
      float dz = sqrtf(radius*radius - dx*dx - dy*dy);
      *n = dz/sqrtf(radius*radius);
      return -INF;
    }
  }
};

这里我们注意到在这个结构体中定义了一个方法hit(float ox,float oy,float *n)。对于来自(ox,oy)处的像素光线，这个方法将计算光线是否与球面相交，那么这个方法将计算光线是否与这个球面相交。如果光线与球面相交那么这个方法将计算从相机到光线命中球面处的距离。我们需要这个信息，原因在前面已经介绍过了：当光先命中多个球面时只有最接近的球面才会被看到。
main函数遵循前面示例大致相同的代码结构：

#include "cuda.h"