对相邻数组元素使用共享内存

Use shared memory for neighboring array elements?

本文关键字：共享内存数组元素更新时间：2023-10-16

我想用 CUDA 处理图像。每个像素的新值是根据一行中的两个相邻像素计算的。对像素值使用__shared__内存是否有意义，因为每个值只会使用两次？瓷砖不是也是错误的方法吗，因为它不适合问题结构？我的方法是在每个像素上运行一个线程，并每次为每个线程加载相邻的像素值。

所有当前支持的 CUDA 架构都有 caches.
从计算能力 3.5 开始，这些对于只读数据特别有效（由于读写数据仅缓存在 L2 中，因此 L1 缓存仅限于只读数据）。如果将指向输入数据的指针标记为const __restrict__，编译器很可能会通过 L1 纹理缓存加载它。您还可以通过使用内置的__ldg()来强制执行此操作。

虽然可以通过共享内存显式管理来自相邻像素的数据重用，但您可能会发现这比仅依赖缓存没有任何好处。

当然，无论是否使用共享内存，您都希望在 x 方向上最大化块大小，并使用 1 blockDim.y 来实现最佳访问局部。

结合使用共享内存和利用合并的内存访问。您需要做的就是确保图像按行存储。每个块将处理一个线性数组块。由于数据重用（除了第一个和最后一个像素之外的每个像素都会参与处理三次），如果在内核开始时将要处理的所有像素的值复制到共享内存中，那将是有益的。