From da9c0771c7f3c254f7790d479aee1004ed8132fa Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Tue, 9 Dec 2025 20:56:13 +0000 Subject: [PATCH 1/6] Initial plan From 8e24725e067b0a6c0895118682c0be8a08d2a772 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Tue, 9 Dec 2025 21:01:56 +0000 Subject: [PATCH 2/6] Make ProgramOptions._as_bytes() public with type annotation and docstring Co-authored-by: leofang <5534781+leofang@users.noreply.github.com> --- cuda_core/cuda/core/experimental/_program.py | 23 ++++++++++- cuda_core/tests/test_program.py | 42 ++++++++++++++++++++ 2 files changed, 63 insertions(+), 2 deletions(-) diff --git a/cuda_core/cuda/core/experimental/_program.py b/cuda_core/cuda/core/experimental/_program.py index cdef7c3be6..15bc9c0917 100644 --- a/cuda_core/cuda/core/experimental/_program.py +++ b/cuda_core/cuda/core/experimental/_program.py @@ -422,7 +422,26 @@ def __post_init__(self): if self.numba_debug: self._formatted_options.append("--numba-debug") - def _as_bytes(self): + def as_bytes(self) -> list[bytes]: + """Convert the formatted program options to a list of byte strings. + + This method encodes each of the formatted options stored in this + `ProgramOptions` instance into byte strings, suitable for passing + to C libraries that perform NVRTC compilation. + + Returns + ------- + list[bytes] + A list of byte-encoded option strings. Each element represents + a single compilation option in the format expected by NVRTC. + + Examples + -------- + >>> options = ProgramOptions(arch="sm_80", debug=True) + >>> byte_options = options.as_bytes() + >>> print(byte_options) + [b'-arch=sm_80', b'--device-debug'] + """ # TODO: allow tuples once NVIDIA/cuda-python#72 is resolved return list(o.encode() for o in self._formatted_options) @@ -609,7 +628,7 @@ def compile(self, target_type, name_expressions=(), logs=None): nvrtc.nvrtcAddNameExpression(self._mnff.handle, n.encode()), handle=self._mnff.handle, ) - options = self._options._as_bytes() + options = self._options.as_bytes() handle_return( nvrtc.nvrtcCompileProgram(self._mnff.handle, len(options), options), handle=self._mnff.handle, diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index 8a6526fcc2..eb1c39025d 100644 --- a/cuda_core/tests/test_program.py +++ b/cuda_core/tests/test_program.py @@ -411,3 +411,45 @@ def test_nvvm_program_options(init_cuda, nvvm_ir, options): assert ".visible .entry simple(" in ptx_text program.close() + + +def test_program_options_as_bytes(): + """Test that ProgramOptions.as_bytes() returns correct byte strings""" + # Test with various options + options = ProgramOptions( + arch="sm_80", + debug=True, + lineinfo=True, + max_register_count=32, + ftz=True, + use_fast_math=True, + ) + + byte_options = options.as_bytes() + + # Verify the return type + assert isinstance(byte_options, list) + assert all(isinstance(opt, bytes) for opt in byte_options) + + # Verify specific options are present in byte format + assert b"-arch=sm_80" in byte_options + assert b"--device-debug" in byte_options + assert b"--generate-line-info" in byte_options + assert b"--maxrregcount=32" in byte_options + assert b"--ftz=true" in byte_options + assert b"--use_fast_math" in byte_options + + +def test_program_options_as_bytes_empty(): + """Test that ProgramOptions.as_bytes() works with minimal options""" + # Test with minimal options (only defaults) + options = ProgramOptions() + + byte_options = options.as_bytes() + + # Should at least have arch option (set to device default) + assert isinstance(byte_options, list) + assert len(byte_options) > 0 + assert all(isinstance(opt, bytes) for opt in byte_options) + # The arch option should be present (even if it's the default) + assert any(b"-arch=" in opt for opt in byte_options) From d2f144a4a5c3b6fe566a2d4a70c1aadf0bca4fec Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Tue, 9 Dec 2025 21:03:58 +0000 Subject: [PATCH 3/6] Address code review comment - improve test comment clarity Co-authored-by: leofang <5534781+leofang@users.noreply.github.com> --- cuda_core/tests/test_program.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index eb1c39025d..747b35e2b1 100644 --- a/cuda_core/tests/test_program.py +++ b/cuda_core/tests/test_program.py @@ -447,9 +447,9 @@ def test_program_options_as_bytes_empty(): byte_options = options.as_bytes() - # Should at least have arch option (set to device default) + # Should at least have arch option (automatically set based on Device if not provided) assert isinstance(byte_options, list) assert len(byte_options) > 0 assert all(isinstance(opt, bytes) for opt in byte_options) - # The arch option should be present (even if it's the default) + # The arch option should be present (automatically determined from current device) assert any(b"-arch=" in opt for opt in byte_options) From 523a246eb7bd4095b58fc9041b8b53d4c8ef4899 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Tue, 9 Dec 2025 21:33:09 +0000 Subject: [PATCH 4/6] Update docstring to refer to underlying compiler library instead of NVRTC Co-authored-by: leofang <5534781+leofang@users.noreply.github.com> --- cuda_core/cuda/core/experimental/_program.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cuda_core/cuda/core/experimental/_program.py b/cuda_core/cuda/core/experimental/_program.py index 15bc9c0917..5a25be821a 100644 --- a/cuda_core/cuda/core/experimental/_program.py +++ b/cuda_core/cuda/core/experimental/_program.py @@ -427,13 +427,13 @@ def as_bytes(self) -> list[bytes]: This method encodes each of the formatted options stored in this `ProgramOptions` instance into byte strings, suitable for passing - to C libraries that perform NVRTC compilation. + to C libraries that calls the underlying compiler library. Returns ------- list[bytes] A list of byte-encoded option strings. Each element represents - a single compilation option in the format expected by NVRTC. + a single compilation option in the format expected by the underlying compiler library. Examples -------- From fa2ce29c170c986bd7a9fe8267e2fe953f12d07d Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Tue, 9 Dec 2025 21:48:13 +0000 Subject: [PATCH 5/6] Add backend parameter to as_bytes() method with support for NVRTC, NVVM, and nvJitLink Co-authored-by: leofang <5534781+leofang@users.noreply.github.com> --- cuda_core/cuda/core/experimental/_program.py | 187 ++++++++++++++++++- cuda_core/tests/test_program.py | 64 +++++++ 2 files changed, 244 insertions(+), 7 deletions(-) diff --git a/cuda_core/cuda/core/experimental/_program.py b/cuda_core/cuda/core/experimental/_program.py index 5a25be821a..690ae16188 100644 --- a/cuda_core/cuda/core/experimental/_program.py +++ b/cuda_core/cuda/core/experimental/_program.py @@ -126,6 +126,116 @@ def _process_define_macro(formatted_options, macro): raise RuntimeError(f"Expected define_macro {union_type}, list[{union_type}], got {macro}") +def _format_options_for_backend(options_dict: dict, backend: str) -> list[str]: + """Format compilation options for a specific backend. + + This helper function converts a dictionary of option names and values into + properly formatted string options for the specified backend. Different backends + (NVRTC, NVVM, nvJitLink) use slightly different option naming conventions and + value formats. + + Parameters + ---------- + options_dict : dict + Dictionary mapping option names to their values. The keys should be + generic option names (e.g., "arch", "debug", "ftz"). + backend : str + The backend to format options for. Must be one of "NVRTC", "NVVM", or "nvJitLink". + + Returns + ------- + list[str] + List of formatted option strings suitable for the specified backend. + + Raises + ------ + ValueError + If an unsupported backend is specified. + + Notes + ----- + - NVRTC uses `--` prefix and "true"/"false" for booleans + - NVVM uses `-` prefix and "1"/"0" for booleans + - nvJitLink uses `-` prefix and "true"/"false" for booleans + """ + if backend not in ("NVRTC", "NVVM", "nvJitLink"): + raise ValueError(f"Unsupported backend '{backend}'. Must be one of: NVRTC, NVVM, nvJitLink") + + formatted = [] + + for key, value in options_dict.items(): + if value is None: + continue + + if backend == "NVRTC": + # NVRTC uses -- prefix + if key == "arch": + formatted.append(f"-arch={value}") + elif key == "debug" and value: + formatted.append("--device-debug") + elif key == "lineinfo" and value: + formatted.append("--generate-line-info") + elif key == "max_register_count": + formatted.append(f"--maxrregcount={value}") + elif key in ("ftz", "prec_sqrt", "prec_div"): + bool_val = "true" if value else "false" + # NVRTC uses hyphens in option names + option_name = key.replace("_", "-") + formatted.append(f"--{option_name}={bool_val}") + elif key == "fma": + bool_val = "true" if value else "false" + formatted.append(f"--fmad={bool_val}") + elif key == "device_code_optimize" and value: + formatted.append("--dopt=on") + elif key == "use_fast_math" and value: + formatted.append("--use_fast_math") + elif key == "link_time_optimization" and value: + formatted.append("--dlink-time-opt") + # Add more NVRTC-specific options as needed + + elif backend == "NVVM": + # NVVM uses - prefix and 1/0 for booleans + if key == "arch": + # NVVM uses compute_ instead of sm_ + arch_val = value + if arch_val.startswith("sm_"): + arch_val = f"compute_{arch_val[3:]}" + formatted.append(f"-arch={arch_val}") + elif key == "debug" and value: + formatted.append("-g") + elif key == "device_code_optimize": + if value is False: + formatted.append("-opt=0") + elif value is True: + formatted.append("-opt=3") + elif key in ("ftz", "prec_sqrt", "prec_div", "fma"): + bool_val = "1" if value else "0" + # NVVM uses hyphens in option names + option_name = key.replace("_", "-") + formatted.append(f"-{option_name}={bool_val}") + + elif backend == "nvJitLink": + # nvJitLink uses - prefix and true/false for booleans + if key == "arch": + formatted.append(f"-arch={value}") + elif key == "debug" and value: + formatted.append("-g") + elif key == "lineinfo" and value: + formatted.append("-lineinfo") + elif key == "max_register_count": + formatted.append(f"-maxrregcount={value}") + elif key in ("ftz", "prec_sqrt", "prec_div", "fma"): + bool_val = "true" if value else "false" + # nvJitLink uses hyphens in option names + option_name = key.replace("_", "-") + formatted.append(f"-{option_name}={bool_val}") + elif key == "link_time_optimization" and value: + formatted.append("-lto") + # Add more nvJitLink-specific options as needed + + return formatted + + @dataclass class ProgramOptions: """Customizable options for configuring `Program`. @@ -422,28 +532,91 @@ def __post_init__(self): if self.numba_debug: self._formatted_options.append("--numba-debug") - def as_bytes(self) -> list[bytes]: + def as_bytes(self, backend: str = "NVRTC") -> list[bytes]: """Convert the formatted program options to a list of byte strings. - This method encodes each of the formatted options stored in this - `ProgramOptions` instance into byte strings, suitable for passing + This method encodes the options stored in this `ProgramOptions` instance + into byte strings formatted for the specified backend, suitable for passing to C libraries that calls the underlying compiler library. + Parameters + ---------- + backend : str, optional + The compiler backend to format options for. Must be one of: + + - "NVRTC" (default): NVIDIA NVRTC compiler, supports all ProgramOptions + - "NVVM": NVIDIA NVVM compiler, supports a subset of options + - "nvJitLink": NVIDIA nvJitLink linker, supports a subset of options + + Different backends use different option naming conventions and support + different option subsets. This method will format and filter options + appropriately for the chosen backend. + Returns ------- list[bytes] A list of byte-encoded option strings. Each element represents a single compilation option in the format expected by the underlying compiler library. + Raises + ------ + ValueError + If an unsupported backend is specified. + Examples -------- >>> options = ProgramOptions(arch="sm_80", debug=True) - >>> byte_options = options.as_bytes() - >>> print(byte_options) + >>> # Get options for NVRTC (default) + >>> nvrtc_options = options.as_bytes() + >>> print(nvrtc_options) [b'-arch=sm_80', b'--device-debug'] + >>> + >>> # Get options for NVVM + >>> nvvm_options = options.as_bytes("NVVM") + >>> print(nvvm_options) + [b'-arch=compute_80', b'-g'] + >>> + >>> # Get options for nvJitLink + >>> nvjitlink_options = options.as_bytes("nvJitLink") + >>> print(nvjitlink_options) + [b'-arch=sm_80', b'-g'] """ - # TODO: allow tuples once NVIDIA/cuda-python#72 is resolved - return list(o.encode() for o in self._formatted_options) + if backend == "NVRTC": + # For NVRTC, use the pre-formatted options (backward compatible) + # TODO: allow tuples once NVIDIA/cuda-python#72 is resolved + return list(o.encode() for o in self._formatted_options) + + elif backend in ("NVVM", "nvJitLink"): + # For NVVM and nvJitLink, extract common options and format appropriately + options_dict = {} + + # Common options supported by multiple backends + if self.arch is not None: + options_dict["arch"] = self.arch + if self.debug is not None: + options_dict["debug"] = self.debug + if self.lineinfo is not None and backend == "nvJitLink": + options_dict["lineinfo"] = self.lineinfo + if self.max_register_count is not None: + options_dict["max_register_count"] = self.max_register_count + if self.ftz is not None: + options_dict["ftz"] = self.ftz + if self.prec_sqrt is not None: + options_dict["prec_sqrt"] = self.prec_sqrt + if self.prec_div is not None: + options_dict["prec_div"] = self.prec_div + if self.fma is not None: + options_dict["fma"] = self.fma + if self.device_code_optimize is not None and backend == "NVVM": + options_dict["device_code_optimize"] = self.device_code_optimize + if self.link_time_optimization is not None and backend == "nvJitLink": + options_dict["link_time_optimization"] = self.link_time_optimization + + formatted_options = _format_options_for_backend(options_dict, backend) + return list(o.encode() for o in formatted_options) + + else: + raise ValueError(f"Unsupported backend '{backend}'. Must be one of: NVRTC, NVVM, nvJitLink") def __repr__(self): # __TODO__ improve this diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index 747b35e2b1..1388d472d3 100644 --- a/cuda_core/tests/test_program.py +++ b/cuda_core/tests/test_program.py @@ -453,3 +453,67 @@ def test_program_options_as_bytes_empty(): assert all(isinstance(opt, bytes) for opt in byte_options) # The arch option should be present (automatically determined from current device) assert any(b"-arch=" in opt for opt in byte_options) + + +def test_program_options_as_bytes_nvvm_backend(): + """Test that ProgramOptions.as_bytes() formats options correctly for NVVM backend""" + options = ProgramOptions( + arch="sm_80", + debug=True, + ftz=True, + prec_sqrt=False, + prec_div=True, + fma=False, + device_code_optimize=True, + ) + + byte_options = options.as_bytes("NVVM") + + # Verify the return type + assert isinstance(byte_options, list) + assert all(isinstance(opt, bytes) for opt in byte_options) + + # NVVM uses compute_ instead of sm_ and 1/0 for booleans, with hyphens in option names + assert b"-arch=compute_80" in byte_options + assert b"-g" in byte_options + assert b"-ftz=1" in byte_options + assert b"-prec-sqrt=0" in byte_options + assert b"-prec-div=1" in byte_options + assert b"-fma=0" in byte_options + assert b"-opt=3" in byte_options + + +def test_program_options_as_bytes_nvjitlink_backend(): + """Test that ProgramOptions.as_bytes() formats options correctly for nvJitLink backend""" + options = ProgramOptions( + arch="sm_80", + debug=True, + lineinfo=True, + max_register_count=32, + ftz=False, + prec_sqrt=True, + link_time_optimization=True, + ) + + byte_options = options.as_bytes("nvJitLink") + + # Verify the return type + assert isinstance(byte_options, list) + assert all(isinstance(opt, bytes) for opt in byte_options) + + # nvJitLink uses - prefix and true/false for booleans, with hyphens in option names + assert b"-arch=sm_80" in byte_options + assert b"-g" in byte_options + assert b"-lineinfo" in byte_options + assert b"-maxrregcount=32" in byte_options + assert b"-ftz=false" in byte_options + assert b"-prec-sqrt=true" in byte_options + assert b"-lto" in byte_options + + +def test_program_options_as_bytes_invalid_backend(): + """Test that ProgramOptions.as_bytes() raises error for invalid backend""" + options = ProgramOptions() + + with pytest.raises(ValueError, match="Unsupported backend 'invalid'"): + options.as_bytes("invalid") From 377707d6093f29ef36dded0f3c181f101604fbfc Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Tue, 9 Dec 2025 21:51:11 +0000 Subject: [PATCH 6/6] Improve backend option filtering - remove backend checks from as_bytes, let helper handle filtering Co-authored-by: leofang <5534781+leofang@users.noreply.github.com> --- cuda_core/cuda/core/experimental/_program.py | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/cuda_core/cuda/core/experimental/_program.py b/cuda_core/cuda/core/experimental/_program.py index 690ae16188..96ce859948 100644 --- a/cuda_core/cuda/core/experimental/_program.py +++ b/cuda_core/cuda/core/experimental/_program.py @@ -204,6 +204,7 @@ def _format_options_for_backend(options_dict: dict, backend: str) -> list[str]: elif key == "debug" and value: formatted.append("-g") elif key == "device_code_optimize": + # NVVM explicitly handles both True and False if value is False: formatted.append("-opt=0") elif value is True: @@ -213,6 +214,7 @@ def _format_options_for_backend(options_dict: dict, backend: str) -> list[str]: # NVVM uses hyphens in option names option_name = key.replace("_", "-") formatted.append(f"-{option_name}={bool_val}") + # lineinfo and link_time_optimization are not supported by NVVM, skip them elif backend == "nvJitLink": # nvJitLink uses - prefix and true/false for booleans @@ -231,7 +233,7 @@ def _format_options_for_backend(options_dict: dict, backend: str) -> list[str]: formatted.append(f"-{option_name}={bool_val}") elif key == "link_time_optimization" and value: formatted.append("-lto") - # Add more nvJitLink-specific options as needed + # device_code_optimize is not supported by nvJitLink, skip it return formatted @@ -595,7 +597,7 @@ def as_bytes(self, backend: str = "NVRTC") -> list[bytes]: options_dict["arch"] = self.arch if self.debug is not None: options_dict["debug"] = self.debug - if self.lineinfo is not None and backend == "nvJitLink": + if self.lineinfo is not None: options_dict["lineinfo"] = self.lineinfo if self.max_register_count is not None: options_dict["max_register_count"] = self.max_register_count @@ -607,9 +609,9 @@ def as_bytes(self, backend: str = "NVRTC") -> list[bytes]: options_dict["prec_div"] = self.prec_div if self.fma is not None: options_dict["fma"] = self.fma - if self.device_code_optimize is not None and backend == "NVVM": + if self.device_code_optimize is not None: options_dict["device_code_optimize"] = self.device_code_optimize - if self.link_time_optimization is not None and backend == "nvJitLink": + if self.link_time_optimization is not None: options_dict["link_time_optimization"] = self.link_time_optimization formatted_options = _format_options_for_backend(options_dict, backend)