OpenCL：输出可变长度的数组答案

【问题标题】：OpenCL: Outputting an array of variable lengthOpenCL：输出可变长度的数组
【发布时间】：2011-12-20 17:24:36
【问题描述】：

我们正在处理 GPGPU 课程的作业。我们选择了一种算法，在 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 个浮点数），我们得到的结果是 31、63 或 95。始终是 16 减 1 的倍数。

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

提前致谢。

【问题讨论】：

您的结果是 16N - 1，因为您在整个 warp 上运行内核。要修复，请将三角形的总数传递给内核。如果global_id 大于三角形计数，则返回。这样，您只需运行与三角形一样多的内核。

标签： c++ arrays opencl gpgpu

【解决方案1】：

我猜这通常是这样处理的：

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

您可能想看看 NVIDIA 的 OpenCL 行进立方体 implementation，其中实现了上述所有三个通道。

最好的，克里斯托夫

【讨论】：