Access to display texture for CUDA interop

Support
1

Hello VTK Community.

Currently I’m working on a 3D streaming application utilizing OpenGL, NVENC and libavformat to capture the 3D context, encode it into H264 frames and send it via libavformat as RTP stream to a media server for client distribution.

Using Native OpenGL code I managed to successfully register a GL_TEXTURE_2D to a cuda resource and transfering this buffer data to the NVEncoder for encoding.

Right now, I want to translate this native OpenGL code into higher level VTK-Code. Therefore, I somehow need to get the correct handle of the GL_TEXTURE_2D of the currently rendered frame to register it with CUDA.

I tried:

auto openGLRenWin = vtkOpenGLRenderWindow::SafeDownCast(renderWindow);
auto fbo = openGLRenWin->GetDisplayFramebuffer();
fbo->SaveCurrentBindingsAndBuffers();

int numColors = fbo->GetNumberOfColorAttachments(); // returns 2 (?)
auto texture = fbo->GetColorAttachmentAsTextureObject(0); // also tried index 1 --> same result

cudaGraphicsGLRegisterImage(&cuda_tex_screen_resource, *tex_screen,
        GL_TEXTURE_2D, cudaGraphicsMapFlagsReadOnly)

There has to be something wrong with my code as the resulting encoded frame is all black - rest of my code is identical to the native OpenGL version, where i create and register the texture as followed:

// create a texture
glGenTextures(1, tex_screen);
glBindTexture(GL_TEXTURE_2D, *tex_screen);

// set basic parameters
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);

// buffer data
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F_ARB, size_x, size_y, 0, GL_RGBA, 
    GL_UNSIGNED_BYTE, NULL);
 
// Register CUDA resource
cudaGraphicsGLRegisterImage(&cuda_tex_screen_resource, *tex_screen,
        GL_TEXTURE_2D, cudaGraphicsMapFlagsReadOnly)

My question is: What is the proper way to get the correct handle of the currently rendered frame (texture) in order to register the CUDA resource?
When processing the CUDA resource, do I have to Bind/Active any VTK buffer objects beforehand?

Thanks in advance,
Philipp

1 Like
2

Ok, I solved the issue myself - or to be said, there was no issue at all.

The registration of the framebuffer texture of the currently displayed frame worked as expected. I made an oopsie when fetching the pixels of the cudaGraphicsResource registered texture.

As you can see above, the InternalFormat of the texture in native OpenGL code was GL_RGBA16F_ARB which is why I wrote the kernel function to account for floating point RGBA values (0.0 - 1.0).

The InternalFormat of the VTK texture on the other hand was GL_RGBA8_EXT - therefore I had to rewrite my kernel function to not compute RGBA from float --> unsigned char but to directly take the RGBA values as is.

1 Like
3

Nice that you got this working :100: Could you post a short working example or update any relevant documentation to help others in the future?

1 Like
4

Hi @pieper

The relevant code is the same as in my initial post:

auto openGLRenWin = vtkOpenGLRenderWindow::SafeDownCast(renderWindow);
auto fbo = openGLRenWin->GetDisplayFramebuffer();
fbo->SaveCurrentBindingsAndBuffers();

int numColors = fbo->GetNumberOfColorAttachments(); // returns 2, because of front + back buffer
auto texture = fbo->GetColorAttachmentAsTextureObject(0); // we want to get the front buffer

cudaGraphicsGLRegisterImage(&cuda_tex_screen_resource, *tex_screen,
        GL_TEXTURE_2D, cudaGraphicsMapFlagsReadOnly)

// call RenderWindow->Render() and fetch rendered image from CUDA resource afterwards

It was working from the beginning on but I fetched the data from the registered CUDA resource in a wrong manner.

After calling RenderWindow->Render() to fetch the registered data from CUDA use following code snippets:

// Map registered resource to get mapped pointer
cudaArray* registered_data;
cudaGraphicsMapResources(1, &cuda_tex_screen_resource, 0);
cudaGraphicsSubResourceGetMappedArray(&registered_data, cuda_tex_screen_resource, 0, 0);

// convert cuda_array into texture object for accessing pixel-data
cudaTextureObject_t inTexObject;

struct cudaChannelFormatDesc desc;
cudaGetChannelDesc(&desc, registered_data);

cudaResourceDesc texRes;
memset(&texRes, 0, sizeof(cudaResourceDesc));

texRes.resType = cudaResourceTypeArray;
texRes.res.array.array = registered_data;

cudaTextureDesc texDescr;
memset(&texDescr, 0, sizeof(cudaTextureDesc));

texDescr.normalizedCoords = false;
texDescr.filterMode = cudaFilterModePoint;
texDescr.addressMode[0] = cudaAddressModeWrap;
texDescr.readMode = cudaReadModeElementType;

cudaCreateTextureObject(&inTexObject, &texRes, &texDescr, NULL);

// launch CUDA kernel function to access pixel-data
// Somewhere in .cu file
__device__ uchar4 getPixel(int x, int y, cudaTextureObject_t inTex)
{
    uchar4 ucres = tex2D<uchar4>(inTex, x, y);
    return ucres;
}
2 Likes
5

in VTK 9, found vtkGetObjectMacro(DisplayFramebuffer.
in VTK 8 i found there is no GetDisplayFramebuffer, etc…do u know the related API in vtk 8.1 ?
Thanks!

Morerover, is it able to write by CUDA to the VTK DisplayFramebuffer? and How?

6

is it able to write to the openGL default ColorAttachment of default BackBuffer which is GL_BACK by cudaSurface_t?

and
cudaGraphicsGLRegisterImage(&cuda_tex_screen_resource, *tex_screen,
GL_TEXTURE_2D, cudaGraphicsMapFlagsReadOnly)
*tex_screen should be texture.texUnitID, right?