如何将图像数据从搜索窗口中移动到本地存储器OpenCL

How to move image data from within a searching window to local memory OpenCL

本文关键字：移动 OpenCL 存储器 窗口 图像 数据 搜索      更新时间：2023-10-16

我已经使用opencl并行实现了详尽的块匹配算法，现在正在尝试通过将搜索窗口移至本地内存来优化算法。到目前为止，我拥有的代码如下：

//loop through whole search space and move to local
    for (int i = -searchWindow-blockSize; i <= searchWindow+blockSize; i++) {
        for (int j = -searchWindow-blockSize; j <= searchWindow+blockSize; j++) {
            tgid = (cache[lid].x + i) + (cache[lid].y + j) * imWidth;
            nlid = (blockSize+searchWindow + i) + (blockSize+searchWindow + j) * ((searchWindow+blockSize) * 2 + 1);
            prevCache[nlid] = prevFrame[tgid]; 
            nextCache[nlid] = nextFrame[tgid];
            barrier(CLK_LOCAL_MEM_FENCE);
        }
    }

嵌套在搜索窗口上的循环循环，并且由于图像数据为1维，我需要通过执行x y *宽度来将其转换为1D标识符，以获取TGID和NLID。缓存[LID]是float2类型，该类型存储了整个图像中内核位置的当前参考点的X和Y坐标。LID是get_local_id(0(的本地工作项ID，因此，我正在从整个图像中的参考点附近的搜索窗口中获取数据，并将其移至本地" PrevCache"answers" NextCache"，然后可以执行这些数据我的块匹配算法。

我遇到的问题是，正确的数据并不总是分配给缓存。为了进行测试，我制作了PrevCache和NextCache仅从NextFrame中存储相同的数据，如以下代码：

//loop through whole search space and move to local
    for (int i = -searchWindow-blockSize; i <= searchWindow+blockSize; i++) {
        for (int j = -searchWindow-blockSize; j <= searchWindow+blockSize; j++) {
            tgid = (cache[lid].x + i) + (cache[lid].y + j) * imWidth;
            nlid = (blockSize+searchWindow + i) + (blockSize+searchWindow + j) * ((searchWindow+blockSize) * 2 + 1);
            prevCache[nlid] = nextFrame[tgid]; 
            nextCache[nlid] = nextFrame[tgid];
            barrier(CLK_LOCAL_MEM_FENCE);
        }
    }
    if (prevCache[220] != nextCache[220])
        printf("test");

此如果底部的语句应始终为false，因为PrevCache和NextCache应包含相同的数据，但是，当我运行代码时，"测试"将几次打印到控制台。在我看来，这是一个同步问题或某些问题，就像我将缓存[lid] .x和cache [lid] .y更改为固定的数字时，"测试"从未打印出来，但此时我毫无头绪。任何帮助将不胜感激！

  1. 关于本地记忆以进行优化，让我在这里粘贴一些论文，以便您考虑：

本地内存较小，但比全局内存更快。每个计算 单元具有自己的本地记忆，每个内存都可以访问 该计算单元中的处理元素。内核可以触发 从全局内存到本地的数据同步数据的批量加载 记忆。这允许快速有效地进行处理 大部分数据的元素。这可以在某些情况下证明 对运行时有益。让我们假设批量的时间 将大小s的大量数据加载到本地内存中，需要T_S时间 单位。每个工作单位都会有k随机访问读或写操作 在这大部分数据中。一个读/写动作所需的时间 对于本地记忆，表示为T_L，以及一个读取/所需的时间 来自全局内存缓存的写作动作表示为t_g。我们猜测 关系t_l＆lt;T_G保持。加载大量的全球记忆 如果读取的数量/ 写动作K足够大以满足：

k·t_g> t_s k·t_l

  1. 一般而言，最好使用OpenCl提供的同步散装操作Async_work_group_copy来复制从全局到本地内存，而不是用Word手动复制Word。

  2. 关于您的代码和问题：

工作单元中的处理元素写入相同的本地内存片，但从全局内存的不同点开始，有效地覆盖了数据。证明：

如果您的工作组大小至少为2个，则有一个lid = 0的元素，一个lid = 1。

。

如果cache [0] .x！= cache1.x或cache [0] .y！= cache1.y，那么TGID对于这两个处理元素具有不同的值。

，只要(i，j(在处理元素中同步(您保证使用本地内存围栏(，您的索引NLID对于所有处理元素都是恒定的。

。
  1. 最后，关于嵌套环的注释。如果您真的想这样做，请尝试一种读取一口气结合的内存的方法。假设您的形象是这样组织的：
   1    2    3 ...  w
 w+1  w+2  w+3 ... 2w
2w+1 2w+2 2w+3 ... 3w
....

如果您想读取从左上角开始的3 x-3窗口，则应像这样构建循环

for(int row = start_pos.y; row < start_pos.y+3; row++){
   for(int col = start_pos.x; col < start_pos.x+3; col++){
      tmp = data[col+w*row];
   }
   barrier(CLK_LOCAL_MEM_FENCE);
}

这样，您以此顺序阅读1,2,3，W 1，W 2，W 3,2W 1,2W 1,2W 2,2W 3

以这种方式，OpenCL将缓存全局内存，然后您的处理元素将从Cache中读取1,2,3，OpenCL将缓存下一个全局内存的块，而您的内核将读取下一个单词，依此类推。

>

您编写的方式将以此顺序读取数据：1，W 1,2W 1,2，W 2,2W 2,3，W 3,2W 3

对于此操作，将再次缓存全局内存，您的处理元素读取了他们的第一个数据字，并且OpenCL将再次开始缓存，然后您的PES将读取其第二个数据字，依此类推。这将导致将相同的三块全局内存加载到全局内存缓存中，当加载每个内存一次就足够了。

相关文章：