J/Ph<:ddlZddlmZddlmZmZmZmZddlm Z ddlm Z ddl m Z m Z ddlZddlZddlmZmZddlmZefd Ze d Gd d e Ze d Gd de ZedkrejdSdS)N) namedtuple)voidint32float32float64) guvectorize)cuda)skip_on_cudasim CUDATestCase)NumbaPerformanceWarning TypingError)override_configc tt|ddddf|ddddf|ddddfgddd}|S)Nz(m,n),(n,p)->(m,p)r targetc |j\}}|j\}}t|D]R}t|D]@}d|||f<t|D]'}|||fxx|||f|||fzz cc<(ASdSNrshaperange) ABCmnpijks c/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_gufunc.py matmulcorez*_get_matmulcore_gufunc..matmulcoresw1w1q 1 1A1XX 1 1!Q$q11AadGGGqAw1a400GGGG1 1 1 1)rr)dtyper!s r _get_matmulcore_gufuncr$s|$uQQQT{E!!!QQQ$Kqqq!!!t==>%   11  1 r"z&ufunc API unsupported in the simulatorceZdZdZdZdZdZdZdZdZ dZ d Z d Z d Z d Zd ZdZdZdZdZdZdZdZdZdZdZdS)TestCUDAGufuncct}d}tj|dzdztj|dd}tj|dzdztj|dd}|||}tj||}|tj||dSNr#r$nparangerreshapematmul assertTrueallcloseselfgufunc matrix_ctrrrGolds r test_gufunc_smallz TestCUDAGufunc.test_gufunc_small!s')) Ii!ma'rz : : : B B9aCD F F Ii!ma'rz : : : B B9aCD F F F1aLLyA  At,,-----r"ct}d}tj|dzdztj|dd}tj|dzdztj|dd}t j|}|||}tj||}| tj ||dSr() r$r.r/rr0r to_device copy_to_hostr1r2r3)r5r6r7rrdBrr8s r test_gufunc_auto_transferz(TestCUDAGufunc.test_gufunc_auto_transfer/s')) Ii!ma'rz : : : B B9aCD F F Ii!ma'rz : : : B B9aCD F F^A   F1bMM & & ( (yA  At,,-----r"ct}d}tj|dzdztj|dd}tj|dzdztj|dd}|||}tj||}|tj||dS)Nr)r*r+r,r-r4s r test_gufunczTestCUDAGufunc.test_gufunc?s')) Ii!ma'rz : : : B B9aCD F F Ii!ma'rz : : : B B9aCD F F F1aLLyA  At,,-----r"ct}d}tj|dzdztjdddd}tj|dzdztjdddd}|||}tj||}|tj||dS)Ndr)r*r+r,r-r4s r test_gufunc_hidimz TestCUDAGufunc.test_gufunc_hidimMs')) Ii!ma'rz : : : B B1b!Q O O Ii!ma'rz : : : B B1b!Q O O F1aLLyA  At,,-----r"ctt}tjddd}tjdd}tj||}|||}tj||||tj|d}tj||dS)Nr+ )rGrI) r$rr.randomrandnr1testingassert_allclosetile)r5r6XYgoldres1res2s r test_gufunc_new_axisz#TestCUDAGufunc.test_gufunc_new_axisYs'g666 IOOB1 % % IOOAq ! !yAva|| ""4...vaJ//00 ""4.....r"ct}d}tj|dzdztj|dd}tj|dzdztj|dd}t j}t j||}t j||}t jd|j |}|||||}| | } | tj ||} | tj| | dS) Nr@r)r*r+r,)r@r)r,)rr#stream)outrV)rV)r$r.r/rr0r rVr; device_arrayr#r< synchronizer1r2r3) r5r6r7rrrVdAr=dCrr8s r test_gufunc_streamz!TestCUDAGufunc.test_gufunc_streamhs<')) Ii!ma'rz : : : B B9aCD F F Ii!ma'rz : : : B B9aCD F F ^Av & & ^Av & &  \ P P P VB6 2 2 2 OO6O * *yA  At,,-----r"cVtttddtddgddd}tjdtjdz}tj|}|||tj||dS) N(x)->(x)r rcJt|jD] }||||<dSNrsizerrrs r copyz&TestCUDAGufunc.test_copy..copy416]]  t!  r"rGr+rIrWrrrr.r/ zeros_likerLrMr5rdrrs r test_copyzTestCUDAGufunc.test_copys d7111:wqqqz223" $ $ $   $ $  Ib + + +a / M!   QA ""1a(((((r"cNttddtddfgddd}tjdtjdz}tj|}||||tj||dS) Nr^r rcJt|jD] }||||<dSr`rarcs r rdz9TestCUDAGufunc.test_copy_unspecified_return..copyrer"rGr+rIrf)rrr.r/rhr2r3ris r test_copy_unspecified_returnz+TestCUDAGufunc.test_copy_unspecified_returns wqqqz7111:./" $ $ $   $ $  Ib + + +a / M!   QA  Aq))*****r"cftttddtddgddd}tjdtjdz}tj|}||||tj||dS) Nr^r rcJt|jD] }||||<dSr`rarcs r rdz*TestCUDAGufunc.test_copy_odd..copyrer" r+rIrf)rrrr.r/rhr2r3ris r test_copy_oddzTestCUDAGufunc.test_copy_odds d7111:wqqqz223" $ $ $   $ $  Ib + + +a / M!   QA  Aq))*****r"c tttddddftddddfgddd}tjdtjddd z}tj|}||| |tj||dS) Nz(x, y)->(x, y)r rct|jdD].}t|jdD]}|||f|||f</dS)NrrIrr)rrxys r copy2dz*TestCUDAGufunc.test_copy2d..copy2dsd171:&& & &qwqz**&&A1gAadGG& & &r"r+r,rIrf) rrrr.r/r0rhr2r3)r5rwrrs r test_copy2dzTestCUDAGufunc.test_copy2ds d7111aaa4='!!!QQQ$-889%" $ $ $ & & $ $ & Ib + + + 3 3Aq 9 9A = M!  qa  Aq))*****r"ctttddtddgdddtjfd}t jdd}t j|}d}| t|5|d ||ddddS#1swxYwYdS) Nz(n)->(n)r rcVt|jdD] }||||<dSrrt)rbrs r gufunc_copyzDTestCUDAGufunc.test_not_supported_call_from_jit..gufunc_copys:171:&&  t!  r"c||Sr`)rr}r~s r cuda_jitzATestCUDAGufunc.test_not_supported_call_from_jit..cuda_jits;q!$$ $r"rz#Untyped global name 'gufunc_copy'.*)rIrI) rrrr jitr.r/astyperhassertRaisesRegexr )r5rrr}msgr~s @r test_not_supported_call_from_jitz/TestCUDAGufunc.test_not_supported_call_from_jits< d58U111X../ 0 0 0   0 0   % % % %  % Ii ' ' 0 0 M!  3  # #K 5 5 ! ! HTN1a  ! ! ! ! ! ! ! ! ! ! ! ! ! ! ! ! ! !s5CCCc 4tdgddd}tjdd}tjdd}tj|jdd}td d 5tj d 5}||||| |dj t| d t|dj| dt|djdddn #1swxYwYddddS#1swxYwYdS)N(void(float32[:], float32[:], float32[:]) (n),(n)->(n)r rcl|jd}t|D]}||||z||<dSrrar}distlenrs r numba_dist_cudazMTestCUDAGufunc.test_inefficient_launch_configuration..numba_dist_cudaE'!*C3ZZ & &A$1+Q & &r"rrrCUDA_LOW_OCCUPANCY_WARNINGSrITrecordz Grid sizez low occupancy)rr.rJrandrzerosrrwarningscatch_warnings assertEqualcategoryr assertInstrmessage)r5rrr}rws r %test_inefficient_launch_configurationz4TestCUDAGufunc.test_inefficient_launch_configurations @A#F 4 4 4 & & 4 4 & INN9 % % , ,Y 7 7 INN9 % % , ,Y 7 7x ##**955 :A > > B B(555 B1d+++  10GHHH k3qt|+<+<=== os1Q4.numba_dist_cuda2rr"ir)ir)rrIrr) rr.rJrrr0rhrrrrr)r5rrr}rrs r #test_efficient_launch_configurationz2TestCUDAGufunc.test_efficient_launch_configurations @A#d6 C C C & & C C & INN: & & - -i 8 8 GK  INN: & & - -i 8 8 GK  }Q :A > > , ,(555 ,  At,,,  Q+++ , , , , , , , , , , , , , , , , , , , , , , , , , , , , , , , , ,s6 D7"1D D7D# #D7&D# 'D77D;>D;c d}tttddtddgddd||t5}tttddtddgddd|dddn #1swxYwY|dt |jdS)NcdSr`rrrs r fooz.TestCUDAGufunc.test_nopython_flag..foo Dr"r^r T)rrFznopython flag must be True)rrr assertRaises TypeErrorrr exception)r5rraisess r test_nopython_flagz!TestCUDAGufunc.test_nopython_flags?    # T'!!!*gaaaj112Jv! # # ##& ( ( (  y ) ) .foorr"r^r TF)rwhat1ever2z/The following target options are not supported:,c8g|]}|dS)z'" )strip).0rs r z5TestCUDAGufunc.test_invalid_flags..s"000Q000r"rr) rrrrrrrrrrsplitset)r5rrheadritemss r test_invalid_flagsz!TestCUDAGufunc.test_invalid_flagss     y ) ) EV @Kgaaaj'!!!*556 %T @ @ @@C E E E E E E E E E E E E E E E E E EA&"## Zc$iiZ$///CIIJJ%%''--c2200%000 gw/00#e**=====sAA11A58A5ctttddtddgddd}tjdtjx}}|t 5}||||dddn #1swxYwYd}|t|j |dS) Nr^r rcdSr`r)inprWs r rz2TestCUDAGufunc.test_duplicated_output..foos Dr"rGr+rfz()r rczd}tt|D]}|||||zz }||d<dSr)rr)rurvrsrs r gu_reducez1TestCUDAGufunc.check_tuple_arg..gu_reducesIA3q66]] ! !QqTAaD[ AaDDDr"rI)axis)rrr.sumasarrayrL assert_equal)r5rr}rrexpecteds r check_tuple_argzTestCUDAGufunc.check_tuple_args wqqqz7111:wqqqz:;]" $ $ $   $ $  IaOO6"*Q--"*Q--7a@@@ !,,,,,r"c:d}d}|||dS)N)?@@@@@)?@ @@@@)rr5rr}s r test_tuple_of_tuple_argz&TestCUDAGufunc.test_tuple_of_tuple_arg+s,   Q"""""r"ctdd}|ddd|dddf}|d d d |d d df}|||dS)NPointrurvzrrrrrrrrrrrr)rr)r5rrr}s r test_tuple_of_namedtuple_argz+TestCUDAGufunc.test_tuple_of_namedtuple_arg2s7O44 USC3 ' ' ' USC3 ' ' ' ) USC3 ' ' ' USC3 ' ' ' ) Q"""""r"ctjdtjdf}tjdtjdf}|||dS)Nrrrr)r.rrrs r test_tuple_of_array_argz&TestCUDAGufunc.test_tuple_of_array_arg:s` Z ( ( Z ( ( * Z ( ( Z ( ( * Q"""""r"cXt}||jddS)Nr!)r$r__name__)r5r6s r test_gufunc_namezTestCUDAGufunc.test_gufunc_nameAs+')) ,77777r"c r|t5}tttddtddgddd}dddn #1swxYwYt |j}|d||d|dS)Nz(m)->(m)r rcdSr`r)rurvs r fz.TestCUDAGufunc.test_bad_return_type..fGsr"z+guvectorized functions cannot return valueszspecifies int32 return type)rrrrrrr)r5terrs r test_bad_return_typez#TestCUDAGufunc.test_bad_return_typeEs   y ) ) R %aaa%(334j P P P  Q P                 ", CSIII 3S99999sAA,,A03A0c ttddtddtddfgddd}tjd}|t 5}||dddn #1swxYwYt |j}|d||d||d||t 5}|||||dddn #1swxYwYt |j}|d||d||d |dS) Nz (m),(m)->(m)r rcdSr`rrs r rz;TestCUDAGufunc.test_incorrect_number_of_pos_args..fP  Dr"r,%gufunc accepts 2 positional argumentszor 3 positional argumentsGot 1 positional argument.zGot 4 positional arguments. rrr.r/rrrrrr5rarrrrs r !test_incorrect_number_of_pos_argsz0TestCUDAGufunc.test_incorrect_number_of_pos_argsOs uQQQxqqq5845#F 4 4 4   4 4 ill  y ) ) R AcFFF               ", =sCCC 13777 2C888  y ) ) "R Ac3S ! ! ! " " " " " " " " " " " " " " "", =sCCC 13777 3S99999s$5 B  BBD##D'*D'N)r __module__ __qualname__r9r>rArErTr\rjrmrqrzrrrrrrrrrrrrrrr"r r&r&ss . . .... . . . . . . / / /...0 ) ) ) + + + + + + + + +!!!(BBB&,,,& N N N > > > 5 5 5 - - -#########888::::::::r"r&c,eZdZdZdZdZdZdZdS)TestMultipleOutputsc tttddtddtddgddd}tjdtjdz}tj|}tj|}||||tj||tj||dS)N (x)->(x),(x)r rc`t|jD]}||||<||||<dSr`rarrrrs r rdzKTestMultipleOutputs.test_multiple_outputs_same_type_passed_in..copymsA16]]  t!t!  r"rGr+rIrg)r5rdrrrs r )test_multiple_outputs_same_type_passed_inz=TestMultipleOutputs.test_multiple_outputs_same_type_passed_inls d7111:wqqqz7111:>>?#" $ $ $   $ $  Ib + + +a / M!   M!   Q1  ""1a((( ""1a(((((r"c tttddtddtddgddd}tjdtjdz}tj|}tj|}||||tj||tj|dz|dS) Nrr rcft|jD]}||||<||dz||<dSNr)rars r copy_and_doublezRTestMultipleOutputs.test_multiple_outputs_distinct_values..copy_and_double~E16]]  t!tax!  r"rGr+rIr)rgr5rrrrs r %test_multiple_outputs_distinct_valuesz9TestMultipleOutputs.test_multiple_outputs_distinct_values|s d7111:wqqqz7111:>>?#" $ $ $   $ $ Ib + + +a / M!   M!  1a    ""1a((( ""1q5!,,,,,r"c tttddtddtddgddd}tjdtjdz}||\}}tj||tj|dz|dS) Nrr rcft|jD]}||||<||dz||<dSrrars r rzLTestMultipleOutputs.test_multiple_output_allocation..copy_and_doublerr"rGr+rIr))rrrr.r/rLrMrs r test_multiple_output_allocationz3TestMultipleOutputs.test_multiple_output_allocations d7111:wqqqz7111:>>?#" $ $ $   $ $ Ib + + +a /q!!1 ""1a((( ""1q5!,,,,,r"c tttddtddtddgddd}t jdtjdz}t j|}t j|tj}||||tj||tj|t jdz|dS) Nrr rcft|jD]}||||<||dz||<dS)Nrrars r copy_and_multiplyzJTestMultipleOutputs.test_multiple_output_dtypes..copy_and_multiplysE16]] " "t!tcz! " "r"rGr+rIr) rrrrr.r/rhrLrM)r5r rrrs r test_multiple_output_dtypesz/TestMultipleOutputs.test_multiple_output_dtypess d58U111Xwqqqz::;#" $ $ $ " " $ $ " Ib ) ) )A - M!   M!2: . . .!Q""" ""1a((( ""1rz##6:::::r"c*ttddtddtddtddfgddd}tjd}|t 5}||dddn #1swxYwYt |j}|d||d||d||t 5}||||||dddn #1swxYwYt |j}|d||d||d |dS) Nz(m),(m)->(m),(m)r rcdSr`r)rurvrrs r rz@TestMultipleOutputs.test_incorrect_number_of_pos_args..frr"r,rzor 4 positional argumentsrzGot 5 positional arguments.rrs r rz5TestMultipleOutputs.test_incorrect_number_of_pos_argss  uQQQxqqq58U111X>?' 8 8 8   8 8 ill  y ) ) R AcFFF               ", =sCCC 13777 2C888  y ) ) 'R Ac3S# & & & ' ' ' ' ' ' ' ' ' ' ' ' ' ' '", =sCCC 13777 3S99999s$ BB"BD22D69D6N)rrrrrr rrrr"r rrjs_))) ---" - - -;;;":::::r"r__main__)numpyr. collectionsrnumbarrrrrr numba.cuda.testingr r unittestrnumba.core.errorsr r numba.tests.supportrr$r&rrmainrr"r rs""""""////////////<<<<<<<<BBBBBBBB//////")     9::H:H:H:H:H:\H:H:;:H:V 9::Y:Y:Y:Y:Y:,Y:Y:;:Y:x zHMOOOOOr"