Repository navigation
Adds CUDA FP16 arithmetic operators - #7885
Merged
Merged
Conversation
Contributor
Author
|
gpuci run tests |
Member
|
gpuci run tests |
gmarkall
reviewed
Apr 11, 2022
gmarkall
left a comment
gmarkall
left a comment
Member
There was a problem hiding this comment.
With this branch, the following:
from numba import cuda
from numba.types import float16, float32
def f(x, y):
return x + y
ptx_16, resty_16 = cuda.compile_ptx(f, (float16, float16), device=True)
ptx_32, resty_32 = cuda.compile_ptx(f, (float32, float32), device=True)
print(ptx_16)
print(resty_16)
print(ptx_32)
print(resty_32)When run with
numba --annotate-html scalar.html scalar.py
Seems to give the following typing for the addition of float16s (in the HTML output):
x = arg(0, name=x) :: float16
y = arg(1, name=y) :: float16
$6binary_add.2 = x + y :: float32
and corresponding PTX:
// begin inline asm
cvt.f32.f16 %f1, %rs1;
// end inline asm
// begin inline asm
cvt.f32.f16 %f2, %rs2;
// end inline asm
add.f32 %f3, %f1, %f2;
st.f32 [%rd1], %f3;
Is it expected with this PR that float16 + float16 should be typed as returning a float16?
(Does this reproduce for you? Perhaps I am missing something / running something incorrectly?)
gmarkall
added
4 - Waiting on author
Contributor
Author
@gmarkall this reproduces for me and you definitely found a issue. Thanks. |
isVoid
reviewed
Apr 13, 2022
sklam
requested changes
Apr 14, 2022
sklam
left a comment
sklam
left a comment
Member
There was a problem hiding this comment.
It turns out the template should be using @register_global for the operator.*.
You can check it without CUDA hardware with:
import operator
from numba.core import types, typing
from numba.cuda.descriptor import cuda_target
fp16 = types.float16
typingctx = cuda_target.typing_context
typingctx.refresh()
fnty = typingctx.resolve_value_type(operator.add)
out = typingctx.resolve_function_type(fnty, (fp16, fp16), {})
print(out)
if out != typing.signature(fp16, fp16, fp16):
raise AssertionError(out)The above code is what the compiler does for function resolution. I'd recommend adding tests like this to numba/cuda/tests/nocuda.
…s test cases for inplace operators
Contributor
Author
|
gpuci run tests |
gmarkall
added
4 - Waiting on reviewer
…e being generated
Contributor
Author
|
gpuci run tests |
Contributor
Author
|
gpuci run tests |
Contributor
Author
|
@gmarkall I added the skip_unless_cc_53 back to the testcase. |
Member
|
@esc Could this have another buildfarm run please? |
Member
|
Member
So, I tested `` on the Anaconda build farm and still received the following errors on Windows. @testhound could you take a look, please? |
Contributor
Author
|
@esc thank you for the feedback, looking into the failures. |
Contributor
Author
|
gpuci run tests |
Contributor
Author
|
@esc can I request another buildfarm run? I fixed the failing tests by skipping if the compute capability requirements are not met. I am confused as to why these tests just started failing, since this code has been in main for quite some time. |
Contributor
Author
|
@gmarkall FYI: I updated the PR to fix tests that should be skipped. |
Member
It suggests that there is something wrong with these tests that we never picked up on before. |
Comment on lines
+20
to
+21
| if out != typing.signature(fp16, fp16, fp16): | ||
| raise AssertionError(out) |
Contributor
There was a problem hiding this comment.
Should this be self.assertEqual(out, typing.signature(fp16, fp16, fp16), msg=str(out)) or similar? (same for the other test in this class).
Contributor
Author
|
gpuci run tests |
Contributor
Author
|
@stuartarchibald and @gmarkall code update per your request to use |
Contributor
Author
|
gpuci run tests |
gmarkall
approved these changes
Sep 13, 2022
gmarkall
left a comment
gmarkall
left a comment
Member
There was a problem hiding this comment.
Looks good to me. Could this have another buildfarm run please @stuartarchibald ?
stuartarchibald
approved these changes
Sep 13, 2022
stuartarchibald
left a comment
stuartarchibald
left a comment
Contributor
There was a problem hiding this comment.
Thanks for the patch and fixes.
stuartarchibald
removed
the
4 - Waiting on author
Member
|
I ran this on the farm as |
esc
added
BuildFarm Passed
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 pull request add operator support for the following fp 16 operations: +, -, *, neg, abs. This PR partly addresses the meta issue "add fp16 support": #4402.
This PR builds on: #7460 which added intrinsic support for the above operators.