Skip to content

CUDA: Deprecate eager compilation of device functions - #7114

Merged
sklam merged 3 commits into
numba:masterfrom
gmarkall:deprecate-eager-device
Jun 22, 2021
Merged

sklam merged 3 commits into
numba:masterfrom
gmarkall:deprecate-eager-device

Conversation

@gmarkall

Copy link
Copy Markdown
Member

In future versions of Numba, the device kwarg to the @cuda.jit decorator will be obviated, and whether a device function or global kernel is compiled will be inferred from the context. With respect to kernel / device functions and lazy / eager compilation, four cases are presently handled:

  1. device=True, eager compilation with a signature provided
  2. device=False, eager compilation with a signature provided
  3. device=True, lazy compilation with no signature
  4. device=False, lazy compilation with no signature

The latter two cases can be differentiated without the device kwarg, because it can be inferred from the calling context - if the call is from the host, then a global kernel should be compiled, and if the call is from a kernel or another device function, then a device function should be compiled.

The first two cases cannot be differentiated in the absence of the device kwarg - without it, it will not be clear from a signature alone whether a device function or global kernel should be compiled. In order to resolve, this, support for eager compilation of device functions will be removed. Eager compilation with the @cuda.jit decorator will in future always imply the immediate
compilation of a global kernel.

This PR is made for Numba 0.54, in order to prepare for the changes in #7107 for 0.55.

In future versions of Numba, the ``device`` kwarg to the ``@cuda.jit`` decorator
will be obviated, and whether a device function or global kernel is compiled will
be inferred from the context. With respect to kernel / device functions and lazy
/ eager compilation, four cases are presently handled:

1. ``device=True``, eager compilation with a signature provided
2. ``device=False``, eager compilation with a signature provided
3. ``device=True``, lazy compilation with no signature
4. ``device=False``, lazy compilation with no signature

The latter two cases can be differentiated without the ``device`` kwarg, because
it can be inferred from the calling context - if the call is from the host, then
a global kernel should be compiled, and if the call is from a kernel or another
device function, then a device function should be compiled.

The first two cases cannot be differentiated in the absence of the ``device``
kwarg - without it, it will not be clear from a signature alone whether a device
function or global kernel should be compiled. In order to resolve, this, support
for eager compilation of device functions will be removed. Eager compilation
with the ``@cuda.jit`` decorator will in future always imply the immediate
compilation of a global kernel.
@gmarkall
gmarkall force-pushed the deprecate-eager-device branch from b36fbe8 to 384d2f6 Compare June 14, 2021 12:24
@gmarkall gmarkall added 3 - Ready for Review CUDA CUDA related issue/PR Effort - short Short size effort needed labels Jun 14, 2021
@gmarkall gmarkall added this to the Numba 0.54 RC milestone Jun 14, 2021
@gmarkall gmarkall changed the title CUDA: Deprecate eager deprecation of device functions CUDA: Deprecate eager compilation of device functions Jun 14, 2021
@sklam sklam modified the milestones: Numba 0.54 RC, Numba 0.55 RC Jun 14, 2021

@stuartarchibald stuartarchibald left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Thanks for the patch, couple of minor things to resolve else looks good.

Comment thread docs/source/reference/deprecation.rst Outdated
Comment thread numba/cuda/decorators.py Outdated
@stuartarchibald stuartarchibald added 4 - Waiting on author Waiting for author to respond to review and removed 3 - Ready for Review labels Jun 16, 2021
Co-authored-by: stuartarchibald <stuartarchibald@users.noreply.github.com>
@gmarkall gmarkall added 4 - Waiting on reviewer Waiting for reviewer to respond to author and removed 4 - Waiting on author Waiting for author to respond to review labels Jun 16, 2021

@stuartarchibald stuartarchibald left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Thanks for the patch and fixes.

@stuartarchibald stuartarchibald added 4 - Waiting on CI Review etc done, waiting for CI to finish Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm and removed 4 - Waiting on reviewer Waiting for reviewer to respond to author labels Jun 16, 2021
@stuartarchibald

stuartarchibald commented Jun 16, 2021 •

Copy link
Copy Markdown
Contributor

Buildfarm ID: numba_smoketest_cuda_yaml_73.

@stuartarchibald

Copy link
Copy Markdown
Contributor

Buildfarm ID: numba_smoketest_cuda_yaml_73.

Passed.

@stuartarchibald stuartarchibald added BuildFarm Passed For PRs that have been through the buildfarm and passed and removed 4 - Waiting on CI Review etc done, waiting for CI to finish Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm labels Jun 21, 2021
@stuartarchibald stuartarchibald added the 5 - Ready to merge Review and testing done, is ready to merge label Jun 21, 2021
@sklam
sklam merged commit 9d2f36a into numba:master Jun 22, 2021
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

5 - Ready to merge Review and testing done, is ready to merge BuildFarm Passed For PRs that have been through the buildfarm and passed CUDA CUDA related issue/PR Effort - short Short size effort needed

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants