cuda - カーネルからストリーム内のスラスト関数を呼び出す方法は?

Question

デバイスカーネルで呼び出して非同期にしたいthrust::scatter（別のホストスレッドで呼び出すこともできます）。thrust::cuda::par.on(stream)デバイスカーネルから呼び出すことができないホスト関数です。次のコードは、Turing アーキテクチャの CUDA 10.1 で試しました。

__global__ void async_scatter_kernel(float* first,
    float* last,
    int* map,
    float* output)
{
    cudaStream_t stream;
    cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
    thrust::scatter(thrust::cuda::par.on(stream), first, last, map, output);
    cudaDeviceSynchronize();
    cudaStreamDestroy(stream);
}

デバイスから呼び出されたときに、スラストが動的並列処理を使用してカーネルを起動することは知っていますが、ストリームを指定する方法が見つかりませんでした。