Cythonize launch & LaunchConfig more

[leofang]

Description

closes#1078

this PR:

In [1]: %timeitlaunch(s, config, kernel, 0, 0)
2.6μs ± 0.656nsperloop (mean ± std. dev. of7runs, 100,000loopseach)

main branch (commit 027ba10)

In [1]: %timeitlaunch(s, config, kernel, 0, 0)
4.02μs ± 5.8nsperloop (mean ± std. dev. of7runs, 100,000loopseach)

cuda.core v0.4.2 (regression):

In [1]: %timeitlaunch(s, config, kernel, 0, 0)
4.42μs ± 6.69nsperloop (mean ± std. dev. of7runs, 100,000loopseach)

cuda.core v0.3.2:

In [1]: %timeitlaunch(s, config, kernel, 0, 0)
4.18μs ± 4.13nsperloop (mean ± std. dev. of7runs, 100,000loopseach)

script:

fromcuda.core.experimentalimportDevice, Program, launch, LaunchConfigdev=Device()
dev.set_current()
prog=Program(
r""" extern "C" __global__ void empty_kernel(const float* a, const double* b) { } """, code_type="c++"
)
obj=prog.compile(target_type="cubin")
kernel=obj.get_kernel("empty_kernel")
config=LaunchConfig(grid=64, block=256, shmem_size=0)
s=dev.default_streamlaunch(s, config, kernel, 0, 0)
dev.sync()

Execute it via ipython -i launch_perf.py to warm up and then repeat the launch via %timeit launch(s, config, kernel, 0, 0).

Checklist

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

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[leofang]

/ok to test 51635bf

[leofang]

For comparison, this is with cuda.bindings:

In [16]: %timeitdriver.cuLaunchKernel(kernel._handle, *config.grid, *config.block, 0, s.handle, ((0, 0), (ctypes.c_void_p, ctypes.c_void_p)), 0)
3.98μs ± 27.6nsperloop (mean ± std. dev. of7runs, 100,000loopseach)

[rparolin]

generally lgtm, a left a few comments for you to review.

[leofang]

Thanks, Rob!

cc @kkraus14@cpcloud@emcastillo for vis

[leofang]

Leaving a quick note here in case I forget.

I wasn't super happy about this PR and it's why I was on the fence of pushing this forward: Our design of LaunchConfig allows reusing the Python object across multiple launches. But, it is too flexible that we pay the price of maintaining the public attributes (both readable and writable in Python) associated with Python overhead. If we don't think this is reasonable, we should break it before GA, and turn LaunchConfig into an immutable object.

[kkraus14]

In theory we could store the native types and have getter / setter properties that translate to/from tuples as needed? I agree we should avoid paying the Python overhead here if possible.

[leofang]

Yes. Though in the case of launch config here, we essentially are wrapping a struct with a flexible array member (the attributes array can be arbitrarily long), which is annoying.

Last night I was thinking about thread safety. But in the present case we still do not offer thread safety anyway (ex: when two threads set grid member at the same time, for example). Something to think about in #1389.

[github-actions]

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

[Andy-Jost]

LGTM