Streams and Events

Stream 0 (default): |--copy H2D--|--kernel--|--copy D2H--|
Stream 1:                   |--copy H2D--|--kernel--|--copy D2H--|
                                                         ^^^ overlap

cudaStream_t stream;
cudaStreamCreate(&stream);
// ... use stream ...
cudaStreamDestroy(stream);

// With flags
cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
// cudaStreamDefault (0) — blocks on the default stream
// cudaStreamNonBlocking — does not block on default stream operations

// With priority (lower number = higher priority)
int loPri, hiPri;
cudaDeviceGetStreamPriorityRange(&loPri, &hiPri);
cudaStreamCreateWithPriority(&stream, cudaStreamNonBlocking, hiPri);

// Enqueue to a stream
cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, stream);
myKernel<<<grid, block, 0, stream>>>(d_in, d_out, n);
cudaMemcpyAsync(h_out, d_out, bytes, cudaMemcpyDeviceToHost, stream);

// Synchronize
cudaStreamSynchronize(stream);   // block CPU until stream empty
cudaDeviceSynchronize();         // block CPU until ALL streams empty

// Non-blocking query
cudaError_t s = cudaStreamQuery(stream);
// cudaSuccess → complete; cudaErrorNotReady → still running

// Enable per-thread default stream globally:
// nvcc --default-stream per-thread myfile.cu
// Or per translation unit:
#define CUDA_API_PER_THREAD_DEFAULT_STREAM
#include <cuda_runtime.h>

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

// Record events into stream
cudaEventRecord(start, stream);
myKernel<<<grid, block, 0, stream>>>(args);
cudaEventRecord(stop, stream);

// Wait on CPU for stop to be reached
cudaEventSynchronize(stop);

float ms;
cudaEventElapsedTime(&ms, start, stop);   // milliseconds
printf("Kernel time: %.3f ms\n", ms);

cudaEventDestroy(start);
cudaEventDestroy(stop);

cudaEventCreateWithFlags(&ev, cudaEventDefault);
cudaEventCreateWithFlags(&ev, cudaEventBlockingSync);  // CPU sleeps instead of polls
cudaEventCreateWithFlags(&ev, cudaEventDisableTiming); // no timing — lowest overhead
cudaEventCreateWithFlags(&ev, cudaEventInterprocess);  // for IPC

cudaEvent_t event;
cudaEventCreateWithFlags(&event, cudaEventDisableTiming);

// Stream A: record event after kernel finishes
kernelA<<<g, b, 0, streamA>>>(d_a);
cudaEventRecord(event, streamA);

// Stream B: wait for that event before starting
cudaStreamWaitEvent(streamB, event, 0);   // 0 = no flags
kernelB<<<g, b, 0, streamB>>>(d_b);

const int CHUNKS = 4;
size_t chunkBytes = totalBytes / CHUNKS;

cudaStream_t streams[2];
cudaStreamCreate(&streams[0]);
cudaStreamCreate(&streams[1]);

float *d_in[2], *d_out[2];
cudaMalloc(&d_in[0],  chunkBytes);  cudaMalloc(&d_out[0], chunkBytes);
cudaMalloc(&d_in[1],  chunkBytes);  cudaMalloc(&d_out[1], chunkBytes);

float *h_in, *h_out;
cudaMallocHost(&h_in,  totalBytes);   // must be pinned for async
cudaMallocHost(&h_out, totalBytes);

for (int i = 0; i < CHUNKS; i++) {
    int s = i & 1;   // ping-pong between 0 and 1
    cudaMemcpyAsync(d_in[s], h_in + i * chunkBytes / sizeof(float),
                    chunkBytes, cudaMemcpyHostToDevice, streams[s]);
    process<<<grid, block, 0, streams[s]>>>(d_in[s], d_out[s], chunkBytes / sizeof(float));
    cudaMemcpyAsync(h_out + i * chunkBytes / sizeof(float), d_out[s],
                    chunkBytes, cudaMemcpyDeviceToHost, streams[s]);
}

cudaDeviceSynchronize();

cudaFreeHost(h_in); cudaFreeHost(h_out);
for (int s = 0; s < 2; s++) {
    cudaFree(d_in[s]); cudaFree(d_out[s]);
    cudaStreamDestroy(streams[s]);
}

cudaGraph_t     graph;
cudaGraphExec_t instance;

// Begin capture
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

// Record operations (enqueued but NOT yet executed)
cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, stream);
myKernel<<<grid, block, 0, stream>>>(d_in, d_out, n);
cudaMemcpyAsync(h_out, d_out, bytes, cudaMemcpyDeviceToHost, stream);

// End capture → graph
cudaStreamEndCapture(stream, &graph);

// Instantiate (compile) and launch
cudaGraphInstantiate(&instance, graph, nullptr, nullptr, 0);

for (int iter = 0; iter < 1000; iter++)
    cudaGraphLaunch(instance, stream);   // very low CPU overhead

cudaStreamSynchronize(stream);

cudaGraphExecDestroy(instance);
cudaGraphDestroy(graph);

cudaGraphExecKernelNodeSetParams(instance, node, &newParams);
cudaGraphExecMemcpyNodeSetParams(instance, node, &newMemcpyParams);
// After updating nodes, optionally call:
cudaGraphExecUpdate(instance, graph, nullptr, nullptr);

void CUDART_CB myCallback(cudaStream_t stream, cudaError_t status, void* data) {
    printf("Stream done: %p\n", data);
}

cudaStreamAddCallback(stream, myCallback, myData, 0);
// 0 = no flags (currently the only option)

auto hostFn = [](void* data) {
    int* val = (int*)data;
    *val = 42;   // runs on CPU in stream order
};

cudaLaunchHostFunc(stream, hostFn, myDataPtr);

// Process A — export
cudaIpcMemHandle_t memHandle;
cudaIpcGetMemHandle(&memHandle, d_ptr);
// Send memHandle via IPC mechanism (pipe, shared mem, etc.)

cudaIpcEventHandle_t evHandle;
cudaIpcGetEventHandle(&evHandle, event);
// Send evHandle

// Process B — import
float* d_shared;
cudaIpcOpenMemHandle((void**)&d_shared, memHandle, cudaIpcMemLazyEnablePeerAccess);
cudaEvent_t sharedEvent;
cudaIpcOpenEventHandle(&sharedEvent, evHandle);

cudaStreamWaitEvent(myStream, sharedEvent, 0);
useKernel<<<g, b, 0, myStream>>>(d_shared);

cudaIpcCloseMemHandle(d_shared);