OpenCL：输出可变长度数组

我们正在为GPGPU课程进行分配。我们选择了一种算法，在CPU上实现，现在将其转换为OpenCL。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.

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

在此先感谢。

来源

2011-12-20 knight666

您的结果是16N-1，因为你在整个warp上运行内核。要修复，请将三角形的总数传递给内核。如果'global_id'大于三角形数量，则返回。这样，您只能运行与三角形一样多的内核。 – 2011-12-20 18:33:58

我猜想，这通常解决如下：

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

您可能想看看NVIDIA的OpenCL行军立方体implementation，上面提到的所有三个通道都已实现。

Best，Christoph

来源

2011-12-20 09:24:36 cheind

OpenCL：输出可变长度数组

回答

相关问题