ai-content-maker/.venv/Lib/site-packages/numba/cuda/tests/cudapy/test_libdevice.py

import numpy as np
from numba.core import types
from numba.cuda.testing import skip_on_cudasim, unittest, CUDATestCase
from numba import cuda
from numba.cuda import libdevice, compile_ptx
from numba.cuda.libdevicefuncs import functions, create_signature


def use_sincos(s, c, x):
    i = cuda.grid(1)

    if i < len(x):
        sr, cr = libdevice.sincos(x[i])
        s[i] = sr
        c[i] = cr


def use_frexp(frac, exp, x):
    i = cuda.grid(1)

    if i < len(x):
        fracr, expr = libdevice.frexp(x[i])
        frac[i] = fracr
        exp[i] = expr


def use_sad(r, x, y, z):
    i = cuda.grid(1)

    if i < len(x):
        r[i] = libdevice.sad(x[i], y[i], z[i])


@skip_on_cudasim('Libdevice functions are not supported on cudasim')
class TestLibdevice(CUDATestCase):
    """
    Some tests of libdevice function wrappers that check the returned values.

    These are mainly to check that the generation of the implementations
    results in correct typing and lowering for each type of function return
    (e.g. scalar return, UniTuple return, Tuple return, etc.).
    """

    def test_sincos(self):
        # Tests return of a UniTuple from a libdevice function
        arr = np.arange(100, dtype=np.float64)
        sres = np.zeros_like(arr)
        cres = np.zeros_like(arr)

        cufunc = cuda.jit(use_sincos)
        cufunc[4, 32](sres, cres, arr)

        np.testing.assert_allclose(np.cos(arr), cres)
        np.testing.assert_allclose(np.sin(arr), sres)

    def test_frexp(self):
        # Test return of a Tuple from a libdevice function
        arr = np.linspace(start=1.0, stop=10.0, num=100, dtype=np.float64)
        fracres = np.zeros_like(arr)
        expres = np.zeros(shape=arr.shape, dtype=np.int32)

        cufunc = cuda.jit(use_frexp)
        cufunc[4, 32](fracres, expres, arr)

        frac_expect, exp_expect = np.frexp(arr)

        np.testing.assert_array_equal(frac_expect, fracres)
        np.testing.assert_array_equal(exp_expect, expres)

    def test_sad(self):
        # Test return of a scalar from a libdevice function
        x = np.arange(0, 200, 2)
        y = np.arange(50, 150)
        z = np.arange(15, 115)
        r = np.zeros_like(x)

        cufunc = cuda.jit(use_sad)
        cufunc[4, 32](r, x, y, z)

        np.testing.assert_array_equal(np.abs(x - y) + z, r)


# A template for generating tests of compiling calls to libdevice functions.
# The purpose of the call and assignment of the return variables is to ensure
# the actual function implementations are not thrown away resulting in a PTX
# implementation that only contains the ret instruction - this may hide certain
# errors.
function_template = """\
from numba.cuda import libdevice

def pyfunc(%(pyargs)s):
    ret = libdevice.%(func)s(%(funcargs)s)
    %(retvars)s = ret
"""


def make_test_call(libname):
    """
    Generates a test function for each libdevice function.
    """

    def _test_call_functions(self):
        # Strip off '__nv_' from libdevice name to get Python name
        apiname = libname[5:]
        apifunc = getattr(libdevice, apiname)
        retty, args = functions[libname]
        sig = create_signature(retty, args)

        # Construct arguments to the libdevice function. These are all
        # non-pointer arguments to the underlying bitcode function.
        funcargs = ", ".join(['a%d' % i for i, arg in enumerate(args) if not
                              arg.is_ptr])

        # Arguments to the Python function (`pyfunc` in the template above) are
        # the arguments to the libdevice function, plus as many extra arguments
        # as there are in the return type of the libdevice function - one for
        # scalar-valued returns, or the length of the tuple for tuple-valued
        # returns.
        if isinstance(sig.return_type, (types.Tuple, types.UniTuple)):
            # Start with the parameters for the return values
            pyargs = ", ".join(['r%d' % i for i in
                                range(len(sig.return_type))])
            # Add the parameters for the argument values
            pyargs += ", " + funcargs
            # Generate the unpacking of the return value from the libdevice
            # function into the Python function return values (`r0`, `r1`,
            # etc.).
            retvars = ", ".join(['r%d[0]' % i for i in
                                 range(len(sig.return_type))])
        else:
            # Scalar return is a more straightforward case
            pyargs = "r0, " + funcargs
            retvars = "r0[0]"

        # Create the string containing the function to compile
        d = { 'func': apiname,
              'pyargs': pyargs,
              'funcargs': funcargs,
              'retvars': retvars }
        code = function_template % d

        # Convert the string to a Python function
        locals = {}
        exec(code, globals(), locals)
        pyfunc = locals['pyfunc']

        # Compute the signature for compilation. This mirrors the creation of
        # arguments to the Python function above.
        pyargs = [ arg.ty for arg in args if not arg.is_ptr ]
        if isinstance(sig.return_type, (types.Tuple, types.UniTuple)):
            pyreturns = [ret[::1] for ret in sig.return_type]
            pyargs = pyreturns + pyargs
        else:
            pyargs.insert(0, sig.return_type[::1])

        pyargs = tuple(pyargs)
        ptx, resty = compile_ptx(pyfunc, pyargs)

        # If the function body was discarded by optimization (therefore making
        # the test a bit weak), there won't be any loading of parameters -
        # ensure that a load from parameters occurs somewhere in the PTX
        self.assertIn('ld.param', ptx)

        # Returning the result (through a passed-in array) should also require
        # a store to global memory, so check for at least one of those too.
        self.assertIn('st.global', ptx)

    return _test_call_functions


@skip_on_cudasim('Compilation to PTX is not supported on cudasim')
class TestLibdeviceCompilation(unittest.TestCase):
    """
    Class for holding all tests of compiling calls to libdevice functions. We
    generate the actual tests in this class (as opposed to using subTest and
    one test within this class) because there are a lot of tests, and it makes
    the test suite appear frozen to test them all as subTests in one test.
    """


for libname in functions:
    setattr(TestLibdeviceCompilation, 'test_%s' % libname,
            make_test_call(libname))


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