Extra stream / event entry points#

Defined in code/simulator.h.

This page documents 12 symbols: 1 macros, 11 functions.

Macros#

Signature

Description

Availability

Location

Example

CUDART_CB

Host callbacks are deferred like any other stream op, so they observe exactly the work that precedes them in their stream.

sim only

simulator.h:2418

Functions#

Signature

Description

Availability

Location

Example

inline cudaError_t cudaDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority)

sim only

simulator.h:2391

View

inline cudaError_t cudaEventCreateWithFlags(cudaEvent_t* e, unsigned )

sim only

simulator.h:2430

View

inline cudaError_t cudaEventQuery(cudaEvent_t e)

Query synchronises then reports ready, so a while (cudaEventQuery(e) ...) spin terminates (the sim has no background execution to poll).

sim only

simulator.h:2433

inline cudaError_t cudaLaunchHostFunc(cudaStream_t s, cudaHostFn_t fn, void* userData)

sim only

simulator.h:2426

inline cudaError_t cudaStreamAddCallback(cudaStream_t s, cudaStreamCallback_t cb, void* userData, unsigned )

sim only

simulator.h:2421

View

inline cudaError_t cudaStreamCreateWithFlags(cudaStream_t* s, unsigned )

sim only

simulator.h:2389

View

inline cudaError_t cudaStreamCreateWithPriority(cudaStream_t* s, unsigned , int )

sim only

simulator.h:2390

View

inline cudaError_t cudaStreamQuery(cudaStream_t s)

Query is treated as a synchronise-then-report-ready, so the common while (cudaStreamQuery(s) != cudaSuccess) {} spin terminates.

sim only

simulator.h:2412

inline cudaError_t cudaStreamWaitEvent(cudaStream_t s, cudaEvent_t e, unsigned = 0)

A real ordering edge: everything queued on s after this point waits until e has been recorded by whichever stream records it.

sim only

simulator.h:2398

View

typedef void (*cudaStreamCallback_t)(cudaStream_t, cudaError_t, void*)

sim only

simulator.h:2420

View

typedef void (*cudaHostFn_t)(void*)

sim only

simulator.h:2425

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.

cudaDeviceGetStreamPriorityRange (test_web_cs_streampri.cu:72)
62    }
63    cudaDeviceProp prop{}; cudaGetDeviceProperties(&prop, 0);
64    std::printf("Device 0: %s (compute %d.%d)\n", prop.name, prop.major, prop.minor);
65
66    // create streams with highest and lowest available priorities
67    cudaStream_t st_low;
68    cudaStream_t st_hi;
69#ifdef __CUDACC__
70    // real GPU: query the priority range and make prioritized non-blocking streams
71    int priority_low, priority_hi;
72    checkCudaErrors(cudaDeviceGetStreamPriorityRange(&priority_low, &priority_hi));
73    std::printf("CUDA stream priority range: LOW: %d to HIGH: %d\n", priority_low, priority_hi);
74    checkCudaErrors(cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low));
75    checkCudaErrors(cudaStreamCreateWithPriority(&st_hi, cudaStreamNonBlocking, priority_hi));
76#else
77    // [SIM] fallback: no priority API in the sim — plain streams (priority is a scheduling
78    // hint only; correctness is identical).
79    std::printf("CUDA stream priority range: (unavailable in sim; using plain streams)\n");
80    checkCudaErrors(cudaStreamCreate(&st_low));
81    checkCudaErrors(cudaStreamCreate(&st_hi));
82#endif
cudaEventCreateWithFlags (test_web_olcf_hw13_graph.cu:75)
65        std::printf("No CUDA device available (error %d: %s).\n"
66                    "Compilation and toolkit are fine; kernels need an NVIDIA GPU.\n",
67                    static_cast<int>(e), cudaGetErrorString(e));
68        return 0;
69    }
70    cudaDeviceProp prop{};
71    cudaGetDeviceProperties(&prop, 0);
72    std::printf("Device 0: %s (compute %d.%d)\n", prop.name, prop.major, prop.minor);
73
74    cudaEvent_t event1, event2;
75    cudaEventCreateWithFlags(&event1, cudaEventDisableTiming);
76    cudaEventCreateWithFlags(&event2, cudaEventDisableTiming);
77
78    const int    num_streams = 2;
79    cudaStream_t streams[num_streams];
80    for (int i = 0; i < num_streams; ++i)
81        cudaStreamCreateWithFlags(&streams[i], cudaStreamNonBlocking);
82
83    // [SIM] std::vector host arrays (MSVC /O2 memcpy bug — see
84    // code\test_web_reduce_idle.cu).
85    std::vector<float> h_x(N), h_y(N);
86    for (int i = 0; i < N; ++i) { h_x[i] = (float)i; h_y[i] = (float)i; }
cudaStreamAddCallback (test_web_pcu_callback.cu:110)
100    CallbackData cbdata[NSTREAM];
101
102    for (int i = 0; i < n_streams; i++) {
103        int offset = i * streamSize;
104        cbdata[i].stream_id = i;
105        cbdata[i].fired_count = &fired_count;
106
107        LAUNCH_STREAM(addKernel, streamSize / block, block, streams[i],
108                      d_a + offset, d_b + offset, d_c + offset, streamSize);
109
110        CHECK(cudaStreamAddCallback(streams[i], my_callback, (void *)&cbdata[i], 0));
111    }
112
113    CHECK(cudaDeviceSynchronize());
114    CHECK(cudaMemcpy(c.data(), d_c, n * sizeof(float), cudaMemcpyDeviceToHost));
115
116    for (int i = 0; i < n_streams; i++)
117        CHECK(cudaStreamDestroy(streams[i]));
118    CHECK(cudaFree(d_a)); CHECK(cudaFree(d_b)); CHECK(cudaFree(d_c));
119
120    // Verify BOTH: every callback fired, and c == a + b.
121    bool result_ok = true;
cudaStreamCreateWithFlags (test_web_cs_awbarrier.cu:174)
164    check(cudaMallocHost(&vecA, sizeof(float) * size), "cudaMallocHost vecA");
165    check(cudaMallocHost(&vecB, sizeof(float) * size), "cudaMallocHost vecB");
166    check(cudaMalloc(reinterpret_cast<void**>(&d_vecA), sizeof(float) * size), "cudaMalloc vecA");
167    check(cudaMalloc(reinterpret_cast<void**>(&d_vecB), sizeof(float) * size), "cudaMalloc vecB");
168
169    const float baseVal = 2.0f;
170    for (int i = 0; i < size; i++) vecA[i] = vecB[i] = baseVal;
171
172    cudaStream_t stream;
173    check(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking), "cudaStreamCreateWithFlags");
174    check(cudaMemcpyAsync(d_vecA, vecA, sizeof(float) * size, cudaMemcpyHostToDevice, stream), "H2D A");
175    check(cudaMemcpyAsync(d_vecB, vecB, sizeof(float) * size, cudaMemcpyHostToDevice, stream), "H2D B");
176
177    // [SIM] fixed geometry instead of the occupancy query: a cooperative launch
178    // needs every block resident, and the sim schedules all of them as fibers.
179    int numBlocks = 8, blockSize = 128;
180    int smemSize  = ((blockSize / 32) + 1) * (int)sizeof(double);
181    check(cudaMalloc(reinterpret_cast<void**>(&d_partialResults), numBlocks * sizeof(double)),
182          "cudaMalloc partialResults");
183    check(cudaStreamSynchronize(stream), "cudaStreamSynchronize (H2D)");
cudaStreamCreateWithPriority (test_web_cs_streampri.cu:74)
64    std::printf("Device 0: %s (compute %d.%d)\n", prop.name, prop.major, prop.minor);
65
66    // create streams with highest and lowest available priorities
67    cudaStream_t st_low;
68    cudaStream_t st_hi;
69#ifdef __CUDACC__
70    // real GPU: query the priority range and make prioritized non-blocking streams
71    int priority_low, priority_hi;
72    checkCudaErrors(cudaDeviceGetStreamPriorityRange(&priority_low, &priority_hi));
73    std::printf("CUDA stream priority range: LOW: %d to HIGH: %d\n", priority_low, priority_hi);
74    checkCudaErrors(cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low));
75    checkCudaErrors(cudaStreamCreateWithPriority(&st_hi, cudaStreamNonBlocking, priority_hi));
76#else
77    // [SIM] fallback: no priority API in the sim — plain streams (priority is a scheduling
78    // hint only; correctness is identical).
79    std::printf("CUDA stream priority range: (unavailable in sim; using plain streams)\n");
80    checkCudaErrors(cudaStreamCreate(&st_low));
81    checkCudaErrors(cudaStreamCreate(&st_hi));
82#endif
83
84    size_t size = TOTAL_SIZE;
cudaStreamWaitEvent (test_web_olcf_hw13_graph.cu:115)
105    const int nReplays = 3;                            // [SIM] was 100
106    for (int i = 0; i < nReplays; ++i) {
107        if (!graphCreated) {
108            cudaStreamBeginCapture(streams[0], cudaStreamCaptureModeGlobal);
109            check("Stream begin capture failed");
110
111            LAUNCH_STREAM(kernel_a, blocks, threads, streams[0], d_x, d_y);
112            cudaEventRecord(event1, streams[0]);
113            LAUNCH_STREAM(kernel_b, blocks, threads, streams[0], d_x, d_y);
114
115            cudaStreamWaitEvent(streams[1], event1);   // streams[1] joins the capture
116            LAUNCH_STREAM(kernel_c, blocks, threads, streams[1], d_x, d_y);
117            cudaEventRecord(event2, streams[1]);
118
119            cudaStreamWaitEvent(streams[0], event2);
120            LAUNCH_STREAM(kernel_d, blocks, threads, streams[0], d_x, d_y);
121
122            cudaStreamEndCapture(streams[0], &graph);
123            check("Stream end capture failed");
124
125            cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
126            check("instantiating graph failed");
void (test_web_cs_cppoverload.cu:129)
119    CHECK(cudaMalloc(&dOutput, sizeof(int) * N));
120
121    std::vector<int> hInput(N * 2), hOutput(N, 0);   // [SIM] std::vector host arrays
122    for (int i = 0; i < N * 2; i++) hInput[i] = i;
123
124    CHECK(cudaMemcpy(dInput, hInput.data(), sizeof(int) * N * 2, cudaMemcpyHostToDevice));
125
126    bool testResult = true, funcResult = true;
127    int a = 1;
128
129    void (*func1)(const int *, int *, int);
130    void (*func2)(const int2 *, int *, int);
131    void (*func3)(const int *, const int *, int *, int);
132    struct cudaFuncAttributes attr;
133
134    // overload 1: simple_kernel(const int*, int*, int)
135    func1 = simple_kernel;
136    memset(&attr, 0, sizeof(attr));
137    CHECK(cudaFuncSetCacheConfig(*func1, cudaFuncCachePreferShared));
138    CHECK(cudaFuncGetAttributes(&attr, *func1));
139    OUTPUT_ATTR(attr);
140    LAUNCH(func1, DIV_UP(N, THREAD_N), THREAD_N, dInput, dOutput, a);
void (test_web_cs_cppoverload.cu:129)
119    CHECK(cudaMalloc(&dOutput, sizeof(int) * N));
120
121    std::vector<int> hInput(N * 2), hOutput(N, 0);   // [SIM] std::vector host arrays
122    for (int i = 0; i < N * 2; i++) hInput[i] = i;
123
124    CHECK(cudaMemcpy(dInput, hInput.data(), sizeof(int) * N * 2, cudaMemcpyHostToDevice));
125
126    bool testResult = true, funcResult = true;
127    int a = 1;
128
129    void (*func1)(const int *, int *, int);
130    void (*func2)(const int2 *, int *, int);
131    void (*func3)(const int *, const int *, int *, int);
132    struct cudaFuncAttributes attr;
133
134    // overload 1: simple_kernel(const int*, int*, int)
135    func1 = simple_kernel;
136    memset(&attr, 0, sizeof(attr));
137    CHECK(cudaFuncSetCacheConfig(*func1, cudaFuncCachePreferShared));
138    CHECK(cudaFuncGetAttributes(&attr, *func1));
139    OUTPUT_ATTR(attr);
140    LAUNCH(func1, DIV_UP(N, THREAD_N), THREAD_N, dInput, dOutput, a);
On this page