我想用CUDA处理图像。基于一行中的两个相邻像素来计算每个像素的新值。因为每个值只使用两次，所以使用__shared__内存存储像素值是否有意义？瓷砖不也是错误的方法吗，因为它不适合问题结构？我的方法是在每个像素上运行一个线程，并每次为每个线程加载相邻的像素值。
__shared__
nbnkbykc1#
所有当前支持的CUDA架构都有缓存。从计算能力3.5开始，这些对于只读数据特别有效（因为读写数据仅缓存在L2中，所以L1缓存仅限于只读数据）。如果将指向输入数据的指针标记为const __restrict__，则编译器很可能会通过L1纹理缓存加载它。您也可以通过显式使用__ldg()内置函数来强制执行此操作。虽然可以通过共享内存显式地管理相邻像素的数据重用，但您可能会发现这与仅依赖该高速缓存相比没有任何好处。当然，无论您是否使用共享内存，您都希望最大化x方向的块大小，并使用blockDim.y为1以获得最佳访问局部性。
const __restrict__
__ldg()
blockDim.y
6l7fqoea2#
合并使用共享内存和利用合并内存访问。您需要做的就是确保图像按行存储。每个块将处理一个线性数组块。由于数据重用（除了第一个和最后一个像素之外的每个像素都将参与处理三次），如果在内核开始时将所有将被处理的像素的值复制到共享内存中，这将是有益的。
2条答案
按热度按时间nbnkbykc1#
所有当前支持的CUDA架构都有缓存。
从计算能力3.5开始，这些对于只读数据特别有效（因为读写数据仅缓存在L2中，所以L1缓存仅限于只读数据）。如果将指向输入数据的指针标记为
const __restrict__，则编译器很可能会通过L1纹理缓存加载它。您也可以通过显式使用__ldg()内置函数来强制执行此操作。虽然可以通过共享内存显式地管理相邻像素的数据重用，但您可能会发现这与仅依赖该高速缓存相比没有任何好处。
当然，无论您是否使用共享内存，您都希望最大化x方向的块大小，并使用
blockDim.y为1以获得最佳访问局部性。6l7fqoea2#
合并使用共享内存和利用合并内存访问。您需要做的就是确保图像按行存储。每个块将处理一个线性数组块。由于数据重用（除了第一个和最后一个像素之外的每个像素都将参与处理三次），如果在内核开始时将所有将被处理的像素的值复制到共享内存中，这将是有益的。