A stream is an ordered queue of device work. Operations in the same stream run in order; operations in different streams may run concurrently. Everything goes into the default stream unless told otherwise, which is why a naive program serialises transfers and kernels that could have overlapped.

cudaStream_t s;
cudaStreamCreate(&s);
cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, s);
kernel<<<blocks, threads, 0, s>>>(d_a);
cudaMemcpyAsync(h_a, d_a, bytes, cudaMemcpyDeviceToHost, s);
cudaStreamSynchronize(s);
cudaStreamDestroy(s);

The main use is overlapping transfers with computation. A GPU has separate copy engines and compute units, so while one chunk is being processed another can be moving across PCIe. Splitting the work into chunks and cycling through a few streams hides most of the transfer cost:

for (int c = 0; c < nchunks; c++) {
    int s = c % nstreams;
    cudaMemcpyAsync(d + off, h + off, csize, cudaMemcpyHostToDevice, stream[s]);
    kernel<<<b, t, 0, stream[s]>>>(d + off);
    cudaMemcpyAsync(h + off, d + off, csize, cudaMemcpyDeviceToHost, stream[s]);
}

With enough chunks this approaches the time of the compute alone, since every transfer except the first and last is hidden behind a kernel.

Asynchronous copies require pinned host memory. With pageable memory cudaMemcpyAsync silently falls back to synchronous behaviour and nothing overlaps, which is a common reason streaming code shows no speedup.

Streams also let independent small kernels run at the same time. A kernel that only fills part of the GPU leaves the rest idle, and issuing several such kernels on different streams fills the gap.

For dependencies across streams, use events rather than a full device synchronise:

cudaEventRecord(ev, stream1);
cudaStreamWaitEvent(stream2, ev, 0);   // stream2 waits for that point in stream1