1

I was playing around with the new OptiX-7 version. To display the launch results (simple mono-colour written from raygen) I create an OpenGL texture, map it and obtain the mapped array via cuGraphicsSubResourceGetMappedArray, and finally create a CUDA surface with the array as backing. However, writing to the surface from the OptiX kernel yields neither a change in the (displayed) texture nor an error, whereas an almost identical regular CUDA kernel has visible effects.

The CUDA code:

__global__ void test_kernel(::CUsurfObject surface) {
    const uint2 imageIndex{
        threadIdx.x + blockDim.x * blockIdx.x,
        threadIdx.y + blockDim.y * blockIdx.y
    };
    if(imageIndex.x < 800 && imageIndex.y < 100)
        ::surf2Dwrite(::make_float4(1.f, 0.f, 1.f, 1.f), surface,
                      imageIndex.x * sizeof(float4), imageIndex.y);
}

The OptiX kernel:

struct Params {
    ::CUsurfObject image;
};

extern "C" {
    __constant__ Params params;
}

extern "C" __global__ void __raygen__solid_color() {
    const uint3 launchIndex = ::optixGetLaunchIndex();
    const auto imageIndex = ::make_uint2(launchIndex.x, launchIndex.y);

    ::surf2Dwrite(::make_float4(1.f, 0.f, 1.f, 1.f), params.image,
                imageIndex.x * sizeof(float4), imageIndex.y);
}

The surface handle is passed via the launch parameters:

cuMemAlloc(&devParamBuf, sizeof(Params));
cuMemcpyHtoD(devParamBuf, &hostParamBuf);
...
optixLaunch(pipeline, stream, devParamBuf, sizeof(Params), &sbt, 800, 600, 1);

The kernel itself gets executed just fine. The OptiX samples contain an example of texture display and simply use the cudaTextureObject_t. However, they also use plain buffers and PBOs for displaying the rendered image. Is it simply not possible to directly write to a CUDA surface?

EDIT: to clarify how the surface is created: Taking an existing OpenGL texture, I register and map it with 

GLuint texture{};
glGenTextures(...);
...
cuGraphicsGLRegisterImage(&resource, texture, GL_TEXTURE_2D, CU_GRAPHICS_REGISTER_FLAGS_SURFACE_LDST);
cuGraphicsMapResources(1u, &resource, stream);
CUarray surfaceArray{};
cuGraphicsSubResourceGetMappedArray(&surfaceArray, resource, 0u, 0u);
CUsurfObject surface{};
CUDA_RESOURCE_DESC surfDesc{};
surfDesc.resType = CU_RESOURCE_TYPE_ARRAY;
surfDesc.res.array.hArray = surfaceArray;
cuSurfObjectCreate(&surface, &surfDesc);

All calls are wrapped with error checking omitted for brevity.

IGarFieldI
  • 515
  • 3
  • 12
  • 1
    It is possible I'm doing it with directx12. If you do a mapping, there should not be any *alloc., but a `cudaResourceDesc` properly filled than will do the mapping between your mapped array and `cudaCreateSurfaceObject`. So you're actually changing a surface that is independant from your GL texture, hence the result. But also I'm not certain the cuGraphicsSubResourceGetMappedArray is the right way to go. More code would be helpful. – Soleil Nov 02 '19 at 12:02
  • I added more code how the surface is created. I tried both unmapping the resource before using the surface as well as leaving it mapped until rendering is done. Tbh this is not exactly a pressing issue - I required atomic operations for the framebuffer anyway, something that surfaces AFAIK do not provide. – IGarFieldI Nov 05 '19 at 05:40
  • Check this; I did that with D3D12 and it will probably be same with GL: https://stackoverflow.com/questions/58626926/cuda-directx-12-texture2d-in-1d-array-interop/58682990#58682990 – Soleil Nov 05 '19 at 10:23
  • Thanks for the link. Unfortunately it is not the answer. The mapped surface _does_ work with plain CUDA already. It is only from an OptiX launch that the write does not become visible. – IGarFieldI Nov 05 '19 at 16:58
  • I also checked the generated PTX, and aside from differences in getting the thread and launch indices the surface store is the exact same, yet appears to not get executed from the OptiX raygen program. – IGarFieldI Nov 05 '19 at 18:57
  • However I see you're mapping directly the array, I think this is wrong. In my sample, I export the memory, then I map it to a mipmaparray, then to an array then I create the surface, then I call the kernel. I don't see the same pipeline in your code. – Soleil Nov 05 '19 at 20:07
  • There is no way to import an OpenGL texture as exported memory - the CUDA API only supports it for a handful of types, of which two are DirectX12-related. This makes sense since CUDA has a dedicated interop API for OpenGL and DirectX up until 11. I also tried mapping it as CUmipmappedArray and then extracting the CUarray of first level, to no avail. Once again, it works perfectly fine for a regular CUDA launch. There must be something in the OptiX pipeline that changes the behaviour. – IGarFieldI Nov 06 '19 at 07:44

0 Answers0