Implement Kernel.num_arguments, and Kernel.arguments_info #612

oleksandr-pavlyk · 2025-05-06T13:25:16Z

Description

closes #568

This PR implements Kernel.num_arguments and Kernel.argument_info properties:

Kernel.num_arguments returns the number of arguments in the kernel instance
Kernel.argument_info returns a list of tuples (offset, size) which describes layout of the struct containing all kernel arguments,

Checklist

New or existing tests cover these changes.
The documentation is up to date with these changes.

copy-pr-bot · 2025-05-06T13:25:19Z

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

leofang

Thanks, Sasha! Looks great except for the docstrings, see my comment below.

We also need to add a feature entry to cuda_core/docs/source/release/0.3.0-notes.rst.

cuda_core/cuda/core/experimental/_module.py

…erties Also parametrize test to check arguments_info to check with all int, and all short arguments.

copy-pr-bot · 2025-05-08T13:58:20Z

Auto-sync is disabled for ready for review pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

oleksandr-pavlyk · 2025-05-08T15:21:29Z

/ok to test

rwgk

LGTM although I'd use a dataclass for better ergonomics.

cuda_core/cuda/core/experimental/_module.py

oleksandr-pavlyk · 2025-05-08T20:06:53Z

I fixed pre-commit, and added a line to "0.3.0-notes.rst".

rwgk

Two minor suggestions, please see here for background:

https://chatgpt.com/share/681d1a68-97d0-8008-bad5-40ae287f7d55

(When I wrote the first prompt I had "module scope" in mind, but I think class scope is better.)

cuda_core/cuda/core/experimental/_module.py

Used modern namedtuple instance constructor.

oleksandr-pavlyk · 2025-05-08T22:00:42Z

This perplexes me:

(dev-cuda-python) opavlyk@ee09c48-lcedt:~/repos/cuda-python/cuda_core$ ipython
Python 3.12.10 | packaged by conda-forge | (main, Apr 10 2025, 22:21:13) [GCC 13.3.0]
Type 'copyright', 'credits' or 'license' for more information
IPython 9.2.0 -- An enhanced Interactive Python. Type '?' for help.
Tip: Use `object?` to see the help on `object`, `object??` to view its source

In [1]: import cuda.core.experimental as cc

In [2]: o1 = cc.Program("__global__ void bar(double *p, int n, double x) { *p = n * x; }", code_type="c++").compile("cubin", name_expressions=("bar",))

In [3]: k1 = o1.get_kernel("bar")

In [4]: k1.arguments_info
Out[4]: []

In [5]: k1.num_arguments
Out[5]: 0

rwgk

This looks perfect to me. Caveat: I cannot explain the ipython observation.

leofang · 2025-05-09T02:52:28Z

This perplexes me

Let me look into it tmr

leofang · 2025-05-13T14:10:35Z

nit: as much as I love credit attribution (thanks), let's not do @someone in the commit messages. GitHub still handles them badly (after all these years) and I end up receiving a ton of notifications every time someone forks/rebases/pushes 😞

leofang · 2025-05-13T14:11:35Z

/ok to test 2ca9704

cuda_core/cuda/core/experimental/_module.py

This required moving ParamInfo definition from class scope to module scope, since referencing Kernel.ParamInfo from annotations of methods of the Kernel class results in error that Kernel class does not yet exist.

cuda_core/tests/test_module.py

leofang · 2025-05-14T01:25:38Z

/ok to test 55f6d31

oleksandr-pavlyk · 2025-05-14T02:46:53Z

Test that num_arguments should raise if CUDA has not been initialized failed:

_________________________ test_num_args_error_handling _________________________

deinit_cuda = None, cuda12_prerequisite_check = True

    @skipif_testing_with_compute_sanitizer
    def test_num_args_error_handling(deinit_cuda, cuda12_prerequisite_check):
        if not cuda12_prerequisite_check:
            pytest.skip("Test requires CUDA 12")
        src = "__global__ void foo(int a) { }"
        prog = Program(src, code_type="c++")
        mod = prog.compile(
            "cubin",
            name_expressions=("foo",),
        )
        krn = mod.get_kernel("foo")
>       with pytest.raises(CUDAError):
E       Failed: DID NOT RAISE <class 'cuda.core.experimental._utils.cuda_utils.CUDAError'>

tests/test_module.py:246: Failed

Did I misunderstand the intent of deinit_cuda fixture?

leofang · 2025-05-14T02:58:56Z

I feel deinit_cuda might be a confusing name. The purpose of it is to clean up the per-thread device instances + ensure no CUDA context is set to current after the test finishes. Did you mean to do this before the test starts?

oleksandr-pavlyk · 2025-05-14T03:02:54Z

I meant it to ensure that cuda is not initialized at the start of this test. Kind of like opposite of cuda_init fixture

leofang · 2025-05-14T03:20:05Z

Yes feel free to implement any fixture you need. No need to refactor the existing fixtures if you don't want to sweat in this PR -- but then we should create an issue to track it (to the very least, we should rename deinit_cuda to something like deinit_device_at_teardown).

To make sure we understand "initialized" in the same way -- You don't mean driver.cuInit(0), do you? Note that it is a once-per-process call, effectively as if std::call_once() is used (I dunno how it's implemented exactly in the driver), so any subsequent calls are cheap no-ops, and there's no way to undo cuInit(0).

Use in test_module.py::test_num_args_error_handling Add comments

oleksandr-pavlyk · 2025-05-14T15:12:27Z

/ok to test

…_error_handling

oleksandr-pavlyk · 2025-05-14T17:20:20Z

/ok to test

cuda_core/tests/test_module.py

1. Changed fixture to provide a function that empties the stack of contexts. The function has hidden max_iters bound. If exceeded, a RuntimeError is raised 2. Modified _device_unset_current utility function to return a boolean. True is returned is a context was popped, False if the stack was already empty.

oleksandr-pavlyk · 2025-05-14T18:27:35Z

/ok to test

rwgk

Just a couple cosmetic suggestions. Optional.

rwgk · 2025-05-14T23:00:40Z

cuda_core/cuda/core/experimental/_module.py

+            raise NotImplementedError("New backend is required")
+        arg_pos = 0
+        param_info_data = []
+        while True:


Maybe (sorry I overlooked this before):

for arg_pos in itertools.count():

Then you don't need arg_pos = 0 above and arg_pos = arg_pos + 1 below.

rwgk · 2025-05-14T23:02:32Z

cuda_core/tests/conftest.py

+            if _device_unset_current():
+                # context was popped, continue until stack is empty
+                continue
+            # no active context, we are ready
+            break


Maybe shorter (replace 5 lines with 2):

if not _device_unset_current(): break

github-actions · 2025-05-14T23:30:04Z

Doc Preview CI
Preview removed because the pull request was closed or merged.

Implement Kernel.num_arguments, and Kernel.arguments_info

e22b891

oleksandr-pavlyk requested review from leofang and ksimpson-work May 6, 2025 13:25

leofang reviewed May 6, 2025

View reviewed changes

cuda_core/cuda/core/experimental/_module.py Outdated Show resolved Hide resolved

cuda_core/cuda/core/experimental/_module.py Outdated Show resolved Hide resolved

leofang assigned oleksandr-pavlyk May 6, 2025

leofang added P1 Medium priority - Should do feature New feature or request cuda.core Everything related to the cuda.core module labels May 6, 2025

leofang added this to the cuda.core beta 4 milestone May 6, 2025

rwgk reviewed May 6, 2025

View reviewed changes

cuda_core/cuda/core/experimental/_module.py Outdated Show resolved Hide resolved

Factor out common logic between num_arguments and arguments_info prop…

6de62d0

…erties Also parametrize test to check arguments_info to check with all int, and all short arguments.

oleksandr-pavlyk marked this pull request as ready for review May 8, 2025 13:58

This comment has been minimized.

Sign in to view

rwgk previously approved these changes May 8, 2025

View reviewed changes

cuda_core/cuda/core/experimental/_module.py Outdated Show resolved Hide resolved

oleksandr-pavlyk added 2 commits May 8, 2025 14:49

Use namedtuple turn to pack offset/size for individual kernel argument

fb1344b

Add feature entry to 0.3.0 release notes

96f60d8

oleksandr-pavlyk dismissed rwgk’s stale review via 96f60d8 May 8, 2025 20:05

oleksandr-pavlyk requested review from rwgk and leofang May 8, 2025 20:07

rwgk previously approved these changes May 8, 2025

View reviewed changes

cuda_core/cuda/core/experimental/_module.py Outdated Show resolved Hide resolved

cuda_core/cuda/core/experimental/_module.py Outdated Show resolved Hide resolved

Move ParamInfo to class context

9193901

Used modern namedtuple instance constructor.

oleksandr-pavlyk dismissed rwgk’s stale review via 9193901 May 8, 2025 21:51

oleksandr-pavlyk requested a review from rwgk May 8, 2025 22:25

rwgk previously approved these changes May 8, 2025

View reviewed changes

oleksandr-pavlyk requested a review from leofang May 13, 2025 14:01

leofang reviewed May 13, 2025

View reviewed changes

cuda_core/cuda/core/experimental/_module.py Outdated Show resolved Hide resolved

leofang previously approved these changes May 13, 2025

View reviewed changes

Use ParamInfo instead of NamedTuple is annotation

55f6d31

This required moving ParamInfo definition from class scope to module scope, since referencing Kernel.ParamInfo from annotations of methods of the Kernel class results in error that Kernel class does not yet exist.

oleksandr-pavlyk dismissed leofang’s stale review via 55f6d31 May 13, 2025 17:12

rwgk previously approved these changes May 13, 2025

View reviewed changes

cuda_core/tests/test_module.py Show resolved Hide resolved

leofang previously approved these changes May 14, 2025

View reviewed changes

leofang enabled auto-merge (squash) May 14, 2025 01:25

Add fixture deinit_context_function

912ae11

Use in test_module.py::test_num_args_error_handling Add comments

oleksandr-pavlyk dismissed stale reviews from leofang and rwgk via 912ae11 May 14, 2025 12:46

Keep popping current context until stack is drained for test_num_args…

7e45902

…_error_handling

rwgk reviewed May 14, 2025

View reviewed changes

cuda_core/tests/test_module.py Outdated Show resolved Hide resolved

oleksandr-pavlyk requested review from leofang and rwgk May 14, 2025 21:02

rwgk approved these changes May 14, 2025

View reviewed changes

leofang merged commit 71762af into NVIDIA:main May 14, 2025
75 checks passed

oleksandr-pavlyk deleted the add-num-arguments branch May 15, 2025 01:06

Implement Kernel.num_arguments, and Kernel.arguments_info #612

Implement Kernel.num_arguments, and Kernel.arguments_info #612

Uh oh!

Conversation

oleksandr-pavlyk commented May 6, 2025

Description

Checklist

Uh oh!

copy-pr-bot bot commented May 6, 2025

Uh oh!

leofang left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

copy-pr-bot bot commented May 8, 2025

Uh oh!

oleksandr-pavlyk commented May 8, 2025

Uh oh!

This comment has been minimized.

rwgk left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

oleksandr-pavlyk commented May 8, 2025

Uh oh!

rwgk left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

oleksandr-pavlyk commented May 8, 2025

Uh oh!

rwgk left a comment

Choose a reason for hiding this comment

Uh oh!

leofang commented May 9, 2025

Uh oh!

leofang commented May 13, 2025

Uh oh!

leofang commented May 13, 2025

Uh oh!

Uh oh!

Uh oh!

leofang commented May 14, 2025

Uh oh!

oleksandr-pavlyk commented May 14, 2025

Uh oh!

leofang commented May 14, 2025

Uh oh!

oleksandr-pavlyk commented May 14, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

leofang commented May 14, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

oleksandr-pavlyk commented May 14, 2025

Uh oh!

oleksandr-pavlyk commented May 14, 2025

Uh oh!

Uh oh!

oleksandr-pavlyk commented May 14, 2025

Uh oh!

rwgk left a comment

Choose a reason for hiding this comment

Uh oh!

rwgk May 14, 2025

Choose a reason for hiding this comment

Uh oh!

rwgk May 14, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

github-actions bot commented May 14, 2025

Uh oh!

Uh oh!

oleksandr-pavlyk commented May 14, 2025 •

edited

Loading

leofang commented May 14, 2025 •

edited

Loading