CUDA中一个简单Z缓冲区的实现
Implementation of a simple Z-Buffer in CUDA
我有一个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的最后一个值。
以下是我如何通过评论解决的:
如果z值较小,我会使用atomicCAS(将浮点转换为int)写入z缓冲区。当当前线程的z值较大时,我只返回。最后,我同步所有剩余的线程(__syncthreads()),这些线程已写入缓冲区,并检查它们的z值是否为最终值。如果这确实是真的,我将点颜色写入这个位置的像素值。
编辑:我应该只使用atomicMin。。。
相关文章:
- C++字符*缓冲区的大小
- 在c++中用vector填充一个简单的动态数组
- (C++)分析树以计算返回错误值的简单算术表达式
- 为什么msgrcv()将垃圾字符馈送到缓冲区
- 我的简单if-else语句是如何无法访问的代码
- 使用动态分配的数组会导致代码分析发出虚假的C6386缓冲区溢出警告
- ostream过载时的缓冲区冲洗
- 创建一个简单的前向迭代器,该迭代器在循环缓冲区的"end"处自动换行
- 简单情况下的模板缓冲区行为(GL_ALWAYS、GL_LEQUAL)
- 我正在编写一个简单的客户端套接字应用程序,但在连接后服务器收到一个空缓冲区
- 堆缓冲区溢出随机发生。对于一个简单的代码?(我是C++新手)
- CUDA中一个简单Z缓冲区的实现
- 简单协议缓冲区程序与G 编译时可起作用,但不clang
- C++ STL 容器用作简单缓冲区时速度很慢(我做了一个基准测试)
- 使用 std::vector 作为简单缓冲区是否是一种好的做法
- 多个包含简单数据的常量缓冲区 DirectX 11
- 简单的程序集代码会产生缓冲区溢出
- 实现任意类型擦除的小缓冲区优化的简单方法(如std::function.)
- 为浮点值实现一个简单的环形缓冲区
- GzipOutputStream和GzipInputStream与协议缓冲区的简单工作示例