diff --git a/cuda_core/tests/test_cuda_utils.py b/cuda_core/tests/test_cuda_utils.py index 1111e3c8d4..04670b96f2 100644 --- a/cuda_core/tests/test_cuda_utils.py +++ b/cuda_core/tests/test_cuda_utils.py @@ -1,11 +1,14 @@ -# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE +import dataclasses + import pytest from cuda.bindings import driver, runtime from cuda.core._utils import cuda_utils +from cuda.core._utils.clear_error_support import assert_type_str_or_bytes_like, raise_code_path_meant_to_be_unreachable def test_driver_cu_result_explanations_health(): @@ -76,3 +79,60 @@ def test_check_runtime_error(): assert enum_name in msg # Smoke test: We don't want most to be unexpected. assert num_unexpected < len(driver.CUresult) * 0.5 + + +def test_precondition(): + def checker(*args, what=""): + if args[0] < 0: + raise ValueError(f"{what}: negative") + + @cuda_utils.precondition(checker, what="value check") + def my_func(x): + return x * 2 + + assert my_func(5) == 10 + with pytest.raises(ValueError, match="negative"): + my_func(-1) + + +@dataclasses.dataclass +class _DummyOptions: + x: int = 1 + y: str = "hello" + + +def test_check_nvrtc_error_without_handle(): + from cuda.bindings import nvrtc + + assert cuda_utils._check_nvrtc_error(nvrtc.nvrtcResult.NVRTC_SUCCESS) == 0 + with pytest.raises(cuda_utils.NVRTCError): + cuda_utils._check_nvrtc_error(nvrtc.nvrtcResult.NVRTC_ERROR_COMPILATION) + + +def test_check_nvrtc_error_with_handle(init_cuda): + from cuda.bindings import nvrtc + + err, prog = nvrtc.nvrtcCreateProgram(b"invalid code!@#$", b"test.cu", 0, [], []) + assert err == nvrtc.nvrtcResult.NVRTC_SUCCESS + try: + (compile_result,) = nvrtc.nvrtcCompileProgram(prog, 0, []) + assert compile_result != nvrtc.nvrtcResult.NVRTC_SUCCESS + with pytest.raises(cuda_utils.NVRTCError, match="compilation log"): + cuda_utils._check_nvrtc_error(compile_result, handle=prog) + finally: + nvrtc.nvrtcDestroyProgram(prog) + + +def test_check_or_create_options_invalid_type(): + with pytest.raises(TypeError, match="must be provided as an object"): + cuda_utils.check_or_create_options(_DummyOptions, 12345, options_description="test options") + + +def test_assert_type_str_or_bytes_like_rejects_non_str_bytes(): + with pytest.raises(TypeError, match="Expected type str or bytes or bytearray"): + assert_type_str_or_bytes_like(12345) + + +def test_raise_code_path_meant_to_be_unreachable(): + with pytest.raises(RuntimeError, match="This code path is meant to be unreachable"): + raise_code_path_meant_to_be_unreachable() diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index 23affe756d..b9b47ec8df 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2024-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 import ctypes @@ -382,3 +382,10 @@ def test_launch_with_buffers_allocated_by_memory_resource(init_cuda, memory_reso # Verify buffer is properly closed assert buffer.handle == 0, f"{name} buffer should be closed" + + +def test_kernel_arg_unsupported_type(): + from cuda.core._kernel_arg_handler import ParamHolder + + with pytest.raises(TypeError, match="unsupported type"): + ParamHolder(["not_a_valid_kernel_arg"]) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 302a3172d0..c36903ab32 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2024-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE @@ -207,3 +207,17 @@ def test_linker_options_as_bytes_driver_not_supported(): options = LinkerOptions(arch="sm_80") with pytest.raises(RuntimeError, match="as_bytes\\(\\) only supports 'nvjitlink' backend"): options.as_bytes("driver") + + +def test_linker_logs_cached_after_link(compile_ptx_functions): + """After a successful link(), get_error_log/get_info_log should return cached strings.""" + options = LinkerOptions(arch=ARCH) + linker = Linker(*compile_ptx_functions, options=options) + linker.link("cubin") + err_log = linker.get_error_log() + info_log = linker.get_info_log() + assert isinstance(err_log, str) + assert isinstance(info_log, str) + # Calling again should return the same observable values. + assert linker.get_error_log() == err_log + assert linker.get_info_log() == info_log diff --git a/cuda_core/tests/test_memory.py b/cuda_core/tests/test_memory.py index 0473d2d183..8005d3ce6c 100644 --- a/cuda_core/tests/test_memory.py +++ b/cuda_core/tests/test_memory.py @@ -1501,6 +1501,18 @@ def test_graph_memory_resource_object(init_cuda): assert gmr1 == gmr2 == gmr3 +@pytest.mark.skipif(np is None, reason="numpy is not installed") +def test_strided_memory_view_dlpack_errors(): + arr = np.zeros(64, dtype=np.uint8) + smv = StridedMemoryView.from_any_interface(arr, stream_ptr=-1) + with pytest.raises(BufferError, match="dl_device other than None"): + smv.__dlpack__(dl_device=()) + with pytest.raises(BufferError, match="copy=True"): + smv.__dlpack__(copy=True) + with pytest.raises(BufferError, match="Expected max_version"): + smv.__dlpack__(max_version=(9, 8, 7)) + + def test_memory_resource_alloc_zero_bytes(init_cuda, memory_resource_factory): MR, MROps = memory_resource_factory diff --git a/cuda_core/tests/test_multiprocessing_warning.py b/cuda_core/tests/test_multiprocessing_warning.py index 2248cd2cd6..0eb47daaee 100644 --- a/cuda_core/tests/test_multiprocessing_warning.py +++ b/cuda_core/tests/test_multiprocessing_warning.py @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 """ @@ -16,7 +16,7 @@ from cuda.core._event import _reduce_event from cuda.core._memory._device_memory_resource import _deep_reduce_device_memory_resource from cuda.core._memory._ipc import _reduce_allocation_handle -from cuda.core._utils.cuda_utils import reset_fork_warning +from cuda.core._utils.cuda_utils import check_multiprocessing_start_method, reset_fork_warning def test_warn_on_fork_method_device_memory_resource(ipc_device): @@ -145,3 +145,19 @@ def test_warning_emitted_only_once(ipc_device): mr1.close() mr2.close() + + +def test_warn_on_unset_start_method_linux(): + """Test warning when get_start_method raises RuntimeError on Linux (unset start method).""" + with ( + patch("multiprocessing.get_start_method", side_effect=RuntimeError), + patch("platform.system", return_value="Linux"), + warnings.catch_warnings(record=True) as w, + ): + warnings.simplefilter("always") + reset_fork_warning() + check_multiprocessing_start_method() + + fork_warnings = [x for x in w if "fork" in str(x.message).lower()] + assert len(fork_warnings) == 1 + assert "not set" in str(fork_warnings[0].message).lower() diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index 565305a1e3..4bc6fbaf9f 100644 --- a/cuda_core/tests/test_program.py +++ b/cuda_core/tests/test_program.py @@ -427,6 +427,19 @@ def test_program_compile_invalid_target_type(): program.compile("invalid_target") +def test_nvrtc_compile_invalid_code(init_cuda): + """Compiling invalid C++ exercises the HANDLE_RETURN_NVRTC error path with compilation log.""" + from cuda.core._utils.cuda_utils import NVRTCError + + code = 'extern "C" __global__ void bad_kernel() { this_symbol_is_undefined(); }' + program = Program(code, "c++") + try: + with pytest.raises(NVRTCError, match="compilation log"): + program.compile("ptx") + finally: + program.close() + + def test_program_backend_property(): code = 'extern "C" __global__ void my_kernel() {}' program = Program(code, "c++") @@ -481,6 +494,20 @@ def test_nvvm_compile_invalid_target(nvvm_ir): program.close() +@nvvm_available +def test_nvvm_compile_invalid_ir(): + """Compiling invalid NVVM IR exercises the HANDLE_RETURN_NVVM error path.""" + from cuda.bindings.nvvm import nvvmError + + bad_ir = "this is not valid NVVM IR" + program = Program(bad_ir, "nvvm") + try: + with pytest.raises(nvvmError): + program.compile("ptx") + finally: + program.close() + + @nvvm_available @pytest.mark.parametrize("target_type", ["ptx", "ltoir"]) @pytest.mark.parametrize(