OpenCL中的深度缓冲区

Question

我想在OpenCL中进行一些自定义渲染。我需要的是一个深度缓冲区。我已经搜索了相当多，但由于编译器优化，许多解决方案都不再有效。

下面是绘制点的基本内核。如果两个点在同一位置，我只想画离相机更近的那个点。

__kernel void render_points(__global const uint4 *points, __global float *zbuffer, __global uint *img)
{
    int i = get_global_id(0);
    uint4 point = points[i];
    int pos = point.y * WIDTH + point.x;

    if (point.z < zbuffer[pos]) {   
        zbuffer[pos] = point.z;
        img[pos] = point.w;
    }
}

然而，这种简单的方法并不起作用，因为由于并行性和缓冲的写回，zbuffer不会立即为所有线程更新。

我正在使用OpenCL 1.2，包括32位原子扩展。

问题

如何实现深度缓冲？

Answer 1

我怀疑问题不在于并行性和原子性，而在于不利的数据类型。

深度缓冲区应为浮点格式，即float (32位)或half (16位)。在3D到2D转换算法中，生成的x和y位置应为int或short类型(uchar会将分辨率限制为256x256)，而生成的z位置应为浮点。这样，当两个点有相似的z位置，比如3.4和3.2，接近的一个仍然可以正确绘制，而如果你使用整数数据类型，这两个点的深度都是3，最后绘制的将决定像素颜色。

另外，对于img缓冲区，我建议使用uint数据类型，以便使用完整的32位颜色。

除了数据类型之外，算法的其余部分也应该有效。下面是我的OpenCL方法的实现，它绘制一个像素并检查z缓冲区。出于性能原因，我对z-buffer使用数据类型half。

void __attribute__((always_inline)) draw(const int x, const int y, const float z, const uint color, global uint* bitmap, global half* zbuffer) {
    if(x<0||x>=2*def_sceen_w||y<0||y>=2*def_sceen_h) return; // cancel drawing if point is off screen
    const int index = y*2*def_sceen_w+x;
    if(z<=vload_half(index,zbuffer)) return; // cancel drawing if point is behind zbuffer
    vstore_half_rte(z,index,zbuffer);
    bitmap[index] = color;
}

Answer 2

一条很长的评论：

对用户函数没有原子性支持。

对于单个整数值，有一个atomic_min函数，它将两个值中较小的一个(1在目标地址中，1在参数中)写入目标地址：

https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/atomic_min.html

int atomic_min (volatile __local int *p ,   int val)

如果将浮点类型的深度值乘以1,000或1000000，则可以使用此基于整数的函数将像素值钳制到后面最接近的多边形。但atomics在全球访问方面速度很慢。如果每个像素没有被许多多边形访问，可能就不会那么慢。但仍然存在对主存的无序访问，这对gpu内存控制器来说是不好的，因为每次访问多个数据的概率很低。要部分修复这个问题，您可以对opencl工作项进行排序(根据它们的线程id和z索引(不是zbuffer，而是屏幕2D位置))，以便“可能”编译器或硬件“连接”(不知道是否可能)多个并行(以及独立的、连续的，如1、2、3、4)原子进行读/写。

但是，由于您希望切换两个32位的值，而不仅仅是深度，因此您应该处理64位整数(64位原子仅在OpenCL 2.0+上)，其最重要的部分是深度(以生成atomic_min)，而另一半(Img)将自动切换。为此，您需要"64位原子“：

https://www.khronos.org/registry/OpenCL/sdk/2.0/docs/man/xhtml/atom_min.html

long atom_min (volatile __local long *p, long val)

只要你保持这个值的前一半(最重要的)作为深度，它应该总是在深度上执行"min“操作，但对于完全相同的两个深度值，它也会用"img”值进行切换。当两个z值完全相同时，这将导致红色、蓝色或绿色中的一个值始终位于zbuffer的顶部。至少这可以解决闪烁问题，如果你还没有使用深度偏移的话。

如果数以千计的多边形位于同一像素后面，原子可能会变得非常慢。最好的情况是每个多边形(或它渲染的任何点)1个像素，并且仅使用功索引的z顺序变换。根据深度对img进行排序应该会使其在性能上成为一个“稳定”的算法，无论场景拓扑如何，您都会看到相同的性能。

但是，如果你以某种方式将屏幕划分为256个方块，并为每个方块分配一个opencl工作组(它将在本地内存上工作，而不是全局内存上)，那么它们将很快使用原子函数。一些架构(如Nvidia Pascal)具有非常好的本地原子性能。他们甚至可能不需要原子，但需要某种波前同步来实现高/可接受的性能。你能把场景分成正方形，然后给每个正方形分配多边形吗？看起来像是直方图(将正方形上的多边形打包)的问题，但你也许可以找到一个更好的解决方案。如果我使用的是GPU或多核，在实现硬件上的东西时，我会尝试分而治之的方法。