xKg<ddlZddlmZddlmZmZmZmZddlm Z ddlm Z ddl m Z m Z ddlZddlZddlmZmZddlmZefd Ze d Gd d e Ze d Gd de Zedk(rej2yy)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]E}t|D]5}d|||f<t|D]}|||fxx|||f|||fzz cc< 7GyNrshaperange) ABCmnpijks g/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_gufunc.py matmulcorez*_get_matmulcore_gufunc..matmulcores|ww1ww1qA1X!Q$qAadGqAw1a400G")rr)dtyper!s r _get_matmulcore_gufuncr$sM$uQT{E!Q$Kq!t=>% 1 1 r"z&ufunc API unsupported in the simulatorceZdZdZdZdZdZdZdZdZ dZ d Z d Z d Z d Zd ZdZdZdZdZdZdZdZdZdZdZy)TestCUDAGufuncct}d}tj|dzdztjj |dd}tj|dzdztjj |dd}|||}tj ||}|j tj||yNr#r$nparangerreshapematmul assertTrueallcloseselfgufunc matrix_ctrrrGolds r test_gufunc_smallz TestCUDAGufunc.test_gufunc_small!s') IIi!ma'rzz : B B9aCD F IIi!ma'rzz : B B9aCD F 1aLyyA  At,-r"ct}d}tj|dzdztjj |dd}tj|dzdztjj |dd}t j |}|||j}tj||}|jtj||yr() r$r.r/rr0r to_device copy_to_hostr1r2r3)r5r6r7rrdBrr8s r test_gufunc_auto_transferz(TestCUDAGufunc.test_gufunc_auto_transfer/s') IIi!ma'rzz : B B9aCD F IIi!ma'rzz : B B9aCD F^^A  1bM & & (yyA  At,-r"ct}d}tj|dzdztjj |dd}tj|dzdztjj |dd}|||}tj ||}|j tj||y)Nr)r*r+r,r-r4s r test_gufunczTestCUDAGufunc.test_gufunc?s') IIi!ma'rzz : B B9aCD F IIi!ma'rzz : B B9aCD F 1aLyyA  At,-r"ct}d}tj|dzdztjj dddd}tj|dzdztjj dddd}|||}tj ||}|j tj||y)Ndr)r*r+r,r-r4s r test_gufunc_hidimz TestCUDAGufunc.test_gufunc_hidimMs') IIi!ma'rzz : B B1b!Q O IIi!ma'rzz : B B1b!Q O 1aLyyA  At,-r"ctt}tjj ddd}tjj dd}tj ||}|||}tj j||||tj|d}tj j||y)Nr+ )rGrI) r$rr.randomrandnr1testingassert_allclosetile)r5r6XYgoldres1res2s r test_gufunc_new_axisz#TestCUDAGufunc.test_gufunc_new_axisYs'g6 IIOOB1 % IIOOAq !yyAa| ""4.aJ/0 ""4.r"ct}d}tj|dzdztjj |dd}tj|dzdztjj |dd}t j }t j||}t j||}t jd|j|}|||||}|j| } |jtj||} |jtj| | y) 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') IIi!ma'rzz : B B9aCD F IIi!ma'rzz : B B9aCD F ^^Av & ^^Av &   \ P B6 2 OO6O *yyA  At,-r"c,tttddtddgddd}tjdtjdz}tj |}|||tj j||y) N(x)->(x)r rcHt|jD] }||||< yNrsizerrrs r copyz&TestCUDAGufunc.test_copy..copy#166]t!#r"rGr+rIrWrrrr.r/ zeros_likerLrMr5rdrrs r test_copyzTestCUDAGufunc.test_copysy d71:wqz23" $  $  IIb +a / MM!  QA ""1a(r"c&ttddtddfgddd}tjdtjdz}tj|}||||j tj ||y) Nr^r rcHt|jD] }||||< yr`rarcs r rdz9TestCUDAGufunc.test_copy_unspecified_return..copyrer"rGr+rIrf)rrr.r/rhr2r3ris r test_copy_unspecified_returnz+TestCUDAGufunc.test_copy_unspecified_returns| wqz71:./" $  $  IIb +a / MM!  QA  Aq)*r"c6tttddtddgddd}tjdtjdz}tj |}||||j tj||y) Nr^r rcHt|jD] }||||< yr`rarcs r rdz*TestCUDAGufunc.test_copy_odd..copyrer" r+rIrf)rrrr.r/rhr2r3ris r test_copy_oddzTestCUDAGufunc.test_copy_odds| d71:wqz23" $  $  IIb +a / MM!  QA  Aq)*r"c ntttddddftddddfgddd}tjdtjj ddd z}tj |}||| |jtj||y) Nz(x, y)->(x, y)r rct|jdD]+}t|jdD]}|||f|||f<-y)NrrIrr)rrxys r copy2dz*TestCUDAGufunc.test_copy2d..copy2dsG1771:&qwwqz*A1gAadG+'r"r+r,rIrf) rrrr.r/r0rhr2r3)r5rwrrs r test_copy2dzTestCUDAGufunc.test_copy2ds d71a4='!Q$-89%" $ & $ & IIb + 3 3Aq 9A = MM! qa  Aq)*r"cttttddtddgdddtjfd}t j djd}t j|}d}|jt|5|d ||dddy#1swYyxYw) Nz(n)->(n)r rcNt|jdD] }||||< yrrt)rbrs r gufunc_copyzDTestCUDAGufunc.test_not_supported_call_from_jit..gufunc_copys(1771:&t!'r"c||Sr`)rr}r~s r cuda_jitzATestCUDAGufunc.test_not_supported_call_from_jit..cuda_jitsq!$ $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 d58U1X./ 0  0   %  % IIi ' ' 0 MM! 3  # #K 5 HTN1a 6 5 5s  B..B7c tdgddd}tjjdj d}tjjdj d}tj |j dj d}td d 5tjd 5}|||||j|djt|jd t|dj|jdt|djddddddy#1swYxYw#1swYyxYw)N(void(float32[:], float32[:], float32[:]) (n),(n)->(n)r rc^|jd}t|D]}||||z||<yrrar}distlenrs r numba_dist_cudazMTestCUDAGufunc.test_inefficient_launch_configuration..numba_dist_cuda4''!*C3ZA$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 & IINN9 % , ,Y 7 IINN9 % , ,Y 7xx #**95 :A >((51d+  10GH k3qt||+<= os1Q4<55? >s%/E!A>EE!E E!!E*c0tdgdddd}tjjdj dj d }tjjdj dj d }tj |}td d 5tjd 5}|||||jt|d ddddddy#1swYxYw#1swYyxYw)NrrTr )nopythonrc^|jd}t|D]}||||z||<yrrrs r numba_dist_cuda2zLTestCUDAGufunc.test_efficient_launch_configuration..numba_dist_cuda2rr"ir)ir)rrIrr) rr.rJrrr0rhrrrrr)r5rrr}rrs r #test_efficient_launch_configurationz2TestCUDAGufunc.test_efficient_launch_configurations @A#d6 C & C & IINN: & - -i 8 GK  IINN: & - -i 8 GK  }}Q :A >((5 At,  Q+6? >55? >s$2D  &D/D D D  Dc `d}tttddtddgddd||jt5}tttddtddgddd|ddd|j dt jy#1swY/xYw)Ncyr`rrrs r fooz.TestCUDAGufunc.test_nopython_flag..foo r"r^r T)rrFznopython flag must be True)rrr assertRaises TypeErrorrr exception)r5rraisess r test_nopython_flagz!TestCUDAGufunc.test_nopython_flags  # T'!*gaj12Jv! ##& (  y )V 7Kgaj'!*56 % 77: <* 5s6;K;K7LM* )s /B$$B-c d}|jt5}tttddtddgdddd|dddd}t j }|j|dt|||t|djjd}|Dcgc]}|jd }}|jtd d gt|y#1swYxYwcc}w) Ncyr`rrs r rz.TestCUDAGufunc.test_invalid_flags..foorr"r^r TF)rwhat1ever2z/The following target options are not supported:,z'" rr) rrrrrrrrrstripsplitset)r5rrheadritemsrs r test_invalid_flagsz!TestCUDAGufunc.test_invalid_flagss   y )V @Kgaj'!*56 %T @@C E*A&""# Zc$i$/CIJ%%'--c2*/0%Q%0 gw/0#e*=* )1s0C65D6C?cdtttddtddgddd}tjdtjx}}|j t 5}||||dddd}|jtj|y#1swY1xYw) Nr^r rcyr`r)inprWs r rz2TestCUDAGufunc.test_duplicated_output..foos r"rGr+rfz()r rc`d}tt|D]}|||||zz }||d<yr)rr)rurvrsrs r gu_reducez1TestCUDAGufunc.check_tuple_arg..gu_reduces:A3q6]QqTAaD[ #AaDr"rI)axis)rrr.sumasarrayrL assert_equal)r5rr}rrexpecteds r check_tuple_argzTestCUDAGufunc.check_tuple_args| wqz71:wqz:;]" $  $  aO66"**Q-"**Q-7a@ !,r"c0d}d}|j||y)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}|j||y)NPointrurvzrrrrrrrrrrrr)rr)r5rrr}s r test_tuple_of_namedtuple_argz+TestCUDAGufunc.test_tuple_of_namedtuple_arg2s\7O4 SC3 ' SC3 ' ) SC3 ' SC3 ' ) Q"r"ctjdtjdf}tjdtjdf}|j||y)Nrrrr)r.rrrs r test_tuple_of_array_argz&TestCUDAGufunc.test_tuple_of_array_arg:sP ZZ ( ZZ ( * ZZ ( ZZ ( * Q"r"cPt}|j|jdy)Nr!)r$r__name__)r5r6s r test_gufunc_namezTestCUDAGufunc.test_gufunc_nameAs') ,7r"c $|jt5}tttddtddgddd}dddt j }|j d||j d|y#1swYCxYw)Nz(m)->(m)r rcyr`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 %a%(34j P Q * ",, CSI 3S9* )s .BBcttddtddtddfgddd}tjd}|j t 5}||dddt j}|jd||jd||jd||j t 5}|||||dddt |j}|jd||jd||jd |y#1swYxYw#1swYaxYw) Nz (m),(m)->(m)r rcyr`rrs r rz;TestCUDAGufunc.test_incorrect_number_of_pos_args..fP r"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 uQxq5845#F 4  4 iil  y )R cF*",, =sC 137 2C8  y )R c3S !*",, =sC 137 3S9* )* )s D) D5)D25D>N)r __module__ __qualname__r9r>rArErTr\rjrmrqrzrrrrrrrrrrrrrrr"r r&r&sw .. . . /.0 ) + + +!(B&,& N > 5 -###8::r"r&c*eZdZdZdZdZdZdZy)TestMultipleOutputsc tttddtddtddgddd}tjdtjdz}tj |}tj |}||||tj j||tj j||y)N (x)->(x),(x)r rcXt|jD]}||||<||||<yr`rarrrrs r rdzKTestMultipleOutputs.test_multiple_outputs_same_type_passed_in..copyms0166]t!t!#r"rGr+rIrg)r5rdrrrs r )test_multiple_outputs_same_type_passed_inz=TestMultipleOutputs.test_multiple_outputs_same_type_passed_inls d71:wqz71:>?#" $  $  IIb +a / MM!  MM!  Q1  ""1a( ""1a(r"c tttddtddtddgddd}tjdtjdz}tj |}tj |}||||tj j||tj j|dz|y) Nrr rc^t|jD]}||||<||dz||<yNr)rars r copy_and_doublezRTestMultipleOutputs.test_multiple_outputs_distinct_values..copy_and_double~4166]t!tax!#r"rGr+rIr)rgr5rrrrs r %test_multiple_outputs_distinct_valuesz9TestMultipleOutputs.test_multiple_outputs_distinct_values|s d71:wqz71:>?#" $  $ IIb +a / MM!  MM! 1a  ""1a( ""1q5!,r"c Ztttddtddtddgddd}tjdtjdz}||\}}tj j ||tj j |dz|y) Nrr rc^t|jD]}||||<||dz||<yrrars r rzLTestMultipleOutputs.test_multiple_output_allocation..copy_and_doublerr"rGr+rIr))rrrr.r/rLrMrs r test_multiple_output_allocationz3TestMultipleOutputs.test_multiple_output_allocations d71:wqz71:>?#" $  $ IIb +a /q!1 ""1a( ""1q5!,r"c tttddtddtddgddd}t j dtjdz}t j |}t j |tj}||||tjj||tjj|t jdz|y) Nrr rc^t|jD]}||||<||dz||<y)Nrrars r copy_and_multiplyzJTestMultipleOutputs.test_multiple_output_dtypes..copy_and_multiplys4166]t!tcz!#r"rGr+rIr) rrrrr.r/rhrLrM)r5r rrrs r test_multiple_output_dtypesz/TestMultipleOutputs.test_multiple_output_dtypess d58U1Xwqz:;#" $ " $ " IIb )A - MM!  MM!2:: .!Q" ""1a( ""1rzz##6:r"cttddtddtddtddfgddd}tjd}|j t 5}||dddt j}|jd||jd||jd||j t 5}||||||dddt |j}|jd||jd||jd |y#1swYxYw#1swYaxYw) Nz(m),(m)->(m),(m)r rcyr`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 uQxq58U1X>?' 8  8 iil  y )R cF*",, =sC 137 2C8  y )R c3S# &*",, =sC 137 3S9* )* )s D2 D>2D;>EN)rrrrrrr rrr"r rrjs) -" -;":r"r__main__)numpyr. collectionsrnumbarrrrrr numba.cuda.testingr r unittestrnumba.core.errorsr r numba.tests.supportrr$r&rrmainrr"r rs"//<B/")  9:H:\H:;H:V 9:Y:,Y:;Y:x zHMMOr"