Auditing CUDA build artifacts on DGX Spark: targets, SASS and PTX

If a CUDA project compiles on DGX Spark but fails at kernel launch with no kernel image is available for execution on the device, first check which GPU architectures the build artifacts support rather than reinstalling the entire environment. A subtler problem is an application that has always relied on PTX just-in-time compilation, while its release records incorrectly claim that native device code is included. This article uses a small, standalone program to bring target detection, artifact inspection, and testing of both execution paths into a single acceptance workflow.

1. Distinguish the host architecture from the GPU target

DGX Spark uses NVIDIA GB10 Grace Blackwell. Its host architecture is ARM64, and it runs NVIDIA DGX OS, based on Ubuntu 24.04, with the CUDA stack preinstalled. After payment, retrieve the connection credentials from the console and connect over SSH.

ARM64 describes the CPU-side executable architecture; CUDA compute capability describes the target for GPU device code. These are separate compatibility requirements. An AArch64 result from file confirms only that the host-side format is correct, not that the binary contains usable kernels.

SASS is machine code for a target GPU, while PTX is an intermediate representation that the driver can compile just in time. A release package can include both, but each must be checked separately. A successful run under normal conditions does not verify both paths.

The following experiment creates source files and executables only in a new directory. It does not modify the driver, system toolchain, or application code. The example results are for compatibility acceptance testing, not performance measurement.

2. Query the actual compute capability instead of guessing the target

Start by creating a separate working directory and recording the system and toolchain information. The CUDA version shown by nvidia-smi indicates what the driver supports; it is not the toolkit version used by the current nvcc. Keep both outputs.

mkdir -p "$HOME/cuda-artifact-check"
cd "$HOME/cuda-artifact-check"
uname -m > host-arch.txt
nvidia-smi > driver.txt
nvcc --version > toolkit.txt
nvcc --list-gpu-code > supported-targets.txt

Save the following as probe.cu. The program uses the CUDA Runtime to query the compute capability of device zero. On failure, it returns a nonzero exit code rather than printing a target identifier that could be used incorrectly.

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

int main() {
    cudaDeviceProp p{};
    cudaError_t e = cudaGetDeviceProperties(&p, 0);
    if (e != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(e));
        return 1;
    }
    std::printf("%d%d\n", p.major, p.minor);
    return 0;
}

nvcc probe.cu -o probe
CC=$(./probe) || exit 1
printf '%s\n' "$CC" > compute-capability.txt
grep -Fx "sm_${CC}" supported-targets.txt || exit 1

The GB10 target identifier is deliberately not hard-coded here. If the toolkit's list does not include the detected sm_ target, check the toolkit's supported targets first. Do not substitute an older identifier just to proceed with the release. If the probe fails, resolve device-access or runtime-library issues before moving on to compilation.

3. Explicitly generate and inspect SASS and PTX

Save the following as check.cu. It launches just one thread to write an integer, then copies the value back to the host and checks it. Every CUDA call is checked for errors so that a failed kernel launch cannot be followed by a misleading success message.

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

#define CUDA_OK(call) do { \
    cudaError_t e = (call); \
    if (e != cudaSuccess) { \
        std::fprintf(stderr, "%s\n", cudaGetErrorString(e)); \
        return 1; \
    } \
} while (0)

__global__ void write_value(int* p) {
    *p = 7;
}

int main() {
    int* device = nullptr;
    int host = 0;
    CUDA_OK(cudaMalloc(&device, sizeof(int)));
    write_value<<<1, 1>>>(device);
    CUDA_OK(cudaGetLastError());
    CUDA_OK(cudaDeviceSynchronize());
    CUDA_OK(cudaMemcpy(&host, device, sizeof(int),
                       cudaMemcpyDeviceToHost));
    CUDA_OK(cudaFree(device));
    if (host != 7) return 2;
    std::puts("PASS");
    return 0;
}

Run the following in the same working directory, explicitly requesting both native code and PTX:

CC=$(cat compute-capability.txt)
nvcc -O2 \
  -gencode "arch=compute_${CC},code=[sm_${CC},compute_${CC}]" \
  check.cu -o check

file check > binary-format.txt
cuobjdump --list-elf ./check > elf-list.txt
cuobjdump --dump-sass ./check > sass.txt
cuobjdump --dump-ptx ./check > ptx.txt

Check that the ELF listing and SASS output contain the expected target and kernel, and that the PTX output includes .version, .target, and the kernel entry point. Empty files, a mismatched target, or a missing expected kernel must block acceptance. Checking command exit codes alone is not enough.

Inspect the final artifact, not intermediate object files

A real project may involve static-library linking, device linking, or packaging steps. Finding PTX in an object file does not prove that it survives in the final deliverable. Repeat the checks on the final executable and the shared libraries actually loaded, and record their file hashes.

4. Verify native execution and just-in-time compilation separately

First disable PTX just-in-time compilation so that the current artifact must run using native device code. Then force the PTX path and disable the on-disk compilation cache. Use a new process for each test to prevent modules already loaded in a process from affecting the result.

env -u CUDA_FORCE_PTX_JIT \
  CUDA_DISABLE_PTX_JIT=1 \
  ./check > native.log 2>&1
native_rc=$?

env -u CUDA_DISABLE_PTX_JIT \
  CUDA_FORCE_PTX_JIT=1 CUDA_CACHE_DISABLE=1 \
  ./check > ptx-jit.log 2>&1
ptx_rc=$?

printf 'native=%s\nptx_jit=%s\n' \
  "$native_rc" "$ptx_rc" > execution-status.txt

cat execution-status.txt native.log ptx-jit.log
test "$native_rc" -eq 0 && test "$ptx_rc" -eq 0

Both logs should contain PASS, and both exit codes should be zero. The test variables apply only to the individual commands. Do not add them to the global environment and leave application processes permanently forced into JIT compilation.

If native execution fails but PTX succeeds, first revisit the native target and final linking step. If native execution succeeds but PTX fails, check whether PTX is present and whether the driver supports its version. If both fail, use the specific errors to investigate the toolchain, runtime libraries, and device state. Repeatedly switching compiler options is no substitute for diagnosing the cause.

5. Turn the test results into a release gate

Passing the small test program proves only that the basic compilation toolchain and this particular kernel work in the current environment. The application may also load extension libraries, choose different kernels based on input, or generate code at runtime. Tests through the actual application entry points are therefore required, covering the main data types and branches.

Keep four sets of evidence for each release:

• Host architecture, driver output, toolkit version, and the actual compilation commands.
• GPU compute capability and the SASS and PTX inspection results for the final artifacts.
• Exit codes and logs for both execution paths, along with records of application-result checks.
• Source revision, SHA-256 hashes of the final artifacts, and a list of dynamic dependencies.

sha256sum ./check identifies the exact artifact accepted in this test. Any subsequent relinking requires a new acceptance test, even if the filename stays the same. If you use containers, inspect the toolchain and final artifacts inside the actual delivery image, and verify execution in the actual runtime environment.

PTX is a compatibility path, not a guarantee of support across all drivers. Nor does one successful test establish that the program will run on other devices. The release gate should explicitly state: “Passed with the current driver, current artifacts, and tested entry points.” Regenerate the evidence whenever the architecture or toolchain changes rather than reusing old logs.

Frequently asked questions