np/NP_future.h at master · simoncblyth/np · GitHub

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
#pragma once

/**
NP_future.h
============

General pattern of async functions::

    // Inside every async function:
    cudaStreamWaitEvent(my_stream, input.ready, 0);   // consume input
    // ... do work ...
    cudaEventRecord(output.ready, my_stream);         // produce output
    cudaEventDestroy(input.ready);                    // safe cleanup
    return output;                                    // ONE event


**/


struct NP_future
{
    NP*          arr   = nullptr;
    cudaEvent_t  ready = nullptr;

    void wait(cudaStream_t stream) const;
};

inline void NP_future::wait(cudaStream_t stream) const
{
    if (ready && stream) {
        cudaStreamWaitEvent(stream, ready, 0);
    }
}