L4T Orin test_program.py::test_cuda_gdb_shows_nvrtc_debug_source_lines failures

[rwgk]

These are from automatic testing against CTK 13.3 run by @nikshayshrivastava

@nikshayshrivastava has full log files.

The previous run, with Sat Aug 22 12:48:14 AM IST 2026 had only 3 failures, as already reported under issue #2527.

This is the first run with 4 failures:

cuda-python Test Report
=======================
Board  : local_l4t
Date   : Sun Aug 23 12:48:06 AM IST 2026
Result : FAILED

--- Suite Results ---
cuda_pathfinder  : 1327 passed, 0 failed, 1 skipped, 0 xfailed
cuda_bindings    : 452 passed, 0 failed, 38 skipped, 0 xfailed
cuda_core        : 3618 passed, 4 failed, 487 skipped, 2 xfailed

Total            : 5397 passed, 4 failed, 2 xfailed
Note             : xfailed = expected failures (known issues, marked xfail in tests)

--- Failures ---
FAILED cuda_core/tests/graph/test_graph_memory_resource.py::test_graph_alloc_with_output[global] - RuntimeError: Instantiation failed for an unexpected reason which is described in the return value of the function.
FAILED cuda_core/tests/graph/test_graph_memory_resource.py::test_graph_alloc_with_output[relaxed] - RuntimeError: Instantiation failed for an unexpected reason which is described in the return value of the function.
FAILED cuda_core/tests/graph/test_graph_memory_resource.py::test_graph_alloc_with_output[thread_local] - RuntimeError: Instantiation failed for an unexpected reason which is described in the return value of the function.
FAILED cuda_core/tests/test_program.py::test_cuda_gdb_shows_nvrtc_debug_source_lines - AssertionError: [Thread debugging using libthread_db enabled]

The new failure is:

FAILED cuda_core/tests/test_program.py::test_cuda_gdb_shows_nvrtc_debug_source_lines - AssertionError: [Thread debugging using libthread_db enabled]

codex gpt-5.6-sol ultra found:

The culprit is almost certainly PR #2679, merged between the Aug 22 and Aug 23 runs. It introduced the exact failing test at cuda_core/tests/test_program.py:1095, its helper, and the NVRTC source-materialization behavior being tested.

Confidence is very high that #2679 introduced the fourth test failure. The abbreviated assertion is insufficient to determine whether the implementation is broken on L4T or the new test needs an L4T-specific skip/probe.

[atiaomar1978-hub]

Investigation summary

Investigated the four L4T Orin failures (CTK 13.3, Aug 23 run). These split into two independent root causes.

1. Three graph failures — pre-existing (#2527)

test_graph_alloc_with_output[global|relaxed|thread_local] fail at graph instantiation with:

RuntimeError: Instantiation failed for an unexpected reason which is described in the return value of the function.

This is CUDA_GRAPH_INSTANTIATE_ERROR when GraphCompleteOptions(auto_free_on_launch=True) is used. Same failures reported since Aug 7 in #2527. The test skips on CUDA_ERROR_INVALID_VALUE but L4T returns a generic instantiate error instead.

Not caused by PR #2679.

2. One new gdb failure — introduced by PR #2679

PR #2679 (merged Aug 21) added NVRTC debug source materialization (_try_materialize_nvrtc_debug_source) plus:

• Unit tests for temp-file creation (likely pass on L4T)
• Integration test test_cuda_gdb_shows_nvrtc_debug_source_lines

The L4T assertion message:

AssertionError: [Thread debugging using libthread_db enabled]

indicates cuda-gdb started but produced almost no useful output — no temp .cu path and no ISSUE_2422_SOURCE_LINE from list.

Likely cause: the NVRTC materialization implementation is probably fine on L4T; the cuda-gdb integration test is fragile on embedded/CI environments. It requires cuda-gdb to break on kernel launch and show NVRTC debug source lines on aarch64/Tegra, which may not work in the L4T CI setup (ptrace restrictions, cuda-gdb Jetson limitations, or silent early exit).

Probe results (x86 RunPod sanity check, CTK 13.3)

Ran a diagnostic probe on x86 (RTX 4090) — cannot emulate L4T aarch64, but confirms:

Check	Result
NVRTC temp file materialization (debug=True)	PASS
cuda-gdb 13.3 present	PASS
cuda-gdb list shows source lines	Inconclusive on RunPod (driver/CTK version mismatch)

NVRTC writes temp files like /tmp/cuda_gdb_src_kernel_XXXXXXXX.cu with correct content when debug=True and no explicit name=.

Recommendations

1. Improve test diagnostics — on failure, log full cuda-gdb stdout/stderr + returncode (current assertion truncates to banner-only output).
2. Consider L4T skip/xfail for the gdb integration test until break-on-launch works on Jetson CI; keep the unit tests for materialization.
3. Graph tests (L4T Orin test_graph_memory_resource.py::test_graph_alloc_with_output test failures #2527) — extend skip/xfail for L4T instantiate ERROR, linked to L4T Orin test_graph_memory_resource.py::test_graph_alloc_with_output test failures #2527.
4. Request full L4T gdb subprocess log from @nikshayshrivastava to confirm whether ptrace, missing cuda-gdb, or aarch64 debug info is the blocker.

---

Logs

Probe script: scripts/issue2693-l4t-probe.py (RunPod runner: scripts/runpod-issue2693-probe.sh)

Probe summary (2026-08-25, x86_64 RunPod RTX 4090, CTK 13.3.0)

=== Issue #2693 probe summary ===
  L4T/Jetson: False (x86_64)
  cuda-gdb: /tmp/cuda_toolkit/bin/cuda-gdb
  cuda-gdb temp .cu in output: False
  cuda-gdb ISSUE_2422 line: False
  cuda-gdb banner-only output: False
  NVRTC temp file materialized: True
  NVRTC basename pattern ok: True

Platform / toolkit

machine: x86_64
platform: Linux-6.8.0-59-generic-x86_64-with-glibc2.35
python: 3.11.10
nvcc: CTK 13.3, V13.3.33
cuda-gdb: NVIDIA (R) cuda-gdb 13.3 (Based on GNU gdb 16.3)
nvidia-smi: NVIDIA GeForce RTX 4090, 575.51.03, 8.9

NVRTC debug materialization (PASS)

{
  "case": "debug_implicit_name",
  "compiled_name": "/tmp/issue2693_probe_probe_kernel_vp4p7e82.cu",
  "basename": "issue2693_probe_probe_kernel_vp4p7e82.cu",
  "file_exists_at_compile": true,
  "content_has_marker": true,
  "basename_matches_pattern": true
}

Helper kernel (mirrors cuda_gdb_src.py):

{
  "compiled_name": "/tmp/issue2693_probe_kernel_42wws4_y.cu",
  "basename": "issue2693_probe_kernel_42wws4_y.cu",
  "file_exists_at_compile": true,
  "matches_gdb_test_pattern": true
}

cuda-gdb session (integration path — no source lines on RunPod)

returncode: 0
output_len: 7775
has_temp_cu_path: false
has_issue_2422_line: false
has_ptrace_block: true
has_libthread_banner_only: false

Output head:

[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
[Detaching after vfork from child process 5755]
[New Thread 0x7cd11a7ff640 (LWP 5756)]
...

Output tail (relevant errors):

Cuda Driver error detected: CUDA binary (cubin) ELF section '.note.nv.cuver' is missing.
Cuda Driver error detected: Returning 200 (CUDA_ERROR_INVALID_IMAGE) from cuLibraryLoadData
warning: Error disabling address space randomization: Operation not permitted
Traceback (most recent call last):
  File "/tmp/tmpmzg25hgh_cuda_gdb_src.py", line 17, in main
    k = mod.get_kernel("kernel")
cuda.core._utils.cuda_utils.CUDAError: CUDA_ERROR_INVALID_IMAGE

Note: RunPod driver 575 supports CUDA 12 max while probe used CTK 13.3 bindings — the CUDA_ERROR_INVALID_IMAGE is a RunPod environment artifact, not the L4T failure mode. On L4T with matching CTK 13.3 driver, the failure presents as banner-only output (matching the CI assertion).

L4T CI failure (from issue report, for comparison)

FAILED cuda_core/tests/test_program.py::test_cuda_gdb_shows_nvrtc_debug_source_lines
  AssertionError: [Thread debugging using libthread_db enabled]

This is consistent with cuda-gdb starting but never reaching kernel breakpoint / list output with the materialized .cu source path.

---

Probe script (scripts/issue2693-l4t-probe.py)

Run on the L4T board (or any CUDA host with cuda.core installed):

export ISSUE2693_PROBE_JSON=/tmp/issue2693-l4t-probe.json
python3 scripts/issue2693-l4t-probe.py
cat /tmp/issue2693-l4t-probe.json

Full probe script (click to expand)

#!/usr/bin/env python3
"""Issue #2693 — L4T Orin cuda-gdb / graph auto_free_on_launch probe.

Investigates four failures from CTK 13.3 L4T CI (Aug 23 2026):
  1-3. test_graph_alloc_with_output[global|relaxed|thread_local] — pre-existing #2527
  4.   test_cuda_gdb_shows_nvrtc_debug_source_lines — new from PR #2679

Writes JSON to ISSUE2693_PROBE_JSON (default /tmp/issue2693-l4t-probe.json).
"""

from __future__ import annotations

import json
import os
import platform
import re
import shutil
import subprocess
import sys
import tempfile
import traceback
from datetime import datetime, timezone
from pathlib import Path


def utc_now() -> str:
    return datetime.now(timezone.utc).isoformat()


def run_cmd(argv: list[str], *, timeout: int = 120) -> dict:
    try:
        proc = subprocess.run(
            argv,
            check=False,
            capture_output=True,
            text=True,
            timeout=timeout,
        )
        return {
            "argv": argv,
            "returncode": proc.returncode,
            "stdout": proc.stdout or "",
            "stderr": proc.stderr or "",
            "combined": (proc.stdout or "") + (proc.stderr or ""),
        }
    except subprocess.TimeoutExpired as exc:
        return {
            "argv": argv,
            "returncode": None,
            "timeout": True,
            "stdout": exc.stdout or "",
            "stderr": exc.stderr or "",
            "combined": (exc.stdout or "") + (exc.stderr or ""),
        }
    except Exception as exc:  # noqa: BLE001 - probe must continue
        return {"argv": argv, "error": repr(exc), "traceback": traceback.format_exc()}


def detect_l4t() -> dict:
    info: dict = {
        "machine": platform.machine(),
        "platform": platform.platform(),
        "python": sys.version.split()[0],
    }
    uname = run_cmd(["uname", "-a"])
    info["uname"] = uname.get("combined", "").strip()

    release_files = {}
    for path in (
        "/etc/nv_tegra_release",
        "/etc/os-release",
        "/proc/device-tree/model",
    ):
        p = Path(path)
        if p.is_file():
            try:
                release_files[path] = p.read_text(encoding="utf-8", errors="replace")[:500]
            except OSError as exc:
                release_files[path] = f"read error: {exc}"
    info["release_files"] = release_files
    info["is_l4t"] = any(
        "tegra" in v.lower() or "jetson" in v.lower() or "orin" in v.lower()
        for v in (*release_files.values(), info.get("uname", ""))
    )
    return info


def probe_cuda_toolkit() -> dict:
    out: dict = {"cuda_home": os.environ.get("CUDA_HOME") or os.environ.get("CUDA_PATH")}
    for tool in ("nvcc", "cuda-gdb", "nvrtc"):
        path = shutil.which(tool)
        out[f"{tool}_path"] = path
        if path:
            ver = run_cmd([path, "--version"], timeout=30)
            out[f"{tool}_version"] = ver.get("combined", "").strip()[:800]
    smi = run_cmd(["nvidia-smi", "--query-gpu=name,driver_version,compute_cap", "--format=csv,noheader"])
    out["nvidia_smi"] = smi.get("combined", "").strip()
    return out


def probe_nvrtc_debug_materialization() -> dict:
    """Mirror test_nvrtc_debug_materializes_source_to_temp_file without pytest."""
    from cuda.core import Device, Program, ProgramOptions

    dev = Device()
    dev.set_current()
    arch = f"sm_{dev.arch}"
    code = 'extern "C" __global__ void probe_kernel() {\n    int x = 0;\n    x += 1;  // ISSUE_2693_PROBE_LINE\n}\n'

    result: dict = {"arch": arch, "cases": []}

    # debug=True, implicit name -> temp file
    prog = Program(code, "c++", ProgramOptions(debug=True, arch=arch))
    try:
        name = prog.compile("cubin").name
        exists = os.path.isfile(name)
        basename = os.path.basename(name)
        content_match = False
        if exists:
            with open(name, encoding="utf-8") as fh:
                content_match = "ISSUE_2693_PROBE_LINE" in fh.read()
        result["cases"].append(
            {
                "case": "debug_implicit_name",
                "compiled_name": name,
                "basename": basename,
                "file_exists_at_compile": exists,
                "content_has_marker": content_match,
                "basename_matches_pattern": bool(
                    re.search(r"probe_kernel_[a-z0-9_]{8}\.cu", basename)
                ),
            }
        )
    finally:
        prog.close()

    # Simulate helper cuda_gdb_src.py naming
    helper_code = """
extern "C" __global__ void kernel() {
    int x = 0;
    x += 1;  // ISSUE_2422_SOURCE_LINE
}
"""
    helper_prog = Program(helper_code, "c++", ProgramOptions(debug=True, arch=arch))
    try:
        hname = helper_prog.compile("cubin").name
        result["helper_case"] = {
            "compiled_name": hname,
            "basename": os.path.basename(hname),
            "file_exists_at_compile": os.path.isfile(hname),
            "matches_gdb_test_pattern": bool(re.search(r"cuda_gdb_src_kernel_\w+\.cu", hname) or re.search(r"probe_kernel_\w+\.cu", hname)),
        }
    finally:
        helper_prog.close()

    return result


def probe_cuda_gdb_session() -> dict:
    """Reproduce test_cuda_gdb_shows_nvrtc_debug_source_lines with verbose logging."""
    cuda_gdb = shutil.which("cuda-gdb")
    if cuda_gdb is None:
        return {"skipped": True, "reason": "cuda-gdb not on PATH"}

    # Write inline helper (same as upstream cuda_gdb_src.py)
    helper = tempfile.NamedTemporaryFile(
        mode="w",
        suffix="_cuda_gdb_src.py",
        delete=False,
        encoding="utf-8",
    )
    helper.write(
        '''"""Probe helper for issue #2693."""
from cuda.core import Device, LaunchConfig, Program, ProgramOptions, launch

CODE = """
extern \\"C\\" __global__ void kernel() {
    int x = 0;
    x += 1;  // ISSUE_2422_SOURCE_LINE
}
"""

def main() -> None:
    dev = Device()
    dev.set_current()
    stream = dev.create_stream()
    prog = Program(CODE, "c++", ProgramOptions(debug=True, arch=f"sm_{dev.arch}"))
    mod = prog.compile("cubin")
    k = mod.get_kernel("kernel")
    launch(stream, LaunchConfig(grid=1, block=1), k)
    stream.sync()

if __name__ == "__main__":
    main()
'''
    )
    helper.flush()
    helper_path = helper.name
    helper.close()

    argv = [
        cuda_gdb,
        "--batch",
        "-ex",
        "set cuda break_on_launch application",
        "-ex",
        "run",
        "-ex",
        "list",
        "--args",
        sys.executable,
        "-u",
        helper_path,
    ]
    proc = run_cmd(argv, timeout=120)
    output = proc.get("combined", "")

    analysis = {
        "cuda_gdb": cuda_gdb,
        "helper_path": helper_path,
        "returncode": proc.get("returncode"),
        "timeout": proc.get("timeout", False),
        "output_len": len(output),
        "output_head": output[:2000],
        "output_tail": output[-2000:] if len(output) > 2000 else "",
        "has_libthread_banner_only": output.strip() == "[Thread debugging using libthread_db enabled]",
        "has_temp_cu_path": bool(re.search(r"cuda_gdb_src_kernel_\w+\.cu", output)),
        "has_issue_2422_line": "ISSUE_2422_SOURCE_LINE" in output,
        "has_no_such_file": "No such file or directory" in output,
        "has_ptrace_block": "Operation not permitted" in output or "ptrace" in output.lower(),
        "has_cuda_break": "CUDA" in output or "kernel" in output.lower(),
    }

    # Also try with additional gdb verbosity
    verbose_argv = [
        cuda_gdb,
        "--batch",
        "-ex",
        "set verbose on",
        "-ex",
        "set cuda break_on_launch application",
        "-ex",
        "run",
        "-ex",
        "info cuda kernels",
        "-ex",
        "list",
        "--args",
        sys.executable,
        "-u",
        helper_path,
    ]
    verbose = run_cmd(verbose_argv, timeout=120)
    analysis["verbose_output_head"] = verbose.get("combined", "")[:3000]

    try:
        os.unlink(helper_path)
    except OSError:
        pass

    return analysis


def probe_graph_auto_free_on_launch() -> dict:
    """Reproduce test_graph_alloc_with_output instantiation path."""
    from cuda.bindings import driver
    from cuda.core import Buffer, Device, DeviceMemoryResource, LaunchConfig, Program, ProgramOptions, launch
    from cuda.core.graph import GraphCompleteOptions, GraphMemoryResource
    from cuda.core._utils.cuda_utils import CUDAError

    device = Device()
    device.set_current()
    stream = device.create_stream()
    dmr = DeviceMemoryResource(device)

    code = """
    __global__ void add_one(char *a, size_t nbytes) {
        size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
        size_t stride = blockDim.x * gridDim.x;
        for (size_t i = idx; i < nbytes; i += stride) {
            a[i] += 1;
        }
    }
    """
    mod = Program(code, "c++", ProgramOptions(std="c++17", arch=f"sm_{device.arch}")).compile(
        "cubin", name_expressions=("add_one",)
    )
    add_one = mod.get_kernel("add_one")

    results = {"modes": []}
    nbytes = 64
    for mode in ("global", "thread_local", "relaxed"):
        entry: dict = {"mode": mode}
        try:
            gb = device.create_graph_builder().begin_building(mode)
            gmr = GraphMemoryResource(device)
            in_buf = Buffer(nbytes, dmr)
            out = gmr.allocate(nbytes, stream=gb)
            out.copy_from(in_buf, stream=gb)
            launch(gb, LaunchConfig(grid=1, block=1), add_one, out, nbytes)
            options = GraphCompleteOptions(auto_free_on_launch=True)
            try:
                graph = gb.end_building().complete(options)
                entry["instantiate"] = "success"
                graph.upload(stream)
                graph.launch(stream)
                stream.sync()
                entry["launch"] = "success"
            except CUDAError as exc:
                entry["instantiate"] = "failed"
                entry["error"] = repr(exc)
                msg = str(exc)
                if "CUDA_ERROR_INVALID_VALUE" in msg:
                    entry["likely_skip"] = "auto_free_on_launch not supported"
            except RuntimeError as exc:
                entry["instantiate"] = "failed"
                entry["error"] = repr(exc)
                if "Instantiation failed for an unexpected reason" in str(exc):
                    entry["likely_issue"] = "#2527 graph instantiate ERROR on this platform"
        except Exception as exc:  # noqa: BLE001
            entry["setup_error"] = repr(exc)
            entry["traceback"] = traceback.format_exc()
        results["modes"].append(entry)

    # Raw instantiate flags probe
    try:
        results["instantiate_flag_value"] = int(
            driver.CUgraphInstantiate_flags.CUDA_GRAPH_INSTANTIATE_FLAG_AUTO_FREE_ON_LAUNCH
        )
    except Exception as exc:  # noqa: BLE001
        results["instantiate_flag_error"] = repr(exc)

    return results


def main() -> int:
    report: dict = {
        "issue": 2693,
        "timestamp": utc_now(),
        "platform": detect_l4t(),
        "toolkit": probe_cuda_toolkit(),
    }

    try:
        import cuda.core  # noqa: F401

        report["cuda_core_version"] = cuda.core.__version__
    except Exception as exc:  # noqa: BLE001
        report["cuda_core_import_error"] = repr(exc)
        _write(report)
        return 1

    for name, fn in (
        ("nvrtc_debug", probe_nvrtc_debug_materialization),
        ("cuda_gdb", probe_cuda_gdb_session),
        ("graph_auto_free", probe_graph_auto_free_on_launch),
    ):
        try:
            report[name] = fn()
        except Exception as exc:  # noqa: BLE001
            report[name] = {"error": repr(exc), "traceback": traceback.format_exc()}

    _write(report)
    _print_summary(report)
    return 0


def _write(report: dict) -> None:
    path = os.environ.get("ISSUE2693_PROBE_JSON", "/tmp/issue2693-l4t-probe.json")
    with open(path, "w", encoding="utf-8") as fh:
        json.dump(report, fh, indent=2)
    print(f"Wrote probe JSON: {path}")


def _print_summary(report: dict) -> None:
    print("\n=== Issue #2693 probe summary ===")
    plat = report.get("platform", {})
    print(f"  L4T/Jetson: {plat.get('is_l4t')} ({plat.get('machine')})")
    tk = report.get("toolkit", {})
    print(f"  cuda-gdb: {tk.get('cuda-gdb_path') or 'NOT FOUND'}")
    gdb = report.get("cuda_gdb", {})
    if gdb.get("skipped"):
        print(f"  cuda-gdb session: SKIPPED ({gdb.get('reason')})")
    else:
        print(f"  cuda-gdb temp .cu in output: {gdb.get('has_temp_cu_path')}")
        print(f"  cuda-gdb ISSUE_2422 line: {gdb.get('has_issue_2422_line')}")
        print(f"  cuda-gdb banner-only output: {gdb.get('has_libthread_banner_only')}")
    nvrtc = report.get("nvrtc_debug", {})
    cases = nvrtc.get("cases", [])
    if cases:
        c0 = cases[0]
        print(f"  NVRTC temp file materialized: {c0.get('file_exists_at_compile')}")
        print(f"  NVRTC basename pattern ok: {c0.get('basename_matches_pattern')}")
    graph = report.get("graph_auto_free", {})
    for m in graph.get("modes", []):
        print(f"  graph auto_free [{m.get('mode')}]: {m.get('instantiate', m.get('setup_error', '?'))}")


if __name__ == "__main__":
    sys.exit(main())

[lijinf2]

Thank you @rwgk ! I will take a look.