CUDA中一个简单Z缓冲区的实现

Question

我有一个3D点云，我将像素投影到图像平面。由于一些3D点被映射到同一个像素，我只想要相机中Z值最低的像素。我使用Z-Buffer(一个浮点数组)来跟踪我的深度值。这里有一些伪代码：

// Initialize z-Buffer with max depth (99999.9f)
// Go through every point of point cloud...
// Project Point (x,y,z) to image plane (u,v) 
int newIndex = v*imgWidth+u;
float oldDepth = zbuffer[newIndex];
if (z < oldDepth){
zbuffer[newIndex] = z; // put z value in buffer
outputImg[newIndex] = pointColor[i]; // put pixel in resulting image
}

我有一个完美工作的单核CPU版本。

cuda版本看起来很好，速度也很快，但只有z测试起作用的区域非常"条纹"，我认为这意味着一些背景点覆盖了前景像素。此外，当我查看彩色图像时，我会看到图像中不存在的随机彩色条纹。

CUDA版本看起来更像这样：

//Initialize, kernel, project, new coordinates...
const float oldDepth = outputDepth[v * outputMaxWidth + u];
if (z < oldDepth){
outputDepth[v * outputMaxWidth + u] = z;
const int inputColorIndex = yIndex * inputImageStep + 3*xIndex;
const int outputColorIndex = yIndex * outputImageStep + 3*xIndex;
outputImage[outputColorIndex] = inputImage[inputColorIndex]; //B
outputImage[outputColorIndex + 1] = inputImage[inputColorIndex + 1]; //G
outputImage[outputColorIndex + 2] = inputImage[inputColorIndex + 2]; //R
}

我认为并发性是一个问题。一个线程可以将该像素的最接近的z值写入z缓冲区，但就在同一时间，另一个线程读取旧值并覆盖正确的值。

如何在CUDA中防止这种情况？

第1版：将块大小从(16,16)减小到(1,1)将减少条纹图案，但看起来像1个像素的洞。

第2版：这里有一个最小的例子：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
cudaError_t insertToZBuffer(int *z, const int *a, unsigned int size);
__global__ void zbufferKernel(int *z, const int *a)
{
int i = threadIdx.x;
if (a[i] < z[0]){
z[0] = a[i]; //  all mapped to pixel index 0        
}    
}
int main(){
for (int i = 0; i < 20; ++i){
const int arraySize = 5;
const int a[arraySize] = { 1, 7, 3, 40, 5 }; // some depth values which get mapped all to index 0
int z[arraySize] = { 999 }; // large depth value
insertToZBuffer(z, a, arraySize);
printf("{%d,%d,%d,%d,%d}n", z[0], z[1], z[2], z[3], z[4]);
cudaDeviceReset();
}   
return 0;
}
cudaError_t insertToZBuffer(int *z, const int *a, unsigned int size){
int *dev_a = 0;
int *dev_z = 0;
cudaError_t cudaStatus;
cudaStatus = cudaSetDevice(0);
cudaStatus = cudaMalloc((void**)&dev_z, size * sizeof(int));
cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
cudaStatus = cudaMemcpy(dev_z, z, size * sizeof(int), cudaMemcpyHostToDevice);
zbufferKernel<<<1, size >>>(dev_z, dev_a);
cudaStatus = cudaGetLastError();
cudaStatus = cudaDeviceSynchronize();   
cudaStatus = cudaMemcpy(z, dev_z, size * sizeof(int), cudaMemcpyDeviceToHost);
cudaFree(dev_z);
cudaFree(dev_a);
return cudaStatus;
}

索引0处z的值应该是1，因为它是最低值，但它是5，这是a的最后一个值。

Answer 1

以下是我如何通过评论解决的：

如果z值较小，我会使用atomicCAS(将浮点转换为int)写入z缓冲区。当当前线程的z值较大时，我只返回。最后，我同步所有剩余的线程(__syncthreads())，这些线程已写入缓冲区，并检查它们的z值是否为最终值。如果这确实是真的，我将点颜色写入这个位置的像素值。

编辑：我应该只使用atomicMin。。。