@@ -21,4 +21,5 @@ cdef class Program:
         str _code_type          # Normalised code_type ("c++", "ptx", "nvvm")
         str _pch_status         # PCH creation outcome after compile
         bytes _nvrtc_name       # Source filepath given to NVRTC; a real path for debug builds
+        bytes _debug_source     # Temp file this Program wrote for cuda-gdb, or None
         list _extra_options     # NVRTC options Program adds on top of ProgramOptions
@@ -104,12 +104,17 @@ cdef class Program:
         self._cleanup_debug_source()
 
     def _cleanup_debug_source(self):
-        # Only a temp file this Program wrote may be removed, and the name having
-        # moved off options.name is what says one was written. Without that test
-        # the caller's own file is deleted whenever options.name happens to match
-        # something on disk, since the unredirected name is just that path.
-        if self._nvrtc_name is not None and self._nvrtc_name != self._options._name:
-            self._unlink_debug_source(self._nvrtc_name.decode())
+        # Only a temp file this Program wrote may be removed, so its path is kept
+        # in a dedicated attribute rather than derived from _nvrtc_name and
+        # _options. When a Program dies inside a reference cycle, the cyclic
+        # collector clears the object attributes of everything in the cycle
+        # before __dealloc__ runs: _options may already be None, or the
+        # ProgramOptions may already have lost its fields (#2876). A bytes
+        # attribute is left alone by that clearing.
+        path = self._debug_source
+        if path is not None:
+            self._debug_source = None
+            self._unlink_debug_source(path.decode())
 
     def _unlink_debug_source(self, path: str) -> None:
         try:
@@ -883,7 +888,7 @@ cdef inline int Program_init(Program self, object code, str code_type, object op
         if (options.debug or options.lineinfo) and options.name == "default_program":
             debug_path = self._try_materialize_nvrtc_debug_source(code)
             if debug_path is not None:
-                self._nvrtc_name = debug_path.encode()
+                self._nvrtc_name = self._debug_source = debug_path.encode()
                 # NVRTC resolves #include "..." against the directory of the name it
                 # is given, so moving the name into the temp dir would otherwise stop
                 # every quoted include from resolving where it did before.
 
@@ -34,6 +34,13 @@ New features
 Fixes and enhancements
 ----------------------
 
+- A :class:`Program` that is collected as part of a reference cycle no longer
+  raises ``AttributeError`` from its destructor, and the source file it wrote
+  for a ``debug`` or ``lineinfo`` build is removed. The cyclic collector
+  clears object attributes before the destructor runs, and the destructor
+  read the program options to locate that file.
+  (`#2876 <https://github.com/NVIDIA/cuda-python/issues/2876>`__)
+
 - ``Graph.__getitem__`` now declares an overload for each node type that has
   an executable view, so type checkers and editors see the precise view type:
   indexing with a :class:`~graph.KernelNode` yields an
 
@@ -2,10 +2,13 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
+import gc
+import os
 import re
 import shutil
 import subprocess
 import sys
+import tempfile
 import warnings
 
 import pytest
@@ -406,6 +409,55 @@ def test_program_init_invalid_code_format():
         Program(code, "c++")
 
 
+@pytest.mark.agent_authored(model="claude-fable-5-1")
+@pytest.mark.thread_unsafe(reason="replaces the process-global sys.unraisablehook and tempfile.mkstemp")
+@pytest.mark.parametrize("cleared_first", ["options", "program"])
+def test_program_collected_in_reference_cycle_is_quiet(cleared_first, monkeypatch):
+    """A Program collected as part of a reference cycle is destroyed quietly and removes its debug source (#2876).
+
+    The cyclic collector clears the attributes of every object in the cycle
+    before it runs __dealloc__, in allocation order: either the ProgramOptions
+    loses its fields first, or the Program loses its _options first. Neither
+    may be consulted in __dealloc__. A failure there is only visible through
+    sys.unraisablehook, so the hook is replaced to observe it.
+    """
+    written = []
+    real_mkstemp = tempfile.mkstemp
+
+    def recording_mkstemp(*args, **kwargs):
+        fd, path = real_mkstemp(*args, **kwargs)
+        written.append(path)
+        return fd, path
+
+    monkeypatch.setattr(tempfile, "mkstemp", recording_mkstemp)
+    unraisable = []
+    monkeypatch.setattr(sys, "unraisablehook", unraisable.append)
+
+    code = 'extern "C" __global__ void my_kernel() {}'
+    # arch is passed explicitly so the current device is not queried; debug=True
+    # makes the Program write its source to a temp file for cuda-gdb.
+    options = ProgramOptions(arch="sm_90", debug=True)
+    if cleared_first == "options":
+        # The cycle is allocated between the options and the Program, so the
+        # collector clears the ProgramOptions before it frees the Program.
+        cycle = []
+        cycle.append(cycle)
+        program = Program(code, "c++", options)
+        cycle.append(program)
+    else:
+        # The cycle is allocated after the Program, so the collector clears the
+        # Program's own attributes, including _options, before it frees it.
+        program = Program(code, "c++", options)
+        cycle = [program]
+        cycle.append(cycle)
+    assert written and all(os.path.exists(path) for path in written)
+    del options, program, cycle
+    gc.collect()
+
+    assert unraisable == []
+    assert not any(os.path.exists(path) for path in written)
+
+
 # arch is passed explicitly so the current device is not queried.
 @pytest.mark.agent_authored(model="claude-opus-5")
 @pytest.mark.parametrize("name", [None, "my_program"])