Improved fastmath code generation for trig, log, and exp/pow.

[testhound]

This pull request adds fastmath code generation support for the following functions:

sin, cos, tan, log, log2, log10, exp, and pow when the fastmath option is used with 32-bit types.

This pull request address poor code generation identified in: #6183

[esc]

@testhound thank you for submitting this! I have added it to the queue for review.

[stuartarchibald]

Thanks for the patch, great to see these fast math functions added. I've made a couple of suggestions, in general looks good. I think this probably needs documenting, both in the CUDA docs and also the cuda.jit doc strings. Thanks again!

[stuartarchibald]

Suggested change

	if ty == float32 and context.fastmath is True:
	if ty == float32 and context.fastmath:

[testhound]

@stuartarchibald thanks for the review, I have made the code changes locally. Can you point me to the documentation files that need to be changed?

[stuartarchibald]

Think the 3x libfunc_impl = context.get_function... can be pulled out as all paths use it, then the branch is just on 32bit float with fastmath and switching the library libfunc = getattr(libdevice, fast_replacement), something like the (entirely untested!) code below perhaps, what do you think?

if ty == float32 and context.fastmath:
    fast_replacement = unarys_fastmath.get(libfunc.__name__)
    if fast_replacement is not None:
        libfunc = getattr(libdevice, fast_replacement)

libfunc_impl = context.get_function(libfunc,
                                    typing.signature(ty, ty))

[stuartarchibald]

Same comment/refactor suggestion as above.

[stuartarchibald]

RE #6619 (comment): I've taken a look through the current CUDA docs and I can't seem to find a section talking about the various options available to the cuda.jit decorator and the impact they have (might be a nice in future to have an FAQ entry https://numba.readthedocs.io/en/latest/cuda/faq.html or something to add as kernel declaration options here https://numba.readthedocs.io/en/latest/cuda/kernels.html#kernel-declaration). For now, I think it's just the fastmath part of this section needs an update: https://numba.readthedocs.io/en/latest/cuda-reference/kernel.html#kernel-declaration which I think comes from here:

numba/numba/cuda/decorators.py

Lines 45 to 48 in 8a4099a

	:param fastmath: If true, enables flush-to-zero and fused-multiply-add,
	disables precise division and square root. This parameter has no effect
	on device function, whose fastmath setting depends on the kernel function
	from which they are called.

Having reviewed the Numba documentation, I've also looked at the nvcc docs for "fast math" and think that what this PR is providing is a deviation from this behaviour. As a result, one concern I now have with this PR is over whether it would cause regressions in existing code bases due to a behavioural change in an existing configuration i.e. fastmath=True. It makes me wonder if low precision/fast math intrinsic functions ought to be activated via a different option, or that fastmath should be a dictionary like:

{
'ftz': <bool>,
'prec-div': <bool>,
'prec-sqrt': <bool>,
'fmad': <bool>,
'fast_libm': <bool>, # this being the thing that switches on the fast math intrinsic functions etc as wired in in this PR
}

On the CPU target, fastmath is a similar to this, it can be True/False but can also be specifically configured.

numba/numba/core/cpu_options.py

Lines 6 to 40 in 1976c66

	class FastMathOptions(object):
	"""
	Options for controlling fast math optimization.
	"""
	def __init__(self, value):
	# https://releases.llvm.org/7.0.0/docs/LangRef.html#fast-math-flags
	valid_flags = {
	'fast',
	'nnan', 'ninf', 'nsz', 'arcp',
	'contract', 'afn', 'reassoc',
	}
	if value is True:
	self.flags = {'fast'}
	elif value is False:
	self.flags = set()
	elif isinstance(value, set):
	invalid = value - valid_flags
	if invalid:
	raise ValueError("Unrecognized fastmath flags: %s" % invalid)
	self.flags = value
	elif isinstance(value, dict):
	invalid = set(value.keys()) - valid_flags
	if invalid:
	raise ValueError("Unrecognized fastmath flags: %s" % invalid)
	self.flags = {v for v, enable in value.items() if enable}
	else:
	msg = "Expected fastmath option(s) to be either a bool, dict or set"
	raise ValueError(msg)
	def __bool__(self):
	return bool(self.flags)
	__nonzero__ = __bool__

What do you think?

[testhound]

RE #6619 (comment): I've taken a look through the current CUDA docs and I can't seem to find a section talking about the various options available to the cuda.jit decorator and the impact they have (might be a nice in future to have an FAQ entry https://numba.readthedocs.io/en/latest/cuda/faq.html or something to add as kernel declaration options here https://numba.readthedocs.io/en/latest/cuda/kernels.html#kernel-declaration). For now, I think it's just the fastmath part of this section needs an update: https://numba.readthedocs.io/en/latest/cuda-reference/kernel.html#kernel-declaration which I think comes from here:

numba/numba/cuda/decorators.py

Lines 45 to 48 in 8a4099a

	:param fastmath: If true, enables flush-to-zero and fused-multiply-add,
	disables precise division and square root. This parameter has no effect
	on device function, whose fastmath setting depends on the kernel function
	from which they are called.

Having reviewed the Numba documentation, I've also looked at the nvcc docs for "fast math" and think that what this PR is providing is a deviation from this behaviour. As a result, one concern I now have with this PR is over whether it would cause regressions in existing code bases due to a behavioural change in an existing configuration i.e. fastmath=True. It makes me wonder if low precision/fast math intrinsic functions ought to be activated via a different option, or that fastmath should be a dictionary like:

{
'ftz': <bool>,
'prec-div': <bool>,
'prec-sqrt': <bool>,
'fmad': <bool>,
'fast_libm': <bool>, # this being the thing that switches on the fast math intrinsic functions etc as wired in in this PR
}

On the CPU target, fastmath is a similar to this, it can be True/False but can also be specifically configured.

numba/numba/core/cpu_options.py

Lines 6 to 40 in 1976c66

	class FastMathOptions(object):
	"""
	Options for controlling fast math optimization.
	"""
	def __init__(self, value):
	# https://releases.llvm.org/7.0.0/docs/LangRef.html#fast-math-flags
	valid_flags = {
	'fast',
	'nnan', 'ninf', 'nsz', 'arcp',
	'contract', 'afn', 'reassoc',
	}
	if value is True:
	self.flags = {'fast'}
	elif value is False:
	self.flags = set()
	elif isinstance(value, set):
	invalid = value - valid_flags
	if invalid:
	raise ValueError("Unrecognized fastmath flags: %s" % invalid)
	self.flags = value
	elif isinstance(value, dict):
	invalid = set(value.keys()) - valid_flags
	if invalid:
	raise ValueError("Unrecognized fastmath flags: %s" % invalid)
	self.flags = {v for v, enable in value.items() if enable}
	else:
	msg = "Expected fastmath option(s) to be either a bool, dict or set"
	raise ValueError(msg)
	def __bool__(self):
	return bool(self.flags)
	__nonzero__ = __bool__

What do you think?

@stuartarchibald interesting suggestion. Let me digest and respond later.

[gmarkall]

Having reviewed the Numba documentation, I've also looked at the nvcc docs for "fast math" and think that what this PR is providing is a deviation from this behaviour.

As discussed out-of-band earlier, I think this is bringing Numba's behaviour into line with NVCC. An example of NVCC using a less precise cos implementation with the fast math flag can be seen in: https://github.com/gmarkall/nvcc-fastmath

As a result, one concern I now have with this PR is over whether it would cause regressions in existing code bases due to a behavioural change in an existing configuration i.e. fastmath=True.

I think any code that was passing fastmath=True should not be relying on the usual precision guarantees - so if it does break some existing code, then that code was probably not safe to use with the fastmath flag.

It makes me wonder if low precision/fast math intrinsic functions ought to be activated via a different option, or that fastmath should be a dictionary like:

In light of the above, would you agree that the fastmath dictionary suggestion wouldn't be required now?

[gmarkall]

This is looking great so far!

• The code changes look good to me.

• Tests run fine for me locally, after merging in master, with both NVVM70 (CUDA >= 11.2) and NVVM34 (CUDA < 11.2)

• For the documentation, I'd suggest two changes:

  1. Adding a section on "fast math" underneath the "Numba for CUDA GPUs" section: https://numba.readthedocs.io/en/stable/cuda/index.html - this section should refer to the CPU fastmath section (https://numba.readthedocs.io/en/stable/user/performance-tips.html#fastmath) and note the following differences for the CUDA target:
     • The flag can only be True or False for CUDA
     • The flag also enables some fast approximate trigonometric functions, and provide the list of them.
  2. Updating the jit decorator documentation to explain that fastmath also substitutes the fast versions of some functions: https://numba.readthedocs.io/en/stable/cuda-reference/kernel.html#numba.cuda.jit, and refer to the new section on "fast math" for CUDA GPUs.

[gmarkall]

Noting here that this checks that the fast version doesn't have any FMA instructions in it. I noticed that the fast version makes use of ex2.approx.ftz.f32, which could be checked for instead, but I also think the test as-is is sufficient, because there shouldn't be any FMA instructions in the fast version.

[testhound]

@gmarkall I have made and verified the first of the two changes you requested. I am unclear how to make the change to numba/cuda/decorators.py that will update the documentation; specifically I updated 'decorators.py' but after generating the documentation with 'make html', I do not get a updated local page for 'CUDA Kernel API. Is there another command to update this portion of the documentation or did I update the wrong file?

[gmarkall]

I think sphinx isn't very good at detecting when docstrings change. Every time I change something documentation related, I always run make clean html, which is a bit inconvenient because it takes a while to build from scratch, but I haven't found a better way.

If you do this, does your change now show up?

[gmarkall]

(I guess this test is taking the approach that could have been used for exp)

[gmarkall]

Many thanks for the documentation additions - I have some further suggestions for the docs, which I've pushed in the commit gmarkall@f6e735f. To summarise the suggestions, they are:

• Fixing the docs build by moving the cuda-fast-math label ahead of the section title,
• Add links to documentation for the underlying libdevice functions and NVVM optimizations,
• Explicitly state the math module functions affected by the transformation,
• Refer to the docs in the docstring of the jit decorator for fastmath - it's got to the point where it does a bit much to easily summarise in the docstring.

Feel free to pull in any / all of the changes in the linked commit, or do let me know what you think of the suggestions.

Following resolution of the documentation suggestions, I think this will be all looking good!

[gmarkall]

Many thanks for the update! This looks good to me!

@esc could this have a buildfarm run please? (the typeguard fail is a general issue not related to this branch)

[esc]

Running on Farm as: numba_smoketest_cuda_yaml_13

[esc]

Build farm was fine: numba_smoketest_cuda_yaml_13.

[gmarkall]

@stuartarchibald @sklam Are you happy with this going RTM?

[stuartarchibald]

most of the :func:``math.FOO`` link to the developer/autogen_math_listing references, and a couple to the intended python doc. Not sure what can be done apart for explicit linking. See rendered doc: https://numba--6619.org.readthedocs.build/en/6619/cuda/fastmath.html

[stuartarchibald]

This is a Numba doc problem, not for fixing in this PR. Deferring to ticket #6737.

[stuartarchibald]

@stuartarchibald @sklam Are you happy with this going RTM?

Once #6619 (comment) is resolved, patch looks good. Thanks for working on this @testhound, thanks for reviewing @gmarkall.

[stuartarchibald]

/AzurePipelines run

[azure-pipelines]

Azure Pipelines successfully started running 1 pipeline(s).

[stuartarchibald]

@testhound please could you resolve the conflicts when you have a moment? Many Thanks.

[testhound]

@testhound please could you resolve the conflicts when you have a moment? Many Thanks.

@stuartarchibald I just resolved the conflicts.

[stuartarchibald]

Thanks for resolving conflicts, looks good.