如何为点云赋予纹理?
How to give texture to a point cloud?
我有一张使用 OpenCV 获得的视差图像。我能够使用 OpenGL 和 CUDA 显示一个 640 x 360 的点网格,并为每个点赋予一个对应于视差值的 Z 值。结果:
现在我想为点云中的每个点赋予一个颜色值,该颜色值对应于具有相同大小 (640 x 380) 的立体相机左图像中的像素。
这是我的 CUDA 内核和调用函数,其中我有两个图像、视差和包含颜色信息的左图像(灰度):
__global__ void simple_vbo_kernel(float4 *pos, unsigned int width, unsigned int height, float time,
uchar* disp, int stepDisp)
{
const int x = threadIdx.x + blockIdx.x * blockDim.x;
const int y = threadIdx.y + blockIdx.y * blockDim.y;
if ((x < width) && (y < height))
{
float u = x / (float) width;
float v = y / (float) height;
u = u*2.0f - 1.0f;
v = v*2.0f - 1.0f;
const int depth = y * stepDisp + x;
float w = static_cast<float>(disp[depth]);
w/=10;
pos[y*width+x] = make_float4(u, w, v, 1.0f);
}
}
extern "C"
void launch_kernel(dim3 grid, dim3 block, float4 *pos, unsigned int mesh_width,unsigned int mesh_height,
float time, cv::Mat disp, cv::Mat left)
{
if(!disp.empty() && !left.empty()){
uchar* d_image;
int dsize = disp.rows * disp.step;
cudaMalloc((void**)&d_image, dsize);
cudaMemcpy(d_image, disp.ptr(), dsize, cudaMemcpyHostToDevice);
simple_vbo_kernel<<< grid, block>>>(pos, mesh_width, mesh_height, time, d_image, disp.step);
cudaDeviceSynchronize();
cudaFree(d_image);
}
}
我的问题是,为点云中的点赋予纹理或颜色的最简单方法是什么,在这种情况下使用 cv::Mat left
?我已经从 CUDA 示例中看到了一些其他示例,但我没有找到如何操作。
对于稀疏点云(如您的点云),通常最好的做法是向顶点添加颜色属性并为其分配所需的颜色。使绘图着色器将该属性的值传递给发出的片段颜色。
我找到了解决办法。在内核部分,我有一个新的 uchar4 颜色映射输入,两个输入指针必须具有相同的大小,才能在点和颜色之间建立正确的关系(float4 *pos 和 uchar4 *color)。由于d_image是灰度图(视差图),而赋予纹理的是彩色图,所以访问每个数组的方式要考虑到这一点。 Opencv 是 BGR 顺序,gstreamer 是 RGB:
__global__ void simple_vbo_kernel(float4 *pos, uchar4 *color, unsigned int width, unsigned int height, uchar* d_image, int stepDisp, uchar* d_color, int stepLeft)
{
const int x = threadIdx.x + blockIdx.x * blockDim.x;
const int y = threadIdx.y + blockIdx.y * blockDim.y;
if ((x < width) && (y < height)){
float u = x / (float) width;
float v = y / (float) height;
u = u*2.0f - 1.0f;
v = v*2.0f - 1.0f;
const int depthi = y * stepDisp + x;
float w = static_cast<float>(d_image[depthi]);
w/=25;
pos[y*width+x] = make_float4(u, w, v, 1.0f);
uchar b = static_cast<uchar>(d_color[y*width*3+x*3 + 0]);
uchar g = static_cast<uchar>(d_color[y*width*3+x*3 + 1]);
uchar r = static_cast<uchar>(d_color[y*width*3+x*3 + 2]);
color[y*width+x] = make_uchar4( r, g, b, 1.0f);
}
}
extern "C"
void launch_kernel(dim3 grid, dim3 block, unsigned int mesh_width, unsigned int mesh_height, float4 *pos, cv::Mat disp, uchar4 *color, cv::Mat left)
{
if(!disp.empty() && !left.empty()){
uchar* d_image;
int dsize = disp.rows * disp.step;
cudaMalloc((void**)&d_image, dsize);
cudaMemcpy(d_image, disp.ptr(), dsize, cudaMemcpyHostToDevice);
uchar* d_color;
int dcolorsize = left.rows * left.step;
cudaMalloc((void**)&d_color, dcolorsize);
cudaMemcpy(d_color, left.ptr(), dcolorsize, cudaMemcpyHostToDevice);
simple_vbo_kernel<<< grid, block>>>(pos, color, mesh_width, mesh_height, d_image, disp.step, d_color, left.step);
cudaDeviceSynchronize();
cudaFree(d_image);
cudaFree(d_color);
}
}
调用内核:
void runCuda(struct cudaGraphicsResource **vbo_resource, Mat dsp, struct cudaGraphicsResource **vbo_resource_color, Mat lft)
{
float4 *dptr_pos;
checkCudaErrors(cudaGraphicsMapResources(1, vbo_resource, 0));
size_t num_bytes;
checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void **)&dptr_pos, &num_bytes, *vbo_resource));
uchar4 *dptr_color;
checkCudaErrors(cudaGraphicsMapResources(1, vbo_resource_color, 0));
size_t num_bytes_color;
checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void **)&dptr_color, &num_bytes_color, *vbo_resource_color));
dim3 block(8, 8, 1);
dim3 grid(mesh_width / block.x, mesh_height / block.y, 1);
launch_kernel(grid,block, mesh_width, mesh_height, dptr_pos, dsp, dptr_color, lft);
checkCudaErrors(cudaGraphicsUnmapResources(1, vbo_resource, 0));
checkCudaErrors(cudaGraphicsUnmapResources(1, vbo_resource_color, 0));
}
在 OpenGl 的初始化中也很重要,添加:
glEnable( GL_TEXTURE_2D );
并启用深度测试。我之前禁用过并且有奇怪的图像重叠:
glEnable(GL_DEPTH_TEST);
以及必须显示或更新图像的部分,在我的例子中,我在 glutDisplayFunc(display)
函数中有它:
glBindBuffer(GL_ARRAY_BUFFER, pos_vbo);
glVertexPointer(4, GL_FLOAT, 0, 0);
glEnableClientState(GL_VERTEX_ARRAY);
glBindBuffer(GL_ARRAY_BUFFER, color_vbo);
glColorPointer(4, GL_UNSIGNED_BYTE, 0, 0);
glEnableClientState(GL_COLOR_ARRAY);
glPointSize(3.0);
glDrawArrays(GL_POINTS, 0, mesh_width * mesh_height);
glDisableClientState(GL_VERTEX_ARRAY);
glDisableClientState(GL_COLOR_ARRAY);
glBindBuffer(GL_ARRAY_BUFFER, 0);
glutSwapBuffers();
glFlush();
glDeleteTextures(1, &color_vbo);
我的图像是灰度图像,因为我仍在尝试使用彩色 gstreamer 0.10 捕获我的 rtsp 流
http://postimg.org/image/ub8ds4zw5/
我新开了一个post试图解决颜色问题:
我有一张使用 OpenCV 获得的视差图像。我能够使用 OpenGL 和 CUDA 显示一个 640 x 360 的点网格,并为每个点赋予一个对应于视差值的 Z 值。结果:
现在我想为点云中的每个点赋予一个颜色值,该颜色值对应于具有相同大小 (640 x 380) 的立体相机左图像中的像素。
这是我的 CUDA 内核和调用函数,其中我有两个图像、视差和包含颜色信息的左图像(灰度):
__global__ void simple_vbo_kernel(float4 *pos, unsigned int width, unsigned int height, float time,
uchar* disp, int stepDisp)
{
const int x = threadIdx.x + blockIdx.x * blockDim.x;
const int y = threadIdx.y + blockIdx.y * blockDim.y;
if ((x < width) && (y < height))
{
float u = x / (float) width;
float v = y / (float) height;
u = u*2.0f - 1.0f;
v = v*2.0f - 1.0f;
const int depth = y * stepDisp + x;
float w = static_cast<float>(disp[depth]);
w/=10;
pos[y*width+x] = make_float4(u, w, v, 1.0f);
}
}
extern "C"
void launch_kernel(dim3 grid, dim3 block, float4 *pos, unsigned int mesh_width,unsigned int mesh_height,
float time, cv::Mat disp, cv::Mat left)
{
if(!disp.empty() && !left.empty()){
uchar* d_image;
int dsize = disp.rows * disp.step;
cudaMalloc((void**)&d_image, dsize);
cudaMemcpy(d_image, disp.ptr(), dsize, cudaMemcpyHostToDevice);
simple_vbo_kernel<<< grid, block>>>(pos, mesh_width, mesh_height, time, d_image, disp.step);
cudaDeviceSynchronize();
cudaFree(d_image);
}
}
我的问题是,为点云中的点赋予纹理或颜色的最简单方法是什么,在这种情况下使用 cv::Mat left
?我已经从 CUDA 示例中看到了一些其他示例,但我没有找到如何操作。
对于稀疏点云(如您的点云),通常最好的做法是向顶点添加颜色属性并为其分配所需的颜色。使绘图着色器将该属性的值传递给发出的片段颜色。
我找到了解决办法。在内核部分,我有一个新的 uchar4 颜色映射输入,两个输入指针必须具有相同的大小,才能在点和颜色之间建立正确的关系(float4 *pos 和 uchar4 *color)。由于d_image是灰度图(视差图),而赋予纹理的是彩色图,所以访问每个数组的方式要考虑到这一点。 Opencv 是 BGR 顺序,gstreamer 是 RGB:
__global__ void simple_vbo_kernel(float4 *pos, uchar4 *color, unsigned int width, unsigned int height, uchar* d_image, int stepDisp, uchar* d_color, int stepLeft)
{
const int x = threadIdx.x + blockIdx.x * blockDim.x;
const int y = threadIdx.y + blockIdx.y * blockDim.y;
if ((x < width) && (y < height)){
float u = x / (float) width;
float v = y / (float) height;
u = u*2.0f - 1.0f;
v = v*2.0f - 1.0f;
const int depthi = y * stepDisp + x;
float w = static_cast<float>(d_image[depthi]);
w/=25;
pos[y*width+x] = make_float4(u, w, v, 1.0f);
uchar b = static_cast<uchar>(d_color[y*width*3+x*3 + 0]);
uchar g = static_cast<uchar>(d_color[y*width*3+x*3 + 1]);
uchar r = static_cast<uchar>(d_color[y*width*3+x*3 + 2]);
color[y*width+x] = make_uchar4( r, g, b, 1.0f);
}
}
extern "C"
void launch_kernel(dim3 grid, dim3 block, unsigned int mesh_width, unsigned int mesh_height, float4 *pos, cv::Mat disp, uchar4 *color, cv::Mat left)
{
if(!disp.empty() && !left.empty()){
uchar* d_image;
int dsize = disp.rows * disp.step;
cudaMalloc((void**)&d_image, dsize);
cudaMemcpy(d_image, disp.ptr(), dsize, cudaMemcpyHostToDevice);
uchar* d_color;
int dcolorsize = left.rows * left.step;
cudaMalloc((void**)&d_color, dcolorsize);
cudaMemcpy(d_color, left.ptr(), dcolorsize, cudaMemcpyHostToDevice);
simple_vbo_kernel<<< grid, block>>>(pos, color, mesh_width, mesh_height, d_image, disp.step, d_color, left.step);
cudaDeviceSynchronize();
cudaFree(d_image);
cudaFree(d_color);
}
}
调用内核:
void runCuda(struct cudaGraphicsResource **vbo_resource, Mat dsp, struct cudaGraphicsResource **vbo_resource_color, Mat lft)
{
float4 *dptr_pos;
checkCudaErrors(cudaGraphicsMapResources(1, vbo_resource, 0));
size_t num_bytes;
checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void **)&dptr_pos, &num_bytes, *vbo_resource));
uchar4 *dptr_color;
checkCudaErrors(cudaGraphicsMapResources(1, vbo_resource_color, 0));
size_t num_bytes_color;
checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void **)&dptr_color, &num_bytes_color, *vbo_resource_color));
dim3 block(8, 8, 1);
dim3 grid(mesh_width / block.x, mesh_height / block.y, 1);
launch_kernel(grid,block, mesh_width, mesh_height, dptr_pos, dsp, dptr_color, lft);
checkCudaErrors(cudaGraphicsUnmapResources(1, vbo_resource, 0));
checkCudaErrors(cudaGraphicsUnmapResources(1, vbo_resource_color, 0));
}
在 OpenGl 的初始化中也很重要,添加:
glEnable( GL_TEXTURE_2D );
并启用深度测试。我之前禁用过并且有奇怪的图像重叠:
glEnable(GL_DEPTH_TEST);
以及必须显示或更新图像的部分,在我的例子中,我在 glutDisplayFunc(display)
函数中有它:
glBindBuffer(GL_ARRAY_BUFFER, pos_vbo);
glVertexPointer(4, GL_FLOAT, 0, 0);
glEnableClientState(GL_VERTEX_ARRAY);
glBindBuffer(GL_ARRAY_BUFFER, color_vbo);
glColorPointer(4, GL_UNSIGNED_BYTE, 0, 0);
glEnableClientState(GL_COLOR_ARRAY);
glPointSize(3.0);
glDrawArrays(GL_POINTS, 0, mesh_width * mesh_height);
glDisableClientState(GL_VERTEX_ARRAY);
glDisableClientState(GL_COLOR_ARRAY);
glBindBuffer(GL_ARRAY_BUFFER, 0);
glutSwapBuffers();
glFlush();
glDeleteTextures(1, &color_vbo);
我的图像是灰度图像,因为我仍在尝试使用彩色 gstreamer 0.10 捕获我的 rtsp 流 http://postimg.org/image/ub8ds4zw5/
我新开了一个post试图解决颜色问题: