Add tests for ObjectCode.from_fatbin() using nvfatbin bindings (#1875)

* Add tests for ObjectCode.from_fatbin() using nvfatbin bindings

Add availability detection for nvfatbin bindings and tests for loading
fatbin code both from memory (bytes) and from file (str path). The fatbin
fixture creates a multi-arch fatbin containing a cubin for the current
device arch and PTX for a second arch, exercising the nvfatbin API
(create, add_cubin, add_ptx, size, get, destroy).

Partially addresses #663.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

* Fix fatbin fixture: compile PTX targeting second_arch

The PTX was being compiled for the current device's arch but labeled
as a different arch in nvfatbin, which could produce an invalid fatbin.
Now compile PTX with ProgramOptions(arch=f"sm_{second_arch}") so the
PTX actually targets the intended architecture.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

* Make cubin compilation arch explicit in fatbin fixture

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

* [pre-commit.ci] auto code formatting

---------

Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

Author: 3 people

cuda_core/tests/test_module.py

@@ -33,6 +33,20 @@
 """
 
 
+def _is_nvfatbin_available():
+    """Check if nvfatbin bindings are available."""
+    try:
+        from cuda.bindings import nvfatbin
+
+        nvfatbin.version()
+        return True
+    except Exception:
+        return False
+
+
+nvfatbin_available = pytest.mark.skipif(not _is_nvfatbin_available(), reason="nvfatbin bindings not available")
+
+
 @pytest.fixture(scope="module")
 def cuda12_4_prerequisite_check():
     return binding_version() >= (12, 0, 0) and driver_version() >= (12, 4, 0)
@@ -90,6 +104,44 @@ def get_saxpy_kernel_ltoir(init_cuda):
     return mod
 
 
+@pytest.fixture
+def get_saxpy_fatbin(init_cuda):
+    from cuda.bindings import nvfatbin
+
+    dev = Device()
+    arch = dev.arch
+
+    # Pick a second arch different from the current device
+    second_arch = "75" if arch != "75" else "80"
+
+    # Compile to cubin for current device arch
+    prog = Program(SAXPY_KERNEL, code_type="c++", options=ProgramOptions(arch=f"sm_{arch}"))
+    mod = prog.compile(
+        "cubin",
+        name_expressions=("saxpy<float>", "saxpy<double>"),
+    )
+    cubin = mod.code
+    sym_map = mod.symbol_mapping
+
+    # Compile to PTX targeting the second arch
+    ptx_mod = Program(SAXPY_KERNEL, code_type="c++", options=ProgramOptions(arch=f"sm_{second_arch}")).compile(
+        "ptx",
+        name_expressions=("saxpy<float>", "saxpy<double>"),
+    )
+    ptx = ptx_mod.code
+
+    # Create fatbin with both cubin + PTX
+    handle = nvfatbin.create([], 0)
+    nvfatbin.add_cubin(handle, cubin, len(cubin), arch, "saxpy")
+    nvfatbin.add_ptx(handle, ptx, len(ptx), second_arch, "saxpy", f"-arch=sm_{second_arch}")
+    fatbin_size = nvfatbin.size(handle)
+    fatbin = bytearray(fatbin_size)
+    nvfatbin.get(handle, fatbin)
+    nvfatbin.destroy(handle)
+
+    return bytes(fatbin), sym_map
+
+
 def test_get_kernel(init_cuda):
     kernel = """extern "C" __global__ void ABC() { }"""
 
@@ -220,6 +272,28 @@ def test_object_code_load_ltoir_from_file(get_saxpy_kernel_ltoir, tmp_path):
     # ltoir doesn't support kernel retrieval directly as it's used for linking
 
 
+@nvfatbin_available
+def test_object_code_load_fatbin(get_saxpy_fatbin):
+    fatbin, sym_map = get_saxpy_fatbin
+    assert isinstance(fatbin, bytes)
+    mod_obj = ObjectCode.from_fatbin(fatbin, symbol_mapping=sym_map)
+    assert mod_obj.code == fatbin
+    assert mod_obj.code_type == "fatbin"
+    mod_obj.get_kernel("saxpy<double>")  # force loading
+
+
+@nvfatbin_available
+def test_object_code_load_fatbin_from_file(get_saxpy_fatbin, tmp_path):
+    fatbin, sym_map = get_saxpy_fatbin
+    assert isinstance(fatbin, bytes)
+    fatbin_file = tmp_path / "test.fatbin"
+    fatbin_file.write_bytes(fatbin)
+    mod_obj = ObjectCode.from_fatbin(str(fatbin_file), symbol_mapping=sym_map)
+    assert mod_obj.code == str(fatbin_file)
+    assert mod_obj.code_type == "fatbin"
+    mod_obj.get_kernel("saxpy<double>")  # force loading
+
+
 def test_saxpy_arguments(get_saxpy_kernel_cubin, cuda12_4_prerequisite_check):
     krn, _ = get_saxpy_kernel_cubin