This project is read-only.

CUDA OpenGL Image Interop

Apr 8, 2014 at 2:07 PM
Edited Apr 8, 2014 at 2:08 PM
Hello,

I have some strange thing with the CudaOpenGLImageInteropResource.

Its like:
cuda_resource = new CudaGraphicsInteropResourceCollection();

cuda_resource.Add(new CudaOpenGLImageInteropResource(tex_ID, CUGraphicsRegisterFlags.None, CudaOpenGLImageInteropResource.OpenGLImageTarget.GL_TEXTURE_2D));
Run:
cuda_resource.MapAllResources();

cuda_tex_ptr_1 = cuda_resource[0].GetMappedPointer();

cuda_integ_kernel.Run(cuda_tex_ptr_1, ... );

cuda_resource.UnmapAllResources();
This is more or less the same way how I map VBOs to CUDA - which works fine. Why is there a problem with mapping a texture like this?
Apr 8, 2014 at 2:25 PM
Hi,

can you tell what problem you encounter? Is it an error (with error message) or just a strange looking image?

My first guess is the GetMappedPointer(), if I recall right, you cannot map a texture to a pointer. But you can access the texture using a CudaArray.

Also check that the mapping and register flags are set right.

-Michael
Apr 8, 2014 at 2:43 PM
Hey,

The Error is: ErrorNotMappedAsPointer.

If I can not write directly to OpenGL Texture Memory (I don't know if the name makes sense, I mean the stuff one created with TexImage2D calls) then I have to get the Array, right?

Get the Array (2D ?) and copy it to device and copy back?

Thank you!

Cheers
MutterOberin
Apr 8, 2014 at 2:52 PM
Yes, bind the openGL texture to a CudaArray2D (I suppose it is 2D). Then you can either copy the data to an flat array (=normal CudaDeviceVariable) or bind the CudaArray to a CudaSurface to write directly to it.

Cheers,
Michael
Apr 8, 2014 at 3:28 PM
How do I get a DevicePointer out of a Surface?

This is the first time I try the surface thing...
cuda_resource.MapAllResources();

CudaArray2D arr = cuda_resource[0].GetMappedArray2D(0, 0);

CudaSurfObject surf = new CudaSurfObject(arr);
Apr 8, 2014 at 3:42 PM
You can't get a device pointer to a surface, you read / write directly to the cuda surface using surf2Dread / surf2Dwrite in your kernel. A surface is kind of a writable texture reference in CUDA...
And there're also two different APIs for that: either the old reference based API or the new object based API which only works on newer device (compute capability >= 3.0)
Apr 8, 2014 at 3:47 PM
Edited Apr 8, 2014 at 3:53 PM
Thank you!

Can you please give a short example?

So host looks like this:
cuda_resource.MapAllResources();

CudaArray2D arr = cuda_resource[0].GetMappedArray2D(0, 0);

CudaSurfObject surf = new CudaSurfObject(arr);

cuda_integ_kernel.Run(surf, this.imageSize, base.fieldMinBound, base.fieldMaxBound,
                                      base.fieldMagnitude, this.row);

cuda_resource.UnmapAllResources();
Device looks like this:
__global__ void CUDA_Kernel(cudaSurfaceObject_t tex, ... )
{
    ...

    surf2Dwrite(act_col, tex, i * 4, j);
}
Apr 8, 2014 at 4:02 PM
It's exactly how you described it. Only you pass surf.SurfObject to your kernel and not surf itself...
Apr 8, 2014 at 4:10 PM
If I do it with surf.SurfObject I get a "ErrorInvalidValue", indicating that one parameter is not within an acceptable range.

I will look to the stuff again....

Thank you very much for the discussion.

Cheers
MutterOberin
Apr 8, 2014 at 4:57 PM
It looks like the mentioned error is produced by the constructor of CudaSurfObject.
Apr 8, 2014 at 7:43 PM
Can you check that you set all necessary flags for resource mapping correctly, i.e. that you correctly bind the OpenGL-texture to a CudaArray which is bound to a surface object with write privilege?
Apr 9, 2014 at 9:03 AM
So the texture is created as follows:
GL.BindTexture(TextureTarget.Texture2D, this.tex_ID);

// Tex Parameters...
... 

GL.TexImage2D(TextureTarget.Texture2D, 0, PixelInternalFormat.Rgba, (int)imageSize.X, (int)imageSize.Y, 0, OpenTK.Graphics.OpenGL.PixelFormat.Rgba, PixelType.Float, IntPtr.Zero);
Then we have the resource
cuda_resource = new CudaGraphicsInteropResourceCollection();

cuda_resource.Add(new CudaOpenGLImageInteropResource(tex_ID, CUGraphicsRegisterFlags.SurfaceLDST, CudaOpenGLImageInteropResource.OpenGLImageTarget.GL_TEXTURE_2D));
Finally the Call
cuda_resource.MapAllResources();

// Get The 2D Array from the Mapped Resource
CudaArray2D arr = cuda_resource[0].GetMappedArray2D(0, 0);

// Create a CUDA Surface out of the Array
CudaSurfObject surf = new CudaSurfObject(arr);
                
// Pass the Surface and Run the Kernel
cuda_integ_kernel.Run(surf.SurfObject, ... );

cuda_resource.UnmapAllResources();
The Kernel takes cudaSurfaceObject_t and writes float4 via surf2Dwrite(data, tex, i * sizeof(float), j) where data is float4

What do I wrong?
Apr 9, 2014 at 1:32 PM
GL.TexImage2D(TextureTarget.Texture2D, 0, PixelInternalFormat.Rgba, (int)imageSize.X, (int)imageSize.Y, 0, OpenTK.Graphics.OpenGL.PixelFormat.Rgba, PixelType.Float, IntPtr.Zero);
is not a float4 texture, this creates an uchar4 typed texture. To get openGL to create a texture of float4s you need PixelInternalFormat.Rgba32f.
Apr 10, 2014 at 9:53 AM
Thanks for pointing this out. Anyway it doesn't help ....

For now I will stick with the solution to compute it, copy the result back and upload it as texture. Takes more time but it works.