在第 5 章中，我们研究了一个简单的算法，用于解决三维点列表中许多点的最近邻问题。我们在第 6 章中遵循了这一点，在那里我们详细检查了共享内存。我们现在将应用这些新知识来大大提高第 5 章中的程序效率。
共享内存最基本的用途之一就是阻塞。阻塞允许更有效地使用全局内存总线，并且还允许在共享内存和寄存器上执行大部分计算，而不是依赖全局内存的 L1 和 L2 缓存来提高速度。它的工作原理是这样的:我们将全局内存块复制到共享内存中，而不是让所有线程重复地读写全局内存。线程处理数据的共享内存副本，完成后，它们从全局内存加载另一个块。这样，线程几乎用共享内存完成所有操作，最大限度地减少了对全局内存的读写。
| 
| 提示:阻塞的另一个好处是，当数据被复制到共享内存时，它为程序员提供了一个改变存储数据格式的机会。例如，我们可以从全局内存中复制三维点，并将它们作为四维点存储在共享内存中。这种特殊的操作可能会也可能不会对性能产生影响，但它肯定会带来许多可能性。 |
最近邻的以下版本(清单 7.1)是我们之前在第 5 章中看到的优化版本。它使用我刚才描述的阻塞技术。线程块的线程各自从全局内存复制一个点到共享内存。它们都检查自己点附近的邻居是否在共享内存中的点列表中。一旦完成，它们会从全局内存中复制更多的点，从而大大提高算法的速度。清单 7.1 中的代码可以替换FindClosestGPU.cu文件中的代码。
// FindClosestGPU.cu
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include "FindClosestGPU.h" // Not required!
// Constant block size
__device__ const int blockSize = 128;
// Find nearest neighbor using shared memory blocking
__global__ void FindClosestGPU(float3* points, int* indices, int count) {
// Shared memory for block
__shared__ float3 OtherBlocksPoints[128];
if(count <= 1) return;
// Calculate a unique idx
int idx = threadIdx.x + blockIdx.x * blockSize;
// Assume the closest if points[-1]
int indexOfClosest = -1;
// This thread's point
float3 thisPoint;
// Read in this block's points
if(idx < count) {
thisPoint = points[idx];
}
// Assume distance to nearest is float.max
float distanceToClosest = 3.40282e38f;
// Read in blocks of other points
for(int currentBlockOfPoints = 0; currentBlockOfPoints < gridDim.x; currentBlockOfPoints++) {
// Read in a block of points to the OtherBlocksPoints array
if(threadIdx.x + currentBlockOfPoints * blockSize < count)
OtherBlocksPoints[threadIdx.x] = points[threadIdx.x + (currentBlockOfPoints * blockSize)];
// Wait until blocks read from global into shared memory
__syncthreads();
if(idx < count) {
// Use pointer for faster addressing
float* ptr = &OtherBlocksPoints[0].x;
// For each point in shared memory block:
for(int i = 0; i < blockSize; i++) {
// Calculate distance
float dist = (thisPoint.x - ptr[0]) * (thisPoint.x - ptr[0]);
dist += (thisPoint.y - ptr[1])*(thisPoint.y - ptr[1]);
dist += (thisPoint.z - ptr[2])*(thisPoint.z - ptr[2]);
ptr+=3;
// If this point is within the list and nearer than the
// current closest, update nearest with this one:
if(dist<distanceToClosest &&(i+currentBlockOfPoints*blockSize)
< count && (i+currentBlockOfPoints*blockSize) != idx) {
distanceToClosest = dist;
indexOfClosest = (i + currentBlockOfPoints * blockSize);
}
}
}
__syncthreads();
}
if(idx < count)
indices[idx] = indexOfClosest;
}
清单 7.1:共享内存的最近邻居
如您所见，清单 7.1 中的内核比第 5 章中的前一个版本快得多。这个版本要求网格用 128 个线程的块来启动。在列表的顶部，我使用了一个名为blockSize的常量int；这将成为编译代码中的文字。文字通常比变量快，甚至像blockDim.x这样的自动变量。
每个线程计算一个唯一的idx值，并从点列表中读取相应的点。每个线程都会找到离自己最近的邻居。这与上一版本在第 5 章中执行的操作相同。
接下来我们有一个for循环，使用变量currentBlockOfPoints进行计数。每次循环迭代时，线程都会将全局内存中的一组点读入属于其块的共享内存中。注意我如何在将点从全局内存读入共享内存块后使用__syncthreads，确保在将整个块复制到共享内存之前，没有线程开始检查点到点的距离。
一旦这些点在共享内存中，线程就遍历数据，看看是否有最近的邻居。一旦整个全局内存阵列通过这些块中的共享内存，计算就完成了。注意在for循环之后还有另一个__syncthreads调用。这确保了所有线程在读取新块之前已经检查完共享内存块中的点。
我已经手动使用了一个指针来寻址共享内存点的值(这在名为ptr的变量中)。这通常(但可能不总是)比使用结构寻址语法更快。当执行复杂的计算时，使用小步骤(如清单 7.1 中dist变量的计算)可能也是一个好主意。这对于编译器优化来说更有效，而不是将整个计算放在一行长代码上。当然，像任何代码一样，总是值得尝试看看某个东西是否有效。
这个版本的算法(清单 7.1)的运行速度是第 5 章中介绍的原始 CUDA 版本的 3 倍，是原始 CPU 实现速度的 14 倍(这些时间将根据您的特定 GPU 和 CPU 之间的性能比而有所不同)。很有可能这段代码可以进一步优化，但是对于当前的示例，它已经足够快了。