diff --git a/cuda_core/tests/graph/test_graph_definition_errors.py b/cuda_core/tests/graph/test_graph_definition_errors.py index 5069dacf4b8..14e48c22501 100644 --- a/cuda_core/tests/graph/test_graph_definition_errors.py +++ b/cuda_core/tests/graph/test_graph_definition_errors.py @@ -6,13 +6,16 @@ import ctypes import pytest +from helpers.contexts import current_context_handle, use_context from helpers.graph_kernels import compile_common_kernels from helpers.misc import try_create_condition -from cuda.core import Device, LaunchConfig -from cuda.core._utils.cuda_utils import CUDAError +from cuda.core import ContextOptions, Device, LaunchConfig, SMResourceOptions +from cuda.core._utils.cuda_utils import CUDAError, driver, handle_return +from cuda.core._utils.version import driver_version from cuda.core.graph import ( EmptyNode, + GraphCompleteOptions, GraphCondition, GraphDefinition, ) @@ -150,6 +153,118 @@ def test_multiple_instantiation(init_cuda): assert g1 is not g2 +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_instantiate_rejects_cycle_and_accepts_repaired_graph(init_cuda): + """Removing a cycle makes an invalid graph executable again.""" + definition = GraphDefinition() + called = [] + first = definition.empty() + last = first.callback(lambda: called.append(True)) + last.succ.add(first) + stream = init_cuda.create_stream() + executable = None + try: + with pytest.raises(RuntimeError, match="invalid structure, such as cycles"): + definition.instantiate() + assert called == [] + last.succ.discard(first) + executable = definition.instantiate() + executable.launch(stream) + stream.sync() + assert called == [True] + finally: + if executable is not None: + executable.close() + last.succ.discard(first) + last.destroy() + first.destroy() + stream.close() + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_device_launch_rejects_host_callback(init_cuda): + """Device-launch graphs reject host callbacks before execution.""" + definition = GraphDefinition() + called = [] + node = definition.callback(lambda: called.append(True)) + try: + with pytest.raises( + RuntimeError, match="device launch failed because the graph contained an unsupported operation" + ): + definition.instantiate(GraphCompleteOptions(device_launch=True)) + assert called == [] + finally: + node.destroy() + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_device_launch_rejects_nodes_from_different_contexts(init_cuda): + """Device-launch instantiation rejects mixed contexts and accepts a repaired graph.""" + if driver_version() < (12, 4, 0): + pytest.skip("green contexts require CUDA driver 12.4 or newer") + if not init_cuda.properties.unified_addressing: + pytest.skip("device launch requires unified addressing") + original_context = current_context_handle() + definition = GraphDefinition() + green_context = first = foreign = executable = stream = None + try: + groups, _ = init_cuda.resources.sm.split(SMResourceOptions(count=None)) + green_context = init_cuda.create_context(ContextOptions(resources=[groups[0]])) + module = compile_common_kernels() + kernel = module.get_kernel("empty_kernel") + config = LaunchConfig(grid=1, block=1) + first = definition.launch(config, kernel) + with use_context(init_cuda, green_context): + foreign = definition.launch(config, kernel) + + # CUkernel handles are context-independent; check each node's resolved context. + first_params = handle_return(driver.cuGraphKernelNodeGetParams(first.handle)) + foreign_params = handle_return(driver.cuGraphKernelNodeGetParams(foreign.handle)) + assert int(first_params.ctx) == original_context + assert int(foreign_params.ctx) == int(green_context.handle) + assert int(foreign_params.ctx) != original_context + assert current_context_handle() == original_context + with pytest.raises(RuntimeError, match="nodes belonging to different contexts"): + definition.instantiate(GraphCompleteOptions(device_launch=True)) + assert current_context_handle() == original_context + + # A single-context graph remains valid after the failed instantiation. + foreign.destroy() + foreign = None + executable = definition.instantiate(GraphCompleteOptions(device_launch=True)) + stream = init_cuda.create_stream() + executable.launch(stream) + stream.sync() + assert current_context_handle() == original_context + finally: + if executable is not None: + executable.close() + if foreign is not None: + foreign.destroy() + if first is not None: + first.destroy() + if stream is not None: + stream.close() + if green_context is not None: + green_context.close() + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_instantiate_rejects_unassociated_condition(init_cuda): + """Instantiation rejects a condition that no conditional node uses.""" + if driver_version() < (12, 8, 0): + pytest.skip("Unused graph-condition diagnostics require CUDA driver 12.8 or newer") + definition = GraphDefinition() + # Deliberately leave this graph-owned condition unassociated with a node. + definition.create_condition(default_value=1) + node = definition.empty() + try: + with pytest.raises(RuntimeError, match="conditional handles are not associated with conditional builders"): + definition.instantiate() + finally: + node.destroy() + + def test_unmatched_alloc_succeeds(init_cuda): """Alloc without corresponding free is valid (graph-scoped lifetime).""" _skip_if_no_mempool() diff --git a/cuda_core/tests/test_device.py b/cuda_core/tests/test_device.py index 44f74620fc2..20dc601b147 100644 --- a/cuda_core/tests/test_device.py +++ b/cuda_core/tests/test_device.py @@ -2,12 +2,14 @@ # SPDX-License-Identifier: Apache-2.0 import contextlib +import ctypes import pickle import subprocess import sys from concurrent.futures import ThreadPoolExecutor import pytest +from helpers.buffers import DummyDeviceMemoryResource from helpers.contexts import ( assert_device_operations_use_bound_context, current_context_handle, @@ -18,7 +20,7 @@ import cuda.core from cuda.bindings import driver, runtime from cuda.core import Device, StreamOptions -from cuda.core._utils.cuda_utils import ComputeCapability, handle_return +from cuda.core._utils.cuda_utils import ComputeCapability, CUDAError, handle_return from cuda.core._utils.version import driver_version @@ -58,6 +60,95 @@ def test_device_set_current(deinit_cuda): assert handle_return(driver.cuCtxGetCurrent()) is not None +@pytest.mark.parametrize("ordinal_kind", ["negative", "one-past-end"]) +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_device_rejects_out_of_range_ordinal(init_cuda, ordinal_kind): + """Invalid ordinals preserve the current context and cached Device.""" + device_count = len(Device.get_all_devices()) + ordinal = -1 if ordinal_kind == "negative" else device_count + original_context = current_context_handle() + with pytest.raises(ValueError): + Device(ordinal) + assert current_context_handle() == original_context + assert Device(init_cuda.device_id) is init_cuda + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_device_context_requires_initialization_on_new_thread(init_cuda): + """Initialization on one thread does not initialize another thread's Device.""" + device_id = init_cuda.device_id + original_context = current_context_handle() + + def worker(): + device = Device(device_id) + assert current_context_handle() == 0 + with pytest.raises(CUDAError): + _ = device.context + with pytest.raises(CUDAError): + device.create_stream() + device.set_current() + stream = None + try: + stream = device.create_stream() + assert int(stream.context.handle) == int(device.context.handle) + with device.allocate(4, stream=stream) as buffer: + stream.sync() + handle_return(driver.cuMemsetD8(int(buffer.handle), 0x29, buffer.size)) + result = (ctypes.c_ubyte * 4)() + handle_return(driver.cuMemcpyDtoH(ctypes.addressof(result), int(buffer.handle), 4)) + assert bytes(result) == b"\x29" * 4 + stream.sync() + finally: + if stream is not None: + stream.close() + handle_return(driver.cuCtxPopCurrent()) + + with ThreadPoolExecutor(max_workers=1) as executor: + executor.submit(worker).result() + assert current_context_handle() == original_context + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_device_memory_resource_controls_allocation(init_cuda): + """Device.allocate uses the public memory resource assigned to the device.""" + device = init_cuda + original_resource = device.memory_resource + replacement = DummyDeviceMemoryResource(device) + try: + device.memory_resource = replacement + assert device.memory_resource is replacement + with device.allocate(16, stream=device.default_stream) as buffer: + assert buffer.memory_resource is replacement + handle_return(driver.cuMemsetD8(int(buffer.handle), 0x5C, buffer.size)) + result = (ctypes.c_ubyte * 16)() + handle_return(driver.cuMemcpyDtoH(ctypes.addressof(result), int(buffer.handle), 16)) + assert bytes(result) == b"\x5c" * 16 + finally: + device.memory_resource = original_resource + assert device.memory_resource is original_resource + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_invalid_memory_resource_assignment_preserves_resource(init_cuda): + """A rejected replacement leaves the existing allocator operational.""" + device = init_cuda + original_resource = device.memory_resource + try: + with pytest.raises(TypeError): + device.memory_resource = object() + assert device.memory_resource is original_resource + with device.allocate(4, stream=device.default_stream) as buffer: + device.default_stream.sync() + assert buffer.memory_resource is original_resource + handle_return(driver.cuMemsetD8(int(buffer.handle), 0x3A, buffer.size)) + result = (ctypes.c_ubyte * 4)() + handle_return(driver.cuMemcpyDtoH(ctypes.addressof(result), int(buffer.handle), 4)) + assert bytes(result) == b"\x3a" * 4 + device.default_stream.sync() + finally: + device.memory_resource = original_resource + + @pytest.mark.agent_authored(model="gpt-5.6-sol") def test_primary_context_cleanup_is_safe_during_python_shutdown(init_cuda, tmp_path): result = subprocess.run( diff --git a/cuda_core/tests/test_green_context.py b/cuda_core/tests/test_green_context.py index a51e8e0bc2b..f9503707701 100644 --- a/cuda_core/tests/test_green_context.py +++ b/cuda_core/tests/test_green_context.py @@ -4,7 +4,7 @@ import numpy as np import pytest -from helpers.contexts import assert_device_operations_use_bound_context, use_context +from helpers.contexts import assert_device_operations_use_bound_context, current_context_handle, use_context from cuda.core import ( ContextOptions, @@ -208,6 +208,24 @@ def test_create_context_requires_resources(init_cuda): init_cuda.create_context(object()) +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_create_context_rejects_empty_resources(init_cuda): + """An empty resource list cannot create a green context.""" + original_context = current_context_handle() + with pytest.raises(ValueError): + init_cuda.create_context(ContextOptions(resources=[])) + assert current_context_handle() == original_context + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_create_context_rejects_unsupported_resource(init_cuda): + """A resource of an unsupported type leaves the current context unchanged.""" + original_context = current_context_handle() + with pytest.raises(TypeError): + init_cuda.create_context(ContextOptions(resources=[object()])) + assert current_context_handle() == original_context + + @pytest.mark.agent_authored(model="claude-opus-4.8") def test_context_handle_alias_and_closed_queries(init_cuda, sm_resource): """``Context._handle`` mirrors ``.handle``; after a (non-current) green diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index a8fe53dab3d..c01bf9103f1 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -1024,3 +1024,64 @@ class SubclassedHandle(driver.CUgraphConditionalHandle): dev = Device() dev.set_current() _assert_kernel_sees_ull(dev, handle, _UINT64_HANDLE_VALUE) + + +@pytest.mark.parametrize("allocation", ["managed", "virtual"]) +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_launch_buffer_subclasses_preserve_kernel_arguments(init_cuda, allocation): + """Managed and virtual buffers marshal their device pointers into kernels.""" + from contextlib import ExitStack + + from cuda.bindings import driver + from cuda.core import ManagedBuffer + from cuda.core._utils.cuda_utils import handle_return + + device = init_cuda + if allocation == "managed" and not device.properties.managed_memory: + pytest.skip("Device does not support managed memory") + if allocation == "virtual" and not device.properties.virtual_memory_management_supported: + pytest.skip("Device does not support CUDA virtual memory") + source = np.arange(16, dtype=np.int32) + actual = np.empty_like(source) + stream = device.create_stream() + buffer = program = managed_pointer = None + try: + if allocation == "managed": + managed_pointer = handle_return( + driver.cuMemAllocManaged(source.nbytes, driver.CUmemAttach_flags.CU_MEM_ATTACH_GLOBAL) + ) + buffer = ManagedBuffer.from_handle(managed_pointer, source.nbytes) + assert isinstance(buffer, ManagedBuffer) + else: + from cuda.core import VirtualMemoryBuffer, VirtualMemoryResource, VirtualMemoryResourceOptions + + resource = VirtualMemoryResource(device, VirtualMemoryResourceOptions(handle_type=None)) + buffer = resource.allocate(source.nbytes, stream=stream) + assert isinstance(buffer, VirtualMemoryBuffer) + handle_return(driver.cuMemcpyHtoDAsync(buffer.handle, source.ctypes.data, source.nbytes, stream.handle)) + program = Program( + 'extern "C" __global__ void update(int* data, int offset) { ' + "int i = threadIdx.x; data[i] = data[i] * 3 + offset; }", + "c++", + ) + module = program.compile("cubin") + kernel = module.get_kernel("update") + launch(stream, LaunchConfig(grid=1, block=16), kernel, buffer, np.int32(7)) + stream.sync() + handle_return(driver.cuMemcpyDtoH(actual.ctypes.data, buffer.handle, actual.nbytes)) + np.testing.assert_array_equal(actual, source * 3 + 7) + buffer.close() + assert buffer.is_closed + with pytest.raises(RuntimeError, match="Buffer has been closed"): + launch(stream, LaunchConfig(grid=1, block=16), kernel, buffer, np.int32(7)) + finally: + # Attempt every cleanup even if synchronization or another cleanup fails. + with ExitStack() as cleanup: + cleanup.callback(stream.close) + if program is not None: + cleanup.callback(program.close) + if managed_pointer is not None: + cleanup.callback(lambda: handle_return(driver.cuMemFree(managed_pointer))) + if buffer is not None: + cleanup.callback(buffer.close) + stream.sync() diff --git a/cuda_core/tests/test_memory.py b/cuda_core/tests/test_memory.py index 6451ae0ba40..c02d92c0e33 100644 --- a/cuda_core/tests/test_memory.py +++ b/cuda_core/tests/test_memory.py @@ -1940,6 +1940,38 @@ def test_vmm_buffers_alive_at_shutdown_are_freed_quietly(init_cuda): assert result.stderr == "" +@pytest.mark.parametrize( + "options", + [ + pytest.param({"addr_align": 3}, id="invalid-alignment"), + pytest.param({"peers": (-1,)}, id="invalid-peer"), + ], +) +@pytest.mark.thread_unsafe(reason="records process-global warnings") +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_failed_vmm_allocation_does_not_block_another_resource(init_cuda, options): + """A failed allocation leaves its config unchanged and another VMM resource usable.""" + device = _vmm_device_or_skip() + invalid_resource = _vmm_resource(device, **options) + error_match = "CUDA_ERROR_INVALID_VALUE" if "addr_align" in options else r"CUDA_ERROR_INVALID_(VALUE|DEVICE)" + with assert_no_cuda_warning(), pytest.raises(CUDAError, match=error_match): + invalid_resource.allocate(4096) + for name, value in options.items(): + assert getattr(invalid_resource.config, name) == value + + # Use a separate resource with valid options; the invalid configuration is unchanged. + resource = _vmm_resource(device) + buffer = resource.allocate(4096) + ptr, size = int(buffer.handle), buffer.size + try: + _vmm_fill(buffer, 0x5C) + assert _vmm_read(buffer, 0, 16) == bytes([0x5C]) * 16 + finally: + with assert_no_cuda_warning(): + buffer.close() + _vmm_assert_released([(ptr, size)]) + + @pytest.mark.agent_authored(model="claude-fable-5-1") @pytest.mark.thread_unsafe(reason="records process-global warnings") def test_vmm_failed_grow_leaves_input_intact(init_cuda): diff --git a/cuda_core/tests/test_module.py b/cuda_core/tests/test_module.py index a3ffa6d4e75..5b8c7a4b791 100644 --- a/cuda_core/tests/test_module.py +++ b/cuda_core/tests/test_module.py @@ -916,3 +916,19 @@ def test_kernel_equality_hash_and_repr(get_saxpy_kernel_cubin): assert kernel != other assert kernel != object() assert repr(kernel).startswith("= 13\.3"): ProgramOptions(arch="sm_80", use_bundled_headers=True) assert calls == [True] + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_nvrtc_debug_close_tolerates_removed_source(init_cuda): + """A temp-file cleaner may remove debug source before Program.close().""" + code = 'extern "C" __global__ void cleanup_kernel() {}' + program = Program(code, "c++", ProgramOptions(debug=True, arch="sm_80")) + try: + object_code = program.compile("ptx") + source_path = object_code.name + assert os.path.isfile(source_path) + with open(source_path, encoding="utf-8") as source: + assert source.read() == code + os.unlink(source_path) + program.close() + program.close() + assert not os.path.exists(source_path) + assert b"cleanup_kernel" in object_code.code + finally: + program.close() + + +@nvrtc_pch_available +@pytest.mark.skipif("NVRTC_DISABLE_PCH" in os.environ, reason="PCH is disabled by NVRTC_DISABLE_PCH") +@pytest.mark.thread_unsafe(reason="creates PCH files in the shared pytest tmp_path") +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_nvrtc_automatic_pch_preserves_compiled_kernel(init_cuda, tmp_path): + """Automatic PCH compilation preserves kernel output and reports creation status.""" + from io import StringIO + + import numpy as np + + from cuda.bindings import driver + from cuda.core import LaunchConfig, launch + + if handle_return(nvrtc.nvrtcGetPCHHeapSize()) == 0: + pytest.skip("PCH is disabled because the NVRTC PCH heap size is zero") + + declarations = "\n".join(f"struct Type{i} {{ int data[{i % 16 + 1}]; }};" for i in range(2048)) + (tmp_path / "pch_header.cuh").write_text( + declarations + "\n__device__ int pch_answer() { return 37; }\n", encoding="utf-8" + ) + code = '#include "pch_header.cuh"\n#pragma nv_hdrstop\nextern "C" __global__ void answer(int* out) { *out = pch_answer(); }' + stream = init_cuda.create_stream() + buffer = None + actual = np.empty(1, dtype=np.int32) + try: + buffer = init_cuda.allocate(4, stream=stream) + for expected_status in (PCHStatusType.CREATED, PCHStatusType.NOT_ATTEMPTED): + program = Program( + code, + "c++", + ProgramOptions(pch=True, pch_dir=str(tmp_path), pch_messages=True, include_path=(str(tmp_path),)), + ) + try: + logs = StringIO() + module = program.compile("cubin", logs=logs) + kernel = module.get_kernel("answer") + handle_return(driver.cuMemsetD32Async(buffer.handle, 0, 1, stream.handle)) + launch(stream, LaunchConfig(grid=1, block=1), kernel, buffer) + stream.sync() + handle_return(driver.cuMemcpyDtoH(actual.ctypes.data, buffer.handle, actual.nbytes)) + assert actual[0] == 37 + if expected_status == PCHStatusType.CREATED and program.pch_status == PCHStatusType.NOT_ATTEMPTED: + pytest.skip("NVRTC did not create an automatic PCH; kernel output was verified") + assert program.pch_status == expected_status, logs.getvalue() + assert list(tmp_path.glob("*.pch")), logs.getvalue() + finally: + program.close() + finally: + if buffer is not None: + buffer.close() + stream.close() + for path in tmp_path.glob("*.pch"): + path.unlink() diff --git a/cuda_core/tests/test_stream.py b/cuda_core/tests/test_stream.py index 02e817cf85b..bb98108bf8f 100644 --- a/cuda_core/tests/test_stream.py +++ b/cuda_core/tests/test_stream.py @@ -1,13 +1,16 @@ # SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 +import ctypes + import pytest +from helpers.contexts import current_context_handle, no_current_context from helpers.misc import StreamWrapper from cuda.core import Device, Stream, StreamOptions from cuda.core._event import Event from cuda.core._stream import LEGACY_DEFAULT_STREAM, PER_THREAD_DEFAULT_STREAM, Stream_accept -from cuda.core._utils.cuda_utils import driver +from cuda.core._utils.cuda_utils import driver, handle_return def test_stream_init_disabled(): @@ -136,6 +139,39 @@ def test_stream_from_handle(): assert isinstance(stream, Stream) +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_borrowed_external_stream_resolves_context_without_taking_ownership(init_cuda): + """A borrowed driver stream resolves its context and remains externally owned.""" + from contextlib import ExitStack + + raw_stream = handle_return(driver.cuStreamCreate(driver.CUstream_flags.CU_STREAM_NON_BLOCKING)) + borrowed = None + try: + borrowed = Stream.from_handle(int(raw_stream)) + with no_current_context(): + assert int(borrowed.context.handle) == int(init_cuda.context.handle) + assert borrowed.device is init_cuda + assert current_context_handle() == 0 + + with init_cuda.allocate(16, stream=borrowed) as buffer: + handle_return(driver.cuMemsetD8Async(int(buffer.handle), 0x6D, buffer.size, raw_stream)) + result = (ctypes.c_ubyte * 16)() + handle_return(driver.cuMemcpyDtoHAsync(ctypes.addressof(result), int(buffer.handle), 16, raw_stream)) + borrowed.sync() + assert bytes(result) == b"\x6d" * 16 + borrowed.sync() + borrowed.close() + assert borrowed.is_closed + handle_return(driver.cuStreamQuery(raw_stream)) + assert int(handle_return(driver.cuStreamGetCtx(raw_stream))) == int(init_cuda.context.handle) + finally: + with ExitStack() as cleanup: + cleanup.callback(lambda: handle_return(driver.cuStreamDestroy(raw_stream))) + if borrowed is not None: + cleanup.callback(borrowed.close) + handle_return(driver.cuStreamSynchronize(raw_stream)) + + def test_legacy_default_stream(): assert isinstance(LEGACY_DEFAULT_STREAM, Stream) diff --git a/cuda_core/tests/test_tensor_map.py b/cuda_core/tests/test_tensor_map.py index 33d24deafa9..de83d3f8db1 100644 --- a/cuda_core/tests/test_tensor_map.py +++ b/cuda_core/tests/test_tensor_map.py @@ -1,13 +1,19 @@ # SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 +import gc +import weakref + import numpy as np import pytest +from helpers.contexts import no_current_context, use_context from helpers.memory import create_managed_memory_resource_or_skip, skip_if_managed_memory_unsupported from cuda.core import ( + ContextOptions, Device, ManagedMemoryResourceOptions, + SMResourceOptions, TensorMapDescriptor, ) from cuda.core._tensor_map import ( @@ -20,6 +26,8 @@ TensorMapSwizzle, _coerce_tensor_map_descriptor_options, ) +from cuda.core._utils.cuda_utils import CUDAError +from cuda.core._utils.version import driver_version from cuda.core.utils import StridedMemoryView @@ -58,6 +66,66 @@ def _as_view(obj): return StridedMemoryView.from_any_interface(obj, stream_ptr=-1) +@pytest.mark.parametrize( + "dtype,data_type", + [ + pytest.param(np.uint8, TensorMapDataType.UINT8, id="uint8"), + pytest.param(np.uint16, TensorMapDataType.UINT16, id="uint16"), + pytest.param(np.uint32, TensorMapDataType.UINT32, id="uint32"), + pytest.param(np.uint64, TensorMapDataType.UINT64, id="uint64"), + pytest.param(np.int32, TensorMapDataType.INT32, id="int32"), + pytest.param(np.int64, TensorMapDataType.INT64, id="int64"), + pytest.param(np.float16, TensorMapDataType.FLOAT16, id="float16"), + pytest.param("bfloat16", TensorMapDataType.BFLOAT16, id="bfloat16"), + pytest.param(np.float32, TensorMapDataType.FLOAT32, id="float32"), + pytest.param(np.float64, TensorMapDataType.FLOAT64, id="float64"), + pytest.param(np.float32, TensorMapDataType.FLOAT32_FTZ, id="float32-ftz"), + pytest.param(np.float32, TensorMapDataType.TFLOAT32, id="tfloat32"), + pytest.param(np.float32, TensorMapDataType.TFLOAT32_FTZ, id="tfloat32-ftz"), + ], +) +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_tiled_tensor_map_encodes_dtypes_with_padded_rows(dev, skip_if_no_tma, dtype, data_type): + """Encode a tiled descriptor for each dtype with non-contiguous row strides.""" + if dtype == "bfloat16": + dtype = pytest.importorskip("ml_dtypes").bfloat16 + dtype = np.dtype(dtype) + stream = dev.create_stream() + source = view = descriptor = None + try: + source = dev.allocate(4 * 32 * dtype.itemsize, stream=stream) + stream.sync() + view = StridedMemoryView.from_buffer( + source, shape=(4, 16), strides=(32, 1), itemsize=dtype.itemsize, dtype=dtype + ) + descriptor = view.as_tensor_map(box_dim=(2, 16), data_type=data_type) + assert isinstance(descriptor, TensorMapDescriptor) + assert descriptor.device == dev + assert ", rank=2," in repr(descriptor) + assert f"dtype={data_type.name}," in repr(descriptor) + finally: + descriptor = view = None + if source is not None: + source.close() + stream.close() + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_as_tensor_map_rejects_misaligned_device_pointer(dev): + """Tensor-map creation rejects a valid device allocation's misaligned view.""" + source = dev.allocate(16 * 4 + 16, stream=dev.default_stream) + tensor = view = None + try: + tensor = _DeviceArray(source, (16,)) + tensor.__cuda_array_interface__["data"] = (int(source.handle) + 4, False) + view = StridedMemoryView.from_any_interface(tensor, stream_ptr=-1) + with pytest.raises(ValueError, match="Global memory address must be 16-byte aligned"): + view.as_tensor_map(box_dim=(16,)) + finally: + tensor = view = None + source.close() + + class TestTensorMapEnums: """Test that enum wrappers expose the expected values.""" @@ -765,3 +833,82 @@ def test_coerce_requires_box_dim_without_options(self): l2_promotion=TensorMapL2Promotion.NONE, oob_fill=TensorMapOOBFill.NONE, ) + + +@pytest.mark.parametrize("context_kind", ["missing", "different"]) +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_tensor_map_replace_address_context_error_preserves_ownership(dev, skip_if_no_tma, context_kind): + """A rejected replacement retains the source; a valid one transfers ownership.""" + if context_kind == "different" and driver_version() < (12, 4, 0): + pytest.skip("green contexts require CUDA driver 12.4 or newer") + + def make_view(allocation): + tensor = _DeviceArray(allocation, (4, 16)) + tensor.__cuda_array_interface__["strides"] = (32 * 4, 4) + return _as_view(tensor), weakref.ref(tensor) + + stream = dev.create_stream() + original = replacement = None + original_view = replacement_view = descriptor = green_context = None + try: + original = dev.allocate(4 * 32 * 4, stream=stream) + replacement = dev.allocate(4 * 32 * 4, stream=stream) + stream.sync() + original_view, original_ref = make_view(original) + descriptor = original_view.as_tensor_map(box_dim=(2, 16)) + original_view = None + gc.collect() + assert original_ref() is not None + + replacement_view, rejected_ref = make_view(replacement) + if context_kind == "missing": + other_context = no_current_context() + else: + groups, _ = dev.resources.sm.split(SMResourceOptions(count=None)) + green_context = dev.create_context(ContextOptions(resources=[groups[0]])) + other_context = use_context(dev, green_context) + error_match = ( + "requires an active CUDA context" + if context_kind == "missing" + else "was created in a different CUDA context" + ) + with other_context, pytest.raises(RuntimeError, match=error_match): + descriptor.replace_address(replacement_view) + replacement_view = None + gc.collect() + assert original_ref() is not None + assert rejected_ref() is None + + replacement_view, replacement_ref = make_view(replacement) + assert descriptor.replace_address(replacement_view) is None + replacement_view = None + gc.collect() + assert original_ref() is None + assert replacement_ref() is not None + assert descriptor.device == dev + + descriptor = None + gc.collect() + assert replacement_ref() is None + finally: + descriptor = original_view = replacement_view = None + if green_context is not None: + green_context.close() + for allocation in (replacement, original): + if allocation is not None: + allocation.close() + stream.close() + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_as_tensor_map_reports_unsupported_hardware_as_cuda_error(dev): + """A valid tensor-map request on unsupported hardware raises CUDAError.""" + if dev.properties.tensor_map_access_supported: + pytest.skip("requires a device without tensor-map access") + with dev.allocate(64 * 4, stream=dev.default_stream) as source: + view = StridedMemoryView.from_buffer(source, shape=(64,), itemsize=4, dtype=np.dtype(np.float32)) + try: + with pytest.raises(CUDAError, match=r"Compute capability 9\.0 or higher is required"): + view.as_tensor_map(box_dim=(16,)) + finally: + view = None diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index 23e1e2976fb..f04f182d7c0 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -691,6 +691,10 @@ def test_resource_descriptor_from_mipmapped_array_rejects_non_mipmap(): def test_texture_object_from_mipmapped_array(init_cuda): + """Mipmap texture creation forwards the requested trilinear-optimization flag.""" + from cuda.bindings import driver + from cuda.core._utils.cuda_utils import handle_return + mip = Device().create_mipmapped_array( MipmappedArrayOptions( shape=(32, 32), @@ -710,11 +714,14 @@ def test_texture_object_from_mipmapped_array(init_cuda): mipmap_level_bias=0.0, min_mipmap_level_clamp=0.0, max_mipmap_level_clamp=float(mip.num_levels - 1), + disable_trilinear_optimization=True, ) tex = Device().create_texture_object(resource=res, options=tex_desc) try: assert tex.handle != 0 assert tex.resource is res + descriptor = handle_return(driver.cuTexObjectGetTextureDesc(tex.handle)) + assert descriptor.flags & driver.CU_TRSF_DISABLE_TRILINEAR_OPTIMIZATION finally: tex.close() finally: @@ -1244,3 +1251,159 @@ def test_array_copy_rejects_non_buffer(init_cuda): stream.close() finally: arr.close() + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_resource_descriptor_public_dimensions_match_backing(init_cuda): + """Linear and pitched resources expose their applicable dimensions.""" + with _alloc_device_buffer(init_cuda, 4096) as buffer: + linear = ResourceDescriptor.from_linear(buffer, format=ArrayFormatType.UINT32, num_channels=1, size_bytes=2048) + pitched = ResourceDescriptor.from_pitch2d( + buffer, format=ArrayFormatType.UINT32, num_channels=1, width=8, height=4, pitch_bytes=64 + ) + assert linear.source is pitched.source is buffer + assert linear.size_bytes == 2048 + assert linear.width is linear.height is linear.pitch_bytes is None + assert pitched.size_bytes is None + assert (pitched.width, pitched.height, pitched.pitch_bytes) == (8, 4, 64) + assert linear.format == pitched.format == ArrayFormatType.UINT32 + assert linear.num_channels == pitched.num_channels == 1 + + +@pytest.mark.parametrize("shape", [(16,), (16, 8, 4)], ids=["1d", "3d"]) +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_mipmapped_array_rank_level_roundtrip(init_cuda, shape): + """Level 1 of one- and three-dimensional mipmaps preserves metadata and data.""" + from contextlib import closing + + options = MipmappedArrayOptions(shape=shape, format=ArrayFormatType.UINT8, num_channels=1, num_levels=3) + with closing(init_cuda.create_stream()) as stream, init_cuda.create_mipmapped_array(options) as mipmap: + text = repr(mipmap) + assert "MipmappedArray(" in text + assert f"shape={shape}" in text + assert "UINT8" in text and "num_channels=1" in text and "num_levels=3" in text + with mipmap.get_level(1) as level: + expected_shape = tuple(extent // 2 for extent in shape) + assert level.shape == expected_shape + assert level.format == ArrayFormatType.UINT8 + assert level.num_channels == 1 + resource = ResourceDescriptor.from_opaque_array(level) + assert resource.source is level + assert "kind='array'" in repr(resource) + source = np.arange(level.size_bytes, dtype=np.uint8) + destination = np.empty_like(source) + level.copy_from(source, stream=stream) + assert level.copy_to(destination, stream=stream) is destination + stream.sync() + np.testing.assert_array_equal(destination, source) + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_surface_invalid_resource_then_live_texture_lifecycle(init_cuda): + """Invalid surface resources leave subsequent array-backed objects usable.""" + from cuda.bindings import driver + from cuda.core._utils.cuda_utils import handle_return + + with pytest.raises(TypeError, match="resource must be a ResourceDescriptor"): + init_cuda.create_surface_object(resource=object()) + with init_cuda.create_opaque_array( + OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=4, is_surface_load_store=True) + ) as array: + resource = ResourceDescriptor.from_opaque_array(array) + assert "kind='array'" in repr(resource) + with ( + init_cuda.create_surface_object(resource=resource) as surface, + init_cuda.create_texture_object( + resource=resource, + options=TextureObjectOptions(read_mode=ReadModeType.NORMALIZED_FLOAT, srgb=True), + ) as texture, + ): + assert surface.device == texture.device == init_cuda + assert surface.handle != 0 and texture.handle != 0 + assert f"handle=0x{int(surface.handle):x}" in repr(surface) + assert f"handle=0x{int(texture.handle):x}" in repr(texture) + descriptor = handle_return(driver.cuTexObjectGetTextureDesc(texture.handle)) + assert descriptor.flags & driver.CU_TRSF_SRGB + assert surface.is_closed and texture.is_closed + assert repr(surface) == "SurfaceObject(handle=0x0)" + assert repr(texture) == "TextureObject(handle=0x0)" + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_texture_border_color_is_sampled_on_device(init_cuda): + """Out-of-bounds texture reads return the configured RGBA border color.""" + from contextlib import closing + + from cuda.bindings import driver + from cuda.core import LaunchConfig, Program, ProgramOptions, launch + from cuda.core._utils.cuda_utils import handle_return + + code = r""" + extern "C" __global__ void sample_border(unsigned long long texture, float* result) { + float coordinate = threadIdx.x == 0 ? -1.f : 0.5f; + float x, y, z, w; + asm("tex.2d.v4.f32.f32 {%0,%1,%2,%3}, [%4,{%5,%6}];" + : "=f"(x), "=f"(y), "=f"(z), "=f"(w) + : "l"(texture), "f"(coordinate), "f"(coordinate)); + result[4 * threadIdx.x] = x; + result[4 * threadIdx.x + 1] = y; + result[4 * threadIdx.x + 2] = z; + result[4 * threadIdx.x + 3] = w; + } + """ + border = (0.25, 0.5, 0.75, 1.0) + texel = np.array([1.0, 2.0, 3.0, 4.0], dtype=np.float32) + source = np.tile(texel, (2, 2, 1)) + actual = np.empty(8, dtype=np.float32) + with ( + closing(Program(code, "c++", options=ProgramOptions(std="c++17", arch=f"sm_{init_cuda.arch}"))) as program, + closing(init_cuda.create_stream()) as stream, + init_cuda.create_opaque_array( + OpaqueArrayOptions(shape=(2, 2), format=ArrayFormatType.FLOAT32, num_channels=4) + ) as array, + init_cuda.allocate(actual.nbytes, stream=stream) as output, + ): + module = program.compile("cubin") + kernel = module.get_kernel("sample_border") + array.copy_from(source, stream=stream) + with init_cuda.create_texture_object( + resource=ResourceDescriptor.from_opaque_array(array), + options=TextureObjectOptions( + address_mode=AddressModeType.BORDER, + read_mode=ReadModeType.NORMALIZED_FLOAT, + border_color=border, + ), + ) as texture: + launch(stream, LaunchConfig(grid=1, block=2), kernel, np.uint64(texture.handle), output) + handle_return(driver.cuMemcpyDtoHAsync(actual.ctypes.data, output.handle, actual.nbytes, stream.handle)) + stream.sync() + np.testing.assert_array_equal(actual[:4], np.asarray(border, dtype=np.float32)) + np.testing.assert_array_equal(actual[4:], texel) + + +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_surface_creation_rejects_foreign_context_resource(init_cuda): + """Rejecting another context's array leaves owner-context surface creation usable.""" + from helpers.contexts import use_context + + from cuda.core import ContextOptions, SMResourceOptions + from cuda.core._utils.version import driver_version + + if driver_version() < (12, 4, 0): + pytest.skip("green contexts require CUDA driver 12.4 or newer") + groups, _ = init_cuda.resources.sm.split(SMResourceOptions(count=None)) + green_context = init_cuda.create_context(ContextOptions(resources=[groups[0]])) + try: + with init_cuda.create_opaque_array( + OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=4, is_surface_load_store=True) + ) as array: + resource = ResourceDescriptor.from_opaque_array(array) + with use_context(init_cuda, green_context): + with pytest.raises(ValueError, match="resource is not compatible with this Device object"): + init_cuda.create_surface_object(resource=resource) + assert current_context_handle() == int(green_context.handle) + with init_cuda.create_surface_object(resource=resource) as surface: + assert surface.resource is resource + assert not surface.is_closed + finally: + green_context.close() diff --git a/cuda_core/tests/test_utils_dlpack.py b/cuda_core/tests/test_utils_dlpack.py index ef1a31e2965..43b135e96ce 100644 --- a/cuda_core/tests/test_utils_dlpack.py +++ b/cuda_core/tests/test_utils_dlpack.py @@ -843,3 +843,20 @@ def test_dlpack_c_exchange_api_managed_tensor_allocator_not_supported(): # Currently sets a Python error when `SetError` isn't passed. api.managed_tensor_allocator(None, ctypes.byref(out), None, None) assert not out.value # set to NULL before the error + + +@pytest.mark.parametrize(("dtype_name", "max_itemsize"), [("longdouble", 8), ("clongdouble", 16)]) +@pytest.mark.agent_authored(model="gpt-6-sol") +def test_dlpack_export_rejects_extended_numpy_width(dtype_name, max_itemsize): + """Supported dtype kinds with extended widths cannot be exported as DLPack.""" + dtype = np.dtype(getattr(np, dtype_name)) + if dtype.itemsize <= max_itemsize: + pytest.skip(f"{dtype_name} has a DLPack-supported width on this platform") + source = np.zeros(4, dtype=dtype) + view = StridedMemoryView.from_array_interface(source) + error_match = "Unsupported floating dtype" if dtype.kind == "f" else "Unsupported complex dtype" + with pytest.raises(BufferError, match=error_match): + view.__dlpack__(max_version=(1, 0)) + assert view.dtype == dtype + assert view.ptr == source.ctypes.data + assert view.exporting_obj is source