从全局内存复制到私有内存时值错误答案

【问题标题】：Wrong value when copying from global to private memory从全局内存复制到私有内存时值错误
【发布时间】：2017-12-10 03:58:34
【问题描述】：

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

__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] 时没有将正确的值复制到私有内存中。

【问题讨论】：

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

标签： opencl gpgpu opencl-c

【解决方案1】：

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

例如，下面代码中的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() 并且这仅适用于每个块（工作组）内部。

【讨论】：

【解决方案2】：

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

【讨论】：