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最后绘制的 on 将决定像素颜色。

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

除了您的数据类型之外，您的算法的其余部分应该可以工作。下面是我对绘制一个像素并检查 z 缓冲区的 OpenCL 方法的实现。出于性能原因，我将数据类型 half 用于 z 缓冲区。

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)

如果您将浮点型深度值乘以 1000 或 1000000，您可以使用这个基于整数的函数将像素值限制在它们后面最近的多边形上。但是原子在全球访问上很慢。如果每个像素不被许多多边形访问，也许不会那么慢。但是仍然会有对主内存的无序访问，这对 gpu 内存控制器不利，因为每次访问多数据的概率很低。要部分解决此问题，您可以对 opencl 工作项进行排序（在它们的线程 id 与它们的 z-index （不是 zbuffer 而是屏幕 2D 位置）上），以便“可能”编译器或硬件“加入”（不知道是否可能）多个并行（和独立的、连续的，如 1、2、3、4）原子一起用于读/写。

但是，由于您想要切换两个 32 位值而不仅仅是深度，您应该使用 64 位整数（64 位原子仅在 OpenCL 2.0+ 上），其最重要的部分是深度（用于生成 atomic_min）和其他half(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)

只要您将此值的前半部分（最重要的一半）作为深度，它应该始终在深度上执行“最小”操作，但对于完全相同的两个深度值，它也会使用“img”进行切换“ 价值。当两个 z 值完全相同时，这将导致红色或蓝色或绿色值之一始终位于 zbuffer 的顶部。如果您还没有使用深度偏差，至少这可以解决闪烁问题。

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

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