For what it’s worth, even with a pure CUDA kernel and access to all the barrier intrinsics, the only reliable way to synchronize all threads, and best advice we have if you want to syncronize all threads is to split it into multiple kernels. You can find multiple forum threads and Stack Overflow threads that all echo the same advice, for example
here: Synchronize all blocks in CUDA - #8 by njuffa
here: cuda - Does __syncthreads() synchronize all threads in the grid? - Stack Overflow
and here: Global thread barrier - #2 by MisterAnderson42
One option would be a wavefront architecture, which is exactly as you describe: storing the path state in a buffer. While it’s not ideal for the reasons you mention, it is not uncommon for people to use a wavefront architecture for GPU rendering, since it can sometimes be easier to scale than the alternative megakernel approach you currently have. One advantage of using a wavefront architecture is that if you have some paths already terminated, and some paths are still bouncing, you have the option to reduce your kernel size with every launch, instead of allowing some threads to become inactive. With your current setup, all the paths that terminate will become inactive threads which waste some time and reduce thread coherence, as fewer and fewer paths are left.
Another option might be to separate into sub-kernels that all share as much state as possible, so you can track most of your state at a global level, and minimize the per-thread state that you would need to store in a buffer.
There is a CUDA 9+ feature called Cooperative Thread Groups that sort-of has some of the syncronization properties you want. It can syncronize all the threads in a kernel, but only if you launch a number of thread blocks that does not exceed the number of SMs. Programming Guide :: CUDA Toolkit Documentation I haven’t used this, but I suspect the block number constraint means the total number of threads cannot exceed your total number of CUDA cores on your GPU, meaning this probably wouldn’t work for you even if you could use it. But, either way, it is not available from an OptiX Launch, so this is not currently an option.
So currently the only option for synchronizing all threads in a ray tracing launch is to use separate kernels. I recommend trying it and measuring the performance before assuming that the cost of the state buffer would slow it down. There is some cost to saving the ray state, of course, but maybe it’s less that you fear, or maybe achieving the synchronization you’re after will provide more benefit than the cost of saving state.
–
David.