surf2Dread in OptiX kernel

Hi,

The following kernel succeeds to be compiled but fails at module compilation.
Is it legal to use CUDA surface object API in OptiX kernel?

struct Params {
    CUsurfObject accumBuffer;
    uint32_t numAccumFrames;
};

extern "C" __constant__ Params params;

extern "C" __global__  void __raygen__test() {
    uint2 launchIndex = make_uint2(optixGetLaunchIndex().x, optixGetLaunchIndex().y);

    float4 accumResult = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
    if (params.numAccumFrames > 1)
        surf2Dread(&accumResult, params.accumBuffer, launchIndex.x * sizeof(float4), launchIndex.y);
    surf2Dwrite(accumResult,
                params.accumBuffer,
                launchIndex.x * sizeof(float4), launchIndex.y);
}

Specifically, I could compile the kernel to a ptx, but failed at optixModuleCreateFromPTX.
The error message is only COMPILE ERROR: Module compilation failed.

I have uploaded the VS project to reproduce the issue:

Thanks,

Windows 10, 1909
NVIDIA Driver: 442.74
RTX 2070

Possibly related to this:
https://forums.developer.nvidia.com/t/bugreport-writing-to-cuda-surfaceobjects-produces-no-result/115109

I’ll file a bug report for investigation.

Looking into the reproducer I could see the compilation error on a different system (Quadro P6000, 441.28) as well, but only when using the Debug executable. The Release executable did not throw a compilation error.

The Debug CUDA compilation settings for the test_kernel.cu have -g and -G for host and device side debug information set.
OptiX never supported debug device code and removing the -G flag solved the compilation error.
(Note that line info would work and is useful for Nsight Compute performance analysis.)

Please try changing the Debug CUDA Compilation settings to match the Release settings.

It could still be that the surface accesses won’t work afterwards as there might still be an issue with those as described inside the linked thread. If that’s the case, feel free to provide a another reproducer, maybe this time with some more descriptions about the pre-requisites for building it. I neither had MSVS 2019 with CUDA 10.1 Visual Studio Integrations nor the CUDA 10.1 and OptiX SDK 7.0.0 installed into their default locations as required by the project.

Thanks for the reply.

I’ve confirmed that removing -G flag resolves the module compilation error.
Is there a note in the document that OptiX disallows that flag?
I remember and can find that is there in OptiX 6 document but couldn’t find in OptiX 7 document.

Surface writes appear to result nothing as you suggest.
But I’m not so sure if my code is correct, so I’m checking now.
I’ll post a reproducer if it seems correct.

Thanks

Thanks. I’ll close the bug about the failed compilation for now and check on status about surface access issues.

Right, the OptiX 6.5.0 programming guide contains the recommendation to avoid the nvcc --device-debug (-G) flag, but I haven’t found anything about it inside the OptiX 7 programming guide.
OptiX behaves like this for over ten years now. Maybe one day it will finally handle it, but I won’t hold my breath. ;-)

I’m back to the issue though I disabled -G flag :(

I created a new reproducer.


This reproducer sets up two surface objects, the one is for surf2Dread, the other is for surf2Dwrite.

Writing via surf2Dwrite seems to have an issue as reported in the topic you mentioned already (and I posted another reproducer there):
[bugreport] Writing to CUDA SurfaceObjects produces no result

The reproducer provides three macro switches to enable/disable surf2D read (instead use a plain buffer to read), surf2D write (instead use a plain buffer to write), OptiX kernel (instead use a pure CUDA kernel).
The issue this time happens when using the following condition:
#define USE_SURFACE_OBJECT_READ
//#define USE_SURFACE_OBJECT_WRITE
#define USE_OPTIX_KERNEL

I get an error which says:

COMPILE ERROR: Module compilation failed

If the reproducer works as expected, it will generates an image (512x256) with red to purple gradient.

I’m not sure this issue shares the same cause with the issue in the topic above.

Thanks

Thanks, I attached the second reproducer to the internal bug report as well.

Would you mind trying this again with the released R450 drivers?
The reproducer you provided should be working in those. Thanks.

Note that the recent OptiX 7.1.0 release strictly requires these R450 drivers.

I tried that already and it works!

Thanks