OpenCL - 多个工作人员可以同时将图像写入同一个坐标吗？

Question

我有一个内核，它可以拍摄图像并根据亮度生成waveform out of it。在这个内核中，工作人员最终不得不将像素写入与其他工作人员完全相同的位置。这安全吗？

在我看来，它看起来不错（英特尔 5000 显卡），但在朋友的带有独立 GPU 的计算机上，它似乎正在泄漏内存或从 GPU 内存中的其他区域获取内存（可能缓存在 GPU 中的随机图像显示在渲染的图像）。这甚至是问题所在吗？开源代码库是GCVideoWaveform。非常感谢您对此的任何想法！

int clampInt(int value, int lowerBound, int upperBound){
    if(value < lowerBound){
        return lowerBound;
    }
    if (value > upperBound){
        return upperBound;
    }
    return value;
}
kernel void waveformLuma(read_only image2d_t input, write_only image2d_t output, uint maxValue) {
    size_t x = get_global_id(0);
    size_t y = get_global_id(1);

    uint4 inputColor = read_imageui(input, sampler, (int2)(x,y));

    float luminance01 = (((float)inputColor.x + (float)inputColor.y + (float)inputColor.z) / 3.0) / (float)maxValue;

    int newY = (get_image_dim(output).y - (luminance01*(float)(get_image_dim(output).y-1)));
    newY = clampInt(newY, 0, get_image_dim(output).y-1);

    write_imageui(output, (int2)(x, newY), (uint4)(maxValue, maxValue, maxValue, maxValue));
}

第一台计算机/集成 GPU

第二台计算机/独立 GPU

Answer 1

write_imageui 和read_imageui 可以被多个线程以相同的参数调用，这没问题。它们与 OpenCL 中的任何其他赋值操作数一样工作。 例如：

__private float b = tanh[(global int *)a[2]]; //like read_image, reads from a global memory
(global int *)a[2] = b; //like write image, writes lo a global memory

注意： 我用了[2]，表示多个WI会写入同一个位置

---

write_imageui 问题不是您遇到的问题，问题是写操作本身。结果完全未知，因为它们严重依赖于 WI 执行顺序。在 CPU 中它提供了很好的结果（并行线程，但一些底层同步），而在 GPU（完全并行设备）上则没有。

但是，由于结果很“奇怪”，我建议您考虑其他问题。我个人并不认为垃圾输出完全是基于这种无序的执行行为。