0

我正在尝试模拟类型的循环方程

s i (t+1) = f[Σ j W ij s j (t) + v i *input(t)]

在 OpenCL 中，其中 f(.) 是一些非线性函数（在下面的代码中，它只是一个具有阈值 th 的阶跃函数），而 s(t) 是一些外部输入。自然地，我为每个 x i实施了一个工人。在每个时间步中，每个工作人员都会计算上面等式的结果，然后将该结果与所有其他工作人员共享。因此，所有工人必须在同一个工作组中。

我当前的 OpenCL 内核看起来像这样

__kernel void part1(__global int* s, __global float* W, __global float* Th, __global float* V, __global float* input, int N, int T)
    {
        unsigned int i = get_global_id(0);

        float value = 0;
        float v = V[i];
        float th = Th[i];  

        for(int t = 0; t < T; t++){
            value = v*input[t];
            for(int j = 0; j < N; j++){
                value = value + W[i*N + j]*s[j];
            }
            barrier(CLK_GLOBAL_MEM_FENCE);
            if (value >= th){
                s[i] = 1;
            } else {
                s[i] = 0;
            }
            barrier(CLK_GLOBAL_MEM_FENCE);
        }
    }

不幸的是，这段代码实际上比等效的 C 实现慢了三倍。此外，我预计工作人员数量的变化不会产生太大的影响（因为新工作人员坐在与其他工作人员并行运行的新线程上），但实际上处理时间随着工作人员数量线性增加。瓶颈似乎是第一个障碍之后的写入操作。消除此操作（但保留屏障）将处理时间减少了 25 倍并消除了线性相关性。

我对 OpenCL 很陌生，如果能帮助我加快这段代码的速度，我将不胜感激！

提前非常感谢！Blue2script

4

1 回答 1

0

正如我在评论中已经说过的，访问全局内存很慢。通常，硬件通过在同一计算单元上运行多个线程子组来隐藏延迟。我所指的子组是 NVIDIA 术语中的扭曲和 AMD 术语中的波前。通常一个工作组由几个子组组成。

因此，同时一个子组等待从全局内存接收数据，另一个已经拥有所有必要资源的子组可以运行。当正在运行的进程因为需要从/向全局内存读取/写入数据而停止运行时，另一个进程可以开始运行，依此类推。

但是，在您的情况下，由于障碍，所有子组中的所有工作人员都必须在能够继续计算之前将其他工作人员写入内存（障碍在工作组级别）。因此，延迟会直接打到你的脸上:)。

现在，改进实现的一种方法是使用本地内存，这次是本地内存级别的屏障（使用标志 CLK_LOCAL_MEM_FENCE）。我刚刚解释的相同原理仍然适用，但访问本地内存要快得多。

据我了解你的代码（很可能我没有得到所有的细节），你的 s 数组有 N 个元素，我猜你也有 N 个工人。因此，您创建了一个包含 N 个元素的本地数组，然后执行以下操作：

kernel void part1(global int* s, global float* W, global float* Th, global float* V, global float* input, local int* local_s, int N, int T)
    {
        unsigned int i = get_global_id(0);
        unsigned int local_i = get_local_id(0);

        float value = 0;
        float v = V[i];
        float th = Th[i];  
        //fetch from global to local and sync before computing
        local_s[local_i] = s[i];
        barrier(CLK_LOCAL_MEM_FENCE);

        for(int t = 0; t < T; t++){
            value = v*input[t];
            for(int j = 0; j < N; j++){
                value = value + W[i*N + j]*local_s[j];
            }
            barrier(CLK_LOCAL_MEM_FENCE);
            if (value >= th){
                local_s[i] = 1;
            } else {
                local_s[i] = 0;
            }
            barrier(CLK_LOCAL_MEM_FENCE);
        }
    //If necessary write some stuff to global (maybe the last s computed?)
    }

现在我必须警告你：

  • 我可能完全误解了您的需求:)
  • 我刚刚在输入此答案时编辑了您的代码，因此很可能存在拼写错误等。
  • 即使使用本地内存，有这么多障碍仍然可能使 opencl 版本比 C 版本慢。

请注意，我删除了前导 __，因为它们不是必需的，而且我认为它更容易阅读。

编辑：关于您对 CLK_LOCAL_MEM_FENCE 与 CLK_GLOBAL_MEM_FENCE 的评论。障碍总是在工作组级别应用，因此工作组内的所有工作人员都必须达到该障碍。作为参数给出的标志是指内存访问。当标志为 CLK_GLOBAL_MEM_FENCE 时，这意味着所有关于全局内存的读/写操作都必须由每个工作人员完成，然后任何工作人员才能继续运行下一条语句。这与 CLK_LOCAL_MEM_FENCE 标志完全相同，但用于本地内存。

于 2013-07-27T10:12:21.653 回答