6.2_常量内存

6.2 常量内存

之前，我们已经介绍了GPU中包含的强大数学处理能力。事实上，正是这种强大的计算优势激发了人们开始研究如何在图形处理器上执行通用计算。由于在GPU上包含有数百个数学计算单元，因此性能瓶颈通常并不在于芯片的数学计算吞吐量，而是在于芯片的内存带宽。由于在图形处理器上包含了非常多的数学逻辑单元（ALU），因此有时输入数据的速率甚至无法维持如此高的计算速率。因此，有必要研究一些手段来减少计算问题时的内存通信量。

到目前为止，我们已经看到了CUDA C程序中可以使用全局内存和共享内存。但是，CUDA C还支持另一种类型的内存，即常量内存。从常量内存的名字就可以看出，常量内存用于保存在核函数执行期间不会发生变化的数据。NVIDIA硬件提供了64KB的常量内存，并且对常量内存采取了不同于标准全局内存的处理方式。在某些情况中，用常量内存来替换全局内存能有效地减少内存带宽。

6.2.1 光线跟踪简介

我们将给出一个简单的光线跟踪（Ray Tracing）应用程序示例，并在这个示例中介绍如何使用常量内存。首先，我们将介绍光线跟踪的一些背景知识。如果你已经熟悉了光线跟踪的一些基本概念，那么可以直接跳到6.2.2节。

简单地说，光线跟踪是从三维对象场景中生成二维图像的一种方式。此时你会奇怪，这不是GPU的设计初衷么？当你玩游戏时，这与OpenGL和 DirectX实现的功能有何不同？没错，GPU确实能解决相同的问题，但它们使用的是一种称之为光栅化（Rasterization）的技术。在许多参考书中都介绍了光栅化技术，因此在这里不会介绍二者的差异。但可以说，它们是解决相同问题的完全不同的方法。

那么，光线跟踪如何从三维场景中生成一张二维图像？原理很简单：在场景中选择一个位置放上一台假想的相机。这台数字相机包含一个光传感器来生成图像，因此我们需要判断哪些光将接触到这个传感器。图像中的每个像素与命中传感器的光线有着相同的颜色和强度。

由于在传感器中命中的光线可能来自场景中的任意位置，因此事实也证明了采用逆向计算或许是更容易实现的。也就是说，不是找出哪些光线将命中某个像素，而是想象从该像素发出一道射线进入场景中。按照这种思路，每个像素的行为都像一只“观察”场景的眼睛。图6.1说明了这些从每个像素投射光纤并进入到场景的过程。


图6.1 一种简单的光线跟踪模式

我们将跟踪从像素中投射出的光线穿过场景，直到光线命中某个物体，然后计算这个像素的颜色。我们说像素都将“看到”这个物体，并根据它所看到物体的颜色来设置它的颜色。光线跟踪中的大部分计算都是光线与场景中物体的相交运算。

在更复杂的光线跟踪模型中，场景中的反光物体能够反射光线，而半透明的物体能够折射光线。这将生成二次射线（Secondary Ray）和三次射线（Tertiary Ray）等等。事实上，这正是光线跟踪最具吸引力的功能之一：实现基本的光线跟踪器是很容易的，如果需要的话，也可以在光线跟踪器中构建更为复杂的成像模型以生成更真实的图像。

6.2.2 在GPU上实现光线跟踪

由于OpenGL和 DirectX等API都不是专门为了实现光线跟踪而设计的，因此我们必须使用CUDA来实现基本的光线跟踪器。我们构造的光线跟踪器非常简单，这样可以将重点放在常量内存的使用上。因此，如果你希望基于这段代码来构建一个功能完备的渲染器，那么是不现实的。我们的光线跟踪器只支持一组包含球状物体的场景，并且相机被固定在Z轴，面向原点。此外，我们将不支持场景中的任何照明，从而避免二次光线带来的复杂性。我们也不计算照明效果，而只是为每个球面分配一个颜色值，然后如果它们是可见的，则通过某个预先计算的值对其着色。

那么，光线跟踪器将实现哪些功能？它将从每个像素发射一道光线，并且跟踪这些光线会命中哪些球面。此外，它还将跟踪每道命中光线的深度。当一道光线穿过多个球面时，只有最接近相机的球面才会被看到。我们的“光线跟踪器”会把相机看不到的球面隐藏起来。

通过一个数据结构对球面建模，在数据结构中包含了球面的中心坐标 (x,y,z)(\mathbf{x},\mathbf{y},\mathbf{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  $=$  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;   
}

我们还将注意到，在这个结构中定义了一个方法hit(float ox, float oy, float *n)。对于来自(ox, oy)处像素的光线，这个方法将计算光线是否与这个球面相交。如果光线与球面相交，那么这个方法将计算从相机到光线命中球面处的距离。我们需要这个信息，原因在前面已经提到了：当光线命中多个球面时，只有最接近相机的球面才会被看见。

main()函数遵循了与前面示例大致相同的代码结构。

include“cuda.h" #include“../common/book.h" #include“../common/cpuBITMAP.h" #definernd(x）(x\*rand()／RAND_MAX) #define SPHERES20 Sphere \*s; int main(void){ //记录起始时间 CUDAEvent_t start，stop; HANDLE_ERROR(cudaEventCreate(&start）);
HANDLE_ERROR(udaEventCreate(&stop) );  
HANDLE_ERROR(udaEventRecord(start,0) );  
CPUBitmap bitmap(DIM, DIM);  
unsigned char *dev_bitmap;  
//在GPU上分配内存以计算输出位图  
HANDLE_ERROR(udaMalloc(void**)&dev_bitmap, bitmap(image_size()));  
//为Sphere数据集分配内存  
HANDLE_ERROR(udaMalloc(void**&s, sizeof(Sphere)*SPHERES))；

我们为输入的数据分配了内存，这些数据是一个构成场景的Sphere数组。由于Sphere数组将在CPU上生成并在GPU上使用，因此我们必须分别调用CUDAAlloc()和malloc()在GPU和CPU上分配内存。我们还需要分配一张位图图像，当在GPU上计算光线跟踪球面时，将用计算得到的像素值来填充这张图像。

在分配输入数据和输出数据的内存后，我们将随机地生成球面的中心坐标，颜色以及半径。

//分配临时内存，对其初始化，并复制到  
//GPU上的内存，然后释放临时内存  
Sphere \*temp_s=(Sphere\*)malloc(sizeof(Sphere)\*SPHERES);for(inti=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;

当前，程序将生成一个包含20个球面的随机数组，但这个数量值是通过一个#define宏指定的，因此可以相应地做出调整。

通过CUDAMemcpy()将这个球面数组复制到GPU，然后释放临时缓冲区。

HANDLE_ERROR(udaMemcpy(s, temp_s, sizeof(Sphere) * SPHERES,udaMemcpyHostToDevice));  
free( temp_s);

现在，输入数据位于GPU上，并且我们已经为输出数据分配好了空间，因此可以启动核函数。

//从球面数据中生成一张位图  
dim3 grids(DIM/16, DIM/16);  
dim3 threads(16, 16);  
kernel<<<grids, threads>>>(dev bitmap);

我们稍后将分析核函数本身，现在你只需知道这个函数将执行光线跟踪计算并且从输入的一组球面中为每个像素计算颜色数据。最后，我们将把输出图像从GPU中复制回来，并显示它。当然，我们还要释放所有已经分配但还未释放的内存。

//将位图从GPU复制回到CPU以显示   
HANDLE_ERROR(udaMemcpy(bitmap.get_ptr(),dev_bitmap, bitmap,image_size(),udaMemcpyDeviceToHost））; bitmap.display_and_exit(); //释放内存   
CUDAFree(dev_bitmap）;   
CUDAFree(s);   
}

现在，所有这些操作对你来说应该很熟悉了。那么，如何实现光线跟踪算法？由于我们已经建立了一个非常简单的光线跟踪模型，因此核函数理解起来就很容易了。每个线程都会为输出影像中的一个像素计算颜色值，因此我们遵循一种惯用的方式，计算每个线程对应的x坐标和y坐标，并且根据这两个坐标来计算输出缓冲区中的偏移。此外，我们还将把图像坐标（x，y）偏移DIM/2，这样z轴将穿过图像的中心。

__global__ void kernel(unsigned char *ptr) {
    // 将threadIdx/BlockIdx映射到像素位置
    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  $\mathbf{r} = 0$ $g = 0$ $b = 0$    
float maxz  $=$  INF;   
for(int  $\mathrm{i} = 0$  .i<SPHERES；i++) { float n; float t  $=$  s[i].hit( ox,oy,&n）; if（t>maxz）{ float fscale  $\equiv$  n;
$\begin{array}{rl} & {\mathrm{r = s[i].r*fscale;}}\\ & {\mathrm{g = s[i].g*fscale;}}\\ & {\mathrm{b = s[i].b*fscale;}} \end{array}$

显然，判断相交计算的大部分代码都包含在for()循环中。对每个输入的球面进行迭代，并调用hit()方法来判断来自像素的光线能否“看到”球面。如果光线命中了当前的球面，那么接着判断命中的位置与相机之间的距离是否比上一次命中的距离更加接近。如果更加接近，那么我们将这个距离保存为新的最接近球面。此外，我们还将保存这个球面的颜色值，这样当循环结束时，线程就会知道与相机最接近的球面的颜色值。由于这就是从像素发出的光线“看到”的颜色值，也就是该像素的颜色值，因此这个值应该保存在输出图像的缓冲区中。

在判断了每个球面的相交情况后，可以将当前的颜色值保存到输出图像中，如下所示：

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;

注意，如果没有命中任何球面，那么保存的颜色值将是变量r，b和g的初始值。在本示例中，r、b和g的初始值都设置为0，因此背景色是黑色。你可以修改这些值以便生成不同颜色的背景。在图6.2中给出了一个输出示例，其中绘制了20个球体，并且背景为黑色。


图6.2 光线跟踪示例的截图

由于我们随机设置了这些球面的位置，颜色和大小，因此，如果你得到的输出与这里的图像并不相同，那么也是正常的。

6.2.3 通过常量内存来实现光线跟踪

你已经注意到，在这个光线跟踪示例中并没有提到常量内存。现在，我们使用常量内存来改进这个示例。由于常量内存是不能修改的，因此显然无法用常量内存来保存输出图像的数据。在这个示例中只有一个输入数据，即球面数组，因此应该将这个数据保存到常量内存中。

常量内存的声明方法与共享内存是类似的。要使用常量内存，那么在代码中将不再像下面这样声明数组：

Sphere \*s;

而是在变量前面加上constant修饰符：

constantSphere[sPHERES];

注意，在最初的示例中，我们声明了一个指针，然后通过CUDAAlloc()来为指针分配GPU内存。当我们将其修改为常量内存时，同样要将这个声明修改为在常量内存中静态地分配空间。我们不再需要对球面数组调用CUDAAlloc()或者CUDAFree()，而是在编译时为这个数组提交一个固定的大小。这对许多应用程序来说是可以接受的，因为常量内存能够带来性能的提升。我们稍后会看到常量内存的优势，但首先来看看如何将main()函数修改为使用常量内存：

int main(void）{CPUBitmap bitmap(DIM，DIM);unsigned char \*dev_bitmap;//在GPU上分配内存以计算输出位图HANDLE_ERROR(cudaMalloc(void\*\*)&dev_bitmap， bitmap(image_size()））；//分配临时内存，对其初始化，并复制到  
//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(udaMemcpyToSymbol(s, temp_s, sizeof(Sphere) * SPHERES)); free(temp_s); //从球面数据中生成一张位图 dim3 grids(DIM/16, DIM/16); dim3 threads(16, 16); kernel<<<grids, threads>>>(dev bitmap); //将位图从GPU复制回到CPU以显示 HANDLE_ERROR(udaMemcpy(bitmap.get_ptr(), dev bitmap, bitmap(image_size(),udaMemcpyDeviceToHost)); bitmap.display_and_exit(); //释放内存 CUDAFree(dev.Bitmap); }

这段代码在很大程度上类似于之前main()的实现。正如在前面提到的，对main()函数的修改之一就是不再需要调用cudaMalloc()为球面数组分配空间。在下面给出了另一处修改：

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

当从主机内存复制到GPU上的常量内存时，我们需要使用这个特殊版本的CUDAMemcpy()。CUDAMemcpyToSymbol()与参数为CUDAMemcpyHostToDevice()的CUDAMemcpy()之间的唯一差异在于，CUDAMemcpyToSymbol()会复制到常量内存，而CUDAMemcpy()会复制到全局内存。

除了__constant__修饰符和对main()的两处修改之外，其他的代码都是相同的。

6.2.4 常量内存带来的性能提升

__constant__将把变量的访问限制为只读。在接受了这种限制后，我们希望获得某种回报。在前面曾提到，与从全局内存中读取数据相比，从常量内存中读取相同的数据可以节约内存带宽，主要有两个原因：

  • 对常量内存的单次读操作可以广播到其他的“邻近（Nearby）”线程，这将节约15次读取操作。
    ·常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生额外的内存通信量。

“邻近”这个词的含义是什么？要回答这个问题，我们需要解释线程束（Warp）的概念。

这里的“Warp”并不是《星际迷航》电影中的曲速引擎（Warp Drive），而是来自纺织（Weaving）领域的概念，这里的线程束与空间旅行速度没有任何关系。线程束可以看成是一组线程通过交织而形成的一个整体。在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“步调一致（Lockstep）”的形式执行。在程序中的每一行，线程束中的每个线程都将在不同的数据上执行相同的指令。

当处理常量内存时，NVIDIA硬件将把单次内存读取操作广播到每个半线程束（Half-Warp）。在半线程束中包含了16个线程，即线程束中线程数量的一半。如果在半线程束中的每个线程都从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求并在随后将数据广播到每个线程。如果从常量内存中读取大量的数据，那么这种方式产生的内存流量只是使用全局内存时的1/16（大约 6%6\% ）。

但在读取常量内存时，所节约的并不仅限于减少了 94%94\% 的带宽。由于这块内存的内容是不会发生变化的，因此硬件将主动把这个常量数据缓存在GPU上。在第一次从常量内存的某个地址上读取后，当其他半线程束请求同一个地址时，那么将命中缓存，这同样减少了额外的内存流量。

在我们的光线跟踪器中，每个线程都要读取球面的相应数据从而计算它与光线的相交情况。在把应用程序修改为将球面数据保存在常量内存后，硬件只需要请求这个数据一次。在缓存数据后，其他每个线程将不会产生内存流量，原因有两个：

·线程将在半线程束的广播中收到这个数据。
·从常量内存缓存中收到数据。

然而，当使用常量内存时，也可能对性能产生负面影响。半线程束广播功能实际上是一把双刃剑。虽然当所有16个线程都读取相同地址时，这个功能可以极大地提升性能，但当所有16个线程分别读取不同的地址时，它实际上会降低性能。

只有当16个线程每次都只需要相同的读取请求时，才值得将这个读取操作广播到16个线程。然而，如果半线程束中的所有16个线程需要访问常量内存中不同的数据，那么这个16次不同的读取操作会被串行化，从而需要16倍的时间来发出请求。但如果从全局内存中读取，那么这些请求会同时发出。在这种情况中，从常量内存读取就慢于从全局内存中读取。

前の章6.1_本章目标
次の章6.3_使用事件来测量性能