L4T Orin test_program.py::test_cuda_gdb_shows_nvrtc_debug_source_lines failures #2693

Activity

1. rwgk

   assigned

   lijinf2

   on Aug 25, 2026
2. rwgk

   added the

   Bug

   issue type on Aug 25, 2026
3. rwgk

   added

   bugSomething isn't workingSomething isn't working

   on Aug 25, 2026
4. atiaomar1978-hub commented on Aug 25, 2026

   

   atiaomar1978-hub

   on Aug 25, 2026

   Last edited by atiaomar1978-hub

   Contributor

   More actions

   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())
5. lijinf2 commented on Aug 25, 2026

   

   lijinf2

   on Aug 25, 2026

   Contributor

   More actions

   Thank you @rwgk ! I will take a look.
6. lijinf2

   added this to the cuda.core 1.2.0 milestone on Aug 25, 2026
7. lijinf2

   added a sub-issue on Aug 27, 2026

   • test_cuda_gdb_shows_nvrtc_debug_source_lines fails on L4T Orin (CUDA debugger back-end unavailable) #2711
8. lijinf2

   closed this as completedin 133bf0aon Aug 27, 2026
9. lijinf2

   added a sub-issue on Aug 27, 2026

   • L4T Orin test_graph_memory_resource.py::test_graph_alloc_with_output test failures #2527
10. STUNSLAP

    added a commit that references this issue on Aug 27, 2026

    fix(cuda.core): xfail cuda-gdb test where GPU debugging is unavailable (

    ...

    This issue will close once commit 133bf0a is merged into the 'main' branch.

    133bf0a
11. rwgk commented on Sep 1, 2026

    

    rwgk

    on Sep 1, 2026

    ContributorAuthor

    More actions

    To close the loop here, this is the first passing job after this fix (there were some other unrelated issues in the meantime that needed to be resolved separately):

    cuda-python Test Report
    =======================
    Board  : local_l4t
    Date   : Tue Sep  1 11:37:20 AM IST 2026
    Result : PASSED
    
    --- Suite Results ---
    cuda_pathfinder  : 1612 passed, 0 failed, 2 skipped, 0 xfailed
    cuda_bindings    : 454 passed, 0 failed, 36 skipped, 0 xfailed
    cuda_core        : 3714 passed, 0 failed, 472 skipped, 3 xfailed
    
    Total            : 5780 passed, 0 failed, 3 xfailed
    Note             : xfailed = expected failures (known issues, marked xfail in tests)
    
    
    Log file  : 20260901_111652.log