Streams: Overlap Copies with Compute

This is a host-side lab: you write the orchestration, not a kernel. The serial version copies the whole array to the device, runs the compute kernel, and copies it back - three phases that each wait for the last. Split the work into chunks and drive them across several CUDA streams so a chunk's copy overlaps the previous chunk's compute. The harness times your version against the serial baseline and gates on the speedup.

Languages: CUDA C++, Rust

Lesson

A serial pipeline leaves the hardware idle: while the host-to-device copy runs, the compute units and the device-to-host copy engine sit waiting, and vice-versa. CUDA streams let independent work run concurrently - copies use the copy engines while a kernel uses the SMs. Chunk the array into pieces and, for each piece on its own stream, issue an async host-to-device copy, the kernel, and an async device-to-host copy. Because the pieces are independent, the GPU can be copying chunk i+1's input while computing chunk i and copying chunk i-1's result out. Two requirements make the overlap real: the host memory must be pinned (the harness gives you a pinned buffer - regular pageable memory can't overlap), and the copies must be the async form (cudaMemcpyAsync on a stream), since the plain cudaMemcpy is synchronous.

Real overlap requires pinned host buffers, independent work and supported copy/compute engines. non_blocking or a second stream alone does not prove overlap; inspect concurrent timeline intervals.

What you will learn

Recognize that a copy -> compute -> copy pipeline runs the three phases serially, leaving engines idle.
Create CUDA streams and drive independent chunks of work across them.
Use cudaMemcpyAsync (not cudaMemcpy) on a stream so copies overlap compute.
Understand why the overlap needs pinned host memory and independent per-chunk work.

Starter code

    // Serial: one big copy in, one launch, one copy back - each phase waits for the
    // last, so the copy engines and the SMs are never busy at the same time.
    cudaMemcpy(d_data, h_data, (size_t)N * sizeof(float), cudaMemcpyHostToDevice);
    streamKernel<<<(N + blockSize - 1) / blockSize, blockSize>>>(d_data, N);
    cudaMemcpy(h_data, d_data, (size_t)N * sizeof(float), cudaMemcpyDeviceToHost);