对相邻的数组元素使用共享内存?
Use shared memory for neighboring array elements?
我想用 CUDA 处理图像。每个像素的新值是根据一行中的两个相邻像素计算的。为像素值使用 __shared__
内存是否有意义,因为每个值只会使用两次?瓷砖不是也是错误的方法吗,因为它不适合问题结构?我的方法是 运行 每个像素上的线程,并每次为每个线程加载相邻像素值。
结合使用共享内存和利用联合内存访问。您需要做的就是确保图像按行存储。每个块将处理一大块线性阵列。由于数据重用(除了第一个和最后一个像素之外的每个像素都将参与处理三次),如果在内核开始时将所有将要处理的像素的值复制到共享内存,那将是有益的。
所有当前支持的 CUDA 架构都有缓存。
从计算能力 3.5 开始,这些对于只读数据特别有效(因为读写数据仅缓存在 L2 中,L1 缓存仅限于只读数据)。如果将指向输入数据的指针标记为 const __restrict__
,编译器很可能会加载它 via the L1 texture cache. You can also force this by explicitly using the __ldg()
builtin.
虽然可以通过共享内存显式管理相邻像素数据的重用,但您可能会发现这与仅依赖缓存相比没有任何好处。
当然,无论您是否使用共享内存,您都希望在 x 方向上最大化块大小并使用 blockSize.y 1 以获得最佳访问位置。
我想用 CUDA 处理图像。每个像素的新值是根据一行中的两个相邻像素计算的。为像素值使用 __shared__
内存是否有意义,因为每个值只会使用两次?瓷砖不是也是错误的方法吗,因为它不适合问题结构?我的方法是 运行 每个像素上的线程,并每次为每个线程加载相邻像素值。
结合使用共享内存和利用联合内存访问。您需要做的就是确保图像按行存储。每个块将处理一大块线性阵列。由于数据重用(除了第一个和最后一个像素之外的每个像素都将参与处理三次),如果在内核开始时将所有将要处理的像素的值复制到共享内存,那将是有益的。
所有当前支持的 CUDA 架构都有缓存。
从计算能力 3.5 开始,这些对于只读数据特别有效(因为读写数据仅缓存在 L2 中,L1 缓存仅限于只读数据)。如果将指向输入数据的指针标记为 const __restrict__
,编译器很可能会加载它 via the L1 texture cache. You can also force this by explicitly using the __ldg()
builtin.
虽然可以通过共享内存显式管理相邻像素数据的重用,但您可能会发现这与仅依赖缓存相比没有任何好处。
当然,无论您是否使用共享内存,您都希望在 x 方向上最大化块大小并使用 blockSize.y 1 以获得最佳访问位置。