Directx11interop Texture

Oct 28, 2015 at 11:06 PM
Hi kunzmi,

I need to create a cube texture from DX11 (sharpdx), and write in it from cuda like "simpleDirectx11Texture", but i'm blocked. after registerresource!, from c++ sample it seems that unlike buffer(vertex) it can't be addressed with only pointer, but need to be seen from cuda like cudarray, can you put me in the right way??.
in C++ ->
    cudaGraphicsD3D11RegisterResource(&g_texture_cube.cudaResource, g_texture_cube.pTexture, cudaGraphicsRegisterFlagsNone);
    getLastCudaError("cudaGraphicsD3D11RegisterResource (g_texture_cube) failed");
    // create the buffer. pixel fmt is DXGI_FORMAT_R8G8B8A8_SNORM
    cudaMallocPitch(&g_texture_cube.cudaLinearMemory, &g_texture_cube.pitch, g_texture_cube.size * 4, g_texture_cube.size);
    getLastCudaError("cudaMallocPitch (g_texture_cube) failed");
    cudaMemset(g_texture_cube.cudaLinearMemory, 1, g_texture_cube.pitch * g_texture_cube.size);
After, it copy memory from device to device -> from array to texture??
// populate the faces of the cube map
for (int face = 0; face < 6; ++face)
    cudaArray *cuArray;
    cudaGraphicsSubResourceGetMappedArray(&cuArray, g_texture_cube.cudaResource, face, 0);
    getLastCudaError("cudaGraphicsSubResourceGetMappedArray (cuda_texture_cube) failed");

    // kick off the kernel and send the staging buffer cudaLinearMemory as an argument to allow the kernel to write to it
    cuda_texture_cube(g_texture_cube.cudaLinearMemory, g_texture_cube.size, g_texture_cube.size, g_texture_cube.pitch, face, t);
    getLastCudaError("cuda_texture_cube failed");

    // then we want to copy cudaLinearMemory to the D3D texture, via its mapped form : cudaArray
        cuArray, // dst array
        0, 0,    // offset
        g_texture_cube.cudaLinearMemory, g_texture_cube.pitch, // src
        g_texture_cube.size*4, g_texture_cube.size,            // extent
        cudaMemcpyDeviceToDevice); // kind
    getLastCudaError("cudaMemcpy2DToArray failed");
I really need help!!! How can i make this in managedcuda??
Thanks a lot
Oct 29, 2015 at 10:34 PM

in general a directX texture maps to a cuda array. In the past, cuda arrays could not be written why one often sees that data is written to an additional normally allocated memory (CudaDeviceVariable in managedCuda). After the data is written to memory, a device to cudaArray copy copies the data to the cuda array.

Today this is no more needed as you can write to cuda arrays through cuda surfaces. Which solution is better in terms of performance, I don't know.

Just a few excerpts of code I have lying around on how to do things in managedCuda (I have all my data already in a CudaDeviceVariable because of other reasons, so no Cuda surfaces here):

Init directX texture and binding:
_texture = new Texture(_device, width, height, 0, Usage.AutoGenerateMipMap, Format.X8R8G8B8, Pool.Default);
CudaDirectXInteropResource resource = new CudaDirectXInteropResource(_texture.ComPointer, CUGraphicsRegisterFlags.None, CudaContext.DirectXVersion.D3D9, CUGraphicsMapResourceFlags.None);
_graphicsres.Add(resource); //_graphicsres is a CudaGraphicsInteropResourceCollection 
then somewhere else:
d3dimage.Lock(); //lock directX
_graphicsres.MapAllResources(); //map the resources
//Get an CudaArray from mapped resource
CudaArray2D arr = _graphicsres[0].GetMappedArray2D(0, 0);
//Copy data to it
//Undo the mapping
_graphicsres[0].UnMap(); //Bug in old managedCuda version: Can't unmap all resources in collection...
Note that only the latest version of managedCuda has the bug fixed, so that _graphicsres.UnMapAllResources() actually works. In older version one must iterate over the collection manually to unmap the resources.

Hope it helps,
Oct 30, 2015 at 10:48 AM
Edited Oct 30, 2015 at 2:49 PM
First of all thanks for the answer! ;)
sure i'm near to achive solution.. but i still have an error on mapping resource,

I Create TextureCube with this code:
                BindFlags bindFlags = BindFlags.ShaderResource;

                Texture2DDescription textureDescription = new Texture2DDescription
                    ArraySize = 6,
                    BindFlags = bindFlags,
                    CpuAccessFlags = CpuAccessFlags.None,
                    Format = SharpDX.DXGI.Format.R8G8B8A8_UNorm,
                    Height = 256,
                    Width = 256,
                    MipLevels = 1,
                    OptionFlags = ResourceOptionFlags.TextureCube|ResourceOptionFlags.Shared,
                    SampleDescription = new SharpDX.DXGI.SampleDescription(1, 0),
                    Usage = ResourceUsage.Default

                newTexture.TextureObject2D = new Texture2D(DXDevice, textureDescription);

                ShaderResourceViewDescription srvViewDesc = new ShaderResourceViewDescription
                    Format = SharpDX.DXGI.Format.R8G8B8A8_UNorm,
                    Dimension = SharpDX.Direct3D.ShaderResourceViewDimension.TextureCube,
                    BufferEx = new ShaderResourceViewDescription.ExtendedBufferResource()
                        Flags = ShaderResourceViewExtendedBufferFlags.None
                    TextureCube = new ShaderResourceViewDescription.TextureCubeResource()
                newTexture.shaderResourceView = new ShaderResourceView(DXDevice, newTexture.TextureObject2D, srvViewDesc);

    return newTexture;
I create cuda context an device like this:
//DXDevice Is sharpdx11 device 
devices = CudaContext.GetDirectXDevices(DXDevice.NativePointer, CUd3dXDeviceList.All, CudaContext.DirectXVersion.D3D11);
                context = new CudaContext(devices[0], DXDevice.NativePointer, CUCtxFlags.SchedAuto, CudaContext.DirectXVersion.D3D11);                
I register resource like this where d3dPointerResource = newTexture.TextureObject2D.NativePointer and this is ok no error.
 resource = new CudaDirectXInteropResource[1 ]
            new CudaDirectXInteropResource(d3dPointerResource, CUGraphicsRegisterFlags.None, CudaContext.DirectXVersion.D3D11, CUGraphicsMapResourceFlags.None)
Routine that will launch kernel:
public void UpdateTexture()
            resource[0].Map(); //<-------------this launch error :(

            for (int i = 0; i < 6; ++i)
               CudaArray2D arr = resource[0].GetMappedArray2D((uint)i, 0);
               arr.CopyFromDeviceToThis<uchar4>(image); // image is a CudaPitchedDeviceVariable<uchar4>
               context.CopyToDevice(CuDevPtrBuffInput, BuffInput);  //frame as vector of bytes grabbed form a video
               context.CopyToDevice(CuDevPtrPar, parametervalues); //some int params
               kernelStep.Run(image.DevicePointer, CuDevPtrPar.Pointer, CuDevPtrBuffInput.Pointer);                    
one difference is that i no lock texture before map as directx11 texture no more expose lock method.
have you idea where is the problem?

Thanks for help,
Oct 30, 2015 at 3:00 PM
Hi some update,

I move mapping from loop to post registration (before any bindings as shaderresource) and now no more error :)
but i only move error from map to:
CudaArray2D arr = resource[0].GetMappedArray2D((uint)i, 0); //ok the properties of arr are all ok.
but the istruction below throw error "ErrorInvalidValue: This indicates that one or more of the parameters passed to the API call is not within an acceptable range of values"
arr.CopyFromDeviceToThis<uchar4>(image); // image = new CudaPitchedDeviceVariable<uchar4>(cubesize,cubesize); 
parhaps i'm misunderstood what image in your post means.


Oct 31, 2015 at 11:21 AM
In managedCuda you have three variants of CUArray: CudaArray1D, 2D and 3D. They all wrap the same CUArray handle but differ mainly in their copy methods and constructors to set internal data. CudaArray2D e.g. calls in constructor cuArrayGetDescriptor to get the array descriptor structure, CudaArray3D calls cuArray3DGetDescriptor. I think that a textureCube in directX maps to three-dimensional texture in Cuda (2D image on six faces is a layered 3D texture). So first check where exactly the error is thrown, also try a debug version of managedCuda where every CUDA-API call is logged on the console. My guess is, that you should use a CudaArray3D for textureCube.

The variable "image" I'm using is just the device memory I'm using to compute some stuff that I want to see on screen in the sample code given and I use it to copy data to the cudaArray. Just ignore it...
  • Michael
Oct 31, 2015 at 4:35 PM
Edited Oct 31, 2015 at 4:36 PM
Hi kunzmi,
finally i solve!
It was an unamap missing in onther cudakernel that share a veretexbuffer... so when i try to map in the loop the texture it gives an error.
if i map the texture at registration, than it was ok, (vertexbuffer not mapped) but it throw an error at launch kernel (buffer mapped). Add umap and map to vertexbuffer solve.

The pass are:

Register resource:
resource = new CudaDirectXInteropResource(pointerres, CUGraphicsRegisterFlags.None, CudaContext.DirectXVersion.D3D11, CUGraphicsMapResourceFlags.None);          
image = new CudaPitchedDeviceVariable<char4>(cubesize, cubesize);
and in the loop
 if (!block && Simulate)
                for (int i = 0; i < 6; ++i)
                    CudaArray2D arr = resource.GetMappedArray2D((uint)i, 0);
                    context.CopyToDevice(CuDevPtrBuffInput, BuffInput);  
                    context.CopyToDevice(CuDevPtrPar, parametervalues);
                    kernelStep.Run(image.DevicePointer, CuDevPtrPar.Pointer, CuDevPtrBuffInput.Pointer);                    
                    arr.CopyFromDeviceToThis(image);//<- Your suggestion! :)
You are the best!