将 async_work_group_copy() 与指针一起使用？

Question

__kernel void kmp(__global char pattern[1*4], __global char* string, 
                  __global int failure[1*4], __global int ret[1], 
                  int g_length, int l_length, int thread_num){
    int pattern_num = 1;
    int pattern_size = 4;
    int gid = get_group_id(0);
    int glid = get_global_id(0);
    int lid = get_local_id(0);
    int i, j, x = 0;

    __local char *tmp_string;
    event_t event;

    if(l_length < pattern_size){
            return;
    }

    event = async_work_group_copy(tmp_string, string+gid*g_length, g_length, 0);
    wait_group_events(1, &event);

这些是我的代码的一部分。

我想在文本中找到匹配的模式。

首先，在 CPU 端初始化所有模式和字符串（我从文本中读取字符串并实验性地仅使用一种模式）。

其次，将它们传输到名为

kmp

的内核 （参数

l_length

和

g_length

是要复制到

lid

和

glid

的字符串的大小。换句话说，就是字符串的片段）。

最后，我想将分割后的字符串复制到本地内存。

但是有一个问题。当我使用

async_work_group_copy()

复制它们时，我无法得到任何有效结果。

当我将

__local char*tmp_string

更改为数组时，问题仍然存在。

我想做的是

1. 分割字符串
2. 将它们复制到每个线程
3. 并计算匹配数。

我想知道这段代码有什么问题。

Answer 1

OpenCL 规范是这样的：

异步复制由工作组中的所有工作项执行，这 因此，内置函数必须被一个中的所有工作项所遇到。 使用相同参数值执行内核的工作组； 否则结果未定义。

因此您不应该提前返回小组中的任何工作项目。无论如何，提前返回更适合CPU。如果这是 GPU，只需使用增强/填充的输入输出缓冲区来计算最后溢出的部分。

否则，您可以提前返回整个组（这应该可以工作，因为没有工作项命中任何异步复制指令）并在 cpu 上完成剩余的工作，除非设备不使用任何工作项（而是专用的秘密管道）进行异步复制操作。

也许您可以将第二个内核（同时在另一个队列中）加入队列，以使用 workgroupsize=remaining_size 来计算剩余的最新项目，而不是使用额外的缓冲区大小或控制逻辑。

---

如果您要向/从其中复制某些内容，则需要初始化/分配

tmp_string

。所以你可能需要它的数组版本。

---

async_work_group_copy

不是同步点，因此在它之前需要一个屏障才能获取本地内存的最新位以用于异步复制到全局。

            __kernel void foo(__global int *a, __global int *b)
            {
                int i=get_global_id(0);
                int g=get_group_id(0);
                int l=get_local_id(0);
                int gs=get_local_size(0);
                __local int tmp[256];
                event_t evt=async_work_group_copy(tmp,&a[g*gs],gs,0);
                // compute foobar here in async to copies
                wait_group_events(1,&evt);

                tmp[l]=tmp[l]+3; // compute foobar2 using local memory
                barrier(CLK_LOCAL_MEM_FENCE);

                event_t evt2=async_work_group_copy(&b[g*gs],tmp,gs,0);
                // compute foobar3 here in async to copies
                wait_group_events(1,&evt2);
            }