Skip to content
Closed
Show file tree
Hide file tree
Changes from 4 commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
23 changes: 21 additions & 2 deletions cuda_core/cuda/core/experimental/_program.py
Original file line number Diff line number Diff line change
Expand Up @@ -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]:

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@copilot it occurs to me that we should take an argument backend and transforms this accordingly. The reason is that which a subset of options is shared by all compiler libraries (nvrtc/nvjitlink/nvvm), the accepted names are slightly different and unfortunately there exists no unique option. So we need to handle this difference for users.

Please inspect the codebase and consolidate a treatment that we can reuse across the Program and Linker classes. The handling of the backend argument can be done in a standalone helper function instead of a class method, if deemed appropriate.

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@copilot keep #1341 in mind when exploring a plan.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Implemented in fa2ce29. The as_bytes() method now accepts a backend parameter ("NVRTC", "NVVM", or "nvJitLink") and formats options appropriately for each backend. Created a reusable helper function _format_options_for_backend() that handles the different option naming conventions and value formats across backends:

  • NVRTC: -- prefix, "true"/"false" booleans
  • NVVM: - prefix, "1"/"0" booleans, compute_ arch prefix
  • nvJitLink: - prefix, "true"/"false" booleans

The method defaults to "NVRTC" for backward compatibility. Added comprehensive tests for all backends.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm unable to access issue #1341 details from this environment. Could you provide the key considerations from that issue that should be incorporated into this implementation? If there are specific changes needed to align with #1341, I can make those adjustments.

"""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 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 the underlying compiler library.

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)

Expand Down Expand Up @@ -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,
Expand Down
42 changes: 42 additions & 0 deletions cuda_core/tests/test_program.py
Original file line number Diff line number Diff line change
Expand Up @@ -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 (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 (automatically determined from current device)
assert any(b"-arch=" in opt for opt in byte_options)