Add support for fp16 comparison native operators - #8010

Why the 2?

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

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.

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?

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).

@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.

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? :-)

@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.

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?

This does explain what conversions we allow based on the type-conversion rules. However, the question in most readers' minds (I imagine) when encountering this typing function, will by why we need an AbstractTemplate with what appears to be a partial re-implementation of Numba's casting logic, instead of just a ConcreteTemplate with a list of cases.

It took me some time to grasp this, and our discussions in the other comments (e.g. above) seem to cover the reasoning. Could you please extend the commentary to explain why we go to all this trouble with the ConcreteTemplate?

Done