CUDA: Lazily add libdevice to compilation units #6795
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
The nvvmLazyAddModuleToProgram function, which will be used in the following commit, was added in CUDA 9.2. This commit drops support for CUDA 9.0 and 9.1 so that the use of this function can be supported without maintaining two different paths.
Modules should be lazily loaded where possible, in order to permit certain optimizations that make compilation more efficient and improve runtime performance. This commit adds support for lazy loading and uses it for libdevice, since it is never the entry point to a compilation unit. This seems to have a noticeable effect on compilation speed. Running: ``` python -m numba.runtests numba.cuda.tests -m ``` prior to this commit gives: ``` Ran 1132 tests in 136.102s ``` on my system. With this commit, it gives: ``` Ran 1132 tests in 86.255s ```
gmarkall
added
3 - Ready for Review
Effort - medium
Contributor
stuartarchibald
left a comment
stuartarchibald
left a comment
There was a problem hiding this comment.
Thanks for the patch, couple of minor things to resolve else looks good.
stuartarchibald
added
4 - Waiting on author
Member
|
Oddly, my OSX CUDA9.2 doesn't have In [1]: p = "/Developer/NVIDIA/CUDA-9.2/nvvm/lib/libnvvm.3.2.0.dylib"
In [2]: import ctypes
In [3]: lib = ctypes.CDLL(p)
In [4]: lib.nvvmLazyAddModuleToProgram
---------------------------------------------------------------------------
AttributeError Traceback (most recent call last)
<ipython-input-4-7fa6d4c0d54f> in <module>
----> 1 lib.nvvmLazyAddModuleToProgram
~/dev/numba/.conda_env_py38/lib/python3.8/ctypes/__init__.py in __getattr__(self, name)
384 if name.startswith('__') and name.endswith('__'):
385 raise AttributeError(name)
--> 386 func = self.__getitem__(name)
387 setattr(self, name, func)
388 return func
~/dev/numba/.conda_env_py38/lib/python3.8/ctypes/__init__.py in __getitem__(self, name_or_ordinal)
389
390 def __getitem__(self, name_or_ordinal):
--> 391 func = self._FuncPtr((name_or_ordinal, self))
392 if not isinstance(name_or_ordinal, int):
393 func.__name__ = name_or_ordinal
AttributeError: dlsym(0x7ff34acdbb90, nvvmLazyAddModuleToProgram): symbol not found
In [5]: lib.nvvmAddModuleToProgram
Out[5]: <_FuncPtr object at 0x1109cc580> |
Member
|
I have just reinstalled CUDA 9.2.148. here's running $ nm libnvvm.3.2.0.dylib | grep nvvm
0000000000012700 T ___nvvmHandle
000000000000f080 T _nvvmAddModuleToProgram
000000000000f290 T _nvvmCompileProgram
000000000000ed30 T _nvvmCreateProgram
000000000000ee90 T _nvvmDestroyProgram
0000000000012380 T _nvvmGetCompiledResult
00000000000122f0 T _nvvmGetCompiledResultSize
00000000000125b0 T _nvvmGetErrorString
00000000000124e0 T _nvvmGetProgramLog
0000000000012440 T _nvvmGetProgramLogSize
000000000000ec80 T _nvvmIRVersion
0000000000011f50 T _nvvmVerifyProgram
000000000000ebf0 T _nvvmVersionThe symbol is not defined. |
Member
Author
From OOB discussion:
So as I understand it, there's no resolution needed here. |
gmarkall
added
4 - Waiting on reviewer
Contributor
Assuming that this is resolved, this PR looks good. Please could you resolve the conflicts and then it can be tested on the build farm? Many Thanks. |
stuartarchibald
added
4 - Waiting on author
gmarkall
added
4 - Waiting on CI
Member
Author
|
@stuartarchibald Thanks for the review - conflicts now resolved, and I understood from @sklam that he was happy with the MacOS situation - this should be good for a BuildFarm run now. |
Contributor
|
Buildfarm ID: |
Member
|
The CUDA 9.0 tests fail on this PR because CUDA 9.0 can not be installed: |
Member
Author
|
Following the OOB conversation - @esc Can we test this with 9.2 as it drops support for 9.0 and 9.1 please? |
Member
|
Will upgrade the build farm to test against 9.2 instead of 9.0 and see if this helps. |
Member
|
Using 9.2 instead yields the following: |
Member
|
Using 9.2 on Windows I get: |
Member
|
@gmarkall thank you for the fixes, |
esc
added
the
BuildFarm Passed
stuartarchibald
approved these changes
Mar 15, 2021
Contributor
stuartarchibald
left a comment
stuartarchibald
left a comment
There was a problem hiding this comment.
Fix looks good, thanks.
stuartarchibald
added
5 - Ready to merge
Member
Author
|
Thanks @stuartarchibald ... just for info, I'm thinking we could remove that fix for 0.55, as by the time it's released CUDA toolkit 10.0 will be over three years old, so it ought to be OK drop support for 9.2 at that point (on the basis of supporting toolkits at least three years back at the time of the release). |
Contributor
No problem. RE 0.55 and dropping 9.2, perhaps add as an issue for https://github.com/numba/numba/milestone/56 ? Seems reasonable to drop it and also review any other existing CUDA <10.0 workarounds. |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Modules should be lazily loaded where possible, in order to permit certain optimizations that make compilation more efficient and improve runtime performance.
This commit adds support for lazy loading and uses it for libdevice, since it is never the entry point to a compilation unit. This seems to have a noticeable effect on compilation speed. Running:
prior to this PR gives:
on my system. With this PR, it gives:
This PR also drops support for CUDA 9.0 and 9.1 on the basis that they will be greater than three years old by the time 0.54 is released (9.1 released in December 2017), and
nvvmLazyAddModuleToProgramwas added in CUDA 9.2 - this avoids supporting two different paths through NVVM. Note that although examination of this PR suggests it would be straightforward to maintain both paths, I expect it to be problematic for future work based on #6769, in particular related to improving support for debug info.