# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
#
# SPDX-License-Identifier: Apache-2.0

import contextlib
import re
import warnings

import pytest

from cuda.core import _linker
from cuda.core._device import Device
from cuda.core._module import Kernel, ObjectCode
from cuda.core._program import Program, ProgramOptions
from cuda.core._utils.cuda_utils import CUDAError, handle_return
from cuda.core.typing import CompilerBackendType, PCHStatusType

pytest_plugins = ("cuda_python_test_helpers.nvvm_bitcode",)

is_culink_backend = _linker._decide_nvjitlink_or_driver()


def _is_nvvm_available():
    """Check if NVVM is available."""
    try:
        from cuda.core._program import _get_nvvm_module

        _get_nvvm_module()
        return True
    except RuntimeError:
        return False


nvvm_available = pytest.mark.skipif(
    not _is_nvvm_available(), reason="NVVM not available (libNVVM not found or cuda-bindings < 12.9.0)"
)

with contextlib.suppress(Exception):
    from cuda.core._utils.cuda_utils import nvrtc


def _get_nvrtc_version_for_tests():
    """
    Get NVRTC version.

    Returns:
        int: Version in format major * 1000 + minor * 100 (e.g., 13200 for CUDA 13.2)
        None: If NVRTC is not available
    """
    try:
        nvrtc_major, nvrtc_minor = handle_return(nvrtc.nvrtcVersion())
        version = nvrtc_major * 1000 + nvrtc_minor * 100
        return version
    except Exception:
        return None


def _has_nvrtc_pch_apis_for_tests():
    required = (
        "nvrtcGetPCHHeapSize",
        "nvrtcSetPCHHeapSize",
        "nvrtcGetPCHCreateStatus",
        "nvrtcGetPCHHeapSizeRequired",
    )
    return all(hasattr(nvrtc, name) for name in required)


nvrtc_pch_available = pytest.mark.skipif(
    (_get_nvrtc_version_for_tests() or 0) < 12800 or not _has_nvrtc_pch_apis_for_tests(),
    reason="PCH runtime APIs require NVRTC >= 12.8 bindings",
)


def _has_check_nvvm_compiler_options():
    try:
        import cuda.bindings.utils as utils
    except ModuleNotFoundError:
        return False
    return hasattr(utils, "check_nvvm_compiler_options")


has_nvvm_option_checker = pytest.mark.skipif(
    not _has_check_nvvm_compiler_options(),
    reason="cuda.bindings.utils.check_nvvm_compiler_options not available (cuda-bindings too old?)",
)


def _check_nvvm_arch(arch: str) -> bool:
    """Check if the given NVVM arch is supported by the installed libNVVM."""
    from cuda.bindings.utils import check_nvvm_compiler_options

    return check_nvvm_compiler_options([f"-arch={arch}"])


def _check_nvvm_supports_numba_debug() -> bool:
    """Check if the installed libNVVM recognizes --numba-debug (CTK 13.2+)."""
    if not _has_check_nvvm_compiler_options():
        return False
    from cuda.bindings.utils import check_nvvm_compiler_options

    return check_nvvm_compiler_options(["--numba-debug"])


@pytest.fixture(scope="session")
def nvvm_ir():
    """Generate working NVVM IR with proper version metadata.
    The try clause here is used for older nvvm modules which
    might not have an ir_version() method. In which case the
    fallback assumes no version metadata will be present in
    the input nvvm ir
    """
    from cuda.core._program import _get_nvvm_module

    nvvm = _get_nvvm_module()
    major, minor, debug_major, debug_minor = nvvm.ir_version()

    nvvm_ir_template = """target triple = "nvptx64-unknown-cuda"
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"

define i32 @ave(i32 %a, i32 %b) {{
entry:
  %add = add nsw i32 %a, %b
  %div = sdiv i32 %add, 2
  ret i32 %div
}}

define void @simple(i32* %data) {{
entry:
  %0 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
  %1 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
  %mul = mul i32 %0, %1
  %2 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  %add = add i32 %mul, %2
  %call = call i32 @ave(i32 %add, i32 %add)
  %idxprom = sext i32 %add to i64
  store i32 %call, i32* %data, align 4
  ret void
}}

declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() nounwind readnone
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x() nounwind readnone
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() nounwind readnone

!nvvm.annotations = !{{!0}}
!0 = !{{void (i32*)* @simple, !"kernel", i32 1}}

!nvvmir.version = !{{!1}}
!1 = !{{i32 {major}, i32 {minor}, i32 {debug_major}, i32 {debug_minor}}}
"""
    return nvvm_ir_template.format(major=major, minor=minor, debug_major=debug_major, debug_minor=debug_minor)


@pytest.fixture(scope="module")
def ptx_code_object():
    code = 'extern "C" __global__ void my_kernel() {}'
    program = Program(code, "c++")
    ptx_object_code = program.compile("ptx")
    return ptx_object_code


@pytest.mark.parametrize(
    "options",
    [
        ProgramOptions(name="abc"),
        ProgramOptions(device_code_optimize=True, debug=True),
        pytest.param(
            ProgramOptions(debug=True, numba_debug=True),
            marks=pytest.mark.skipif(
                (_get_nvrtc_version_for_tests() or 0) < 13200,
                reason="numba_debug requires NVRTC >= 13.2",
            ),
        ),
        ProgramOptions(relocatable_device_code=True, max_register_count=32),
        ProgramOptions(ftz=True, prec_sqrt=False, prec_div=False),
        ProgramOptions(fma=False, use_fast_math=True),
        ProgramOptions(extra_device_vectorization=True),
        ProgramOptions(link_time_optimization=True),
        ProgramOptions(define_macro="MY_MACRO"),
        ProgramOptions(define_macro=("MY_MACRO", "99")),
        ProgramOptions(define_macro=[("MY_MACRO", "99")]),
        ProgramOptions(define_macro=[("MY_MACRO", "99"), ("MY_OTHER_MACRO", "100")]),
        ProgramOptions(undefine_macro=["MY_MACRO", "MY_OTHER_MACRO"]),
        ProgramOptions(undefine_macro="MY_MACRO", include_path="/usr/local/include"),
        ProgramOptions(builtin_initializer_list=False, disable_warnings=True),
        ProgramOptions(restrict=True, device_as_default_execution_space=True),
        ProgramOptions(device_int128=True, optimization_info="inline"),
        ProgramOptions(no_display_error_number=True),
        ProgramOptions(diag_error=1234, diag_suppress=1234),
        ProgramOptions(diag_error=[1234, 1223], diag_suppress=(1234, 1223)),
        ProgramOptions(diag_warn=1000),
        ProgramOptions(std="c++11", ptxas_options=["-v"]),
        ProgramOptions(std="c++11", ptxas_options=["-v", "-O2"]),
        ProgramOptions(brief_diagnostics=True),
        ProgramOptions(builtin_move_forward=False),
        ProgramOptions(extensible_whole_program=True),
        ProgramOptions(fdevice_syntax_only=True),
        ProgramOptions(gen_opt_lto=True),
        ProgramOptions(minimal=True),
        ProgramOptions(no_source_include=True),
        # TODO: Add test for pre_include once we have a suitable header in the test environment
        # ProgramOptions(pre_include="cuda_runtime.h"),
        ProgramOptions(no_cache=True),
        pytest.param(
            ProgramOptions(arch="sm_100", device_float128=True),
            marks=pytest.mark.skipif(
                Device().compute_capability < (100, 0),
                reason="device_float128 requires sm_100 or later",
            ),
        ),
        ProgramOptions(frandom_seed="12345"),
        ProgramOptions(ofast_compile="min"),
        pytest.param(
            ProgramOptions(pch=True),
            marks=pytest.mark.skipif(
                (_get_nvrtc_version_for_tests() or 0) < 12800,
                reason="PCH requires NVRTC >= 12.8",
            ),
        ),
        # TODO: pch_dir requires actual PCH directory to exist - needs integration test
        # pytest.param(
        #     ProgramOptions(pch_dir="/tmp/pch"),
        #     marks=pytest.mark.skipif(
        #         (_get_nvrtc_version_for_tests() or 0) < 12800,
        #         reason="PCH requires NVRTC >= 12.8",
        #     ),
        # ),
        pytest.param(
            ProgramOptions(pch_verbose=True),
            marks=pytest.mark.skipif(
                (_get_nvrtc_version_for_tests() or 0) < 12800,
                reason="PCH requires NVRTC >= 12.8",
            ),
        ),
        pytest.param(
            ProgramOptions(pch_messages=False),
            marks=pytest.mark.skipif(
                (_get_nvrtc_version_for_tests() or 0) < 12800,
                reason="PCH requires NVRTC >= 12.8",
            ),
        ),
        pytest.param(
            ProgramOptions(instantiate_templates_in_pch=True),
            marks=pytest.mark.skipif(
                (_get_nvrtc_version_for_tests() or 0) < 12800,
                reason="PCH requires NVRTC >= 12.8",
            ),
        ),
    ],
)
def test_cpp_program_with_various_options(init_cuda, options):
    code = 'extern "C" __global__ void my_kernel() {}'
    program = Program(code, "c++", options)
    assert program.backend == "NVRTC"
    assert isinstance(program.backend, CompilerBackendType)
    program.compile("ptx")
    program.close()


@pytest.mark.skipif(
    (_get_nvrtc_version_for_tests() or 0) < 13000,
    reason="buggy with NVRTC < 13.0 (File 'trace.json.json' could not be opened)",
)
def test_cpp_program_with_trace_option(init_cuda, tmp_path):
    code = 'extern "C" __global__ void my_kernel() {}'
    path = tmp_path / "trace"
    options = ProgramOptions(fdevice_time_trace=path)
    program = Program(code, "c++", options)
    assert program.backend == "NVRTC"
    program.compile("ptx")
    program.close()


@pytest.mark.skipif((_get_nvrtc_version_for_tests() or 0) < 12800, reason="PCH requires NVRTC >= 12.8")
def test_cpp_program_with_pch_options(init_cuda, tmp_path):
    code = 'extern "C" __global__ void my_kernel() {}'

    path = str(tmp_path / "test.pch")

    for opts in ({"create_pch": path}, {"use_pch": path}):
        options = ProgramOptions(**opts)
        program = Program(code, "c++", options)
        assert program.backend == "NVRTC"
        program.compile("ptx")
        program.close()


@nvrtc_pch_available
def test_cpp_program_pch_auto_creates(init_cuda, tmp_path):
    code = 'extern "C" __global__ void my_kernel() {}'
    pch_path = str(tmp_path / "test.pch")
    program = Program(code, "c++", ProgramOptions(create_pch=pch_path))
    assert program.pch_status is None  # not compiled yet
    program.compile("ptx")
    assert program.pch_status in ("created", "not_attempted", "failed")
    assert isinstance(program.pch_status, PCHStatusType)
    program.close()


def test_cpp_program_pch_status_none_without_pch(init_cuda):
    code = 'extern "C" __global__ void my_kernel() {}'
    program = Program(code, "c++")
    program.compile("ptx")
    assert program.pch_status is None
    program.close()


options = [
    ProgramOptions(max_register_count=32),
    ProgramOptions(debug=True),
    ProgramOptions(lineinfo=True),
    ProgramOptions(ftz=True),
    ProgramOptions(prec_div=True),
    ProgramOptions(prec_sqrt=True),
    ProgramOptions(fma=True),
    # Plumb-through; no-op at link time. See #1287.
    ProgramOptions(debug=True, numba_debug=True),
]
if not is_culink_backend:
    options += [
        ProgramOptions(time=True),
        ProgramOptions(split_compile=True),
    ]


@pytest.mark.parametrize("options", options)
def test_ptx_program_with_various_options(init_cuda, ptx_code_object, options):
    program = Program(ptx_code_object.code.decode(), "ptx", options=options)
    assert program.backend == ("driver" if is_culink_backend else "nvJitLink")
    program.compile("cubin")
    program.close()


def test_program_init_valid_code_type():
    code = 'extern "C" __global__ void my_kernel() {}'
    program = Program(code, "c++")
    assert program.backend == "NVRTC"
    assert program.handle is not None


def test_program_init_invalid_code_type():
    code = "goto 100"
    with pytest.raises(
        RuntimeError, match=r"^Unsupported code_type='fortran' \(supported_code_types=\('c\+\+', 'ptx', 'nvvm'\)\)$"
    ):
        Program(code, "FORTRAN")


def test_program_init_invalid_code_format():
    code = 12345
    with pytest.raises(TypeError):
        Program(code, "c++")


# arch is passed explicitly so the current device is not queried.
@pytest.mark.agent_authored(model="claude-opus-5")
@pytest.mark.parametrize("name", [None, "my_program"])
def test_program_options_name_accepts_none(name):
    options = ProgramOptions(name=name, arch="sm_90")
    expected = "default_program" if name is None else name
    assert options.name == expected
    assert options._name == expected.encode()


# This is tested against the current device's arch
def test_program_compile_valid_target_type(init_cuda):
    code = 'extern "C" __global__ void my_kernel() {}'
    program = Program(code, "c++", options={"name": "42"})

    with warnings.catch_warnings(record=True) as w:
        warnings.simplefilter("always")
        ptx_object_code = program.compile("ptx")
        assert isinstance(ptx_object_code, ObjectCode)
        assert ptx_object_code.name == "42"
        if any("The CUDA driver version is older than the backend version" in str(warning.message) for warning in w):
            pytest.skip("PTX version too new for current driver")
        ptx_kernel = ptx_object_code.get_kernel("my_kernel")
        assert isinstance(ptx_kernel, Kernel)

    program = Program(ptx_object_code.code.decode(), "ptx", options={"name": "24"})
    cubin_object_code = program.compile("cubin")
    assert isinstance(cubin_object_code, ObjectCode)
    assert cubin_object_code.name == "24"
    cubin_kernel = cubin_object_code.get_kernel("my_kernel")
    assert isinstance(cubin_kernel, Kernel)


def test_program_compile_invalid_target_type():
    code = 'extern "C" __global__ void my_kernel() {}'
    program = Program(code, "c++")
    with pytest.raises(ValueError):
        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++")
    assert program.backend == "NVRTC"


def test_program_handle_property():
    code = 'extern "C" __global__ void my_kernel() {}'
    program = Program(code, "c++")
    assert program.handle is not None


def test_program_close():
    code = 'extern "C" __global__ void my_kernel() {}'
    program = Program(code, "c++")
    program.close()
    # close() is idempotent
    program.close()


@nvvm_available
def test_nvvm_deferred_import():
    """Test that our deferred NVVM import works correctly"""
    from cuda.core._program import _get_nvvm_module

    nvvm = _get_nvvm_module()
    assert nvvm is not None


@nvvm_available
def test_nvvm_program_creation_compilation(nvvm_ir):
    """Test basic NVVM program creation"""
    program = Program(nvvm_ir, "nvvm")
    assert program.backend == "NVVM"
    assert program.handle is not None
    obj = program.compile("ptx")
    try:
        ker = obj.get_kernel("simple")
    except CUDAError as e:
        if re.search(r"CUDA_ERROR_UNSUPPORTED_PTX_VERSION", str(e)):
            pytest.xfail("PTX version not supported by current CUDA Driver")
        raise
    program.close()


@nvvm_available
def test_nvvm_compile_invalid_target(nvvm_ir):
    """Test that NVVM programs reject invalid compilation targets"""
    program = Program(nvvm_ir, "nvvm")
    with pytest.raises(ValueError, match='Unsupported target_type="cubin" for NVVM'):
        program.compile("cubin")
    program.close()


@nvvm_available
def test_nvvm_accepts_bytearray_input(nvvm_ir):
    """Program(..., 'nvvm') must accept bytearray input.

    Regression for a bug where the NVVM init branch retained the coerced
    ``self._code`` as bytes but still cast the original ``code`` object to
    ``<bytes>`` for the C pointer -- tripping a runtime type error for
    bytearray inputs before nvvmAddModuleToProgram was called.
    """
    program = Program(bytearray(nvvm_ir, "utf-8"), "nvvm")
    try:
        assert program.backend == "NVVM"
        assert program.handle is not None
    finally:
        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(
    "options",
    [
        ProgramOptions(name="test1", arch="sm_90", device_code_optimize=False),
        ProgramOptions(name="test2", arch="sm_100", device_code_optimize=False),
        ProgramOptions(name="test3", arch="sm_100", link_time_optimization=True),
        ProgramOptions(
            name="test4",
            arch="sm_90",
            ftz=True,
            prec_sqrt=False,
            prec_div=False,
            fma=True,
            device_code_optimize=True,
            link_time_optimization=True,
        ),
        pytest.param(
            ProgramOptions(name="test_sm110_1", arch="sm_110", device_code_optimize=False),
            marks=[
                has_nvvm_option_checker,
                pytest.mark.skipif(
                    _has_check_nvvm_compiler_options() and not _check_nvvm_arch("compute_110"),
                    reason="Compute capability 110 not supported by installed libNVVM",
                ),
            ],
        ),
        pytest.param(
            ProgramOptions(
                name="test_sm110_2",
                arch="sm_110",
                ftz=True,
                prec_sqrt=False,
                prec_div=False,
                fma=True,
                device_code_optimize=True,
            ),
            marks=[
                has_nvvm_option_checker,
                pytest.mark.skipif(
                    _has_check_nvvm_compiler_options() and not _check_nvvm_arch("compute_110"),
                    reason="Compute capability 110 not supported by installed libNVVM",
                ),
            ],
        ),
        pytest.param(
            ProgramOptions(name="test_sm110_3", arch="sm_110", link_time_optimization=True),
            marks=[
                has_nvvm_option_checker,
                pytest.mark.skipif(
                    _has_check_nvvm_compiler_options() and not _check_nvvm_arch("compute_110"),
                    reason="Compute capability 110 not supported by installed libNVVM",
                ),
            ],
        ),
    ],
)
def test_nvvm_program_options(init_cuda, nvvm_ir, options, target_type):
    """Test NVVM programs with different options and target types (ptx/ltoir)"""
    program = Program(nvvm_ir, "nvvm", options)
    assert program.backend == "NVVM"

    result = program.compile(target_type)
    assert isinstance(result, ObjectCode)
    assert result.name == options.name

    code_content = result.code
    assert len(code_content) > 0

    if target_type == "ptx":
        ptx_text = code_content.decode() if isinstance(code_content, bytes) else str(code_content)
        assert ".visible .entry simple(" in ptx_text

    program.close()


@nvvm_available
def test_nvvm_program_with_single_extra_source(nvvm_ir):
    """Test NVVM program with a single extra source"""
    from cuda.core._program import _get_nvvm_module

    nvvm = _get_nvvm_module()
    major, minor, debug_major, debug_minor = nvvm.ir_version()
    helper_nvvmir = f"""target triple = "nvptx64-unknown-cuda"
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"

define i32 @helper_add(i32 %x) {{
entry:
  %result = add i32 %x, 1
  ret i32 %result
}}

!nvvmir.version = !{{!0}}
!0 = !{{i32 {major}, i32 {minor}, i32 {debug_major}, i32 {debug_minor}}}
"""

    options = ProgramOptions(
        name="multi_module_test",
        extra_sources=[
            ("helper", helper_nvvmir),
        ],
    )
    program = Program(nvvm_ir, "nvvm", options)

    assert program.backend == "NVVM"

    ptx_code = program.compile("ptx")
    assert isinstance(ptx_code, ObjectCode)
    assert ptx_code.name == "multi_module_test"

    program.close()


@nvvm_available
def test_nvvm_program_with_multiple_extra_sources():
    """Test NVVM program with multiple extra sources"""
    from cuda.core._program import _get_nvvm_module

    nvvm = _get_nvvm_module()
    major, minor, debug_major, debug_minor = nvvm.ir_version()

    main_nvvm_ir = f"""target triple = "nvptx64-unknown-cuda"
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"

declare i32 @helper_add(i32) nounwind readnone
declare i32 @helper_mul(i32) nounwind readnone

define void @main_kernel(i32* %data) {{
entry:
  %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  %ptr = getelementptr inbounds i32, i32* %data, i32 %tid
  %val = load i32, i32* %ptr, align 4

  %val1 = call i32 @helper_add(i32 %val)
  %val2 = call i32 @helper_mul(i32 %val1)

  store i32 %val2, i32* %ptr, align 4
  ret void
}}

declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() nounwind readnone

!nvvm.annotations = !{{!0}}
!0 = !{{void (i32*)* @main_kernel, !"kernel", i32 1}}

!nvvmir.version = !{{!1}}
!1 = !{{i32 {major}, i32 {minor}, i32 {debug_major}, i32 {debug_minor}}}
"""

    helper1_ir = f"""target triple = "nvptx64-unknown-cuda"
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"

define i32 @helper_add(i32 %x) nounwind readnone {{
entry:
  %result = add i32 %x, 1
  ret i32 %result
}}

!nvvmir.version = !{{!0}}
!0 = !{{i32 {major}, i32 {minor}, i32 {debug_major}, i32 {debug_minor}}}
"""

    helper2_ir = f"""target triple = "nvptx64-unknown-cuda"
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"

define i32 @helper_mul(i32 %x) nounwind readnone {{
entry:
  %result = mul i32 %x, 2
  ret i32 %result
}}

!nvvmir.version = !{{!0}}
!0 = !{{i32 {major}, i32 {minor}, i32 {debug_major}, i32 {debug_minor}}}
"""

    options = ProgramOptions(
        name="nvvm_multi_helper_test",
        extra_sources=[
            ("helper1", helper1_ir),
            ("helper2", helper2_ir),
        ],
    )
    program = Program(main_nvvm_ir, "nvvm", options)

    assert program.backend == "NVVM"
    ptx_code = program.compile("ptx")
    assert isinstance(ptx_code, ObjectCode)
    assert ptx_code.name == "nvvm_multi_helper_test"

    ltoir_code = program.compile("ltoir")
    assert isinstance(ltoir_code, ObjectCode)
    assert ltoir_code.name == "nvvm_multi_helper_test"

    program.close()


@nvvm_available
def test_bitcode_format(minimal_nvvmir):
    from contextlib import ExitStack, closing

    if len(minimal_nvvmir) < 4:
        pytest.skip("Bitcode file is not valid or empty")

    options = ProgramOptions(name="minimal_nvvmir_bitcode_test", arch="sm_90")

    with ExitStack() as stack:
        program = stack.enter_context(closing(Program(minimal_nvvmir, "nvvm", options)))
        assert program.backend == "NVVM"
        ptx_result = program.compile("ptx")
        assert isinstance(ptx_result, ObjectCode)
        assert ptx_result.name == "minimal_nvvmir_bitcode_test"
        assert len(ptx_result.code) > 0

        program_lto = stack.enter_context(closing(Program(minimal_nvvmir, "nvvm", options)))
        ltoir_result = program_lto.compile("ltoir")
        assert isinstance(ltoir_result, ObjectCode)
        assert len(ltoir_result.code) > 0


def test_cpp_program_with_extra_sources():
    # negative test with NVRTC with multiple sources
    code = 'extern "C" __global__ void my_kernel(){}'
    helper = 'extern "C" __global__ void helper(){}'
    options = ProgramOptions(extra_sources=[("helper", helper)])
    with pytest.raises(ValueError, match="extra_sources is not supported by the NVRTC backend"):
        Program(code, "c++", options)


def test_program_options_as_bytes_nvrtc():
    """Test ProgramOptions.as_bytes() for NVRTC backend"""
    options = ProgramOptions(arch="sm_80", debug=True, lineinfo=True, ftz=True)
    nvrtc_options = options.as_bytes(CompilerBackendType.NVRTC)
    assert isinstance(nvrtc_options, list)
    assert all(isinstance(opt, bytes) for opt in nvrtc_options)
    options_str = [opt.decode() for opt in nvrtc_options]
    assert "-arch=sm_80" in options_str
    assert "--device-debug" in options_str
    assert "--generate-line-info" in options_str
    assert "--ftz=true" in options_str


@nvvm_available
def test_program_options_as_bytes_nvvm():
    """Test ProgramOptions.as_bytes() for NVVM backend"""
    options = ProgramOptions(arch="sm_80", debug=True, ftz=True, device_code_optimize=True)
    nvvm_options = options.as_bytes("nvvm")
    assert isinstance(nvvm_options, list)
    assert all(isinstance(opt, bytes) for opt in nvvm_options)
    options_str = [opt.decode() for opt in nvvm_options]
    assert "-arch=compute_80" in options_str
    assert "-g" in options_str
    assert "-ftz=1" in options_str
    assert "-opt=3" in options_str


def test_program_options_as_bytes_invalid_backend():
    """Test ProgramOptions.as_bytes() with invalid backend"""
    options = ProgramOptions(arch="sm_80")
    with pytest.raises(ValueError, match="Unknown backend 'invalid'"):
        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
    (see #1287)."""
    options = ProgramOptions(arch="sm_80", debug=True, numba_debug=True)
    nvvm_bytes = options.as_bytes("nvvm")
    assert b"--numba-debug" in nvvm_bytes
    assert b"-g" in nvvm_bytes


@nvvm_available
@pytest.mark.skipif(
    not _check_nvvm_supports_numba_debug(),
    reason="installed libNVVM does not recognize --numba-debug (needs CTK 13.2+)",
)
def test_nvvm_program_numba_debug(init_cuda, nvvm_ir):
    options = ProgramOptions(arch="sm_80", debug=True, numba_debug=True)
    program = Program(nvvm_ir, "nvvm", options)
    try:
        assert program.backend == "NVVM"
        result = program.compile("ptx")
        assert isinstance(result, ObjectCode)
        assert len(result.code) > 0
    finally:
        program.close()


def test_program_options_repr():
    """ProgramOptions.__repr__ returns a human-readable string."""
    opts = ProgramOptions(name="mykernel", arch="sm_80")
    r = repr(opts)
    assert "ProgramOptions" in r
    assert "mykernel" in r
    assert "sm_80" in r


def test_program_options_bad_define_macro_short_tuple():
    """define_macro with a 1-element tuple raises RuntimeError."""
    opts = ProgramOptions(name="test", arch="sm_80", define_macro=("ONLY_NAME",))
    with pytest.raises(RuntimeError, match="Expected define_macro tuple"):
        opts.as_bytes("nvrtc")


def test_program_options_bad_define_macro_non_str_value():
    """define_macro tuple with a non-string value raises RuntimeError."""
    opts = ProgramOptions(name="test", arch="sm_80", define_macro=("MY_MACRO", 99))
    with pytest.raises(RuntimeError, match="Expected define_macro tuple"):
        opts.as_bytes("nvrtc")


def test_program_options_bad_define_macro_list_non_str():
    """define_macro list containing a non-str/non-tuple item raises RuntimeError."""
    opts = ProgramOptions(name="test", arch="sm_80", define_macro=[42])
    with pytest.raises(RuntimeError, match="Expected define_macro"):
        opts.as_bytes("nvrtc")


def test_program_options_bad_define_macro_list_bad_tuple():
    """define_macro list with a malformed tuple inside raises RuntimeError."""
    opts = ProgramOptions(name="test", arch="sm_80", define_macro=[("ONLY_NAME",)])
    with pytest.raises(RuntimeError, match="Expected define_macro"):
        opts.as_bytes("nvrtc")


def test_ptx_program_extra_sources_unsupported(ptx_code_object):
    """PTX backend raises ValueError when extra_sources is specified."""
    options = ProgramOptions(extra_sources=[("module1", b"data")])
    with pytest.raises(ValueError, match="extra_sources is not supported by the PTX backend"):
        Program(ptx_code_object.code.decode(), "ptx", options)


def test_ptx_program_handle_is_linker_handle(init_cuda, ptx_code_object):
    """Program.handle for the PTX backend delegates to the linker handle."""
    program = Program(ptx_code_object.code.decode(), "ptx")
    handle = program.handle
    assert handle is not None
    assert int(handle) != 0
    program.close()


@nvvm_available
def test_nvvm_program_wrong_code_type():
    """NVVM backend raises TypeError when code is not str/bytes/bytearray."""
    with pytest.raises(TypeError, match="NVVM IR code must be provided as str, bytes, or bytearray"):
        Program(42, "nvvm")


def test_extra_sources_not_sequence():
    """extra_sources must be a sequence; non-sequence raises TypeError."""
    with pytest.raises(TypeError, match="extra_sources must be a sequence of 2-tuples"):
        ProgramOptions(name="test", arch="sm_80", extra_sources=42)


def test_extra_sources_bad_module_not_tuple():
    """extra_sources items must be 2-tuples; non-tuple item raises TypeError."""
    with pytest.raises(TypeError, match="Each extra module must be a 2-tuple"):
        ProgramOptions(name="test", arch="sm_80", extra_sources=["not_a_tuple"])


def test_extra_sources_bad_module_name_not_str():
    """extra_sources module name must be a string; non-str raises TypeError."""
    with pytest.raises(TypeError, match="Module name at index 0 must be a string"):
        ProgramOptions(name="test", arch="sm_80", extra_sources=[(42, b"source")])


def test_extra_sources_bad_module_source_wrong_type():
    """extra_sources module source must be str/bytes/bytearray."""
    with pytest.raises(TypeError, match="Module source at index 0 must be str"):
        ProgramOptions(name="test", arch="sm_80", extra_sources=[("mod", 42)])


def test_extra_sources_empty_source():
    """extra_sources module source cannot be empty bytes."""
    with pytest.raises(ValueError, match="Module source for 'mod'.*cannot be empty"):
        ProgramOptions(name="test", arch="sm_80", extra_sources=[("mod", b"")])


@pytest.mark.parametrize(
    ("extra_sources", "expected"),
    [
        (None, None),
        ([("mod_s", "kernel-as-string")], [(b"mod_s", b"kernel-as-string")]),
        (
            [("mod_ba", bytearray(b"\x00\x01module-as-bytearray"))],
            [(b"mod_ba", b"\x00\x01module-as-bytearray")],
        ),
        ([("mod_b", b"\x00\x01module-as-bytes")], [(b"mod_b", b"\x00\x01module-as-bytes")]),
    ],
    ids=["none", "str", "bytearray", "bytes"],
)
def test_prepare_extra_sources_bytes(extra_sources, expected):
    """_prepare_extra_sources_bytes converts each input type to (bytes, bytes) tuples (None passthrough)."""
    # arch is set to skip __post_init__'s Device() lookup, keeping this a pure unit test.
    opts = ProgramOptions(name="t", arch="sm_80", extra_sources=extra_sources)
    result = opts._prepare_extra_sources_bytes()
    assert result == expected
    # bytearray == bytes by content, so == alone misses type regressions.
    if result is not None:
        for name, source in result:
            assert isinstance(name, bytes), f"name should be bytes, got {type(name).__name__}"
            assert isinstance(source, bytes), f"source should be bytes, got {type(source).__name__}"


def test_find_libdevice_path_delegates_to_pathfinder(monkeypatch):
    """_find_libdevice_path calls cuda.pathfinder.find_bitcode_lib('device') and returns its result."""
    import cuda.pathfinder
    from cuda.core import _program

    captured = []
    sentinel = "/fake/path/libdevice.10.bc"

    def fake_find(name):
        captured.append(name)
        return sentinel

    monkeypatch.setattr(cuda.pathfinder, "find_bitcode_lib", fake_find)
    assert _program._find_libdevice_path() == sentinel
    assert captured == ["device"]


def test_nvrtc_compile_with_logs_capture(init_cuda):
    """Program.compile with logs= exercises the NVRTC program-log reading path."""
    import io

    # #warning generates a non-empty NVRTC program log, ensuring logsize > 1.
    code = '#warning "test log capture"\nextern "C" __global__ void my_kernel() {}'
    program = Program(code, "c++")
    logs = io.StringIO()
    result = program.compile("ptx", logs=logs)
    assert isinstance(result, ObjectCode)
    assert logs.getvalue(), "Expected non-empty compilation log from #warning directive"
    program.close()
