Can host event and stream objects be used in device code

[tomato]

Is it possible to use a cuda stream and event created on the host in the device?

Something like,

__global__ void kernel(..., cudaEvent_t  _event,
                            cudaStream_t _stream) {
    // code 
    cudaEventRecord(event, stream);
    // code
}

int main() {
    cudaStream_t stream {};
    cudaEvent_t  event {};
    CUDA_CHECK(cudaStreamCreate(&stream));
    CUDA_CHECK(cudaEventCreate(&event));
    kernel<<<x, y>>>(..., stream, event);
}

Is this “legal”, as in, is this supposed to work?

I tried it out, and it doesn’t work.

It seems like in-kernel streams and events only work on those created on device itself, i.e; in-device event must use event and stream objects created on the device only. Is this correct? I might’ve missed this in the documentation, please do point me to it, for reference.

Thank you

[Robert_Crovella]

No, the set of all usable streams and events in the host API is fully distinct from the set of all usable streams and events in the device API. The CDP documentation in the programming guide doesn’t state it explicitly this way, but it does state that all streams and events are private to the grid level they were created in/for, and cannot be used at a parent or child level. Therefore passing a stream or event handle to a kernel via a kernel parameter in any setting is unusable.

[tomato]

Ah damn, that’s a bummer. Would’ve been sweet if there was a way somehow

Thank you for the confirmation, and your time!