Find CUDA out-of-bounds accesses on DGX Spark

A CUDA kernel launches, but the error only appears when the results are copied back. The same program runs successfully with a smaller input. Symptoms like these call for checking device memory accesses before changing drivers. On DGX Spark, a small standalone program can establish a complete diagnostic chain: deliberately trigger an out-of-bounds access, locate the source line, fix the index, and then verify that error detection is still active. The process requires no model weights and does not involve a production service.

1. Establish the environment and isolate the test

TcpNode’s DGX Spark uses NVIDIA GB10 Grace Blackwell, the ARM64 architecture, and NVIDIA DGX OS, with 128GB of LPDDR5x unified memory. Sharing memory between the CPU and GPU does not mean pointers can access memory beyond their allocation boundaries. The example below uses cudaMalloc to allocate device memory, making the valid address range explicit.

After payment, retrieve the connection credentials from the console and log in over SSH. Use a separate working directory and run the test as a regular user. Do not inject errors into an inference process that is serving requests.

mkdir -p "$HOME/cuda-access-check"
cd "$HOME/cuda-access-check"
uname -m
nvidia-smi
command -v nvcc
command -v compute-sanitizer
nvcc --version
compute-sanitizer --version

The architecture should be reported as aarch64. Record the tool versions and driver information, and use the same environment for subsequent verification. A preinstalled CUDA stack does not guarantee that every tool is available in the current PATH. If a command is missing, first check the installed toolchain and its installation paths rather than replacing the entire driver stack.

Run the out-of-bounds example only in a separate, disposable process. After an illegal memory access, terminate that process and start a new run. Subsequent results from the same CUDA context are not valid evidence.

2. Build an out-of-bounds reproducer with a single variable

Save the following as repro.cu. It allocates space for just one integer, so zero is the only valid index. The write position is selected at compile time using TEST_INDEX. This keeps the faulty and corrected versions in the same source file, without changing the input, allocation size, and execution path at the same time.

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

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

__global__ void write_value(int* data) {
    data[TEST_INDEX] = 7;
}

int main() {
    int* data = nullptr;
    int result = 0;
    checked(cudaMalloc(reinterpret_cast<void**>(&data), sizeof(int)));
    checked(cudaMemset(data, 0, sizeof(int)));
    write_value<<<1, 1>>>(data);
    checked(cudaGetLastError());
    checked(cudaDeviceSynchronize());
    checked(cudaMemcpy(&result, data, sizeof(int), cudaMemcpyDeviceToHost));
    checked(cudaFree(data));
    std::printf("result=%d\n", result);
    return result == 7 ? EXIT_SUCCESS : EXIT_FAILURE;
}

Compile the two programs separately:

nvcc -arch=native -lineinfo -DTEST_INDEX=1 repro.cu -o repro-bad
nvcc -arch=native -lineinfo -DTEST_INDEX=0 repro.cu -o repro-good
sha256sum repro.cu repro-bad repro-good > build.sha256

-arch=native compiles for the GPU currently visible to the compiler. It is suitable for local verification, not for producing a portable build for different devices. If the compiler does not support the current device, check toolchain compatibility first rather than substituting an arbitrary architecture option. -lineinfo preserves device source-line information. Initial diagnosis does not require full device debugging, which can substantially change execution behavior.

Why check both launch and synchronization errors?

cudaGetLastError checks for launch-related errors, but kernels generally execute asynchronously. cudaDeviceSynchronize waits for preceding device work to finish, allowing errors that occur during execution to be returned there. Check both calls separately; whether the program prints a result is not enough.

3. Turn detection results into an actionable exit status

Run the faulty version first, saving standard output and the diagnostic log separately:

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-bad > bad.stdout 2> bad.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat bad.stdout bad.stderr

Expect a report of an invalid global memory write pointing to the write operation in write_value. The report typically also includes the access size, thread location, and the address’s relation to the allocation. Its exact formatting varies by tool version, so do not use a fixed text fragment as the sole pass/fail criterion.

--error-exitcode 86 makes detected errors affect the automation result. If the tested program itself fails, however, the result may be a different nonzero exit code returned by the program. The acceptance criterion is therefore “a nonzero exit status and a report pointing to the expected out-of-bounds access,” not that every failure must return 86.

Do not pipe the command directly into an ordinary logging pipeline and then check only the status of the final logging command. The redirection and conditional branch above preserve the actual run status and also work in scripts configured to exit on failure.

4. Fix indexing based on the address range, not the symptoms

The example allocates space for one integer but writes at index one, which is already outside the valid range. In a real program, associate the reported address with its allocation, then check the index calculation, element type, stride, and batch offsets.

Common causes include treating allocation sizes expressed in bytes as element counts when indexing; a final thread block with fewer valid elements than a full block; a two-dimensional array whose row stride differs from its logical width; and another task using a pointer after it has been freed. Increasing the allocation may temporarily hide an error, but it does not prove that the algorithm is correct.

For one-dimensional work, the usual approach is to check index < count before accessing memory and ensure that the integer type used for index calculations can hold the product. In kernels that cooperate through shared memory, do not arbitrarily let some threads return early while the remaining threads reach a block-level barrier. That can introduce new synchronization problems.

If the error occurs only in a large workflow, temporarily add synchronization checks after suspect stages to narrow down the location. Add one checkpoint at a time and preserve the first error. Do not keep submitting work after an error and generate a flood of secondary failures.

5. Verify the fix with positive and negative tests

Run the corrected version:

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-good > good.stdout 2> good.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat good.stdout good.stderr

This example must meet all three conditions: a zero exit code, no memory access errors in the diagnostic report, and output of result=7. “No errors reported” is not sufficient. A test can appear to pass even if the kernel never runs or the result is never actually written.

Then run the faulty version again and confirm that it still fails. This negative test can reveal a removed detection command, a test accidentally using an old binary, or a pipeline that swallows failure statuses. Tests for a real kernel should also cover the smallest valid input, full-block boundaries, and a partial final block. If the interface permits empty input, define how it is handled as well.

Finally, archive the source code, build commands, version information, checksums, and logs from both runs. The detector changes execution overhead, so timings measured under detection cannot be used directly as performance conclusions. Passing memcheck also does not rule out data races or synchronization problems. A reliable conclusion leaves independently reviewable evidence for valid memory accesses, correct results, and failure propagation—not just a screenshot of a successful run.

Frequently asked questions