Skip to content
Merged
Show file tree
Hide file tree
Changes from 1 commit
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
22 changes: 21 additions & 1 deletion numba/cuda/cudadecl.py
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@
signature, Registry)
from numba.cuda.types import dim3, grid_group
from numba import cuda

import operator

registry = Registry()
register = registry.register
Expand Down Expand Up @@ -366,6 +366,20 @@ class Cuda_fp16_cmp(ConcreteTemplate):
return Cuda_fp16_cmp


def _genfp16_comparison_operator(l_key):
@register_global(l_key)
class Cuda_fp16_cmp2(AbstractTemplate):

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

Why the 2?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Because it was determined in this PR: #7885
with input from @sklam that operator registration requires using 'AbstractTemplate' while intrinsics need to use 'ConcreteTemplate'.

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

To clarify my thinking - I think the name Cuda_fp16_cmp2 is a little confusing - there is the other class in the preceding method, but there's no name collision so it wouldn't have been necessary to call it ...cmp2 as opposed to ..cmp - the difference between this class and the one from the preceding method is that this one is for typing comparisons using the operator module, so naming it Cuda_fp16_operator_cmp (for example) would be more descriptive.

key = l_key

def generic(self, args, kws):
assert not kws
if len(args) == 2 and \

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

Can you explain the thinking behind the logic here? Some questions that spring to mind when I read this implementation:

  • Why not just explicitly require both arguments to be float16? If they are convertible, won't a cast be inserted by typing?
  • What types are convertible to float16 safely that would result in convertible being True?

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

Further to this, I notice that the following implementation passes all tests:

          def generic(self, args, kws):
              assert not kws  
              if args == (types.float16, types.float16):
                  return signature(types.b1, types.float16, types.float16)

Does this imply that there is some missing test coverage, or that the simple implementation suffices? (my suspicion is that the simple implementation suffices, but I don't think I've understood yet why the more complex implementation could be needed, so I'm not certain either way).

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

@gmarkall requiring both arguments to be float16 will fail, forcing one argument to be float16 ended up generating incorrect mixed type code. As an example comparing a float16 and float32 converted the float32 to float16 and then doing the comparison which is incorrect. The types int16 and int8 can be converted to float16.

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

Thankyou for the information - however, I don't really understand the explanation. I have the following patch applied on top of your commit d9abed6:

diff --git a/numba/cuda/cudadecl.py b/numba/cuda/cudadecl.py
index c0ef42652..a54740efd 100644
--- a/numba/cuda/cudadecl.py
+++ b/numba/cuda/cudadecl.py
@@ -5,7 +5,6 @@ from numba.core.typing.templates import (AttributeTemplate, ConcreteTemplate,
                                          AbstractTemplate, CallableTemplate,
                                          signature, Registry)
 from numba.cuda.types import dim3, grid_group
-from numba.core.typeconv import Conversion
 from numba import cuda
 import operator
 
@@ -374,17 +373,8 @@ def _genfp16_comparison_operator(l_key):
 
         def generic(self, args, kws):
             assert not kws
-
-            if len(args) == 2 and \
-                    (args[0] == types.float16 or args[1] == types.float16):
-                if (args[0] == types.float16):
-                    convertible = self.context.can_convert(args[1], args[0])
-                else:
-                    convertible = self.context.can_convert(args[0], args[1])
-
-                if (convertible == Conversion.exact) or \
-                        (convertible == Conversion.promote):
-                    return signature(types.b1, types.float16, types.float16)
+            if args == (types.float16, types.float16):
+                return signature(types.b1, types.float16, types.float16)
 
 
 Cuda_hadd = _genfp16_binary(cuda.fp16.hadd)

I run the testsuite like this:

python -m numba.runtests numba.cuda.tests  -m

and I get the following summary as a result:

----------------------------------------------------------------------
Ran 1295 tests in 115.148s

OK (skipped=18, expected failures=8)

From this I can see three possibilities:

  • The implementation could just require both arguments to be float16, and everything will still work. If this is the case, then please could you change to the simple implementation outlined above.
  • Requiring both arguments to be float16 is in fact incorrect, but there is no test that demonstrates the need for the more complex implementation in this PR. If this is the case, could you please add a test that will pass with the implementation in this PR, and fail with the simple implementation outlined above.
  • There's something else going on that I don't understand - if this is the case, could you please help me to understand where I'm going wrong? :-)

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

@gmarkall I added two test cases that verify the generated ptx. One of these test cases fails when using your code that requires both arguments to be float16.

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

Thanks - this now makes sense. I now recall that the complexity here is to work around the fact that we can't have multiple concrete templates for a single global. I think what's going on here is going to be quite unclear to anyone coming across this for the first time - can you please add some comments explaining the background and how the workaround works to the Cuda_fp16_operator_cmp class?

(args[0] == types.float16 or args[1] == types.float16):
return signature(types.b1, types.float16, types.float16)

return Cuda_fp16_cmp2
Comment thread
gmarkall marked this conversation as resolved.
Resolved


Cuda_hadd = _genfp16_binary(cuda.fp16.hadd)
Cuda_hsub = _genfp16_binary(cuda.fp16.hsub)
Cuda_hmul = _genfp16_binary(cuda.fp16.hmul)
Expand All @@ -374,11 +388,17 @@ class Cuda_fp16_cmp(ConcreteTemplate):
Cuda_hneg = _genfp16_unary(cuda.fp16.hneg)
Cuda_habs = _genfp16_unary(cuda.fp16.habs)
Cuda_heq = _genfp16_binary_comparison(cuda.fp16.heq)
Cuda_eq = _genfp16_comparison_operator(operator.eq)
Comment thread
gmarkall marked this conversation as resolved.
Resolved
Cuda_hne = _genfp16_binary_comparison(cuda.fp16.hne)
Cuda_ne = _genfp16_comparison_operator(operator.ne)
Cuda_hge = _genfp16_binary_comparison(cuda.fp16.hge)
Cuda_ge = _genfp16_comparison_operator(operator.ge)
Cuda_hgt = _genfp16_binary_comparison(cuda.fp16.hgt)
Cuda_gt = _genfp16_comparison_operator(operator.gt)
Cuda_hle = _genfp16_binary_comparison(cuda.fp16.hle)
Cuda_le = _genfp16_comparison_operator(operator.le)
Cuda_hlt = _genfp16_binary_comparison(cuda.fp16.hlt)
Cuda_lt = _genfp16_comparison_operator(operator.lt)


# generate atomic operations
Expand Down
6 changes: 6 additions & 0 deletions numba/cuda/cudaimpl.py
Original file line number Diff line number Diff line change
Expand Up @@ -545,11 +545,17 @@ def ptx_fp16_comparison(context, builder, sig, args):


lower(stubs.fp16.heq, types.float16, types.float16)(_gen_fp16_cmp('eq'))
lower(operator.eq, types.float16, types.float16)(_gen_fp16_cmp('eq'))
lower(stubs.fp16.hne, types.float16, types.float16)(_gen_fp16_cmp('ne'))
lower(operator.ne, types.float16, types.float16)(_gen_fp16_cmp('ne'))
lower(stubs.fp16.hge, types.float16, types.float16)(_gen_fp16_cmp('ge'))
lower(operator.ge, types.float16, types.float16)(_gen_fp16_cmp('ge'))
lower(stubs.fp16.hgt, types.float16, types.float16)(_gen_fp16_cmp('gt'))
lower(operator.gt, types.float16, types.float16)(_gen_fp16_cmp('gt'))
lower(stubs.fp16.hle, types.float16, types.float16)(_gen_fp16_cmp('le'))
lower(operator.le, types.float16, types.float16)(_gen_fp16_cmp('le'))
lower(stubs.fp16.hlt, types.float16, types.float16)(_gen_fp16_cmp('lt'))
lower(operator.lt, types.float16, types.float16)(_gen_fp16_cmp('lt'))


def lower_fp16_minmax(fn, fname, op):
Expand Down
67 changes: 66 additions & 1 deletion numba/cuda/tests/cudapy/test_operator.py
Original file line number Diff line number Diff line change
@@ -1,7 +1,34 @@
import numpy as np
from numba.cuda.testing import unittest, CUDATestCase
from numba.cuda.testing import (unittest, CUDATestCase, skip_unless_cc_53,
skip_on_cudasim)
from numba import cuda
from numba.core.types import f2,b1
import operator
from numba.cuda import compile_ptx


def simple_fp16_gt(ary, a, b):
ary[0] = a > b


def simple_fp16_ge(ary, a, b):
ary[0] = a >= b


def simple_fp16_lt(ary, a, b):
ary[0] = a < b


def simple_fp16_le(ary, a, b):
ary[0] = a <= b


def simple_fp16_eq(ary, a, b):
ary[0] = a == b


def simple_fp16_ne(ary, a, b):
ary[0] = a != b


class TestOperatorModule(CUDATestCase):
Expand Down Expand Up @@ -36,6 +63,44 @@ def test_truediv(self):
def test_floordiv(self):
self.operator_template(operator.floordiv)

@skip_unless_cc_53
def test_fp16_comparison(self):
functions = (simple_fp16_gt, simple_fp16_ge,
simple_fp16_lt, simple_fp16_le,
simple_fp16_eq, simple_fp16_ne)
ops = (operator.gt, operator.ge, operator.lt, operator.le,
operator.eq, operator.ne)

for fn, op in zip(functions, ops):
with self.subTest(op=op):
kernel = cuda.jit("void(b1[:], f2, f2)")(fn)

expected = np.zeros(1, dtype=np.bool8)
Comment thread
gmarkall marked this conversation as resolved.
Resolved
got = np.zeros(1, dtype=np.bool8)
arg1 = np.random.random(1).astype(np.float16)
arg2 = np.random.random(1).astype(np.float16)

kernel[1, 1](got, arg1[0], arg2[0])
expected = op(arg1, arg2)
self.assertEqual(got[0], expected)

@skip_on_cudasim('Compilation unsupported in the simulator')
def test_fp16_comparison_ptx(self):
functions = (simple_fp16_gt, simple_fp16_ge,
simple_fp16_lt, simple_fp16_le,
simple_fp16_eq, simple_fp16_ne)
ops = (operator.gt, operator.ge, operator.lt, operator.le,
operator.eq, operator.ne)
opstring = ('setp.gt.f16', 'setp.ge.f16',
'setp.lt.f16', 'setp.le.f16',
'setp.eq.f16', 'setp.ne.f16')
args = (b1[:], f2, f2)

for fn, op, s in zip(functions, ops, opstring):
with self.subTest(op=op):
ptx, _ = compile_ptx(fn, args, cc=(5, 3))
self.assertIn(s, ptx)


if __name__ == '__main__':
unittest.main()