一个特定的OpenCL内核在移动设备和PC上表现不同。

Question

我试图在Adreno 630和我的笔记本电脑上运行一个OpenCL内核，结果发现内核在手机上运行得完美无缺，但每次都会导致我的笔记本电脑崩溃。我仍在尝试自己找出原因。这是我的内核。希望您能帮助我解决问题，谢谢。

__kernel void gen_mapxy( __read_only image2d_t _disp, const float offsetX, __write_only image2d_t _mapxy )
{
    const int y = get_global_id(0);

    const int local_y = get_local_id(0);
    __local short temp[24][1080];

    const int imageWidth = get_image_width(_disp);
    for(int x = 0; x < imageWidth; ++x)
        temp[local_y][x] = 0;
    for(int x = imageWidth - 1; x >= 0; --x){
        int tempDisp = read_imagei(_disp, sampler_nearest, (int2)(x, y)).x;
        int newPos = clamp((int)(x + offsetX * (tempDisp) / 255), 0, imageWidth - 1);
        temp[local_y][newPos] = tempDisp;
        write_imagef(_mapxy, (int2)(newPos, y), (float4)(x, y, 0, 0));
}

- wen dai

---

也许你的输入在计算机代码中没有很好地设置。检查一下你的输入是否正确，并且具有正确的尺寸。你可能在手机上使用不同的输入图像/数据，这可能是原因。 - Nikos M.

你漏掉了第二个for循环的闭合}。 - Dithermaster

Answer 1

您正在使用一个很大的本地数组。

最初的回答：

__local short temp[24][1080]

2字节 * 24 * 1080 = 50.6kB。一些台式机GPU（以及其笔记本电脑对应物）具有较少的可用本地内存限制。例如，GTX 1060支持值CL_DEVICE_LOCAL_MEM_SIZE 49152字节。但是Adreno 620可能会在全局数组中模拟局部数组（限于数百兆字节），因此要么默默地忽略数组使用，要么支持更大的局部数组。如果它们确实支持芯片内快速本地内存，则存在更多“忽略”问题或它们确实将上一代Adrenos的本地内存限制加倍。

即使GPU支持确切的值，使用全部会严重降低每个管道的线程级并行性，从而严重降低潜在的性能增益。

如果上一代Adreno GPU相同，

https://compubench.com/device.jsp?benchmark=compu15m&os=Android&api=cs&D=Samsung+Galaxy+S7+%28SM-G930x%29&testgroup=info

此页面说：

CL_DEVICE_LOCAL_MEM_SIZE
32768

CL_DEVICE_LOCAL_MEM_TYPE
CL_LOCAL

它很快，但只有32kB，因此它可能会忽略错误，或者您在其中漏掉了必要的错误捕获逻辑，或两者都有。

"Original Answer"翻译成中文是"最初的回答"

- huseyin tugrul buyukisik

---

我同意这一点，也许在电脑和手机之间存在一些不同的配置。但实际上，在手机上当行大小小于24时，我总是得到正确的输出结果；当我试图继续增加它（例如32），每次在手机上都会崩溃。这仍然很奇怪。 - wen dai

有时候驱动程序会在OpenCL上导致一些奇怪的问题。可能是移动SoC GPU供应商没有像AMD、NVIDIA和INTEL那样给予足够的关注。 - huseyin tugrul buyukisik