How to access launch "params" buffer from a pure CUDA kernel?

Gaming and Visualization Technologies Visualization OptiX
1

Hi,

a question related to OptiX 7.0.0:

What is the appropriate way to access the “params” buffer from a separate CUDA kernel (which is not an OptiX-Kernel) ?

When using this decleration in both (pure CUDA kernel AND OptiX-Kernel)

extern "C"
    {
    __constant__ LaunchParams params;
    }

this error occurs when both are accessed (even when device sync is done around the launches):
[2][ ERROR]: Failed to copy internal launch buffer (CUDA error string: an illegal memory access was encountered, CUDA error code: 700)…

No problems occur, when the pure buffer offset of that device buffer is passed to
the CUDA kernel in a struct:

struct BufferStruct
    {
      LaunchParams* base;
      void* reserved;
    };
    __constant__ __device__  BufferStruct  BufData   = {NULL, NULL};

But that means, that any access to that buffer is an access to a pointer of a pointer. Is there a way to get a direct pointer access to the buffer through such (or a similar ) decleration ?
Its also possible to pass the buffer offset to the kernel as parameter, but also then the access is not directly done as if it would be defined globally (as it is in the OptiX kernels).

Thank you.

My system:
OptiX 7.0.0 SDK GTX 1050 2GB driver 436.48 CUDA 10.1.243 Win10PRO 64bit (version 1809) VS2019 (toolkit v140 of VS2015)

2

Hi m1,

The OptiX launch params declaration is a special OptiX feature. You create a buffer in global memory, put some data there, and then give that pointer to optixLaunch(). What happens under the hood is the global buffer is copied to a different constant memory buffer before launch. This makes accesses to your launch params from OptiX faster than global memory accesses.

If you want to use the same data in a CUDA kernel, you’ll need to use the global memory device pointer that you passed into optixLaunch(). You cannot use the same device side symbol declaration in your CUDA kernel that you’re using in your OptiX shaders. But you can, if you want, copy those launch params into a different constant memory buffer of your own making and reference that buffer in your CUDA kernel.

Does that description make sense?

–
David.

3

Hi David,

thank you very much for the answer.
Yes, that makes sense now.

m1