commit squash: add ProgramOptions to Program

Author: ksimpson-work

cuda_core/cuda/core/experimental/__init__.py

@@ -5,5 +5,5 @@
 from cuda.core.experimental._device import Device
 from cuda.core.experimental._event import EventOptions
 from cuda.core.experimental._launcher import LaunchConfig, launch
-from cuda.core.experimental._program import Program
+from cuda.core.experimental._program import Program, ProgramOptions
 from cuda.core.experimental._stream import Stream, StreamOptions

cuda_core/cuda/core/experimental/_program.py

@@ -211,9 +211,7 @@ def wait(self, event_or_stream: Union[Event, Stream]):
                 try:
                     stream = Stream._init(event_or_stream)
                 except Exception as e:
-                    raise ValueError(
-                        "only an Event, Stream, or object supporting __cuda_stream__ can be waited"
-                    ) from e
+                    raise ValueError("only an Event, Stream, or object supporting __cuda_stream__ can be waited") from e
             else:
                 stream = event_or_stream
             event = handle_return(cuda.cuEventCreate(cuda.CUevent_flags.CU_EVENT_DISABLE_TIMING))

cuda_core/cuda/core/experimental/_stream.py

@@ -87,6 +87,13 @@ def check_or_create_options(cls, options, options_description, *, keep_none=Fals
     return options
 
 
+def _handle_boolean_option(option: bool) -> str:
+    """
+    Convert a boolean option to a string representation.
+    """
+    return str(option).lower()
+
+
 def precondition(checker: Callable[..., None], what: str = "") -> Callable:
     """
     A decorator that adds checks to ensure any preconditions are met.

cuda_core/cuda/core/experimental/_utils.py

@@ -0,0 +1,12 @@
+# `cuda.core` Release notes
+
+Released on <TODO>, 2024
+
+## Hightlights
+- Add ProgramOptions to facilitate the passing of runtime compile options to [Program](#program)
+
+## Limitations
+-<TODO>
+
+## Breaking Changes
+- The `Program.Compile` method no longer accepts an options argument. Instead, you can optionally pass an instance of `ProgramOptions` to the constructor of `Program`.

cuda_core/docs/source/release/0.2.0-notes.md

@@ -6,7 +6,7 @@
 
 import cupy as cp
 
-from cuda.core.experimental import Device, LaunchConfig, Program, launch
+from cuda.core.experimental import Device, LaunchConfig, Program, ProgramOptions, launch
 
 # compute out = a * x + y
 code = """
@@ -29,13 +29,10 @@
 s = dev.create_stream()
 
 # prepare program
-prog = Program(code, code_type="c++")
+program_options = ProgramOptions(std="c++11", gpu_architecture="sm_" + "".join(f"{i}" for i in dev.compute_capability))
+prog = Program(code, code_type="c++", options=program_options)
 mod = prog.compile(
     "cubin",
-    options=(
-        "-std=c++11",
-        "-arch=sm_" + "".join(f"{i}" for i in dev.compute_capability),
-    ),
     logs=sys.stdout,
     name_expressions=("saxpy<float>", "saxpy<double>"),
 )

cuda_core/examples/saxpy.py

@@ -4,7 +4,7 @@
 
 import cupy as cp
 
-from cuda.core.experimental import Device, LaunchConfig, Program, launch
+from cuda.core.experimental import Device, LaunchConfig, Program, ProgramOptions, launch
 
 # compute c = a + b
 code = """
@@ -26,15 +26,9 @@
 s = dev.create_stream()
 
 # prepare program
-prog = Program(code, code_type="c++")
-mod = prog.compile(
-    "cubin",
-    options=(
-        "-std=c++17",
-        "-arch=sm_" + "".join(f"{i}" for i in dev.compute_capability),
-    ),
-    name_expressions=("vector_add<float>",),
-)
+program_options = ProgramOptions(std="c++17", gpu_architecture="sm_" + "".join(f"{i}" for i in dev.compute_capability))
+prog = Program(code, code_type="c++", options=program_options)
+mod = prog.compile("cubin", name_expressions=("vector_add<float>",))
 
 # run in single precision
 ker = mod.get_kernel("vector_add<float>")

cuda_core/examples/vector_add.py

@@ -8,8 +8,45 @@
 
 import pytest
 
-from cuda.core.experimental import Program
 from cuda.core.experimental._module import Kernel, ObjectCode
+from cuda.core.experimental._program import Program, ProgramOptions
+
+
+def test_program_with_various_options(init_cuda):
+    code = 'extern "C" __global__ void my_kernel() {}'
+
+    options_list = [
+        ProgramOptions(ptxas_options="-v"),
+        ProgramOptions(ptxas_options=["-v", "-O3"]),
+        ProgramOptions(device_optimize=True, device_debug=True),
+        ProgramOptions(relocatable_device_code=True, maxrregcount=32),
+        ProgramOptions(ftz=True, prec_sqrt=False, prec_div=False),
+        ProgramOptions(fmad=False, use_fast_math=True),
+        ProgramOptions(extra_device_vectorization=True),
+        ProgramOptions(dlink_time_opt=True, gen_opt_lto=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(pre_include="my_header.h", no_source_include=True),
+        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="5678"),
+        ProgramOptions(diag_warn="91011", brief_diagnostics=True),
+        ProgramOptions(time="compile_time.csv", split_compile=4),
+        ProgramOptions(fdevice_syntax_only=True, minimal=True),
+    ]
+
+    # TODO compile the program once the CI is set up
+    for options in options_list:
+        program = Program(code, "c++", options)
+        assert program.backend == "nvrtc"
+        program.close()
+        assert program.handle is None
 
 
 def test_program_init_valid_code_type():