fix(core): release driver-backend buffers in Linker.close()

Linker.close() reset only the _culink_handle, leaving the retained
option-key/value vectors and the log-buffer bytearrays alive until
Python GC/tp_dealloc. Those buffers exist for cuLinkDestroy's sake,
but cuLinkDestroy has already run at this point via the shared_ptr
deleter, so they can be released immediately.

Cache decoded logs into _info_log/_error_log before releasing the
raw buffers so get_error_log() / get_info_log() remain callable
after close(), including on the failed-link path where link()
never caches them itself. Swap the option vectors with empty
locals to actually free the backing allocation (std::vector::clear
only sets size to 0 and keeps capacity).

Adds two driver-backend regression tests: one that links
successfully, closes + drops, then performs another full
compile + link cycle (prior heap-corruption bugs only surfaced in
the next CUDA op after teardown); another that triggers a link
failure, closes, and checks get_error_log() still returns the
captured diagnostic.

Author: cpcloud

cuda_core/cuda/core/_linker.pyx

@@ -144,10 +144,31 @@ cdef class Linker:
 
     def close(self):
         """Destroy this linker."""
+        cdef vector[cydriver.CUjit_option] empty_keys
+        cdef vector[void*] empty_values
         if self._use_nvjitlink:
             self._nvjitlink_handle.reset()
         else:
+            # link() caches decoded logs into _info_log/_error_log on the
+            # success path only. A failed link leaves them as None with
+            # _drv_log_bufs as the only source, so cache them here before
+            # the raw buffers are released.
+            if self._drv_log_bufs is not None:
+                if self._info_log is None:
+                    self._info_log = self.get_info_log()
+                if self._error_log is None:
+                    self._error_log = self.get_error_log()
+            # .reset() drops the last shared_ptr to the CUlinkState and runs
+            # cuLinkDestroy synchronously via the custom deleter. The driver
+            # is no longer looking at our option arrays or log buffers, so
+            # release the host-side retainers now rather than waiting for
+            # tp_dealloc. Swap with empty vectors to actually free the
+            # backing allocation; vector.clear() only sets size to 0 and
+            # retains capacity.
             self._culink_handle.reset()
+            self._drv_jit_keys.swap(empty_keys)
+            self._drv_jit_values.swap(empty_values)
+            self._drv_log_bufs = None
 
     @property
     def handle(self) -> LinkerHandleT:

cuda_core/tests/test_linker.py

@@ -252,6 +252,66 @@ def test_linker_handle(compile_ptx_functions):
     assert int(handle) != 0
 
 
+def test_driver_linker_lifetime_no_heap_corruption(monkeypatch, compile_ptx_functions):
+    """Driver-backend teardown must not leave cuLinkCreate option arrays or log buffers dangling.
+
+    Two prior bugs corrupted the heap during driver-linker teardown: the log
+    buffer bytearrays were cleared before cuLinkDestroy ran, and the
+    optionValues array was a stack-local vector destroyed when Linker_init
+    returned. Both manifested in the NEXT CUDA operation after the Linker
+    was destroyed, not at destruction itself. This test forces the driver
+    backend, links, closes + drops the Linker, and then performs a full
+    compile + link cycle that would previously segfault.
+    """
+    monkeypatch.setattr(_linker, "_probe_nvjitlink", lambda: None)
+
+    linker = Linker(*compile_ptx_functions, options=LinkerOptions(arch=ARCH))
+    assert linker.backend == "driver"
+    linker.link("cubin")
+    linker.close()
+    del linker
+
+    obj_a = Program(kernel_a, "c++", ProgramOptions(relocatable_device_code=True)).compile("ptx")
+    obj_b = Program(device_function_b, "c++", ProgramOptions(relocatable_device_code=True)).compile("ptx")
+    obj_c = Program(device_function_c, "c++", ProgramOptions(relocatable_device_code=True)).compile("ptx")
+    linker2 = Linker(obj_a, obj_b, obj_c, options=LinkerOptions(arch=ARCH))
+    assert linker2.backend == "driver"
+    linker2.link("cubin")
+    linker2.close()
+    del linker2
+
+
+def test_driver_linker_get_error_log_after_close_on_failed_link(init_cuda, monkeypatch):
+    """close() must preserve get_error_log() output when link() failed.
+
+    link() only caches _info_log/_error_log on the success path, so after
+    a failed cuLinkComplete the driver log buffers are the only source of
+    the error diagnostic. close() releases those buffers, and callers
+    should still be able to read the captured error log afterward.
+    """
+    monkeypatch.setattr(_linker, "_probe_nvjitlink", lambda: None)
+
+    bad_kernel = """
+extern __device__ int Z();
+__global__ void A() { int r = Z(); }
+"""
+    bad_obj = Program(bad_kernel, "c++", ProgramOptions(relocatable_device_code=True)).compile("ptx")
+    linker = Linker(bad_obj, options=LinkerOptions(arch=ARCH))
+    assert linker.backend == "driver"
+    with pytest.raises(CUDAError):
+        linker.link("cubin")
+
+    pre_close_err = linker.get_error_log()
+    assert isinstance(pre_close_err, str)
+    assert pre_close_err  # failed link must have produced a diagnostic
+
+    linker.close()
+    # close() releases the raw driver buffers; the cached decoded logs must
+    # still be readable.
+    assert linker.get_error_log() == pre_close_err
+    assert isinstance(linker.get_info_log(), str)
+
+
 @pytest.mark.skipif(is_culink_backend, reason="nvjitlink options only tested with nvjitlink backend")
 def test_linker_options_nvjitlink_options_as_str():
     """_prepare_nvjitlink_options(as_bytes=False) returns plain strings."""