c++ - OpenCL:输出可变长度的数组

Question

我们正在处理 GPGPU 类(class)的作业。我们选择了一种算法，在 CPU 上实现了它，现在正在将其转换为 OpenCL。

我们选择的算法将模型加载为一组三角形并将它们光栅化为体素。体素被定义为点数据的 VBO。然后，我们使用几何着色器将这些点转换为三角形体素。

因此我们的 OpenCL 程序需要获取三角形列表并输出可变点列表。

而且输出可变长度数组似乎是个问题。

我们找到的解决方案是自动递增一个计数器并将该计数器用作输出数组的索引和数组的最终大小。除了...我们的 GPU 都不支持原子操作的扩展。

这是我们目前所拥有的:

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable

#define POS1      i0 * 3 + 0
#define POS2      i0 * 3 + 1
#define POS3      i0 * 3 + 2

void WritePosition( __global float* OutBuffer, uint inIndex, __global float* inPosition )
{
    OutBuffer[ inIndex * 3 ] = inPosition[0];
    OutBuffer[ inIndex * 3 + 1] = inPosition[1];
    OutBuffer[ inIndex * 3 + 2] = inPosition[2];
}

__kernel void Voxelize( 
    __global float* outPointcloudBuffer, 
    __global float* inTriangleBuffer, 
    __global uint* inoutIndex
)
{
    size_t i0 = get_global_id(0);
    size_t i1 = get_local_id(0);

    WritePosition( outPointcloudBuffer, inIndex[0], &inTriangleBuffer[ i0 ] );

    //atomic_inc(inoutIndex[0]);
    inoutIndex[0] = max(inoutIndex[0], i0);
}

而且这个输出很奇怪。我们正在测试一个非常小的模型(12 个三角形、36 个位置、108 个 float )，我们得到的结果是 31、63 或 95。总是 16 减 1 的倍数。

我们如何获得可变长度输出数组的长度？

提前致谢。

Answer 1

我猜这通常是按如下方式解决的:

• 第一步:使用 scan 在 GPU 上计算所需的数组大小(并行前缀和)原语。以上链接包含来自 Apple 的示例实现。
• 使用扫描算法的结果在主机端分配所需的资源。请注意，扫描算法的结果通常可以用作单个工作项结果的索引提示。
• 第二遍(可选):将数组压缩为第三遍需要考虑的元素。
• 第三遍:重新运行算法，传递目标索引和分配的数组。

您可能想看看 NVIDIA 的 OpenCL 行进立方体 implementation上面提到的所有三个 channel 都实现了。

最好的，克里斯托夫