第四章 CUDA内存处理2.pdf

第4章 CUDA 内存处理2 1、常数存储器 2、纹理存储器 4.1. 常数存储器 常数存储器具有只读和片上缓存的特点，可用来存储在整个核函 数执行期间只读的数据。与全局存储器不同，它拥有缓存。在常 数存储器的数据如果half-warp 中的线程访问同一个地址，那么 GPU就产生一次访问请求随即广播给每个线程，这样整个half- warp访问数据的时间就只需要访问全局存储器1/16的时间，获得 很高的访问效率，有效的提高核函数的执行效率。 但是常数存储器空间较小，GTX480 只提供64KB 的常数存储器 空间，因此，一般用于存储需要频繁访问的只读参数，比如权重 数组、系数矩阵和搜索偏移序列等。有效利用常数存储器的缓存 机制，可以节约带宽、提高访问的速度。 一个简单例子： 定义常数存储器时，需要将其定义在所有函数之外，作用范围为 整个文件，并且对主机端和设备端函数可见。常数存储器的使用 有两种方法： 我们更多的是采用第二种方法。 一个光线跟踪的例子 （使用常数存储器加速）： 光线跟踪如何从三维场景中生成一张二维图像呢？在场景中选择 一个位置放上一台假想的相机，这台数字相机包含一个光传感器 来生成图像，因此我们要判断哪些光接触到这个传感器，图像中 的每个像素与命中传感器的光线有着相同的颜色和强度。 由于在传感器中命中的光线可能来自场景中的任意位置，因此事 实也证明了采用逆向计算或许更容易实现。即不是找出那些光线 将命中某个像素，而是想象从该像素发出一道射线进入场景中。 按照这样的思路，每个像素的行为都像一只 “观察”场景的眼睛。 下图说明了这些从每个像素投射光线并进入到场景的过程。跟踪 从像素中投射出的光线穿过场景，直到光线命中某个物体，再计 算这个像素的颜色。我们说像素都将 “看到”这个物体，并根据 它所看到物体的颜色设置它的颜色。光线跟踪中的大部分计算都 是光线与场景中物体的相交运算。 在GPU上实现光线跟踪： 光线跟踪将实现哪些功能呢？它将从每个像素发射一道光线， 并且跟踪这些光线会命中哪些球面，此外，它还将跟踪每道命中 光线的深度。当一道光线穿过多个球面时，只有最接近相机的球 面才会被看到。 “光线跟踪器”会把相机看不到的球面隐藏起来。 通过一个数据结构对球面建模，在数据结构中包含了数据的中心 坐标 （x,y,z ）、半径radius及颜色值 （r,g,b ） Main函数： 核函数： 如何使用常数存储器进行光线追踪：由于常数存储器是只读的，因 而无法使用它来存放输出图像的数据。在这个例子中，只有一个输 入数据，即球面数组，因此可将这个数据保存至常数存储器中。 在最初的示例中，我们声明一个指针,通过cudaMalloc()为指针分 配空间。当我们修改为常量内存时，将声明修改为在常量内存中静 态分配空间。而是在编译时为这个数组提交固定的大小。 Main() 函数中代码类似上面代码，对其中修改的一个地方就是 不再调用cudaMalloc()为球面数组分配空间，另一处修改如下 （cudaMemcpyToSymbol()直接将数据复制至常数存储器）： 当处理常量内存时，NVIDIA硬件将把单词内存读取操作广播 到每个Half-Warp 。在Half-Warp 中包含了16个线程，如果在 alf-Warp 中每个线程都从常量内存得相同地址上读取数据，那 么GPU只会产生一次读取请求随后将数据广播到每个线程，这 种方式产生的内存通信是全局存储器的1/16 （约6%）。 在光线跟踪中，每个线程都要读取球面的相应数据从而计算 它与光线的相交情况，在把应用程序改为将球面数据保存至常 量内存后，硬件只需请求这个数据一次。在缓存数据后，其它 每个线程不会产生内存通信，原因有两方面： 线程在Half-Warp 的广播中接收到这个数据； 从常量内存缓存中收到数据。 但当使用常量内存时，也有可能对性能产生负面影响，当 Half-Warp 中的16个线程读取相同的地址时，这个可以提高性 能，但当16个线程分别读取不同的地址时，它其实会降低性能。 因为这种情况下，16次不同的操作会被串行化，从而需要16倍 的时间发出请求。但如果从全局存储器中读取数据，这些请求 会同时发出，在这种情况下，从常量内存读取就慢于从全局存 储器中读取。 4.2. 纹理存储器 纹理存储器是一种专为二维访问而优化的