Streams / events (deferred + replayed; see “Asynchronous stream queues”)#

A stream owns a queue of deferred ops; a synchronisation point replays them. Events are a completion flag the replay scheduler honours, so a cross-stream cudaEventRecord/cudaStreamWaitEvent pair is a REAL ordering edge here. Events still measure no time (elapsed time is always 0).

Defined in code/simulator.h.

This page documents 10 symbols: 10 functions.

Functions#

Signature

Description

Availability

Location

Example

inline cudaError_t cudaEventCreate(cudaEvent_t* e)

Events: complete flips false at record and true when the record op replays; that flag is what cudaStreamWaitEvent gates on.

sim only

simulator.h:2162

View

inline cudaError_t cudaEventDestroy(cudaEvent_t e)

sim only

simulator.h:2166

View

inline cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t , cudaEvent_t )

sim only

simulator.h:2182

View

inline cudaError_t cudaEventRecord(cudaEvent_t e, cudaStream_t s = nullptr)

sim only

simulator.h:2167

View

inline cudaError_t cudaEventSynchronize(cudaEvent_t e)

sim only

simulator.h:2178

View

inline cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t size, cudaMemcpyKind , cudaStream_t s = nullptr)

Async copy/set: DEFERRED into the stream (inline on the NULL stream).

sim only

simulator.h:2143

View

inline cudaError_t cudaMemsetAsync(void* devPtr, int value, size_t count, cudaStream_t s = nullptr)

sim only

simulator.h:2151

View

inline cudaError_t cudaStreamCreate(cudaStream_t* s)

sim only

simulator.h:2125

View

inline cudaError_t cudaStreamDestroy(cudaStream_t s)

sim only

simulator.h:2130

View

inline cudaError_t cudaStreamSynchronize(cudaStream_t s)

sim only

simulator.h:2138

View

Examples#

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

cudaEventCreate (test_runtime_api.cu:78)
68    for (int i = 0; i < N; ++i) { a[i] = static_cast<float>(i); b[i] = static_cast<float>(2 * i); }
69
70    // cudaMemset must zero the buffer.
71    CUDA_CHECK(cudaMemset(c, 0, N * sizeof(float)));
72    { bool z = true; for (int i = 0; i < N; ++i) if (c[i] != 0.0f) z = false;
73      all &= report("cudaMemset zeroes buffer", z); }
74
75    // Stream + events around a launch.
76    cudaStream_t stream; CUDA_CHECK(cudaStreamCreate(&stream));
77    cudaEvent_t startE, stopE;
78    CUDA_CHECK(cudaEventCreate(&startE));
79    CUDA_CHECK(cudaEventCreate(&stopE));
80    CUDA_CHECK(cudaEventRecord(startE, stream));
81    LAUNCH(kAdd, blocks, BLOCK, a, b, c, N);
82    CUDA_CHECK(cudaEventRecord(stopE, stream));
83    CUDA_CHECK(cudaEventSynchronize(stopE));
84    float ms = -1.0f;
85    CUDA_CHECK(cudaEventElapsedTime(&ms, startE, stopE));
86    CUDA_CHECK(cudaStreamSynchronize(stream));
87    all &= report("event elapsed-time is defined", ms >= 0.0f);  // 0 in the sim
88
89    // Result of the managed-memory launch, read directly on the host.
cudaEventDestroy (test_runtime_api.cu:114)
104    // Error-code round-trip through the widened table.
105    {
106        bool ok = true;
107        ok &= (std::string(cudaGetErrorName(cudaErrorNotReady)) == "cudaErrorNotReady");
108        ok &= (std::string(cudaGetErrorString(cudaErrorInvalidDevicePointer)) == "invalid device pointer");
109        ok &= (cudaPeekAtLastError() == cudaSuccess);
110        int dev = -1; ok &= (cudaGetDevice(&dev) == cudaSuccess) && dev == 0;
111        all &= report("error/query round-trip", ok);
112    }
113
114    CUDA_CHECK(cudaEventDestroy(startE));
115    CUDA_CHECK(cudaEventDestroy(stopE));
116    CUDA_CHECK(cudaStreamDestroy(stream));
117    CUDA_CHECK(cudaFree(d));
118    CUDA_CHECK(cudaFree(a));
119    CUDA_CHECK(cudaFree(b));
120    CUDA_CHECK(cudaFree(c));
121
122    std::printf("\ntest_runtime_api: %s\n", all ? "ALL PASSED" : "FAILED");
123    return all ? 0 : 1;
124}
cudaEventElapsedTime (test_runtime_api.cu:85)
75    // Stream + events around a launch.
76    cudaStream_t stream; CUDA_CHECK(cudaStreamCreate(&stream));
77    cudaEvent_t startE, stopE;
78    CUDA_CHECK(cudaEventCreate(&startE));
79    CUDA_CHECK(cudaEventCreate(&stopE));
80    CUDA_CHECK(cudaEventRecord(startE, stream));
81    LAUNCH(kAdd, blocks, BLOCK, a, b, c, N);
82    CUDA_CHECK(cudaEventRecord(stopE, stream));
83    CUDA_CHECK(cudaEventSynchronize(stopE));
84    float ms = -1.0f;
85    CUDA_CHECK(cudaEventElapsedTime(&ms, startE, stopE));
86    CUDA_CHECK(cudaStreamSynchronize(stream));
87    all &= report("event elapsed-time is defined", ms >= 0.0f);  // 0 in the sim
88
89    // Result of the managed-memory launch, read directly on the host.
90    { bool ok = true; for (int i = 0; i < N; ++i) if (c[i] != a[i] + b[i]) ok = false;
91      all &= report("managed-memory kernel result", ok); }
92
93    // Async device round-trip on the stream: memsetAsync then D2D copy of c.
94    float* d = nullptr; CUDA_CHECK(cudaMalloc(&d, N * sizeof(float)));
95    CUDA_CHECK(cudaMemsetAsync(d, 0, N * sizeof(float), stream));
96    CUDA_CHECK(cudaMemcpyAsync(d, c, N * sizeof(float), cudaMemcpyDeviceToDevice, stream));
cudaEventRecord (test_runtime_api.cu:80)
70    // cudaMemset must zero the buffer.
71    CUDA_CHECK(cudaMemset(c, 0, N * sizeof(float)));
72    { bool z = true; for (int i = 0; i < N; ++i) if (c[i] != 0.0f) z = false;
73      all &= report("cudaMemset zeroes buffer", z); }
74
75    // Stream + events around a launch.
76    cudaStream_t stream; CUDA_CHECK(cudaStreamCreate(&stream));
77    cudaEvent_t startE, stopE;
78    CUDA_CHECK(cudaEventCreate(&startE));
79    CUDA_CHECK(cudaEventCreate(&stopE));
80    CUDA_CHECK(cudaEventRecord(startE, stream));
81    LAUNCH(kAdd, blocks, BLOCK, a, b, c, N);
82    CUDA_CHECK(cudaEventRecord(stopE, stream));
83    CUDA_CHECK(cudaEventSynchronize(stopE));
84    float ms = -1.0f;
85    CUDA_CHECK(cudaEventElapsedTime(&ms, startE, stopE));
86    CUDA_CHECK(cudaStreamSynchronize(stream));
87    all &= report("event elapsed-time is defined", ms >= 0.0f);  // 0 in the sim
88
89    // Result of the managed-memory launch, read directly on the host.
90    { bool ok = true; for (int i = 0; i < N; ++i) if (c[i] != a[i] + b[i]) ok = false;
91      all &= report("managed-memory kernel result", ok); }
cudaEventSynchronize (test_runtime_api.cu:83)
73      all &= report("cudaMemset zeroes buffer", z); }
74
75    // Stream + events around a launch.
76    cudaStream_t stream; CUDA_CHECK(cudaStreamCreate(&stream));
77    cudaEvent_t startE, stopE;
78    CUDA_CHECK(cudaEventCreate(&startE));
79    CUDA_CHECK(cudaEventCreate(&stopE));
80    CUDA_CHECK(cudaEventRecord(startE, stream));
81    LAUNCH(kAdd, blocks, BLOCK, a, b, c, N);
82    CUDA_CHECK(cudaEventRecord(stopE, stream));
83    CUDA_CHECK(cudaEventSynchronize(stopE));
84    float ms = -1.0f;
85    CUDA_CHECK(cudaEventElapsedTime(&ms, startE, stopE));
86    CUDA_CHECK(cudaStreamSynchronize(stream));
87    all &= report("event elapsed-time is defined", ms >= 0.0f);  // 0 in the sim
88
89    // Result of the managed-memory launch, read directly on the host.
90    { bool ok = true; for (int i = 0; i < N; ++i) if (c[i] != a[i] + b[i]) ok = false;
91      all &= report("managed-memory kernel result", ok); }
92
93    // Async device round-trip on the stream: memsetAsync then D2D copy of c.
94    float* d = nullptr; CUDA_CHECK(cudaMalloc(&d, N * sizeof(float)));
cudaMemcpyAsync (test_runtime_api.cu:96)
 86    CUDA_CHECK(cudaStreamSynchronize(stream));
 87    all &= report("event elapsed-time is defined", ms >= 0.0f);  // 0 in the sim
 88
 89    // Result of the managed-memory launch, read directly on the host.
 90    { bool ok = true; for (int i = 0; i < N; ++i) if (c[i] != a[i] + b[i]) ok = false;
 91      all &= report("managed-memory kernel result", ok); }
 92
 93    // Async device round-trip on the stream: memsetAsync then D2D copy of c.
 94    float* d = nullptr; CUDA_CHECK(cudaMalloc(&d, N * sizeof(float)));
 95    CUDA_CHECK(cudaMemsetAsync(d, 0, N * sizeof(float), stream));
 96    CUDA_CHECK(cudaMemcpyAsync(d, c, N * sizeof(float), cudaMemcpyDeviceToDevice, stream));
 97    CUDA_CHECK(cudaStreamSynchronize(stream));
 98    std::vector<float> back(N);
 99    CUDA_CHECK(cudaMemcpyAsync(back.data(), d, N * sizeof(float), cudaMemcpyDeviceToHost, stream));
100    CUDA_CHECK(cudaStreamSynchronize(stream));
101    { bool ok = true; for (int i = 0; i < N; ++i) if (back[i] != c[i]) ok = false;
102      all &= report("async memset + memcpy round-trip", ok); }
103
104    // Error-code round-trip through the widened table.
105    {
106        bool ok = true;
107        ok &= (std::string(cudaGetErrorName(cudaErrorNotReady)) == "cudaErrorNotReady");
cudaMemsetAsync (test_runtime_api.cu:95)
 85    CUDA_CHECK(cudaEventElapsedTime(&ms, startE, stopE));
 86    CUDA_CHECK(cudaStreamSynchronize(stream));
 87    all &= report("event elapsed-time is defined", ms >= 0.0f);  // 0 in the sim
 88
 89    // Result of the managed-memory launch, read directly on the host.
 90    { bool ok = true; for (int i = 0; i < N; ++i) if (c[i] != a[i] + b[i]) ok = false;
 91      all &= report("managed-memory kernel result", ok); }
 92
 93    // Async device round-trip on the stream: memsetAsync then D2D copy of c.
 94    float* d = nullptr; CUDA_CHECK(cudaMalloc(&d, N * sizeof(float)));
 95    CUDA_CHECK(cudaMemsetAsync(d, 0, N * sizeof(float), stream));
 96    CUDA_CHECK(cudaMemcpyAsync(d, c, N * sizeof(float), cudaMemcpyDeviceToDevice, stream));
 97    CUDA_CHECK(cudaStreamSynchronize(stream));
 98    std::vector<float> back(N);
 99    CUDA_CHECK(cudaMemcpyAsync(back.data(), d, N * sizeof(float), cudaMemcpyDeviceToHost, stream));
100    CUDA_CHECK(cudaStreamSynchronize(stream));
101    { bool ok = true; for (int i = 0; i < N; ++i) if (back[i] != c[i]) ok = false;
102      all &= report("async memset + memcpy round-trip", ok); }
103
104    // Error-code round-trip through the widened table.
105    {
106        bool ok = true;
cudaStreamCreate (test_runtime_api.cu:76)
66    CUDA_CHECK(cudaMallocManaged(&b, N * sizeof(float)));
67    CUDA_CHECK(cudaMallocManaged(&c, N * sizeof(float)));
68    for (int i = 0; i < N; ++i) { a[i] = static_cast<float>(i); b[i] = static_cast<float>(2 * i); }
69
70    // cudaMemset must zero the buffer.
71    CUDA_CHECK(cudaMemset(c, 0, N * sizeof(float)));
72    { bool z = true; for (int i = 0; i < N; ++i) if (c[i] != 0.0f) z = false;
73      all &= report("cudaMemset zeroes buffer", z); }
74
75    // Stream + events around a launch.
76    cudaStream_t stream; CUDA_CHECK(cudaStreamCreate(&stream));
77    cudaEvent_t startE, stopE;
78    CUDA_CHECK(cudaEventCreate(&startE));
79    CUDA_CHECK(cudaEventCreate(&stopE));
80    CUDA_CHECK(cudaEventRecord(startE, stream));
81    LAUNCH(kAdd, blocks, BLOCK, a, b, c, N);
82    CUDA_CHECK(cudaEventRecord(stopE, stream));
83    CUDA_CHECK(cudaEventSynchronize(stopE));
84    float ms = -1.0f;
85    CUDA_CHECK(cudaEventElapsedTime(&ms, startE, stopE));
86    CUDA_CHECK(cudaStreamSynchronize(stream));
87    all &= report("event elapsed-time is defined", ms >= 0.0f);  // 0 in the sim
cudaStreamDestroy (test_runtime_api.cu:116)
106        bool ok = true;
107        ok &= (std::string(cudaGetErrorName(cudaErrorNotReady)) == "cudaErrorNotReady");
108        ok &= (std::string(cudaGetErrorString(cudaErrorInvalidDevicePointer)) == "invalid device pointer");
109        ok &= (cudaPeekAtLastError() == cudaSuccess);
110        int dev = -1; ok &= (cudaGetDevice(&dev) == cudaSuccess) && dev == 0;
111        all &= report("error/query round-trip", ok);
112    }
113
114    CUDA_CHECK(cudaEventDestroy(startE));
115    CUDA_CHECK(cudaEventDestroy(stopE));
116    CUDA_CHECK(cudaStreamDestroy(stream));
117    CUDA_CHECK(cudaFree(d));
118    CUDA_CHECK(cudaFree(a));
119    CUDA_CHECK(cudaFree(b));
120    CUDA_CHECK(cudaFree(c));
121
122    std::printf("\ntest_runtime_api: %s\n", all ? "ALL PASSED" : "FAILED");
123    return all ? 0 : 1;
124}
cudaStreamSynchronize (test_runtime_api.cu:86)
76    cudaStream_t stream; CUDA_CHECK(cudaStreamCreate(&stream));
77    cudaEvent_t startE, stopE;
78    CUDA_CHECK(cudaEventCreate(&startE));
79    CUDA_CHECK(cudaEventCreate(&stopE));
80    CUDA_CHECK(cudaEventRecord(startE, stream));
81    LAUNCH(kAdd, blocks, BLOCK, a, b, c, N);
82    CUDA_CHECK(cudaEventRecord(stopE, stream));
83    CUDA_CHECK(cudaEventSynchronize(stopE));
84    float ms = -1.0f;
85    CUDA_CHECK(cudaEventElapsedTime(&ms, startE, stopE));
86    CUDA_CHECK(cudaStreamSynchronize(stream));
87    all &= report("event elapsed-time is defined", ms >= 0.0f);  // 0 in the sim
88
89    // Result of the managed-memory launch, read directly on the host.
90    { bool ok = true; for (int i = 0; i < N; ++i) if (c[i] != a[i] + b[i]) ok = false;
91      all &= report("managed-memory kernel result", ok); }
92
93    // Async device round-trip on the stream: memsetAsync then D2D copy of c.
94    float* d = nullptr; CUDA_CHECK(cudaMalloc(&d, N * sizeof(float)));
95    CUDA_CHECK(cudaMemsetAsync(d, 0, N * sizeof(float), stream));
96    CUDA_CHECK(cudaMemcpyAsync(d, c, N * sizeof(float), cudaMemcpyDeviceToDevice, stream));
97    CUDA_CHECK(cudaStreamSynchronize(stream));
On this page