VpfE <dZddlmZddlmZmZmZddlmZm Z ddl Z ddl m Z ddl Z ddl mZddl mZdd lmZdd lmZdd lmZdd lmZdd lmZddlmZddlmZddlmZddlmZddlmZddlmZdd l mZ!ddl"m#Z#ddl$m%Z&ddl$m'Z(ddl)m*Z*m+Z+m,Z,m-Z-m.Z.m/Z/ddl0m1Z2ddl1Z3e*e4cZ4Z5e+e6cZ6Z7e!j8Z8e!j9Z9e!j:Z:e!j;Z;e!j<ZZ>e!j?Z?e!j@Z@ejAdZBdeB_CdZDdZEdZFd{dZGd|d#ZHd$ZId%ZJd&ZKd'ZLd}d4ZMeBNeLd~d5ZOeBPeOdd6ZQeQejReB<dd?ZSddDZTddFZUddGZVeVejWeB<ddMZXddOZYeYejZeB<e/ddVZ[ddYZ\ej]dZej^d[d\d]^Z_ddaZ`dddZaejbeBeaddiZcecejdeB<djZeeeejfeB<ddke@e@id\d\dddl ddyZg ddzlhmiZjdS#ek$rdZjYdSwxYw)z-Module for calling pallas functions from JAX.) annotations)CallableIterableSequence)partialreduceN)Any)api_util)lax)ad_util)checkify)config)core)effects) linear_util) tree_util)ad)batching)mlir) partial_eval)uninitialized_value) discharge)utils)safe_mapsafe_zip split_list tuple_insertunzip2weakref_lru_cache pallas_callTcH||J|S|Jtd|D}tj|||}ttjt |tj|tj}tj||S)Nc3VK|]$}tj|tjV%dSdtypeNjnpasarrayint32.0ss [/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/_src/pallas/pallas_call.py z'_maybe_dynamic_slice..J3GGCK333GGGGGG) slice_sizesr$) tupler dynamic_slicenparangelenarraybool_squeeze) start_idx block_shapevalue is_indexingoutput squeeze_dimss r-_maybe_dynamic_slicer@Es     L  GGYGGGGG)  UI; G G G&ry[!1!12228KBD(4L4L4LMNN, V\ * **r0c ||J|S|Jtd|D}tdt|D}tj|||}|j|ksJtj|||S)Nc3VK|]$}tj|tjV%dSr#r&r*s r-r.z._maybe_dynamic_update_slice..Vr/r0c3$K|] \}}||V dSNr+ibs r-r.z._maybe_dynamic_update_slice..Ws>##tq! !#######r0)r2 enumerater broadcast_in_dimshapedynamic_update_slice)r:r;r<updater=broadcast_dimss r-_maybe_dynamic_update_slicerOPs     M  GGYGGGGG)##y'='=#####.   ^ D D&  $ $ $ $ !% ; ;;r0c&tdt|j|D}||jkrYtdt||jD}td|j}t j|||}|S)ayPads values so the shape evenly divides into block dimensions. For example, if values has a shape of (33, 2, 5) with a block_shape of (32, 2, 4), this function will pad the value of shape to (64, 2, 8). Args: value: Array to be padded. block_shape: Block shapes to use for padding. If None, no padding will be performed. Returns: A padded array. c38K|]\}}|dz |zdz|zVdSNrE)r+vrHs r-r.z1_pad_values_to_block_dimension..lsI!%AA!|a1r0c3*K|]\}}d||z fVdS)rNrE)r+arHs r-r.z1_pad_values_to_block_dimension..ps.JJ41aq!A#hJJJJJJr0rErKr%)constant_values)r2ziprKrr%r'pad)r<r; padded_shape pad_width pad_values r-_pad_values_to_block_dimensionr^]s),U[+)F)F,U[  JJ3|U[+I+IJJJJJI#"EK@@@I GE9i @ @ @E ,r0returntuple[jax.Array, ...]cLd|D}td|DS)Nc3>K|]}tj|VdSrDjax_coreraise_to_shapedr+xs r-r.z+_initialize_scratch_vals..vs-FF18+A..FFFFFFr0c3JK|]}t|j|jVdSrD)rrKr%r+rVs r-r.z+_initialize_scratch_vals..ws1LL"17AG44LLLLLLr0)r2) scratch_avalss r-_initialize_scratch_valsrkus2FF FFF- LLmLLL L LLr0block_mappings_outputIterable[BlockMapping]Sequence[jax.Array]cd|D}g}t|D]b\}}||vr"||||+|t|jj|jjc|S)Nci|]\}}|| SrErE)r+krTs r- z+_initialize_output_vals..|s 2 2 2TQAq 2 2 2r0)rIappendrarray_shape_dtyperKr%)rl input_argsinput_output_aliasesoi_map output_valsrGbms r-_initialize_output_valsrzys 3 21 2 2 2&+.//JJeaF{{F1I.////,R-A-G-/-A-GIIJJJJ r0ct|dr4t|jdr|j|jS|S)aConverts logical dtypes into JAX dtypes for interpret mode. This function is used to convert device-specific dtypes that have no corresponding equivalent in JAX/XLA into a type that can be executed by the XLA interpreter (e.g. TPU semaphores -> int32). _rulespallas_interpret_element_aval)hasattrr|r}r%r$s r- _logical_to_interpret_mode_dtypersKeXC el;<<C < 5 5e < < BB ,r0c,t|tjr*t|j}||St|t jr5t|j }t j|j ||j S|S)z*Logical to interpret mode aval conversion.) inner_aval) weak_type) isinstance pallas_coreAbstractMemoryRef$_logical_aval_to_interpret_mode_avalrrMrd ShapedArrayrr%rKr)avalr inner_dtypes r-rrsk344.5doFFJ ;;*; - --h*++S24:>>K   K4> R R RR +r0c2g}d}ttt||D]M\}}tj||dz|}||k}|tj|d|Nt t|S)NTrSr)reversedlistrYr'wherersr2)gridindices next_indicescarrydim_sizeindexrGs r-_get_next_indicesrs, %!$s4'9'9":":;;00oh %E**A ME %A..//// x %% & &&r0cJrDrE)argskwargss r-_pallas_call_implrs,r0jaxprjax_core.Jaxprnamestrdebugboolrvtuple[tuple[int, int], ...] grid_mapping GridMappingcompiler_paramsr c f !"#$~~t|jg\}}t|tfdjDt dJ5tj|d\dddn #1swxYwY|rttj ||}|j #|t#d} |jj} d| D} t!| $g} t#t%j| |jD]\} }t+|jt.jrn|jj}|`t5d|DrG|rt7dt9d| j}t=j| |d|D} | | djD d t# jDtCtD| } | #$t| t|z!tIj%d ftz}rtMtHj'"nd ""fd } !#$f d }t=j(||tIj%d |g| R^}}} | t| t| t|z}g}t#|j D]\}}t+|jt.jry|jj}|kt5d|DrR|rt7dt#|\}}dt#|j)|D}t=j*|||}|j)|j+j)kr(t=j*|d|j,z|j+j)}| ||S)Nc3VK|]#}|tjur|ntV$dSrD)rdynamic_grid_dimnext)r+rVdynamic_grid_args_iters r-r.z._pallas_call_impl_interpret..sT  K0 0 0aa & ' 'r0rEcg|] }|j SrErr+rTs r- z/_pallas_call_impl_interpret..s222a16222r0c3"K|] }|dkV dSrrNrEr+ps r-r.z._pallas_call_impl_interpret..&$B$BQQ&[$B$B$B$B$B$Br0z$Padding with aliasing not supported.rWcg|] }g|dR SrrErs r-rz/_pallas_call_impl_interpret..s""<"<"..s* < <A# # < < < < < .sA  < ..s.DD41aaQDDDDDDr0)r2rYr;)r+iidrys r-rz/_pallas_call_impl_interpret..sZ #rkdd DD3sBN+C+CDDD D Dr0rrSc|^}}|kSrDrE)rrG_num_iterationss r-condz)_pallas_call_impl_interpret..conds EA ~ r0c  |^} }tfdtt D}t|g\}}t j|5 fdjD}dddn #1swxYwYtt| |}t j|5tj tt|ztzksGJtj tt|tftj g||R}dddn #1swxYwYt|j g\}} } tt| || } |dzt g| | RS)Nc3`K|](\}\}}|jvtj||V)dSrD) vmapped_dimsrGridAxis)r+dimidxrHrs r-r.z<_pallas_call_impl_interpret..body..sR C#q l/ / / S!$$ / / / /r0c4g|]}|dn |jgRSrD)compute_start_indices_interpret)r+ryloop_idxscalarss r-rz=_pallas_call_impl_interpret..body..sH111*$$"D""DX"XPW"X"X"X111r0rS)r2rIrYrrgrid_envblock_mappingsmapr@r6invarsrd eval_jaxprnum_index_operandsrOr)rrG carry_blockslocal_grid_envcarry_consts_insscratch start_indicesblocksr out_inout out_scratch out_carryr block_shapesdischarged_constsdischarged_jaxprrris_indexing_dimnum_inout_blocksrscratch_valuess @r-bodyz)_pallas_call_impl_interpret..bodys!&Ax,&s8T':':;;N !+<:J9K L Lg  n - -1111111 /111m111111111111111%}l!?44F  n - - 6 6 !( ) )S\\CKK-G# KK. % & & g,, f++ n    "#35F66$*6-4666f 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6!+02BC!E!EAy+/ $iBBI E$T844 &  &$ & &&s%&BB B >B$E..E25E2c3"K|] }|dkV dSrrErs r-r.z._pallas_call_impl_interpret..rr0cg|] \}}||z  SrErE)r+r,rs r-rz/_pallas_call_impl_interpret.."s BBB41aQBBBr0r)-rnum_dynamic_grid_boundsiterr2rr trace_envstate_dischargedischarge_stateprintrzrlslice_index_opsr6rslice_scratch_opsrkrY itertoolschainrr indexing_moder UnblockedpaddinganyNotImplementedErrorrr%r rZrsrr^extendr'r)rmultiply while_looprKslicertndim)%rrrrvrrrdynamic_grid_argsout block_argsscratch_invarsrjrrgryrr]grid_start_indicesrrrout_out out_nopadopad_lowpad_high limit_indicesrrrrrrrrrrs% ` @@@@@@@@@@r-_pallas_call_impl_interpretrst' \ 1 2T 122    $ $d + + 3 3 3UU*9*I%QS*T*T''UUUUUUUUUUUUUUU   B $&: < <# - .'CLLMM"*< >?.22>222-+M::. %9?:s33\5PQQea""K$9::> (g  $B$B'$B$B$B!B!B   L#$JKK K'b@@@ GAy"<" D39Q<??  ea""K$9::1 (g  $B$B'$B$B$B!B!B   L#$JKK KMBB3qw+A+ABBB Ia- 0 0w"&,,, )Ataf}b&:&@ A Aa Q s<B!!B%(B%c>td|jDS)Nc3hK|]-}tj|jj|jjV.dSrD)rdrrtrKr%rs r-r.z-_pallas_call_abstract_eval..,sV==#B$8$>$&$8$>@@======r0)r2rl)ravalsrs r-_pallas_call_abstract_evalr+s3 ==%;=== = ==r0c :|jrtd|jrt|rtdd|D} d|D}| dg|jzz} t j|d} t j| | g\} } | j| j c}\t|j t||jt|g\}}}}g||||R}g}|j D]j}t|tjr9|||j |j}||k|||}|rt)|t|jt|g\}}g||||R}|||jd z|jd z }t/jg||R||d |||d|d }t|t|d zg\}}||fS) Nz.interpret with dynamic grid bounds unsupportedz JVP with aliasing not supported.cDg|]}t|tj SrE)rr Zeror+ts r-rz)_pallas_call_jvp_rule..:s'HHH!*Q 555HHHr0cHg|]}t|tju| SrE)typer rrs r-rz)_pallas_call_jvp_rule..;s* A A AAT!WWGL%@%@a%@%@%@r0TrE) input_index)rrr num_inputs num_outputs_jvp)rrr interpretrrvr)rrrrrd ClosedJaxprr jvp_jaxprrconstsrrr6rrJaxprInputEffectreplacerr rsrrr pallas_call_pbind)primalstangentsrrrvrrrrnonzero_tangentsnonzero_tangents_with_outputs closed_jaxpr jvp_jaxpr_rr primal_refsprimal_out_refs tangent_refstangent_out_refsreffseffin_bmsout_bmsjvp_bmsjvp_grid_mappingout_flat out_primals out_tangentss r-_pallas_call_jvp_ruler-1s)P N O OO$ B @ A AAHHxHHH A A A A A("2dVl>V5V"V%eR00,,|-JBOO-*a"J$5-)RBLW|'?XOBB>+ .> N[ M< M/ M._block_map_functionqsc!-"?"E"/"?"F)#')))G (%%% nnS'""" >>r0)r;rtr6)rindex_map_grid_avalr6in_avalsrpetrace_to_jaxpr_dynamiclu wrap_initr;rr7rtrrjaxShapeDtypeStructrKr%rdrr)rr.rrr2r: idx_avalsblock_mapping_jaxprrrrKnew_block_shapenew_array_shape_dtypers `` r-_batch_block_mappingrGls} .X1N1WX)66)+)B ())9*6*6&FB666666666666666  #%H O);"5#{/ABBO0]4:   '- //  2F ; ;%   ?1F/4  6 66s.A00A47A4rdims"Sequence[int | batching.NotMapped]Btuple[tuple[jax.Array, ...], tuple[int | batching.NotMapped, ...]]cJt|}t|}|D]e\}}||}d||<|tjur tj|||d||<@|dkrt j|||d||<ft |t |fS)a$Broadcast input/output operands. When we have input/output aliasing, since the output will be mapped, we need to make sure to broadcast the input across that dimension if it is not mapped. If the input is mapped, but on a different axis, we tranpose the input to match the output. r)rrr7 broadcastr'moveaxisr2) rrHrvr.args_dims_r rrs r-_broadcast_input_output_aliasesrPs t**% t**%,CCnk1  CE+ h!!!#-eK.@)QOOeK <[(93BBeK uuU|| ##r0rc N stddtD\ t \ fdjD} df d } tjd | | d } | dt| zfS)aBatch the pallas_call by calling it in loop over the batch size. This function provides a fallback implementation of batching a pallas_call for the cases in which adding a batch dimension to the pallas grid is not supported. This is currently the case when the batched dimension corresponds to a dynamic axis or a scalar prefetch argument. This implementation builds a HLO loop that dynamic_slices the inputs according to the current iteration index and dynamic_updates an (initially empty) output allocation. 'vmapping pallas_call with no arguments.cJh|] \}}|tju|j|!SrErr7rK)r+argrs r- z,_batch_with_explicit_loop..s: #s H' ' ' in ' ' 'r0rvr.c g|];}tjt|jjd|jj.sY  i R17IFF*0222r0 batch_index jax.Arraystatelist[jax.Array]r_c  g}t D]s\}}|tjur||)|t jt j||d||ttj |  d}t|D]1\}}t j ||||d||<2|S)NrS)operand start_index slice_sizeaxisrbrrrrvrrrr) rYrr7rsr'r9rAr dynamic_slice_in_dimrrrIdynamic_update_index_in_dim)rZr\ batch_argsrUr batch_outrGbatch_out_arrayrrrrHrrvrrrs r-rz'_batch_with_explicit_loop..bodys.JdOO  S # # ## K,, +  -     "  !1'   I( 22?44 (   5eAhh Lr0rF)unrollr)rZr[r\r]r_r])rrYrPrlrAr fori_loopr6) rrHrrrrvrrr initial_staterresultr.s ````````` @r-_batch_with_explicit_looprns . I G H HH$oo,9 /  / *$2- &&&&&&&&&&&&&&P 7  Q 4u  M M& F # ##r0c dd} dt||D\} | dkrFt| ||}tj||||||||d } d | Dd t | zfSt ||jg\} }t ||jg\} }td t| | Drt| | | } n8td | Drt| |z| |z||||||| S ~ |j rt ||j g\}}t ||j g\}}tdt||Dr=t| ||}tj gt |z}g||R}g||R}nt||z||z||||||| S|std|j}d|jD}t#|||| \}}t%|dg|jzz}|j }|j}||t ||z }tt+t,|| |||d|}|j|j\}}|rJt4jf|z}t9j|if\}}|| g|jRtA|||d tAd|j!Dz}tjg| |R|d||||||d } | d t | zfS)Nrgr[bdimr1r_cN|tjur|Stj||S)Nrc)rr7r'r9)rgrps r-_maybe_squeeze_out_bdimz;_pallas_call_batching_rule.._maybe_squeeze_out_bdims, x""" h ;qt $ $ $$r0cJh|] \}}|tju|j|!SrErT)r+rgds r-rVz-_pallas_call_batching_rule..s9111tq!H/// ///r0rSrdc8g|]}tj|dSr)r' expand_dimsrfs r-rz._pallas_call_batching_rule..)s$ / / /aCOAq ! ! / / /r0rc3ZK|]&\}}|tjup|j|dkV'dSrRrTr+rUrps r-r.z-_pallas_call_batching_rule..3sR   #t h!!9SYt_%9      r0c32K|]}|tjuVdSrD)rr7)r+rps r-r.z-_pallas_call_batching_rule..:s* I It4x* * I I I I I Ir0) rrHrrrrvrrrc3ZK|]&\}}|tjup|j|dkV'dSrRrTrxs r-r.z-_pallas_call_batching_rule..RsR   C ##;sy!';      r0rRcg|] }|j SrErrs r-rz._pallas_call_batching_rule..ls ( ( (a16 ( ( (r0rWrc3 K|] }|dzV dSrRrEris r-r.z-_pallas_call_batching_rule..s&II!AIIIIIIr0)rrindex_map_avalsindex_map_treerbatched_)rgr[rpr1r_r[)"rYrrrr6rrallrrrnrrr7rrrrPrrnum_scratch_operandsrrGr~ unflattenr}rr;r tree_flattenrrr2r)rrHrrrrvrrrrrr.rrdynamic_grid_dims scalar_args scalar_bdimsbdimsrrall_dimsrravals_to_batchbatched_block_mappingsindex_map_tree_argsindex_map_tree_kwargsbatched_index_map_argsbatched_index_map_avalsbatched_index_map_treebatched_grid_mappings r-_pallas_call_batching_rulers%%%%11s4111*)!^^ &d 3 3D    !1'   C 0 /3 / / /C @@' \ 1 2T' \ 1 2T   ,.?@@    !!24E I I7H I I III  %  %  % !1'     $"4,*I)JKKK$TL,K+LMML%    [,77   4k<PPk)*S-=-==l "{ "T " "d $| $e $ $dd'T!e##3)     I G H HH.. ( (5< ( ( (%/ D';y*$$ZZ1# 88 8(#6%: +SZZ:N-NOP. "L)<<!""# 0.s...A...r0c3>K|]}tj|VdSrDrcrfs r-r.z,pallas_call_checkify_rule..s-NNQX5a88NNNNNNr0rEcnt|g\}}}}}}d|D}g|||||}tjjt|ksJt jjjg|R} t| g\} } t||| D]\} } }|| d<|| d<gS)Ncg|] }|d SrrE)r+err_refs r-rzHpallas_call_checkify_rule..checked_kernel_fn..sCCC' CCCr0r)rr6rrrdrrrY)rr in_error_refsinputsout_error_refsoutputsrinput_error_vals jaxpr_args result_flat output_errorsrin_refout_refrr num_err_valsnum_kernel_inputsnum_kernel_outputs num_scalarss r-checked_kernel_fnz4pallas_call_checkify_rule..checked_kernel_fns  l L*< > ? ?WmV^Wg DC]CCCL#KgKKK7KJ }") * *c*oo = = = =%]1@4>@@@K!+ ~>>M1"%~}#6#6fTlgdmm Ir0c&t|tjr+|j}tjd|jz||jSt|t jrtj |d|jzStj |ggS)N)rSrS)r%r) rrdrr%rKrrAArrayr'reshaper7)rUr%s r-_ensure_2d_error_shapez9pallas_call_checkify_rule.._ensure_2d_error_shapes#x+,, ie  !&39"4E,/M;;;; C # # [fsy0 1 11 Yw  r0cVg|]&}tj|tjj'SrE)rr MemorySpaceERRORr+err_vals r-rz-pallas_call_checkify_rule.. sFOOO29#4 {&,..OOOr0Fcheckify_pallasrr)r}r~r mapped_dimswhatrc32K|]\}}|z|zfVdSrDrE)r+rGrrs r-r.z,pallas_call_checkify_rule..3sMKK+1Aqq~q~&KKKKKKr0c3&K|] }|z|fV dSrDrE)r+rGrs r-r.z,pallas_call_checkify_rule..5s=*7*7q}a*7*7*7*7*7*7r0)rrrrvcg|] }|d SrrErs r-rz-pallas_call_checkify_rule..As 0 0 0gGDM 0 0 0r0),rrrr6rr= close_jaxprr_add_placeholder_effectsrAtreeflattenrr rrr2rrrrrrr flatten_fun_nokwargsr?r@ debug_infor> BlockSpecrtree_flatten_with_pathr$_convert_block_spec_to_block_mappingr}r~rrrrrangerrr)5rrrrrvrrrdynamic_grid_boundsrr_jaxprrrr err_in_treeshaped_err_avals input_avalsshaped_input_avalscheckify_in_avalsclosed_kernel_jaxprerror_out_treerrerror_memref_avalshaped_scalar_avals input_aval output_aval scratch_avalretrace_in_avalsjaxpr_flat_avals jaxpr_in_treewrapped_kernel_with_errout_tree_thunkr final_jaxprerror_block_specs error_pathserror_block_mappingsinput_block_mappingsoutput_block_mappingsgrid_mapping_with_errorinput_output_aliases_with_error new_vals_inrmerrorsresults new_errorrrrrrs5 @@@@@r-pallas_call_checkify_rulers (2 \ 1  , .(($wG +$ii#/&&,>NE<99&!]  ( ( 7 7%(**511(K18<</....+X,NN+NNNNN,(,*,u--#L$5r::NN'/'G^[(N;L(N(N(N$M>1NNNNNNNNNNNNNNN:   /1ABB ' 2 2(OO=MOOO?I;(9;MN@P@P<z; G*G->GG(G+6G9EG$-$:;K$L$L!M,4,Il$%%}-6-6)> -}ne=N P P%#L$5r::::5!15::KAr::::::::::::::: #,T4889C@Pz(_trace_kernel_to_jaxpr..\s)FtT)R)Rr0)rmake_abstract_refrEr)path array_avalr}r~rrr)rnum_constant_operandszIThe kernel function in a pallas_call should return None. Found a PyTree: )!r2rrr rr?r@r=rrr> state_utilshoist_consts_to_refsrr6rIrrrr SequenceKeyrdrrKr%r}r~rrsrrtree_structure ValueError)rrrrrwrapped_kernel_funrrrrrrconst_block_mappingsc_idxcconst_block_mappingkernel_out_trees r-_trace_kernel_to_jaxprr Fsa ,A)++,,L'/'Dl3()()$n -^^UM R R%45G5A5JJE1fb . /RRTTTe"&kk'' 9 9(%)N  !$ - -'..0+AGQW==(8'6"     ##$78888!))N/N,2MNN 5*l5<#N$$/ 06666  -* - - . .. uf $$s;DF  FFf str | NonecJ| t|dr|jr|jnd}|S)N__name__func)r~r)r rs r-_extract_function_namerws- \ J// JAJ J1::FD +r0jax_pallas_use_mosaic_gpuJAX_PALLAS_USE_MOSAIC_GPUFz\If True, lower Pallas kernels to the experimental Mosaic GPU dialect, instead of Trition IR.)defaulthelpplatform Exceptionc(td|dS)Nz&Cannot lower pallas_call on platform: z. To use Pallas on GPU, install jaxlib GPU 0.4.24 or newer. To use Pallas on TPU, install jaxlib TPU and libtpu. See https://jax.readthedocs.io/en/latest/installation.html.r)rs r-_unsupported_lowering_errorrs. AxAAA  r0ctxmlir.LoweringRuleContextc |r/ttfi|}tj|d|g|RSd d}d d}d d }tj|d t |||| dt jg|Rd |i|S)NT)multiple_resultsrrin_nodes'mlir.ir.Value | Sequence[mlir.ir.Value]c td)Nz0Only interpret mode is supported on CPU backend.rrrparamss r- cpu_loweringz+_pallas_call_lowering..cpu_lowerings G H HHr0cXttdtj|g|Ri|S)Ntpu)mosaic_tpu_backendrpallas_call_tpu_lowering_ruler!s r- tpu_loweringz+_pallas_call_lowering..tpu_loweringsJ! ' . ..  ;        r0c tjrddlm}nddlm}n#t $rt dwxYw|j|g|Ri|S)Nrpallas_call_registrationgpu)_PALLAS_USE_MOSAIC_GPUr<jax._src.pallas.mosaic_gpur+jax._src.pallas.triton ImportErrorrpallas_call_lowering)rrr"r+s r- gpu_loweringz+_pallas_call_lowering..gpu_lowerings/  %DGGGGGGGCCCCCC /// ' . ../ 8 # 8        s6r )cpur%cudarocmr)rrrr)rrr lower_funlower_per_platformdictr no_effects)rrrr"implr#r(r2s r-_pallas_call_loweringr;sG . 9 9& 9 9D 64>$ 6 6 6s FX F F FFIIII       m!%,*6+7+7"9"9"9"&!(!3 +#+ + + +,5 +$* + + +r0primjax_core.Primitiver"dict[Any, Any]c ~~dS)Nr rE)r<r"s r-_pallas_custom_str_eqn_compactr@s F r0c|5tj|d|i|cdddS#1swxYwYdS)Nr)rr abstract_eval)rr<r"s r-_pallas_call_typecheck_rulerCs  &   , 06  s 599rE) grid_specrin_specs out_specsrvrrrrCallable[..., None] out_shaperDGridSpec | Noner TupleGridrE BlockSpecTreerFdict[int, int]dict[str, Any] | NoneCallable[..., Any]c ( t  i t|||nJ|rtd||turtd||turtd|~~~t j\ t |trt|}tj |\} t| \dDtj   f d} | S)amInvokes a Pallas kernel on some inputs. See `Pallas Quickstart `_. Args: f: the kernel function, that receives a Ref for each input and output. The shape of the Refs are given by the ``block_shape`` in the corresponding ``in_specs`` and ``out_specs``. out_shape: a PyTree of :class:`jax.ShapeDtypeStruct` describing the shape and dtypes of the outputs. grid_spec: An alternative way to specify ``grid``, ``in_specs``, and ``out_specs``. If given, those other parameters must not be also given. grid: the iteration space, as a tuple of integers. The kernel is executed as many times as ``prod(grid)``. See details at :ref:`pallas_grid`. in_specs: a PyTree of :class:`jax.experimental.pallas.BlockSpec` with a structure matching that of the positional arguments. The default value for ``in_specs`` specifies the whole array for all inputs, e.g., as ``pl.BlockSpec(x.shape, lambda *indices: (0,) * x.ndim)``. See details at :ref:`pallas_blockspec`. out_specs: a PyTree of :class:`jax.experimental.pallas.BlockSpec` with a structure matching that of the outputs. The default value for ``out_specs`` specifies the whole array, e.g., as ``pl.BlockSpec(x.shape, lambda *indices: (0,) * x.ndim)``. See details at :ref:`pallas_blockspec`. input_output_aliases: a dictionary mapping the index of some inputs to the index of the output that aliases them. These indices are in the flattened inputs and outputs. debug: if True, Pallas prints various intermediate forms of the kernel as it is being processed. interpret: runs the ``pallas_call`` as a ``jax.jit`` of a scan over the grid whose body is the kernel lowered as a JAX function. This does not require a TPU or a GPU, and is the only way to run Pallas kernels on CPU. This is useful for debugging. name: TO BE DOCUMENTED. compiler_params: TO BE DOCUMENTED. Returns: A function that can be called on a number of positional array arguments to invoke the Pallas kernel. Nz=If `grid_spec` is specified, then `grid` must be `()`. It is zLIf `grid_spec` is specified, then `in_specs` must be `no_block_spec`. It is zMIf `grid_spec` is specified, then `out_specs` must be `no_block_spec`. It is cLg|]!}tj|j|j"SrE)rArBrKr%rfs r-rzpallas_call..)s9///)!'17;;///r0c tj|\}}t|\}}td|D}tdD}t j||||\}}tj|\} } t|t| | \}} } D]\} }| tt|vr)td| d|d| dt|d |tt|vr)td| d|d |dt|d || }||}|j |j ks|j |j krRtd| d|d tj|| d |d tj|d |d t||jg\}}t#jg|| |R| |td}tj|}|S)Nc3bK|]*}tjtj|V+dSrD)rdreget_avalris r-r.z/pallas_call..wrapped../sL..#283DQ3G3GHH......r0c3TK|]#}tj|j|jV$dSrD)rdrrKr%rs r-r.z/pallas_call..wrapped..1sG55 $/AA555555r0)rz+input_output_aliases contains the mapping ':z' with input index z outside the range [0, )z' with output index z' referring to inputz with abstract value z and to outputz! with a different abstract value .)rrrrrrvr)rrrr2rget_grid_mappingrr itemsrr6rrKr%keystrrrrrtree_unflatten) rflat_args_with_pathsin_treein_paths flat_args flat_in_avalsflat_out_avalsrrflat_kernel_avalsrrri_idxo_idxin_avalout_aval index_args rest_argsr*rrrrr flat_out_shapesrDrvrr out_pathsrs r-wrappedzpallas_call..wrapped+s\$-$DT$J$J!' !566Hi..#,.....M55$355555N"-!=w)"-"-L,)2(>|(L(L%~"8 <011>###L%-224477 u eC ..// / / )% ) )% ) )  ) )}%% ) ) )** * eC//00 0 0 *% * *% * *! * *~&& * * *++ +e$g&h (. ( (GMX^,K,K 6% 6 6% 6 6!*!1(5/!B!B 6 6% 6 6&,Yu-=>> 6 6+3  6 6 677 7-L'y<3R2STTJ !) ))),2)5>))$!"#7#=#=#?#?@@')))H  "8X 6 6C Jr0)rGridSpecr no_block_specrunzip_dynamic_grid_boundsrrr2rrrrAjit)r rHrDrrErFrvrrrrflat_out_shapes_with_pathsrkrrirjrs` ` ````` @@@@r-r r sp 4 ( ($Ox33II $  #  # # $ $$}$$  2'/ 2 2 3 33 %%  3'0 3 3 4 44 Hi#.#H#S#S )  4  !i  I)2)I))T)T&h%&@AA)_//-////7////////////// 7/` .r0r*)r_r`)rlrmr_rn) rrrrrrrvrrrrr )rr)rvrrr ) rrr.r/rr0rr1r2r3r_r3) rrnrHrIrvrr.r/r_rJ)rrnrHrIrrrrrrrvrrrrrrr )rrrrrrrvrrrrrrr )rrrrrrr_r) rrrrrrrrrvrrr) rrrrrrrrrr)r rrr r_r)rrr_r)rrrr)r<r=r"r>r_r)r rGrHr rDrIrrJrErKrFrKrvrLrrrrrr rrMr_rN)l__doc__ __future__rcollections.abcrrr functoolsrrrtypingr rAr r jax._srcr r rrrdrrr?rjax._src.interpretersrrrrr=jax._src.pallasrjax._src.pallas.primitivesrjax._src.staterrrr jax._src.utilrrrrrr jax.numpynumpyr'r4r unsafe_maprY unsafe_zipGridrJrlr3rrrK NoBlockSpecrm Primitiverrr@rOr^rkrzrrrrrdef_implrdef_abstract_evalr-primitive_jvpsrGrPrnrprimitive_batchersrr error_checksr r bool_flagbool_envr-rr;register_loweringr@custom_str_eqn_compact_rulesrCcustom_typechecksr jax._src.pallas.mosaicr+r&r0rEr0r-rs43""""""8888888888%%%%%%%% %%%%%%&&&&&&$$$$$$******&&&&&&444444//////::::::777777//////CZCZ  !  ' %  ! ) % ) #"=11 !%  + + + < < <0MMMM       '''B()))==== :;;;8#8#8#8#r$9- 6666D$$$$:Y$Y$Y$Y$xXXXXv.HM* 0 0 0 0OOOO`(Am$.%.%.%.%` *) FO7 ? ? + -+-+-+-+`}&;<<<# %m4 -H=)"&+,+--1GGGGGG\SSSSSSSSsHHH