Skip to content

Add support for fastmath 32-bit floating point divide #6788

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
sklam merged 3 commits into from
Mar 8, 2021
Merged

Add support for fastmath 32-bit floating point divide #6788

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
99553d3
Add support for fastmath divide
testhound Mar 3, 2021
b7344ba
Add fastmath divide test cases and streamline implementation
testhound Mar 4, 2021
4c4ac66
Merge remote-tracking branch 'upstream/master' into testhound/fastmat…
testhound Mar 4, 2021
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
Prev Previous commit
Next Next commit
Add fastmath divide test cases and streamline implementation
  • Loading branch information
@testhound
testhound committed Mar 4, 2021
commit b7344ba4535ea67453d289346cd15f14eb6adaef
11 changes: 4 additions & 7 deletions numba/cuda/mathimpl.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,16 +73,13 @@ def math_isinf_isnan_int(context, builder, sig, args):

@lower(operator.truediv, types.float32, types.float32)
def maybe_fast_truediv(context, builder, sig, args):
with cgutils.if_zero(builder, args[1]):
context.error_model.fp_zero_division(builder, ("division by zero",))

if context.fastmath:
fast_replacement = 'fast_fdividef'
actual_libfunc = getattr(libdevice, fast_replacement)
sig = typing.signature(float32, float32, float32)
libfunc_impl = context.get_function(actual_libfunc, sig)
return libfunc_impl(builder, args)
impl = context.get_function(libdevice.fast_fdividef, sig)
return impl(builder, args)
else:
with cgutils.if_zero(builder, args[1]):
context.error_model.fp_zero_division(builder, ("division by zero",))
res = builder.fdiv(*args)
return res

Expand Down
24 changes: 23 additions & 1 deletion numba/cuda/tests/cudapy/test_fastmath.py
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,13 +1,14 @@
from numba import cuda, float32
from math import cos, sin, tan, exp, log, log10, log2, pow
import numpy as np
from numba.cuda.testing import CUDATestCase, skip_on_cudasim
import unittest


class TestFastMathOption(CUDATestCase):
@skip_on_cudasim('fast divide not available in CUDASIM')
def test_kernel(self):

@gmarkall gmarkall Mar 3, 2021

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.

For a while I was wondering what this is checking that differs from what test_divf checks for below... Eventually I concluded that it tests that casting an integer to a float32 then using it in a division results in a div.approx instruction being emitted - given how finicky casts and the type system can be, I think it's good to keep this check alongside test_divf - Could you add a comment here explaining that this is testing the cast of an int being used in the divide, to make it a bit easier for future readers to see the difference?

Alternatively - do you have a different view on the difference between this and test_divf? I reached my conclusion by trying to derive meaning from the code (thinking backwards a bit) so maybe there's something else going on here that I haven't followed.


# Test the cast of an int being used in fastmath divide
def foo(arr, val):
i = cuda.grid(1)
if i < arr.size:
Expand Down Expand Up @@ -116,6 +117,27 @@ def f9(r, x, y):
slowver = cuda.jit("void(float32[::1], float32, float32)")(f9)
self.assertIn('div.approx.ftz.f32 ', fastver.ptx)
self.assertNotIn('div.approx.ftz.f32 ', slowver.ptx)
self.assertIn('div.rn.f32', slowver.ptx)
self.assertNotIn('div.rn.f32', fastver.ptx)

@skip_on_cudasim('fast divide not available in CUDASIM')
def test_divf_exception(self):
def f10(r, x, y):
r[0] = x / y

fastver = cuda.jit("void(float32[::1], float32, float32)",
fastmath=True, debug=True)(f10)
slowver = cuda.jit("void(float32[::1], float32, float32)",
debug=True)(f10)
nelem = 10
ary = np.empty(nelem, dtype=np.float32)
with self.assertRaises(ZeroDivisionError):
slowver[1, nelem](ary, 10.0, 0.0)

try:
fastver[1, nelem](ary, 10.0, 0.0)
except ZeroDivisionError:
self.fail("Divide in fastmath should not throw ZeroDivisionError")

def test_device(self):
# fastmath option is ignored for device function
Expand Down