1

我目前正在学习 OpenCL，并且我的内核在直接访问全局数组时工作得很好，但是在私有内存上使用中间值时会给出错误的结果，例如，下面代码中的aux。

__kernel void kernel_cte(__global float *U0,__global float *U1,__constant float *VP0, uint stride, uint nnoi, __constant float *g_W, uint k0, uint k1, float FATMDFX, float FATMDFY, float FATMDFZ) {

uint index = get_global_id(1)*nnoi + get_global_id(0) + k0 * stride;

uint k;
float aux;
aux = U0[index+1];

for(k=k0;k<k1;++k) {
    if(VP0[index] > 0.0f){
      U1[index] = 2.0f * U0[index] - U1[index]
        + FATMDFX * VP0[index] * VP0[index] * (
          + g_W[6] * (U0[index - 6] + U0[index + 6])
          + g_W[5] * (U0[index - 5] + U0[index + 5])
          + g_W[4] * (U0[index - 4] + U0[index + 4])
          + g_W[3] * (U0[index - 3] + U0[index + 3])
          + g_W[2] * (U0[index - 2] + U0[index + 2])
          + g_W[1] * (U0[index - 1] + aux)
          + g_W[0] * U0[index]
        )
        + FATMDFY * VP0[index] * VP0[index] * (
          + g_W[6] * (U0[index - 6 * nnoi] + U0[index + 6 * nnoi])
          + g_W[5] * (U0[index - 5 * nnoi] + U0[index + 5 * nnoi])
          + g_W[4] * (U0[index - 4 * nnoi] + U0[index + 4 * nnoi])
          + g_W[3] * (U0[index - 3 * nnoi] + U0[index + 3 * nnoi])
          + g_W[2] * (U0[index - 2 * nnoi] + U0[index + 2 * nnoi])
          + g_W[1] * (U0[index -     nnoi] + U0[index +     nnoi])
          + g_W[0] * U0[index]
        )
        + FATMDFZ * VP0[index] * VP0[index] * (
          + g_W[6] * (U0[index + 6 * stride] + U0[index - 6 * stride])
          + g_W[5] * (U0[index + 5 * stride] + U0[index - 5 * stride])
          + g_W[4] * (U0[index + 4 * stride] + U0[index - 4 * stride])
          + g_W[3] * (U0[index + 3 * stride] + U0[index - 3 * stride])
          + g_W[2] * (U0[index + 2 * stride] + U0[index - 2 * stride])
          + g_W[1] * (U0[index +     stride] + U0[index -     stride])
          + g_W[0] * U0[index]
        );
    } // end if
    index += stride;
}
}

我想使用向量来执行这些计算，但我不明白为什么当我执行 aux = U0[index+1] 时没有将正确的值复制到私有内存中。

4

2 回答 2

0

如果每个工作项都在自己的数据集上工作，那么他们只需要使用栅栏提交全局内存操作（如果他们正在使用它们并在同一个内核中多次更改它们）。

例如，U1[index]在下面的代码中，如果不打算缓存，则需要提交到全局内存。

mem_fence(CLK_GLOBAL_MEM_FENCE);
if(VP0[index] > 0.0f){
      U1[index] = 2.0f * U0[index] - U1[index]
        + FATMDFX * VP0[index] * VP0[index] * (
          + g_W[6] * (U0[index - 6] + U0[index + 6])
          + g_W[5] * (U0[index - 5] + U0[index + 5])
          + g_W[4] * (U0[index - 4] + U0[index + 4])
          + g_W[3] * (U0[index - 3] + U0[index + 3])
          + g_W[2] * (U0[index - 2] + U0[index + 2])
          + g_W[1] * (U0[index - 1] + aux)
          + g_W[0] * U0[index]
        )
        + FATMDFY * VP0[index] * VP0[index] * (
          + g_W[6] * (U0[index - 6 * nnoi] + U0[index + 6 * nnoi])
          + g_W[5] * (U0[index - 5 * nnoi] + U0[index + 5 * nnoi])
          + g_W[4] * (U0[index - 4 * nnoi] + U0[index + 4 * nnoi])
          + g_W[3] * (U0[index - 3 * nnoi] + U0[index + 3 * nnoi])
          + g_W[2] * (U0[index - 2 * nnoi] + U0[index + 2 * nnoi])
          + g_W[1] * (U0[index -     nnoi] + U0[index +     nnoi])
          + g_W[0] * U0[index]
        )
        + FATMDFZ * VP0[index] * VP0[index] * (
          + g_W[6] * (U0[index + 6 * stride] + U0[index - 6 * stride])
          + g_W[5] * (U0[index + 5 * stride] + U0[index - 5 * stride])
          + g_W[4] * (U0[index + 4 * stride] + U0[index - 4 * stride])
          + g_W[3] * (U0[index + 3 * stride] + U0[index - 3 * stride])
          + g_W[2] * (U0[index + 2 * stride] + U0[index - 2 * stride])
          + g_W[1] * (U0[index +     stride] + U0[index -     stride])
          + g_W[0] * U0[index]
        );
mem_fence(CLK_GLOBAL_MEM_FENCE);

因为 GPU 乱序指令执行能力或编译器可以在不询问的情况下重新排序读/写，而栅栏/屏障正在阻止它们这样做并按照开发人员需要的方式保持顺序。

如果工作项旨在更改彼此的数据区域，则至少需要 barrier() 并且这仅在每个块（工作组）内部有效。

于 2017-12-10T11:31:42.560 回答
0

我发现了问题，很明显，读取 aux = U0[index+1] 应该在 for 循环内执行。

于 2017-12-10T21:25:02.020 回答