From 21f1913ba91ceaeb79d1298947c4ab1d4db35a5c Mon Sep 17 00:00:00 2001 From: Rui Luo Date: Fri, 28 Aug 2026 17:05:56 +0800 Subject: [PATCH] coverage: add cuda.core tests for graph, IPC, launcher, program, and DLPack Signed-off-by: Rui Luo --- cuda_core/tests/graph/test_graph_builder.py | 127 ++++++++++- .../tests/graph/test_graph_definition.py | 30 +++ .../graph/test_graph_definition_lifetime.py | 90 +++++++- .../tests/graph/test_graph_node_update.py | 97 +++++++++ cuda_core/tests/memory_ipc/test_errors.py | 94 +++++++++ cuda_core/tests/test_launcher.py | 139 ++++++++++++ cuda_core/tests/test_memory.py | 60 ++++++ .../tests/test_optional_dependency_imports.py | 23 ++ cuda_core/tests/test_program.py | 165 ++++++++++++++- cuda_core/tests/test_utils.py | 23 ++ cuda_core/tests/test_utils_dlpack.py | 198 ++++++++++++++++++ 11 files changed, 1036 insertions(+), 10 deletions(-) diff --git a/cuda_core/tests/graph/test_graph_builder.py b/cuda_core/tests/graph/test_graph_builder.py index 8889b5c4bf8..b681e19de8a 100644 --- a/cuda_core/tests/graph/test_graph_builder.py +++ b/cuda_core/tests/graph/test_graph_builder.py @@ -17,7 +17,7 @@ import cuda.bindings from cuda.core import Device, LaunchConfig, LegacyPinnedMemoryResource, Program, ProgramOptions, StreamOptions, launch -from cuda.core.graph import GraphBuilder, GraphCompleteOptions, GraphDefinition +from cuda.core.graph import Graph, GraphBuilder, GraphCompleteOptions, GraphDefinition from cuda.core.graph._graph_builder import ( _capture_callback_with_tail_failure_for_testing, ) @@ -32,6 +32,13 @@ def _wait_until(predicate, timeout=5.0): time.sleep(0.02) +def _skip_if_conditional_handles_unsupported(): + from cuda.core._utils.version import binding_version, driver_version + + if driver_version() < (12, 3, 0) or binding_version() < (12, 3, 0): + pytest.skip("conditional handles require CUDA driver and bindings 12.3+") + + def test_graph_is_building(init_cuda): gb = Device().create_graph_builder() assert gb.is_building is False @@ -895,3 +902,121 @@ def test_pdl_same_stream_primary_secondary_overlap_via_graph(init_cuda): "PDL (Programmatic Dependent Launch) graph overlap was not observed. " "If this keeps xfailing in CI, manually re-check on a quiet Hopper+ GPU." ) + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_join_rejects_non_builder(init_cuda): + """join() type-checks its arguments before looking at capture state.""" + gb = Device().create_graph_builder() + with pytest.raises(TypeError, match="All arguments must be GraphBuilder"): + GraphBuilder.join(gb, object()) + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_builder_cuda_stream_protocol(init_cuda): + """The builder exports its underlying stream, and stops doing so once closed.""" + gb = Device().create_graph_builder() + protocol = gb.__cuda_stream__() + assert protocol[0] == 0 + assert int(protocol[1]) == int(gb.stream.handle) + gb.close() + with pytest.raises(RuntimeError, match="GraphBuilder has been closed"): + gb.__cuda_stream__() + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_end_building_requires_active_capture(init_cuda): + """end_building() on a builder that never started capturing is rejected.""" + gb = Device().create_graph_builder() + with pytest.raises(RuntimeError, match="Graph builder is not building"): + gb.end_building() + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_debug_dot_print_requires_finished_build(init_cuda, tmp_path): + """debug_dot_print() needs a completed capture, both before and during building.""" + gb = Device().create_graph_builder() + with pytest.raises(RuntimeError, match="Graph has not finished building"): + gb.debug_dot_print(str(tmp_path / "unfinished.dot")) + gb.begin_building() + try: + with pytest.raises(RuntimeError, match="Graph has not finished building"): + gb.debug_dot_print(str(tmp_path / "capturing.dot")) + finally: + gb.end_building() + gb.debug_dot_print(str(tmp_path / "finished.dot")) + assert (tmp_path / "finished.dot").exists() + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_callback_requires_active_capture(init_cuda): + """callback() is rejected outside an active capture.""" + gb = Device().create_graph_builder() + with pytest.raises(RuntimeError, match="Cannot add callback when graph is not being built"): + gb.callback(lambda: None) + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_create_condition_requires_active_capture(init_cuda): + """create_condition() is rejected outside an active capture.""" + _skip_if_conditional_handles_unsupported() + gb = Device().create_graph_builder() + with pytest.raises(RuntimeError, match="Cannot create a condition when graph is not being built"): + gb.create_condition() + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_embed_requires_finished_child_and_capturing_parent(init_cuda): + """embed() rejects an unfinished child and a parent that is not capturing.""" + parent = Device().create_graph_builder() + # embed() checks the child before the parent, so the child must already be + # ended for the parent guard to be the one that fires here. + child = Device().create_graph_builder().begin_building().end_building() + with pytest.raises(ValueError, match="Parent graph is not being built"): + parent.embed(child) + + unfinished = Device().create_graph_builder().begin_building() + capturing = Device().create_graph_builder().begin_building() + try: + with pytest.raises(ValueError, match="Child graph has not finished building"): + capturing.embed(unfinished) + finally: + capturing.end_building() + unfinished.end_building() + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_graph_builder_and_graph_cannot_be_constructed_directly(): + """Both types are factory-only; the guards run before any CUDA call.""" + with pytest.raises(NotImplementedError, match="directly creating"): + GraphBuilder() + with pytest.raises(RuntimeError, match="directly constructing"): + Graph() + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_graph_builder_close_ends_active_capture(init_cuda): + """close() during capture ends it and hands the stream back usable.""" + empty = compile_common_kernels().get_kernel("empty_kernel") + stream = Device().create_stream() + gb = stream.create_graph_builder().begin_building() + launch(gb, LaunchConfig(grid=1, block=1), empty) + assert gb.is_building + gb.close() + with pytest.raises(RuntimeError, match="has been closed"): + _ = gb.is_building + # Ending capture via close() must leave the stream usable. + launch(stream, LaunchConfig(grid=1, block=1), empty) + stream.sync() + stream.close() + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_if_then_requires_active_capture(init_cuda): + """Conditional nodes cannot be added once capture has ended.""" + _skip_if_conditional_handles_unsupported() + gb = Device().create_graph_builder().begin_building() + condition = try_create_condition(gb) + gb.end_building() + with pytest.raises(RuntimeError, match="Cannot add conditional node when not actively capturing"): + gb.if_then(condition) diff --git a/cuda_core/tests/graph/test_graph_definition.py b/cuda_core/tests/graph/test_graph_definition.py index f8f427567a0..4444bafc5de 100644 --- a/cuda_core/tests/graph/test_graph_definition.py +++ b/cuda_core/tests/graph/test_graph_definition.py @@ -4,7 +4,9 @@ """Tests for GraphDefinition topology, node types, instantiation, and execution.""" import ctypes +import gc import sys +import weakref from collections.abc import Callable from dataclasses import dataclass, field @@ -910,6 +912,34 @@ def test_alloc_peer_access(mempool_device_x2): assert d1.device_id in node.peer_access +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_alloc_memory_type_host(init_cuda): + """HOST graph alloc nodes reconstruct memory_type from the driver, not the Python argument.""" + _skip_if_no_mempool() + from cuda.core._utils.cuda_utils import CUDAError + + g = GraphDefinition() + try: + with xfail_on_graph_mempool_oom(): + node = g.allocate(ALLOC_SIZE, memory_type=GraphMemoryType.HOST) + except CUDAError as e: + if "CUDA_ERROR_NOT_SUPPORTED" in str(e): + pytest.skip("Driver does not support graph alloc memory_type='host'") + raise + + expected_dptr = node.dptr + succ = node.record(Device().create_event()) + node_ref = weakref.ref(node) + del node + gc.collect() + assert node_ref() is None + reconstructed = next(iter(succ.pred)) + assert isinstance(reconstructed, AllocNode) + assert reconstructed.memory_type == GraphMemoryType.HOST + assert reconstructed.dptr == expected_dptr + assert reconstructed.dptr != 0 + + # ============================================================================= # Join API # ============================================================================= diff --git a/cuda_core/tests/graph/test_graph_definition_lifetime.py b/cuda_core/tests/graph/test_graph_definition_lifetime.py index 16d9f8f9869..dfb20e581b7 100644 --- a/cuda_core/tests/graph/test_graph_definition_lifetime.py +++ b/cuda_core/tests/graph/test_graph_definition_lifetime.py @@ -76,15 +76,19 @@ def _wait_until(predicate, timeout=None, interval=0.02): raise AssertionError(f"condition not satisfied within {timeout}s") -from cuda.core import Device, DeviceMemoryResource, EventOptions, Kernel, LaunchConfig +from cuda.core import Device, DeviceMemoryResource, EventOptions, Kernel, LaunchConfig, LegacyPinnedMemoryResource from cuda.core._utils.cuda_utils import CUDAError from cuda.core._utils.version import driver_version from cuda.core.graph import ( ChildGraphNode, ConditionalNode, + EventRecordNode, + EventWaitNode, + FreeNode, GraphDefinition, HostCallbackNode, KernelNode, + MemcpyNode, ) @@ -1213,6 +1217,90 @@ def test_kernel_node_reconstruction_preserves_validity(init_cuda): stream.sync() +def _pred_chain_memcpy(g, bufs): + memory_resource = LegacyPinnedMemoryResource() + src = memory_resource.allocate(8) + dst = memory_resource.allocate(8) + bufs.extend((src, dst)) + node = g.memcpy(dst, src, 8) + src_ptr, dst_ptr, size = node.src, node.dst, node.size + succ = node.record(Device().create_event()) + + def check(reconstructed): + assert isinstance(reconstructed, MemcpyNode) + assert reconstructed.src == src_ptr + assert reconstructed.dst == dst_ptr + assert reconstructed.size == size + + return node, succ, check + + +def _pred_chain_event_record(g, bufs): + event = Device().create_event() + node = g.record(event) + succ = node.wait(Device().create_event()) + + def check(reconstructed): + assert isinstance(reconstructed, EventRecordNode) + assert reconstructed.event.handle == event.handle + + return node, succ, check + + +def _pred_chain_event_wait(g, bufs): + wait_event = Device().create_event() + node = g.wait(wait_event) + succ = node.record(Device().create_event()) + + def check(reconstructed): + assert isinstance(reconstructed, EventWaitNode) + assert reconstructed.event.handle == wait_event.handle + + return node, succ, check + + +def _pred_chain_free(g, bufs): + _skip_if_no_mempool() + with xfail_on_graph_mempool_oom(): + alloc = g.allocate(64) + node = alloc.deallocate(alloc.dptr) + free_dptr = node.dptr + succ = node.record(Device().create_event()) + + def check(reconstructed): + assert isinstance(reconstructed, FreeNode) + assert reconstructed.dptr == free_dptr + + return node, succ, check + + +@pytest.mark.parametrize( + "factory", + [ + _pred_chain_memcpy, + _pred_chain_event_record, + _pred_chain_event_wait, + _pred_chain_free, + ], + ids=["memcpy", "event_record", "event_wait", "free"], +) +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_graph_nodes_reconstructed_via_pred_chain(init_cuda, factory): + """Dropping the original Python node forces `_create_from_driver` on pred walk.""" + g = GraphDefinition() + bufs = [] + try: + node, succ, check = factory(g, bufs) + node_ref = weakref.ref(node) + del node + _wait_until(lambda: node_ref() is None) + reconstructed = next(iter(succ.pred)) + check(reconstructed) + finally: + for buf in bufs: + buf.close() + + # ============================================================================= # Kernel argument lifetime — kernel nodes should keep argument objects alive # ============================================================================= diff --git a/cuda_core/tests/graph/test_graph_node_update.py b/cuda_core/tests/graph/test_graph_node_update.py index f49e547e9cb..4ac20b40f08 100644 --- a/cuda_core/tests/graph/test_graph_node_update.py +++ b/cuda_core/tests/graph/test_graph_node_update.py @@ -22,6 +22,7 @@ ChildGraphNode, EventRecordNode, EventWaitNode, + ExecutableGraphNode, GraphDefinition, HostCallbackNode, KernelNode, @@ -1217,3 +1218,99 @@ def blocking_callback(): stream.sync() _wait_until(lambda: not inflight_weak) + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_memory_node_update_validates_owners_and_noops(init_cuda): + """Memory-node updates validate owners, preserve no-ops, and update geometry.""" + if driver_version() < (12, 2, 0): + pytest.skip("individual graph node updates require CUDA 12.2+") + + memory_resource = LegacyPinnedMemoryResource() + with memory_resource.allocate(16) as src, memory_resource.allocate(16) as dst: + graph_def = GraphDefinition() + memset_node = graph_def.memset(dst, 0x11, 8) + memcpy_node = graph_def.memcpy(dst, src, 8) + + with pytest.raises(ValueError, match=r"^dst_owner requires dst$"): + memset_node.update(dst_owner=dst) + memset_node.update() + assert memset_node.value == 0x11 + assert memset_node.width == 8 + + memset_node.update(width=4, height=2, pitch=8) + assert memset_node.width == 4 + assert memset_node.height == 2 + assert memset_node.pitch == 8 + + with pytest.raises(ValueError, match=r"^dst_owner requires dst$"): + memcpy_node.update(dst_owner=dst) + with pytest.raises(ValueError, match=r"^src_owner requires src$"): + memcpy_node.update(src_owner=src) + memcpy_node.update() + assert memcpy_node.size == 8 + assert memcpy_node.dst == int(dst.handle) + assert memcpy_node.src == int(src.handle) + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_executable_graph_node_cannot_be_constructed_directly(): + """Executable-node views are factory-only and fail before any CUDA call.""" + with pytest.raises(RuntimeError, match=r"^directly constructing an executable graph node is not supported$"): + ExecutableGraphNode() + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_executable_node_repr_reports_graph_and_node(init_cuda): + """An executable-node view reprs its subclass name and both handles.""" + if driver_version() < (12, 2, 0): + pytest.skip("individual graph node updates require CUDA 12.2+") + + kernel = compile_common_kernels().get_kernel("empty_kernel") + graph_def = GraphDefinition() + node = graph_def.launch(LaunchConfig(grid=1, block=1), kernel) + graph = graph_def.instantiate() + + assert repr(graph[node]) == f"" + + +@pytest.mark.parametrize("config_kind", ["clustered", "cooperative"]) +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_executable_kernel_update_rejects_unsupported_config(init_cuda, config_kind): + """Executable kernel updates reject clustered and cooperative launches.""" + if driver_version() < (12, 2, 0): + pytest.skip("individual graph node updates require CUDA 12.2+") + + kernel = compile_common_kernels().get_kernel("empty_kernel") + graph_def = GraphDefinition() + node = graph_def.launch(LaunchConfig(grid=1, block=1), kernel) + view = graph_def.instantiate()[node] + + config = LaunchConfig(grid=1, block=1) + if config_kind == "clustered": + config.cluster = (1, 1, 1) + else: + config.is_cooperative = True + with pytest.raises( + NotImplementedError, + match=r"^updating clustered or cooperative kernel nodes is not supported$", + ): + view.update(config=config, kernel=kernel, args=()) + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_ctypes_host_callback_repr(init_cuda): + """A ctypes host callback repr reports its node and function addresses.""" + callback_type = ctypes.CFUNCTYPE(None, ctypes.c_void_p) + + @callback_type + def host_fn(_user_data): + return None + + graph_def = GraphDefinition() + node = graph_def.callback(host_fn) + assert isinstance(node, HostCallbackNode) + assert node.callback is None + cfunc = ctypes.cast(host_fn, ctypes.c_void_p).value + assert cfunc is not None + assert repr(node) == f"" diff --git a/cuda_core/tests/memory_ipc/test_errors.py b/cuda_core/tests/memory_ipc/test_errors.py index 8038d62570c..a12cf714f21 100644 --- a/cuda_core/tests/memory_ipc/test_errors.py +++ b/cuda_core/tests/memory_ipc/test_errors.py @@ -4,6 +4,7 @@ import multiprocessing import os import pickle +import platform import re import uuid @@ -44,6 +45,15 @@ def test_outer_timeout_marker_is_applied(request): assert marker.args == (expected,), f"unexpected timeout value: {marker.args!r}" +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_ipc_types_cannot_be_constructed_directly(): + """Factory-only IPC types reject direct construction.""" + with pytest.raises(RuntimeError, match=r"^IPCBufferDescriptor objects cannot be instantiated directly\."): + IPCBufferDescriptor() + with pytest.raises(RuntimeError, match=r"^IPCAllocationHandle objects cannot be instantiated directly\."): + IPCAllocationHandle() + + def test_import_truncated_buffer_descriptor(ipc_device, ipc_memory_resource): """Truncated IPC buffer descriptor payload is rejected before driver import.""" desc = IPCBufferDescriptor._init(b"\x00" * 8, NBYTES) @@ -254,3 +264,87 @@ def CHILD_ACTION(self, queue): def ASSERT(self, exc_type, exc_msg): assert exc_type is RuntimeError assert re.match(r"Memory resource [a-z0-9-]+ was not found", exc_msg) + + +@pytest.mark.skipif(platform.system() != "Linux", reason="CUDA mempool IPC is Linux-only") +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_from_allocation_handle_raw_fd_imports_mapped_pool(ipc_device): + """from_allocation_handle accepts a raw int fd and constructs an unregistered mapped MR.""" + from helpers.buffers import PatternGen + + device = ipc_device + stream = device.default_stream + exporter = DeviceMemoryResource(device, DeviceMemoryResourceOptions(max_size=POOL_SIZE, ipc_enabled=True)) + try: + dup_fd = os.dup(exporter.allocation_handle.handle) + try: + imported = DeviceMemoryResource.from_allocation_handle(device, dup_fd) + finally: + # The int overload dups the fd internally, so the imported pool must + # outlive the caller's copy: everything below runs without it. + os.close(dup_fd) + try: + assert imported.is_mapped + assert imported.is_ipc_enabled + assert imported.device_id == device.device_id + # No uuid was supplied, so the pool never entered the registry. + assert imported.uuid is None + # Mapped pools cannot allocate; import a peer buffer instead. + with exporter.allocate(NBYTES, stream=stream) as exported: + descriptor = exported.ipc_descriptor + with Buffer.from_ipc_descriptor(imported, descriptor, stream=stream) as mapped_buf: + pgen = PatternGen(device, NBYTES, stream=stream) + pgen.fill_buffer(mapped_buf, seed=1) + pgen.verify_buffer(exported, seed=1) + finally: + imported.close() + finally: + exporter.close() + + +@pytest.mark.skipif(platform.system() != "Linux", reason="CUDA mempool IPC is Linux-only") +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_allocation_handle_forking_pickler_roundtrip(ipc_device): + """ForkingPickler transfers an IPCAllocationHandle by duplicating its fd.""" + from multiprocessing.reduction import ForkingPickler + + device = ipc_device + mr = DeviceMemoryResource(device, DeviceMemoryResourceOptions(max_size=POOL_SIZE, ipc_enabled=True)) + try: + handle = mr.allocation_handle + restored = ForkingPickler.loads(ForkingPickler.dumps(handle)) + try: + assert isinstance(restored, IPCAllocationHandle) + # DupFd must hand back a real duplicate: a shared fd number would mean + # restored.close() also clobbers the exporter's handle. The number + # itself is unpredictable, so only check validity and distinctness. + assert restored.handle > 0 + assert restored.handle != handle.handle + assert restored.uuid == handle.uuid + finally: + restored.close() + finally: + mr.close() + + +@pytest.mark.skipif(platform.system() != "Linux", reason="CUDA mempool IPC is Linux-only") +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_ipc_registry_dedups_repeated_imports(ipc_device): + """from_allocation_handle registers the mapped pool; later imports hit the cache.""" + device = ipc_device + exporter = DeviceMemoryResource(device, DeviceMemoryResourceOptions(max_size=POOL_SIZE, ipc_enabled=True)) + mapped = None + try: + key = exporter.uuid + mapped = DeviceMemoryResource.from_allocation_handle(device, exporter.allocation_handle) + assert mapped.is_mapped + assert DeviceMemoryResource.from_registry(key) is mapped + mapped2 = DeviceMemoryResource.from_allocation_handle(device, exporter.allocation_handle) + assert mapped2 is mapped + # Registering under a key that is already taken hands back the existing + # entry, so the exporter itself never enters the registry. + assert exporter.register(key) is mapped + finally: + if mapped is not None: + mapped.close() + exporter.close() diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index 2ab766cc2f0..f08f5d2ab6f 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -756,3 +756,142 @@ class MyComplex(complex): holder = ParamHolder([MyBool(1), MyFloat(1.5), MyComplex(1 + 2j)]) assert holder.ptr != 0 + + +_NUMPY_SUBCLASS_FALLBACK_PARAMS = [ + # One case per prepare_numpy_arg isinstance-fallback branch (exact type is + # skipped because type(arg) is the subclass). Values catch width/sign mixups. + (np.bool_, np.bool_, "bool", True), + (np.int8, np.int8, "signed char", -42), + (np.int16, np.int16, "signed short", -1234), + (np.int32, np.int32, "signed int", -123456), + (np.int64, np.int64, "signed long long", -123456789), + (np.uint8, np.uint8, "unsigned char", 200), + (np.uint16, np.uint16, "unsigned short", 60000), + (np.uint32, np.uint32, "unsigned int", 4000000000), + (np.uint64, np.uint64, "unsigned long long", 0x1_0000_0001), + (np.float64, np.float64, "double", 2.718281828), +] +_NUMPY_SUBCLASS_FALLBACK_IDS = [ + "numpy_bool", + "numpy_int8", + "numpy_int16", + "numpy_int32", + "numpy_int64", + "numpy_uint8", + "numpy_uint16", + "numpy_uint32", + "numpy_uint64", + "numpy_float64", +] +if helpers.CCCL_INCLUDE_PATHS is not None: + _NUMPY_SUBCLASS_FALLBACK_PARAMS += [ + (np.float16, np.float16, "half", 0.78), + (np.complex64, np.complex64, "cuda::std::complex", 1 + 2j), + (np.complex128, np.complex128, "cuda::std::complex", -3 - 4j), + ] + _NUMPY_SUBCLASS_FALLBACK_IDS += ["numpy_float16", "numpy_complex64", "numpy_complex128"] + + +@requires_module(np, "2.2.5", reason="need numpy 2.2.5+ (numpy GH #28632)") +@pytest.mark.parametrize( + ("base_type", "np_dtype", "cpp_type", "raw_value"), + _NUMPY_SUBCLASS_FALLBACK_PARAMS, + ids=_NUMPY_SUBCLASS_FALLBACK_IDS, +) +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_launch_numpy_scalar_subclass_fallback(base_type, np_dtype, cpp_type, raw_value): + """Subclassed numpy scalars take prepare_numpy_arg's isinstance fallback and reach the kernel (readback).""" + + class Subclassed(base_type): + pass + + scalar = Subclassed(raw_value) + expected = np_dtype(raw_value) + + dev = Device() + dev.set_current() + + mr = LegacyPinnedMemoryResource() + b = mr.allocate(np.dtype(np_dtype).itemsize) + arr = np.from_dlpack(b).view(np_dtype) + arr[:] = 0 + + code = r""" + template + __global__ void write_scalar(T* arr, T val) { + arr[0] = val; + } + """ + if helpers.CCCL_INCLUDE_PATHS is not None: + code = ( + r""" + #include + #include + """ + + code + ) + + arch = "".join(f"{i}" for i in dev.compute_capability) + pro_opts = ProgramOptions(std="c++17", arch=f"sm_{arch}", include_path=helpers.CCCL_INCLUDE_PATHS) + prog = Program(code, code_type="c++", options=pro_opts) + ker_name = f"write_scalar<{cpp_type}>" + mod = prog.compile("cubin", name_expressions=(ker_name,)) + ker = mod.get_kernel(ker_name) + + stream = dev.default_stream + config = LaunchConfig(grid=1, block=1) + launch(stream, config, ker, arr.ctypes.data, scalar) + stream.sync() + + assert arr[0] == expected + + +# Truncates to 1 if the launcher packs the handle as uint32 instead of uint64. +_UINT64_HANDLE_VALUE = 0x1_0000_0001 + + +def _compile_write_ull_kernel(dev): + code = r""" + extern "C" __global__ void write_ull(unsigned long long *out, unsigned long long val) { + *out = val; + } + """ + arch = "".join(f"{i}" for i in dev.compute_capability) + prog = Program(code, code_type="c++", options=ProgramOptions(std="c++17", arch=f"sm_{arch}")) + return prog.compile("cubin", name_expressions=("write_ull",)).get_kernel("write_ull") + + +def _assert_kernel_sees_ull(dev, kernel_arg, expected): + mr = LegacyPinnedMemoryResource() + buf = mr.allocate(np.dtype(np.uint64).itemsize) + try: + arr = np.from_dlpack(buf).view(np.uint64) + arr[:] = 0 + ker = _compile_write_ull_kernel(dev) + stream = dev.default_stream + launch(stream, LaunchConfig(grid=1, block=1), ker, arr.ctypes.data, kernel_arg) + stream.sync() + assert int(arr[0]) == int(expected) + finally: + buf.close() + + +@pytest.mark.parametrize("use_subclass", [False, True], ids=["exact_type", "subclass_fallback"]) +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_launch_graph_conditional_handle_as_kernel_arg(init_cuda, use_subclass): + """CUgraphConditionalHandle is packed as its uint64 value (readback).""" + from cuda.bindings import driver + + if not hasattr(driver, "CUgraphConditionalHandle"): + pytest.skip("CUgraphConditionalHandle requires cuda-bindings 12.3+") + + class SubclassedHandle(driver.CUgraphConditionalHandle): + pass + + handle_cls = SubclassedHandle if use_subclass else driver.CUgraphConditionalHandle + handle = handle_cls(_UINT64_HANDLE_VALUE) + + dev = Device() + dev.set_current() + _assert_kernel_sees_ull(dev, handle, _UINT64_HANDLE_VALUE) diff --git a/cuda_core/tests/test_memory.py b/cuda_core/tests/test_memory.py index 44227d4b1c5..0a0ab8d1e57 100644 --- a/cuda_core/tests/test_memory.py +++ b/cuda_core/tests/test_memory.py @@ -179,6 +179,51 @@ def test_buffer_initialization(): buffer_initialization(MemoryResource()) +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_buffer_direct_init_forbidden(): + """Buffers must come from a MemoryResource, never from ``Buffer()``.""" + with pytest.raises(RuntimeError, match=r"^Buffer objects cannot be instantiated directly\."): + Buffer() + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_buffer_context_manager_closes_on_exit(): + """``with buf`` yields the buffer, closes it on exit, and does not swallow inner exceptions.""" + device = Device() + device.set_current() + mr = DummyDeviceMemoryResource(device) + buf = mr.allocate(size=64, stream=device.default_stream) + with buf as entered: + assert entered is buf + assert buf.handle != 0 + assert buf.handle == 0 + assert buf.memory_resource is None + + buf = mr.allocate(size=64, stream=device.default_stream) + with pytest.raises(RuntimeError, match="^boom$"), buf: + raise RuntimeError("boom") + assert buf.handle == 0 + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_memory_resource_abstract_stubs(): + """Every abstract MemoryResource member reports itself as unimplemented.""" + device = Device() + device.set_current() + mr = MemoryResource() + stream = device.default_stream + with pytest.raises(TypeError, match=r"^MemoryResource\.allocate must be implemented"): + mr.allocate(1, stream=stream) + with pytest.raises(TypeError, match=r"^MemoryResource\.deallocate must be implemented"): + mr.deallocate(0, 1, stream=stream) + with pytest.raises(TypeError, match=r"^MemoryResource\.is_device_accessible must be implemented"): + _ = mr.is_device_accessible + with pytest.raises(TypeError, match=r"^MemoryResource\.is_host_accessible must be implemented"): + _ = mr.is_host_accessible + with pytest.raises(TypeError, match=r"^MemoryResource\.device_id must be implemented"): + _ = mr.device_id + + def buffer_copy_to(dummy_mr: MemoryResource, device: Device, check=False): src_buffer = dummy_mr.allocate(size=1024) dst_buffer = dummy_mr.allocate(size=1024) @@ -271,6 +316,20 @@ def test_buffer_copy_from_size_mismatch_raises(): src_buffer.close() +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_copy_to_auto_dst_requires_memory_resource(): + """``copy_to()`` cannot mint a destination without a memory resource.""" + device = Device() + device.set_current() + owner = (ctypes.c_byte * 32)() + buf = Buffer.from_handle(ctypes.addressof(owner), 32, owner=owner) + try: + with pytest.raises(ValueError, match="does not have a memory_resource"): + buf.copy_to(stream=device.default_stream) + finally: + buf.close() + + def _bytes_repeat(pattern: bytes, size: int) -> bytes: assert len(pattern) > 0 assert size % len(pattern) == 0 @@ -392,6 +451,7 @@ def test_buffer_external_host(): a = (ctypes.c_byte * 20)() ptr = ctypes.addressof(a) buffer = Buffer.from_handle(ptr, 20, owner=a) + assert buffer.owner is a assert not buffer.is_device_accessible assert buffer.is_host_accessible assert buffer.device_id == -1 diff --git a/cuda_core/tests/test_optional_dependency_imports.py b/cuda_core/tests/test_optional_dependency_imports.py index 9ba7358f9fe..b08b7d344d9 100644 --- a/cuda_core/tests/test_optional_dependency_imports.py +++ b/cuda_core/tests/test_optional_dependency_imports.py @@ -7,6 +7,10 @@ from cuda.core import _linker, _program from cuda.pathfinder import DynamicLibNotFoundError +# The autouse fixture below resets module-level import state for every test in this +# file, so the whole module is thread-unsafe -- not just the tests that monkeypatch. +pytestmark = pytest.mark.thread_unsafe(reason="resets cuda.core._program / _linker optional-import globals") + @pytest.fixture(autouse=True) def restore_optional_import_state(): @@ -31,6 +35,25 @@ def restore_optional_import_state(): _linker._use_nvjitlink_backend = saved_use_nvjitlink +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_get_nvvm_module_rejects_old_bindings(monkeypatch): + """NVVM import requires cuda-bindings >= 12.9.0 and caches a failed attempt.""" + calls = 0 + + def old_binding_version(): + nonlocal calls + calls += 1 + return (12, 8, 0) + + monkeypatch.setattr(_program, "binding_version", old_binding_version) + + with pytest.raises(RuntimeError, match="cuda-bindings >= 12.9.0"): + _program._get_nvvm_module() + with pytest.raises(RuntimeError, match="previous import attempt failed"): + _program._get_nvvm_module() + assert calls == 1 + + def test_get_nvvm_module_reraises_nested_module_not_found(monkeypatch): monkeypatch.setattr(_program, "binding_version", lambda: (12, 9, 0)) diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index 3b280cc48cf..e1bb96a1e11 100644 --- a/cuda_core/tests/test_program.py +++ b/cuda_core/tests/test_program.py @@ -772,14 +772,6 @@ def test_program_options_as_bytes_invalid_backend(): options.as_bytes("invalid") -@nvvm_available -def test_program_options_as_bytes_nvvm_unsupported_option(): - """Test that unsupported options raise CUDAError for NVVM backend""" - options = ProgramOptions(arch="sm_80", lineinfo=True) - with pytest.raises(CUDAError, match="not supported by NVVM backend"): - options.as_bytes("nvvm") - - @nvvm_available def test_nvvm_program_options_as_bytes_numba_debug(): """numba_debug must be plumbed through to libNVVM as -numba-debug @@ -1141,3 +1133,160 @@ def test_nvrtc_compile_with_logs_capture(init_cuda): assert isinstance(result, ObjectCode) assert logs.getvalue(), "Expected non-empty compilation log from #warning directive" program.close() + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_program_options_bad_define_macro_nested_list_invalid_element(): + """Nested define_macro list with a non-processable element raises at the element.""" + # [("MACRO", "1")] makes is_nested_sequence True; 42 fails the inner processor. + opts = ProgramOptions(name="test", arch="sm_80", define_macro=[("MACRO", "1"), 42]) + with pytest.raises(RuntimeError, match=r"Expected define_macro.*got 42"): + opts.as_bytes("nvrtc") + + +@pytest.mark.parametrize( + "kwargs", + [ + {"relocatable_device_code": True}, + {"extensible_whole_program": True}, + {"lineinfo": True}, + {"ptxas_options": "-v"}, + {"max_register_count": 32}, + {"use_fast_math": True}, + {"extra_device_vectorization": True}, + {"gen_opt_lto": True}, + {"define_macro": "M"}, + {"undefine_macro": "M"}, + {"include_path": "include-dir"}, + {"pre_include": "header.h"}, + {"no_source_include": True}, + {"std": "c++17"}, + {"builtin_move_forward": False}, + {"builtin_initializer_list": False}, + {"disable_warnings": True}, + {"restrict": True}, + {"device_as_default_execution_space": True}, + {"device_int128": True}, + {"optimization_info": "inline"}, + {"no_display_error_number": True}, + {"diag_error": 1}, + {"diag_suppress": 1}, + {"diag_warn": 1}, + {"brief_diagnostics": True}, + {"time": "timing.csv"}, + {"split_compile": 2}, + {"fdevice_syntax_only": True}, + {"minimal": True}, + ], + ids=lambda kw: next(iter(kw)), +) +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_nvvm_options_reject_each_unsupported_flag(kwargs): + """Every NVVM-unsupported option is rejected, named, and reported alone.""" + # This table mirrors _prepare_nvvm_options_impl's rejection list one-for-one. + options = ProgramOptions(arch="sm_80", **kwargs) + name = next(iter(kwargs)) + with pytest.raises(CUDAError, match=rf"^The following options are not supported by NVVM backend: {name}$"): + options.as_bytes("nvvm") + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_nvrtc_as_bytes_emits_sequence_and_uncommon_flags(): + """as_bytes emits the NVRTC spellings that compile-option tests do not hit.""" + options = ProgramOptions( + arch="sm_80", + ptxas_options="-v", + pre_include=["a.h", "b.h"], + device_float128=True, + diag_warn=[1000, 1001], + time="timing.csv", + split_compile=2, + pch_dir="pch-cache", + ) + flags = [opt.decode() for opt in options.as_bytes("nvrtc")] + assert "--ptxas-options=-v" in flags + assert "--pre-include=a.h" in flags + assert "--pre-include=b.h" in flags + assert "--device-float128" in flags + assert "--diag-warn=1000" in flags + assert "--diag-warn=1001" in flags + assert "--time=timing.csv" in flags + assert "--split-compile=2" in flags + assert "--pch-dir=pch-cache" in flags + + single_pre = ProgramOptions(arch="sm_80", pre_include="only.h") + assert "--pre-include=only.h" in [opt.decode() for opt in single_pre.as_bytes("nvrtc")] + + +@pytest.mark.thread_unsafe(reason="patches the process-global os.fdopen and tempfile.mkstemp") +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_nvrtc_debug_falls_back_when_temp_file_write_fails(init_cuda, monkeypatch): + """A write failure removes the temporary source and falls back to the default name.""" + import os + + from cuda.core import _program + + real_fdopen = os.fdopen + real_mkstemp = _program.tempfile.mkstemp + temp_paths = [] + + class _FailingWriter: + def write(self, _code): + raise OSError("No space left on device") + + @contextlib.contextmanager + def _write_fails(fd, *args, **kwargs): + with real_fdopen(fd, *args, **kwargs): + yield _FailingWriter() + + def _record_mkstemp(*args, **kwargs): + fd, path = real_mkstemp(*args, **kwargs) + temp_paths.append(path) + return fd, path + + monkeypatch.setattr(_program.os, "fdopen", _write_fails) + monkeypatch.setattr(_program.tempfile, "mkstemp", _record_mkstemp) + + code = 'extern "C" __global__ void matmul() {}' + prog = Program(code, "c++", ProgramOptions(debug=True, arch="sm_80")) + try: + assert len(temp_paths) == 1 + assert not os.path.exists(temp_paths[0]) + assert prog.compile("ptx").name == "default_program" + finally: + prog.close() + + +@nvvm_available +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_nvvm_compile_with_libdevice(nvvm_ir): + """use_libdevice resolves a referenced libdevice function into the generated PTX.""" + store = " store i32 %call, i32* %data, align 4" + declaration = "declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()" + assert store in nvvm_ir and declaration in nvvm_ir + libdevice_ir = nvvm_ir.replace( + store, + """ %arg = sitofp i32 %call to double + %result = call double @__nv_sin(double %arg) + %converted = fptosi double %result to i32 + store i32 %converted, i32* %data, align 4""", + ).replace( + declaration, + """declare double @__nv_sin(double) + +declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()""", + ) + from cuda.pathfinder import BitcodeLibNotFoundError + + program = Program(libdevice_ir, "nvvm", ProgramOptions(use_libdevice=True, arch="sm_80")) + try: + try: + obj = program.compile("ptx") + except BitcodeLibNotFoundError: + pytest.skip("libdevice bitcode not found") + assert isinstance(obj, ObjectCode) + assert obj.code + # Without libdevice, NVVM leaves an external __nv_sin declaration in PTX. + assert not any(b".extern" in line and b"__nv_sin" in line for line in obj.code.splitlines()) + finally: + program.close() diff --git a/cuda_core/tests/test_utils.py b/cuda_core/tests/test_utils.py index c0dffc5a323..2380af2c593 100644 --- a/cuda_core/tests/test_utils.py +++ b/cuda_core/tests/test_utils.py @@ -1023,6 +1023,29 @@ def test_strided_memory_view_proxy_cai_only_has_dlpack_false(): assert proxy.obj is obj +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_strided_memory_view_proxy_cai_view(init_cuda): + """A CAI-only proxy materializes its view through the CAI branch.""" + from cuda.core._memoryview import _StridedMemoryViewProxy + + obj = _make_cuda_array_interface_obj(shape=(2,), strides=None) + view = _StridedMemoryViewProxy(obj).view(-1) + assert view.exporting_obj is obj + assert view.shape == (2,) + assert view.is_device_accessible is True + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_strided_memory_view_view_rejects_dtype_itemsize_mismatch(): + """Changing dtype cannot change the existing layout's element size.""" + view = StridedMemoryView.from_any_interface( + np.arange(4, dtype=np.int16), + stream_ptr=-1, + ) + with pytest.raises(ValueError, match="dtype's itemsize"): + view.view(dtype=np.int8) + + def test_view_as_cai_device_pointer_and_stream_ordering(init_cuda): """``view_as_cai`` on a real device pointer resolves the device ordinal via ``cuPointerGetAttribute`` and takes the cross-stream branch when the CAI diff --git a/cuda_core/tests/test_utils_dlpack.py b/cuda_core/tests/test_utils_dlpack.py index e8796ddd7d6..c3a5994ae1c 100644 --- a/cuda_core/tests/test_utils_dlpack.py +++ b/cuda_core/tests/test_utils_dlpack.py @@ -23,6 +23,10 @@ _PyCapsule_IsValid.argtypes = (ctypes.py_object, ctypes.c_char_p) _PyCapsule_IsValid.restype = ctypes.c_int +_Py_DecRef = ctypes.pythonapi.Py_DecRef +_Py_DecRef.argtypes = (ctypes.c_void_p,) +_Py_DecRef.restype = None + _NUMPY_NATIVE_DLPACK_DTYPES = ( np.uint8, @@ -76,6 +80,31 @@ def test_dlpack_export_roundtrip_special_shapes(shape): _assert_dlpack_export_roundtrip(np.zeros(shape, dtype=np.complex128)) +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_dlpack_export_array_interface_reports_cpu(init_cuda): + """An array-interface view without a DLPack tensor exports as CPU memory.""" + src = np.arange(6, dtype=np.int32) + view = StridedMemoryView.from_array_interface(src) + assert view.is_device_accessible is False + assert view.device_id == init_cuda.device_id + assert view.__dlpack_device__() == (int(DLDeviceType.kDLCPU), 0) + assert np.array_equal(np.from_dlpack(view), src) + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_dlpack_view_of_buffer_reuses_exporting_buffer(init_cuda): + """Re-viewing a Buffer-imported tensor reuses the original Buffer owner.""" + buffer = init_cuda.memory_resource.allocate(16, stream=init_cuda.default_stream) + try: + view = StridedMemoryView.from_dlpack(buffer, stream_ptr=-1) + adjusted = view.view(dtype=np.uint8) + assert adjusted.exporting_obj is buffer + assert adjusted.ptr == int(buffer.handle) + del adjusted, view + finally: + buffer.close() + + def test_dlpack_export_unversioned_capsule_and_deleter(): """``__dlpack__()`` with no ``max_version`` yields an *unversioned* unused DLPack capsule; dropping it unconsumed runs ``_smv_pycapsule_deleter`` on @@ -121,6 +150,21 @@ def __dlpack__(self, **kwargs): StridedMemoryView.from_dlpack(_FakeUnsupportedDevice(), stream_ptr=0) +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_from_dlpack_cuda_stream_none_ambiguous(): + """A CUDA DLPack source requires an explicit consumer stream.""" + + class _FakeCudaDevice: + def __dlpack_device__(self): + return (int(DLDeviceType.kDLCUDA), 0) + + def __dlpack__(self, **kwargs): + raise AssertionError("__dlpack__ must not be reached") + + with pytest.raises(BufferError, match="stream=None is ambiguous"): + StridedMemoryView.from_dlpack(_FakeCudaDevice(), stream_ptr=None) + + class _DLPackNoMaxVersion: """Wraps a StridedMemoryView but rejects the ``max_version`` kwarg, forcing the TypeError fallback in ``view_as_dlpack`` and an *unversioned* capsule import. @@ -166,6 +210,11 @@ def test_from_dlpack_typeerror_fallback_unversioned_import(): # consumer would, exercising the StridedMemoryView exchange-API implementation. # Pointers use PYFUNCTYPE so a failing call raises its real Python exception # (TypeError/RuntimeError/NotImplementedError). +# +# dlpack.h documents every `*_no_sync` entry point as returning "-1 on failure +# with a Python exception set", so every failure below is asserted with +# `pytest.raises`. A test that settles for `assert rc == -1` would be asserting a +# contract violation, not the contract. # --------------------------------------------------------------------------- _PyCapsule_GetPointer = ctypes.pythonapi.PyCapsule_GetPointer @@ -219,6 +268,72 @@ class _DLManagedTensorVersioned(ctypes.Structure): ] +# DLPACK_FLAG_BITMASK_READ_ONLY in dlpack.h. +_FLAG_READ_ONLY = 1 << 0 + + +class _VersionedCapsuleExport: + def __init__(self, base, capsule): + self.base = base + self.capsule = capsule + + def __dlpack_device__(self): + return self.base.__dlpack_device__() + + def __dlpack__(self, **kwargs): + return self.capsule + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_dlpack_versioned_readonly_export_and_import(init_cuda): + """The versioned readonly flag survives a StridedMemoryView round-trip.""" + src = np.arange(4, dtype=np.int32) + src.setflags(write=False) + base = StridedMemoryView.from_array_interface(src) + capsule = base.__dlpack__(max_version=(1, 0)) + dlm = ctypes.cast( + _PyCapsule_GetPointer(capsule, b"dltensor_versioned"), + ctypes.POINTER(_DLManagedTensorVersioned), + ) + assert dlm.contents.flags & _FLAG_READ_ONLY + + imported = StridedMemoryView.from_dlpack( + _VersionedCapsuleExport(base, capsule), + stream_ptr=-1, + ) + assert imported.readonly is True + + +@pytest.mark.parametrize( + ("code", "bits", "lanes", "exception", "match"), + [ + pytest.param(0, 32, 2, NotImplementedError, "vector dtypes", id="lanes"), + pytest.param(1, 24, 1, TypeError, "uint24", id="uint-bits"), + pytest.param(0, 24, 1, TypeError, "int24", id="int-bits"), + pytest.param(2, 8, 1, TypeError, "float8", id="float-bits"), + pytest.param(5, 32, 1, TypeError, "complex32", id="complex-bits"), + pytest.param(6, 1, 1, TypeError, "1-bit bool", id="bool-bits"), + pytest.param(255, 8, 1, TypeError, "Unsupported dtype", id="code"), + ], +) +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_from_dlpack_malformed_dtype_rejected_on_access(code, bits, lanes, exception, match): + """Accessing ``.dtype`` rejects malformed producer dtype metadata.""" + base = StridedMemoryView.from_any_interface(np.arange(4, dtype=np.int32), stream_ptr=-1) + capsule = base.__dlpack__(max_version=(1, 0)) + dlm = ctypes.cast( + _PyCapsule_GetPointer(capsule, b"dltensor_versioned"), + ctypes.POINTER(_DLManagedTensorVersioned), + ) + dlm.contents.dl_tensor.dtype = _DLDataType(code, bits, lanes) + imported = StridedMemoryView.from_dlpack( + _VersionedCapsuleExport(base, capsule), + stream_ptr=-1, + ) + with pytest.raises(exception, match=match): + _ = imported.dtype + + @pytest.mark.agent_authored(model="cursor-grok-4.5") @pytest.mark.parametrize( "max_version, capsule_name, managed_cls", @@ -334,6 +449,14 @@ def test_dlpack_c_exchange_api_current_work_stream(): assert not out.value # set back to NULL +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_dlpack_c_exchange_api_current_work_stream_null_output(): + """``current_work_stream`` rejects a NULL output pointer.""" + api = _get_exchange_api() + with pytest.raises(RuntimeError, match="out_current_stream cannot be NULL"): + api.current_work_stream(int(DLDeviceType.kDLCPU), 0, None) + + def test_dlpack_c_exchange_api_dltensor_from_py_object(): """``dltensor_from_py_object_no_sync`` fills a borrowed DLTensor from a view.""" api = _get_exchange_api() @@ -357,6 +480,27 @@ def test_dlpack_c_exchange_api_dltensor_from_py_object_type_error(): api.dltensor_from_py_object_no_sync(id(not_a_view), ctypes.byref(out)) +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_dlpack_c_exchange_api_dltensor_from_py_object_null_output(): + """``dltensor_from_py_object_no_sync`` rejects a NULL output pointer.""" + api = _get_exchange_api() + view = StridedMemoryView.from_any_interface(np.arange(3), stream_ptr=-1) + with pytest.raises(RuntimeError, match="out cannot be NULL"): + api.dltensor_from_py_object_no_sync(id(view), None) + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_dlpack_c_exchange_api_dltensor_from_py_object_scalar(): + """A borrowed scalar DLTensor has NULL shape and strides pointers.""" + api = _get_exchange_api() + view = StridedMemoryView.from_any_interface(np.array(7, dtype=np.int16), stream_ptr=-1) + out = _DLTensor() + assert api.dltensor_from_py_object_no_sync(id(view), ctypes.byref(out)) == 0 + assert out.ndim == 0 + assert not out.shape + assert not out.strides + + def test_dlpack_c_exchange_api_managed_tensor_roundtrip(): """``managed_tensor_from_py_object_no_sync`` produces a managed tensor that ``managed_tensor_to_py_object_no_sync`` turns back into a StridedMemoryView. @@ -385,6 +529,30 @@ def test_dlpack_c_exchange_api_managed_tensor_roundtrip(): assert imported.ptr == src.ctypes.data +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_dlpack_c_exchange_api_managed_tensor_from_py_object_errors(): + """The managed-tensor producer validates both output and object inputs.""" + api = _get_exchange_api() + view = StridedMemoryView.from_any_interface(np.arange(3), stream_ptr=-1) + with pytest.raises(RuntimeError, match="out cannot be NULL"): + api.managed_tensor_from_py_object_no_sync(id(view), None) + + not_a_view = object() + out = ctypes.c_void_p() + with pytest.raises(TypeError, match="must be a StridedMemoryView"): + api.managed_tensor_from_py_object_no_sync(id(not_a_view), ctypes.byref(out)) + assert not out.value + + +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_dlpack_c_exchange_api_to_py_object_null_output(): + """``managed_tensor_to_py_object_no_sync`` rejects a NULL output pointer.""" + api = _get_exchange_api() + tensor = _DLManagedTensorVersioned() + with pytest.raises(RuntimeError, match="out_py_object cannot be NULL"): + api.managed_tensor_to_py_object_no_sync(ctypes.byref(tensor), None) + + def test_dlpack_c_exchange_api_to_py_object_null_tensor(): """``managed_tensor_to_py_object_no_sync`` rejects a NULL tensor (RuntimeError).""" api = _get_exchange_api() @@ -394,6 +562,36 @@ def test_dlpack_c_exchange_api_to_py_object_null_tensor(): assert not out_obj.value # set to NULL before the error +@pytest.mark.parametrize( + "device_type", + [ + DLDeviceType.kDLCUDA, + DLDeviceType.kDLCUDAHost, + DLDeviceType.kDLCUDAManaged, + ], +) +@pytest.mark.agent_authored(model="gpt-5.6-sol") +def test_dlpack_c_exchange_api_to_py_object_device_accessible(device_type): + """Supported CUDA-family devices produce device-accessible views.""" + api = _get_exchange_api() + tensor = _DLManagedTensorVersioned() + tensor.version = _DLPackVersion(1, 0) + tensor.dl_tensor.device = _DLDevice(int(device_type), 0) + tensor.dl_tensor.dtype = _DLDataType(0, 32, 1) + out_obj = ctypes.c_void_p() + assert api.managed_tensor_to_py_object_no_sync(ctypes.byref(tensor), ctypes.byref(out_obj)) == 0 + assert out_obj.value + try: + imported = ctypes.cast(out_obj, ctypes.py_object).value + assert imported.is_device_accessible is True + assert imported.device_id == 0 + del imported + finally: + # The C API returned a new reference. Release it while the synthetic + # tensor backing the view is still alive -- __dealloc__ dereferences it. + _Py_DecRef(out_obj) + + def test_dlpack_c_exchange_api_managed_tensor_allocator_not_supported(): """Covers the ``managed_tensor_allocator`` entry point, which is unsupported and only ever raises NotImplementedError (StridedMemoryView never allocates)."""