U hd@sddlZddlmZddlmZmZmZddlm Z ddl m Z m Z ddl mZmZddZd d Zd d Zed GdddeZdZddZedGdddejZeD]ZeedeeeqedkredS)N)types)skip_on_cudasimunittest CUDATestCase)cuda) libdevice compile_ptx) functionscreate_signaturecCs<td}|t|kr8t||\}}|||<|||<dSN)rgridlenrZsincos)scxisrcrr_/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_libdevice.py use_sincos s   rcCs<td}|t|kr8t||\}}|||<|||<dSr )rr rrfrexp)fracexprrZfracrexprrrr use_frexps   rcCs8td}|t|kr4t||||||||<dSr )rr rrZsad)rryzrrrruse_sads  r z0Libdevice functions are not supported on cudasimc@s(eZdZdZddZddZddZdS) TestLibdevicea( 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.). cCsjtjdtjd}t|}t|}tt}|d|||tjt ||tjt ||dS)Nd)dtype ) nparangefloat64 zeros_likerjitrtestingZassert_allclosecossin)selfarrsresZcrescufuncrrr test_sincos,s   zTestLibdevice.test_sincoscCsxtjdddtjd}t|}tj|jtjd}tt }|d|||t |\}}tj ||tj ||dS)Ng?g$@r")startstopnumr#)shaper#r$) r'Zlinspacer)r*Zzerosr7Zint32rr+rrr,assert_array_equal)r/r0ZfracresZexpresr2Z frac_expectZ exp_expectrrr test_frexp8s  zTestLibdevice.test_frexpcCsltddd}tdd}tdd}t|}tt}|d||||tjt||||dS) Nr2sr$) r'r(r*rr+r r,r8abs)r/rrrrr2rrrtest_sadFs    zTestLibdevice.test_sadN)__name__ __module__ __qualname____doc__r3r9rArrrrr!"s r!z{from numba.cuda import libdevice def pyfunc(%(pyargs)s): ret = libdevice.%(func)s(%(funcargs)s) %(retvars)s = ret csfdd}|S)z@ Generates a test function for each libdevice function. csfdd}tt|}t\}}t||}dddt|D}t|jtj tj frdddt t |jD}|d|7}dddt t |jD}n d|}d}||||d } t | } i} t| t| | d } d d|D}t|jtj tj frd d|jD} | |}n|d |jdddt|}t| |\}}|d||d|dS)Nz, cSsg|]\}}|jsd|qS)za%d)is_ptr).0rargrrr osz@make_test_call.._test_call_functions..cSsg|] }d|qS)zr%drrHrrrrrJyscSsg|] }d|qS)zr%d[0]rrKrrrrJszr0, zr0[0])funcpyargsfuncargsretvarspyfunccSsg|]}|js|jqSr)rGty)rHrIrrrrJscSsg|]}|dddqSr r)rHretrrrrJsrr zld.paramz st.global)getattrrr r join enumerate isinstance return_typerTupleZUniTuplerangerfunction_templateexecglobalsinserttuplerZassertIn)r/ZapinameZapifuncZrettyargssigrNrMrOdcodelocalsrPZ pyreturnsZptxZrestylibnamerr_test_call_functionsfsB           z,make_test_call.._test_call_functionsr)rerfrrdrmake_test_callas Brgz.Compilation to PTX is not supported on cudasimc@seZdZdZdS)TestLibdeviceCompilationa: 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. N)rBrCrDrErrrrrhsrhztest_%s__main__)numpyr'Z numba.corerZnumba.cuda.testingrrrZnumbarZ numba.cudarrZnumba.cuda.libdevicefuncsr r rrr r!rZrgZTestCaserhresetattrrBmainrrrrs(    5 J