### 如何給點雲紋理？

#### [英]How to give texture to a point cloud?

I have a disparity image obtained using OpenCV. I am able to display a 640 x 360 grid of points using OpenGL and CUDA, and give to each point a Z value which corresponds to the disparity value. The result:

Now I wish to give to each point in the point cloud a colour value corresponding to the pixel in the left image of the stereo camera which have same size (640 x 380).

This is my CUDA kernel and the calling function where I have both images, the disparity and the left image (gray scale) containing the color information:

__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);
cudaFree(d_image);
}
}

My question is, what is the easiest way to give a texture or colour to the points in the point cloud, in this case using the cv::Mat left? I have seen some other examples from the CUDA samples but I did not find how to do it.

## 2 个解决方案

### #1

In case of sparse point clouds (like your's) usually the best course of action is adding a colour attribute to the vertices and assign it the desired color. Make the drawing shaders pass the value of that attribute to the emited fragment color.

### #2

I found a solution. In the kernel part I have a new uchar4 mapped input for the color, both input pointers must have same size for a correct relation between point and color (float4 *pos and uchar4 *color). Since d_image is a gray scale image (disparity image) and the one to give texture is a color image, this should be taken into account for the way to access to each array. Opencv it is in BGR order and gstreamer in 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);
cudaFree(d_image);
cudaFree(d_color);

}
}

Calling to the kernel:

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));
}

Also important in the initialization of the OpenGl add this:

glEnable( GL_TEXTURE_2D );

and also enable depth test. I had Disabled before and I have strange image overlapings:

glEnable(GL_DEPTH_TEST);

And the part where the image must be shown or updated, in my case I have it in my glutDisplayFunc(display) function:

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);

My image is a grayscale image because I am still trying to capture my rtsp stream with gstreamer 0.10 in color http://postimg.org/image/ub8ds4zw5/

I have opened a new post trying to solve the color issue:

GstBuffer to color Mat

GstBuffer為Mat着色