Validate CUDA stream dependencies and buffer lifetimes on DGX Spark

When two CUDA streams share the same device buffer, a program may produce correct results over several runs, then return stale values when the input size changes. The cause is often not the operator’s formula: the consumer stream has started reading before the producer stream has finished writing. This article builds a minimal two-stream test on DGX Spark to validate submission order, execution dependencies, and buffer reuse separately. A single successful run is not enough to establish correctness.

1. Map the data dependencies before creating streams

DGX Spark uses NVIDIA GB10 Grace Blackwell and provides 128GB of LPDDR5x unified memory, with the CPU and GPU sharing the same memory pool. Physically unified memory does not, however, establish cross-stream execution order for the application. Explicit copies, kernel launches, and host accesses must still respect their respective completion requirements.

The experiment uses explicit device allocation and page-locked host memory to make the boundaries easier to observe. The producer stream copies the input to the device and multiplies each element by two. The consumer stream adds one to each element and copies the result back to the host. The full dependency chain is:

Input ready → copy complete → multiplication complete → addition starts → copy back complete → host validation.

Operations within a single stream execute in stream order; dependencies between two streams must be established explicitly. Record an event after the multiplication kernel in the producer stream. The consumer stream must wait for that event before the addition kernel is submitted.

Calling the producer stream first and the consumer stream second on the host establishes only submission order. Cross-stream data dependencies need an explicit synchronization mechanism; they cannot rely on an execution order that happens to hold under a light workload.

2. Build a minimal program that does not rely on the default stream

After paying on TcpNode, retrieve the connection credentials from the console and connect to NVIDIA DGX OS over SSH. First run uname -m, nvidia-smi, and nvcc --version, saving the host architecture, device and driver details, and compiler toolchain information, respectively. The architecture should be aarch64, corresponding to ARM64. Do not infer the compiler version from the driver output.

Save the following as stream_contract.cu. Both streams use the non-blocking flag so that implicit synchronization with the legacy default stream cannot mask missing dependencies. The event is used only for ordering, not for timing.

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>

static void check(cudaError_t result) {
    if (result != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(result));
        std::exit(1);
    }
}

__global__ void scale(float* data, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count) data[i] *= 2.0f;
}

__global__ void add(float* data, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count) data[i] += 1.0f;
}

int main() {
    constexpr int count = 4096;
    const size_t bytes = count * sizeof(float);
    float *input = nullptr, *output = nullptr, *device = nullptr;
    cudaStream_t producer, consumer;
    cudaEvent_t ready;

    check(cudaMallocHost(reinterpret_cast<void**>(&input), bytes));
    check(cudaMallocHost(reinterpret_cast<void**>(&output), bytes));
    check(cudaMalloc(reinterpret_cast<void**>(&device), bytes));
    check(cudaStreamCreateWithFlags(&producer, cudaStreamNonBlocking));
    check(cudaStreamCreateWithFlags(&consumer, cudaStreamNonBlocking));
    check(cudaEventCreateWithFlags(&ready, cudaEventDisableTiming));

    for (int i = 0; i < count; ++i) {
        input[i] = static_cast<float>(i);
        output[i] = -1.0f;
    }

    check(cudaMemcpyAsync(device, input, bytes,
                          cudaMemcpyHostToDevice, producer));
    scale<<<(count + 255) / 256, 256, 0, producer>>>(device, count);
    check(cudaGetLastError());
    check(cudaEventRecord(ready, producer));

    check(cudaStreamWaitEvent(consumer, ready, 0));
    add<<<(count + 255) / 256, 256, 0, consumer>>>(device, count);
    check(cudaGetLastError());
    check(cudaMemcpyAsync(output, device, bytes,
                          cudaMemcpyDeviceToHost, consumer));
    check(cudaStreamSynchronize(consumer));

    int failures = 0;
    for (int i = 0; i < count; ++i) {
        float expected = 2.0f * static_cast<float>(i) + 1.0f;
        if (output[i] != expected) ++failures;
    }
    std::printf("elements=%d mismatches=%d\n", count, failures);

    check(cudaEventDestroy(ready));
    check(cudaStreamDestroy(consumer));
    check(cudaStreamDestroy(producer));
    check(cudaFree(device));
    check(cudaFreeHost(output));
    check(cudaFreeHost(input));
    return failures == 0 ? 0 : 2;
}

3. Validate both the output and the exit code

Compile and run in the same directory:

nvcc -O2 stream_contract.cu -o stream_contract
./stream_contract
printf 'exit=%s\n' "$?"

A passing result is elements=4096 mismatches=0 with an exit code of zero. An API or synchronization failure returns one; a numerical mismatch returns two. Checking only that the process did not crash will miss computations that completed but produced incorrect data.

All values in this example are integers within the range that single precision can represent exactly, so a direct equality comparison is appropriate. Do not carry this criterion over unchanged to a real floating-point model. Define absolute and relative error tolerances, along with rules for handling non-finite values, according to the algorithm.

Next, change the element count to 1, 257, and 4097 in turn, recompiling and running each time. These sizes cover a single element, a partial thread block, and boundaries across multiple blocks. They are test inputs, not performance measurements. For every test, save the source file, compilation command, toolchain information, and program output rather than keeping only a screenshot of a successful run.

Verify that the failure path actually works

Temporarily change the final 1.0f in the expected-value expression to 2.0f. The program should report mismatches and return two. Then restore the original expression. This confirms that result validation is working. Do not use “removing the wait must cause an error” as a negative test: whether a race manifests depends on scheduling, and a program with a missing dependency may still pass for a while.

4. Apply event boundaries to buffer lifetimes

When cudaMemcpyAsync returns, it generally tells you only that the call has been submitted—not that it is safe to overwrite the input. In this example, the input remains unchanged from copy submission until synchronization of the consumer stream has completed. The output is likewise read only after synchronizing the stream that performs the copy back. This is a conservative policy, but its boundaries are clear, making it a useful starting point for correctness.

Page-locked memory helps meet the requirements for asynchronous transfers; it is not a completion notification. Likewise, a device buffer must not be freed while another stream is still accessing it. Having 128GB of unified memory does not mean that large amounts of host memory should all be page-locked. Batch capacity should be constrained by the actual working set.

The producer event covers only operations that precede it. It does not cover the addition or the copy back in the consumer stream. Completion of the producer event therefore confirms that the production stage has finished, but not that the device buffer is ready to be overwritten by the next batch. When extending the program into a pipeline, record a completion event after the copy back in the consumer stream and use it to control when the entire set of buffers can be reclaimed.

Each concurrent batch slot should have its own input, output, device buffer, and completion state. Validate a single slot first, then add more slots. Do not simply increase the number of streams while continuing to share the same pointer without protection.

5. Narrow the synchronization scope during troubleshooting

If results are intermittently incorrect, first check that the wait is placed before the consumer kernel, that the event is recorded after the producer kernel, and that the record and wait operations use the same event. Moving the event record one line earlier can incorrectly weaken “the data has been fully written” to “the copy has completed.”

Next, check for errors after every kernel launch and inspect the return value of the final synchronization call. The former catches some launch problems; the latter reports errors from asynchronous execution. Neither replaces the other. If you check only for launch errors without waiting for completion, an error may not surface until a later API call.

In a diagnostic copy of the program, you can temporarily synchronize the producer stream to determine whether the problem is related to cross-stream ordering. If the error disappears, continue investigating dependencies and buffer ownership rather than treating device-wide synchronization as the final fix. Global synchronization changes scheduling and can hide the original lifetime problem.

The final implementation should retain explicit producer events, consumer waits, and buffer-reclamation events, together with repeatable output validation. Only when both dependency edges and memory ownership are clearly accounted for does a multi-stream program have a sound basis for further concurrency.

Frequently asked questions