具有许多 bin 的 OpenCL 直方图

Question

我正在使用 OpenCL 编程指南第 14 章中提供的代码来计算直方图。它适用于 256 个箱子，但不幸的是我的应用程序需要 65536 个箱子。这会导致问题，如果我使用这种方法，本地数组会变得太大。

local uint tmp_histogram[256 * 256];

因此，程序未构建 (CL_BUILD_PROGRAM_FAILURE)。

您知道如何解决这个问题吗？我考虑过使用多个内核来计算不同 bin 的值（即拆分直方图，以便我首先计算 bin 0-255 的值，然后计算 256-511 等）。但是，在这种情况下，我必须在递增之前检查值是否在该范围内，这意味着我将需要条件...

Answer 1

将您的直方图移至全局存储。如果您的应用程序适合这种大小，另一种解决方案可能是使用 unsigned short。最后你可以运行你的代码两次。第一次为较低的 32000 值，第二次为上半部分。

Answer 2

使用全局内存可以解决问题，但不会产生非常快的内核。我建议创建多个工作组，并使用每个组只计算一个范围内的值。

#define RANGE_SIZE 8192

kernel void histo(__global uint data,__constant int dataSize){
    int wid = get_local_id(0);
    int wSize = get_local_size(0);

    int gid = get_group_id(0);
    int numGroups = get_num_groups(0);

    int rangeStart = gid * RANGE_SIZE / numGroups;
    int rangeEnd = (gid+1) * RANGE_SIZE / numGroups;

    local uint tmp_histogram[RANGE_SIZE];

    uint value;

    for(int i=wid; i< dataSize; i+= wSize){
        value = data[i];
        if(value >= rangeStart && value < rangeEnd){
            atomic_inc(tmp_histogram[value - rangeStart]);
        }
    }
    //barrier...
    //use the local data here
}

假设有 32kb 的本地内存可用。如果你减少 RANGE_SIZE，它不必是 2 的幂，但你确实需要确保你调用内核有足够的工作组来命中所有高达 64k 的值。

具有许多 bin 的 OpenCL 直方图

OpenCL histogram with many bins

opencl

gpu

histogram