[CUDA GRAPH][VPI] capturing VPI internal kernels within a CUDA graph

[Kawa]

My mistake, I was actually not synchronizing VPI within the stream capture.

• If I don’t synchronize, my dummy example finishes with an empty graph.
• If I synchronize my VPI stream before cudaStreamEndCapture(), I get the following error: VPI_ERROR_INTERNAL: (cudaErrorStreamCaptureUnsupported)

I tried several VPI calls and get the same error. My guess is that the whole VPI library don’t support CUDA graphs capture. Is there any plan to make VPI supporting Graph capture?

This could be a useful workaround for the current limitations of VPI when using a common stream for regular CUDA kernels and VPI calls. From VPI library:
CUDA kernels can only be submitted directly to cudaStream_t if it’s guaranteed that all tasks submitted to VPIStream are finished.

This issue is further explained here:

vpiStreamCreateWrapperCUDA with cuda kernels

I want to write some cuda kernels to work with the output of a VPI pipeline. I was thinking of creating a cudaStream_t cuda_stream. Then wrapping it with a vpi stream using vpiStreamCreateWrapperCUDA. Then doing something like this: …vpi pipeline… vpiSubmitStereoDisparityEstimator(vpi_stream, CUDA backend, …) myCudaKernel<<<cuda_stream, …>>>(…) However, in the documentation, it says “CUDA kernels can only be submitted directly to cudaStream_t if it’s guaranteed that all tasks submitted to …

These limitations are quite challenging as they impose many synchronization on complex pipelines, hindering performance of my application.

Best.