5

我是 OpenCL 的新手。

我想编写一个通用内核，以便稍后我可以将其使用扩展到其他内存非合并模式并将其与Rectangular stencil pattern简单起见配对（也避免越界访问）。

该内核控制本地内存 ( __local float ∗lmem) 的使用。

截至目前，我的.cl文件结构如下：

__kernel void kmain (
    __global float ∗in ,
    __global float ∗out ,
    __global float ∗in2 ,
    __local float ∗lmem)
{
    int wg_x = get group id(0);
    int wg_y = get group id(1);
    int wi_x = get local id(0);
    int wi_y = get local id(1);     

    // number of work units each work-item processes
    for (int iter_x = 0; iter_x< NUM_WUS_X-1, iter_x++ ) {
        for (int iter_y = 0; iter_y< NUM_WUS_Y-1; iter_x++) {
            int wu_x, wu_y;

            // The current work unit coordinate (wu_x, wu_y) is computed based on work group ID (wg_x, wg_y), work item ID (wi_x, wi_y) and work unit ID (iter_x, iter_y) :
            (wu_x, wu_y) = func(wg_x, wg_y
                          wi_x, wi_y,
                          iter_x ,iter_y);

        // This is where to cooperatively load
        // a region of <in> to the local memory.
        // barrier (...);

            for (int i = 0; i < N-1, i++) {
                for (int j = 0; j< M-1, j++) {

                // (fo, fi) detemines the home access pattern centered around (idx_o, idx_i). WI(*,*) defines the memory access pattern i.e: (wi_x) = (wi_y) :
                int idx_o = fo(wu_x, wu_y, i, j);
                int idx_i = fi(wu_x, wu_y, i, j);

                // offsets CO's and CI's  determine stencil pattern within each work-item
                ... = in[idx_o + CO_1][idx_i + CI_1];
                ... // context (inner loop body)
                ... = in[idx_o + CO_k][idx_i + CI_k];
                ... // context (inner loop body)
            }
        }
        // barrier (...);
        ... // context (epilogue)
        out[y][x] = ...;
        }
    }
}

有没有人对使用相应的通用主机实现这种模式有任何想法？

4

1 回答 1

0

您可以在 OpenCL 绑定上开发主机端封装，这样，

  • 它需要一些来自用户的通用代码字符串
  • 它在运行时使用用户字符串生成内核，也取决于您选择重新塑造内核的“策略”
  • 内核字符串还定义了一个名为“scratch_pad”的自定义资源
    • 它具有可选的内存类型（本地/全局/寄存器/常量）
    • 由必要的输入自动填充
    • 这受用户最大尺寸的限制
    • 它具有[]针对不同内存类型和数据模式的不同访问模式的大量运算符实现
  • 自动将相关的主机端缓冲区绑定到内核端缓冲区

然后您可以简单地更改参数列表中的标志，以针对本地与全局内存性能对其进行测试，或者为其提供不同的内核字符串，但这并不比简单地编写不同的 cl 文件更容易。如果您只有几个不同的实现，看起来工作量太大。当您不知道结果内核字符串时，调试也会变得更加困难。

抱歉迟了回应。

于 2019-03-22T14:13:17.493 回答