__cudaPopCallConfiguration, neither of which are documented in the CUDA API but they preceed any cudaLaunchKernel() call and push the call configuration on a stack which is then popped when cudaLaunchKernel() is called. Currently the max stack size is 1, an assert will be triggered when this is violated.