Repository navigation
Add FP16 support for CUDA - #7460
Merged
Merged
Conversation
testhound
requested review from
esc,
gmarkall,
sklam and
stuartarchibald
as code owners
October 6, 2021 13:52
Member
|
Looks like the CI failure is because the fp16 intrinsics is not available in cudasim |
Contributor
Author
@sklam pushed a patch for cudasim. |
Member
|
The following: from numba import cuda
import numpy as np
@cuda.jit
def fadd(r, x, y):
r[0] = cuda.fp16.hadd(x[0], y[0])
x = np.ones(1, dtype=np.float16) * 2.5
y = np.ones(1, dtype=np.float16) * 3.7
r = np.zeros_like(x)
fadd[1, 1](r, x, y)yields Am I doing something we don't expect to support yet here? |
Contributor
Author
@gmarkall no I don't see that you are doing anything wrong. Need to dig into this. |
Member
|
I tracked the recent CI error to the ufunc loop selection code. The change in 95727a4#diff-32316606fefc1b740b8e2738aa2ed5fa533d7cec87b1a691bbe65ac5daab12f3R32 allow ufunc loop for diff --git a/numba/np/numpy_support.py b/numba/np/numpy_support.py
index c83dc3cf9..a5530c7da 100644
--- a/numba/np/numpy_support.py
+++ b/numba/np/numpy_support.py
@@ -497,7 +497,9 @@ def ufunc_find_matching_loop(ufunc, arg_types):
# (e.g. float16), try other candidates
continue
else:
- return UFuncLoopSpec(inputs, outputs, candidate)
+ loopspec = UFuncLoopSpec(inputs, outputs, candidate)
+ if supported_ufunc_loop(ufunc, loopspec):
+ return loopspec
return None |
Member
|
Looks like I shouldn't have changed The patch: diff --git a/numba/np/numpy_support.py b/numba/np/numpy_support.py
index a5530c7da..b8ee24135 100644
--- a/numba/np/numpy_support.py
+++ b/numba/np/numpy_support.py
@@ -455,6 +455,9 @@ def ufunc_find_matching_loop(ufunc, arg_types):
for candidate in ufunc.types:
ufunc_inputs = candidate[:ufunc.nin]
ufunc_outputs = candidate[-ufunc.nout:] if ufunc.nout else []
+ if 'e' in ufunc_inputs:
+ # Skip float16 arrays since we don't have implementation for them
+ continue
if 'O' in ufunc_inputs:
# Skip object arrays
continue
@@ -497,9 +500,7 @@ def ufunc_find_matching_loop(ufunc, arg_types):
# (e.g. float16), try other candidates
continue
else:
- loopspec = UFuncLoopSpec(inputs, outputs, candidate)
- if supported_ufunc_loop(ufunc, loopspec):
- return loopspec
+ return UFuncLoopSpec(inputs, outputs, candidate)
return None
|
gmarkall
requested changes
Oct 13, 2021
gmarkall
left a comment
gmarkall
left a comment
Member
There was a problem hiding this comment.
Many thanks for the PR! The big picture looks good - there are some comments on the diff, and a couple of other observations:
- The changes in boxing.py, builtins.py, and cudamath.py appear to be unnecessary to support fp16 in CUDA and the currently-implemented intrinsics. The commit bfa193d undoes these changes (and a whitespace change in compiler.py).
- The logic for lowering could be simplified and shortened - commit 484557f factors some of them and simplifies things a little (in the same way that the typing is already factored in this PR).
The above commits along with Siu's new proposed fix test OK on CI (checked in #7481) and also all passes locally for me with hardware and the simulator - feel free to add those commits to this branch if you'd like.
| 'unsigned long long': 'y', # unsigned __int64 | ||
| '__int128': 'n', | ||
| 'unsigned __int128': 'o', | ||
| 'float16' : 'Dh', |
Member
There was a problem hiding this comment.
Note to any other reviewers: correlates with https://itanium-cxx-abi.github.io/cxx-abi/abi.html#mangling-builtin
…actor/simplify and fix doc errors
gmarkall
requested changes
Oct 18, 2021
gmarkall
left a comment
gmarkall
left a comment
Member
There was a problem hiding this comment.
Many thanks for the updates ! Things are looking good, though there are a couple of comments / questions on the diff.
The additions to the type casting rules and corresponding tests look good, and appear to be consistent with existing type casting rules. I'm not 100% sure I understand why promoting to a wider float from a narrower one is considered unsafe (e.g. the existing f32 -> f64, but I'm assuming it's part of a tradeoff somewhere to try and avoid widening things unnecessarily.
Comment on lines
+444
to
+447
| arg1 = args[0] | ||
| arg2 = args[1] | ||
| arg3 = args[2] | ||
| return builder.call(hfma_inline, [arg1, arg2, arg3]) |
Member
There was a problem hiding this comment.
Suggested change
| arg1 = args[0] | |
| arg2 = args[1] | |
| arg3 = args[2] | |
| return builder.call(hfma_inline, [arg1, arg2, arg3]) | |
| return builder.call(hfma_inline, args) |
|
|
||
|
|
||
| @skip_on_cudasim('CUDA Driver API unsupported in the simulator') | ||
| @unittest.skip |
- Don't raise a lowering error if CUDA toolkit is < 10.2 when lowering habs - use an alternative implementation instead. - Rename the `float16` C type to `half`. - Swap order of some lines that could raise so the raise happens earlier in the function - Simplify return type of a function declaration in `integer_to_float16_cast`.
Member
|
gpuci run tests |
Member
|
gpuci run tests |
stuartarchibald
previously approved these changes
Dec 14, 2021
stuartarchibald
left a comment
stuartarchibald
left a comment
Contributor
There was a problem hiding this comment.
Thanks for the updates to address the review @gmarkall, patch looks good.
stuartarchibald
added
4 - Waiting on CI
Member
|
|
Member
|
Interestingly enough this breaks across all Windows builds with the following errors: |
gmarkall
dismissed stale reviews from stuartarchibald and themself
via
December 14, 2021 16:40
20c1d4f
Member
|
gpuci run tests |
Contributor
|
This shows a limitation of the internal Numba Windows test server, which currently has a |
Member
|
I've pushed a fix that skips the test if the CC is not new enough. I think we should consider dropping support for CC less than at least 5.2 after Numba 0.55. |
stuartarchibald
approved these changes
Dec 14, 2021
stuartarchibald
left a comment
stuartarchibald
left a comment
Contributor
There was a problem hiding this comment.
Thanks for addressing the issue with CC < 5.3.
Member
|
BFID |
sklam
added
5 - Ready to merge
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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
6 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.
This PR begins the process of adding 16-bit floating point support for the cuda target. This PR adds support for the following 16-bit fp operators: add, subtract, multiply, multiply accumulate, negate and absolute value. The divide operator was omitted as it is more involved and will be included in the next PR.