uki4dZddlmZddlmZmZddlmZddlZddlZddl Z ddl Z ddl m Z ddl mcmZddlmZddlmZdd lmZdd lmZdd lmZdd lmZdd lmZddlm ZddlmZddlmZddlmZddlmZddlm Z ddl!m"Z"ddl!m#Z$ddl%m&Z&ddl'm(Z(dd l)mZ*ddl)m+Z,ddl-m.Z/ddl-m0Z0ddl-m1Z2ddl-m3Z4ddl5m6Z6ddlm7Z8e0jrZ9e0jtZ:e jveZdedjd?e6jedd@ddAZeedjdBZfejdCZgegjdDZhdEZieiejeg<dFZkeke"jeg<dGZm ddIZnejendJKZpe/jegdLZqejdMZrerjdNZsdOZtetejer<dPZueue"jer<e/jerdQZvdddddHdR ddSZwdddTdU ddVZxdddWddXZyddYZz d ddZZ{ejd[Z|dHd\d]Z}e|jd^Z~dHd\ dd_Ze6je|edd`Z ddaZejdbZejdcZdJe_ddZee_deZee_dfdg ddhZejdiZdjZe/jeeeje6jedkZejdlZdHe_dJe_ej"j%eddmZejdnZdoZe/jeedpZGdqdrejZ ddsZdtZejduZdHe_ddvZejdwZdxZe/jeeeej>zdzeeej>zdyfzee eej>zfzZejdzZdJe_ ddejHdd{ dd|Zej dd}Z dd~Zeeje<dZe/jeeejdZdJe_ ddJd ddZejdZ ddZeeje<dZe/jeeddZ ddZejdZdJe_ejdZddZy)zPallas-specific JAX primitives.) annotations)CallableSequence)HashableN)Any) tree_util)ad_util)api_util)core)config) debugging)dtypes)typing)effects) linear_util)pretty_printer)state)util)ad) partial_eval)ir)arith)utils) discharge)indexing) primitives)types)mlir)numpy program_idc.tj|S)aReturns the kernel execution position along the given axis of the grid. For example, with a 2D ``grid`` in the kernel execution corresponding to the grid coordinates ``(1, 2)``, ``program_id(axis=0)`` returns ``1`` and ``program_id(axis=1)`` returns ``2``. The returned value is an array of shape ``()`` and dtype ``int32``. Args: axis: the axis of the grid along which to count the program. axis) program_id_pbindr"s U/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/_src/pallas/primitives.pyr r >s     %%c|jd}tj}|r||jStj}|j |}t jjt|dt|SNr#r") pop pallas_corecurrent_grid_envindex axis_framesizejax_core Primitivebind_with_tracer$dict)trace_paramsr#grid_envframes r&program_id_bind_with_tracer:Lsp F $  ) ) +( D>    "% jj!    + +L%Tt_ UUr'c JtjdtjSNr*r1 ShapedArrayjnpint32r6s r&_program_id_abstract_evalrBY   b#)) ,,r' num_programsc.tj|S)z2Returns the size of the grid along the given axis.r")num_programs_pr%r"s r&rDrD_s   $  ''r'c@|jd}tj}|r||jStj}|j|}|tj ur0t jjt|dt|S|Sr)) r+r,r-r0r/dynamic_grid_dimr1r2r3rFr4)r5r6r7r#r8r9r0s r&_num_programs_bind_with_tracerIcs F $  ) ) +( D>    "% D $ [ ) ))    - -neRSW YY +r'c JtjdtjSr<r=rAs r&_num_programs_abstract_evalrKqrCr'c(eZdZdZdZdZdZdZdZdZ y) AtomicOpTypexchgaddmaxminandorxorN) __name__ __module__ __qualname__XCHGADDMAXMINANDORXORr*r'r&rMrMus% $ # # # # " #r'rM atomic_rmwcB~|j|\}}}}|^} } tj|| }|t|tj k(rd} nS|tj k(rtj} n/|tjk(rtj} n t|td| jDr| j} | D cgc]"} t| t xr | j $}} | D cgc] } t| tr | j n| "}} t#d| D}t%j&|||}t#d|D}||}| ||}t%j(|||}t#d|D}||}natd| jDr?|| j}|j*| jj-| ||}nt|fd t/|d z zz|fScc} wcc} w) Nc ||zSNr*)xys r&z,_atomic_rmw_discharge_rule..s !a%r'c3ZK|]#}t|txs |j %ywrb isinstanceSliceshape.0ss r& z-_atomic_rmw_discharge_rule..%D1*Q  -agg+ -D)+c3XK|]"}t|tr |jnd$ywNrhrir0rks r&rnz-_atomic_rmw_discharge_rule.."O!*Q"6A=O(* slice_sizesc3:K|]}|rdn tdywrbslicerlscalars r&rnz-_atomic_rmw_discharge_rule..sRFE$K7R start_indicesc3:K|]}|rdn tdywrNrzr|s r&rnz-_atomic_rmw_discharge_rule..OVt4Or~c3>K|]}t|t ywrbrhrirks r&rnz-_atomic_rmw_discharge_rule.. 9z!U# # 9rbrs) unflattenstate_dischargetransform_arrayNotImplementedErrorrMrYrZr?maximumr[minimumallindicesrhrirjstarttuplelax dynamic_slicedynamic_update_sliceatsetlen)in_avals out_avals args_tree atomic_type args_flatref transformsvalmaskprev_transformsidxmonoidrrm scalar_dims slice_startsrxout_ones val_indexerx_new out_indexerouts r&_atomic_rmw_discharge_rulers(229=#z3$?C''_=#  L$$$ Fl&&& [[Fl&&& [[F k **D DDkkGELMz!U++;AGG ;MKMDKLqz!U3AGG:LLLOwOOK  l LHRkRRK k C h C  $ $S#\ JEO;OOK ; C 9S[[ 99 ckk C FF3;;  # #F3$4 5E  Gs8}q01 13 66NLs 'H>%Hc|j|\}}}}|jtjdk(r,|tjk7rt d|j d|jtjdtjdtjdtjhvr&t d|j d|jdt|d |iS) Nfloat16z`atomic_z` does not support f16.boolint8int16z` does not support .r) rdtyper?rMrY ValueErrorvaluebfloat16_swap_abstract_eval)rr avals_flatrr6s r&_atomic_abstract_evalrs$$Z0,#q!QYY#))I&&;,:J:J+J x 1 122IJ KKYY ii ii ii ll   ;$$%%8 1E  j >I >>r')rctj||d\}}tj||||f\}}t j |||dS)Nr_)rr)spget_ref_and_transformsr tree_flatten atomic_rmw_pr%) x_ref_or_viewrrrrx_refrrrs r& _atomic_rmwrsX//S,%#// C0NO)Y   I; r'c>t||||tjS)zAtomically exchanges the given value with the value at the given index. Args: x_ref_or_view: The ref to operate on. idx: The indexer to use. mask: TO BE DOCUMENTED. Returns: The value at the given index prior to the aupdate. rr)rrMrXrrrrs r& atomic_xchgrs" S#Dl6G6G r'c>t||||tjS)zAtomically computes ``x_ref_or_view[idx] += val``. Args: x_ref_or_view: The ref to operate on. idx: The indexer to use. mask: TO BE DOCUMENTED. Returns: The value at the given index prior to the atomic operation. r)rrMrYrs r& atomic_addr" S#Dl6F6F r'c>t||||tjS)aAtomically computes ``x_ref_or_view[idx] = max(x_ref_or_view[idx], val)``. Args: x_ref_or_view: The ref to operate on. idx: The indexer to use. mask: TO BE DOCUMENTED. Returns: The value at the given index prior to the atomic operation. r)rrMrZrs r& atomic_maxrrr'c>t||||tjS)aAtomically computes ``x_ref_or_view[idx] = min(x_ref_or_view[idx], val)``. Args: x_ref_or_view: The ref to operate on. idx: The indexer to use. mask: TO BE DOCUMENTED. Returns: The value at the given index prior to the atomic operation. r)rrMr[rs r& atomic_minrrr'c>t||||tjS)zAtomically computes ``x_ref_or_view[idx] &= val``. Args: x_ref_or_view: The ref to operate on. idx: The indexer to use. mask: TO BE DOCUMENTED. Returns: The value at the given index prior to the atomic operation. r)rrMr\rs r& atomic_andrrr'c>t||||tjS)zAtomically computes ``x_ref_or_view[idx] |= val``. Args: x_ref_or_view: The ref to operate on. idx: The indexer to use. mask: TO BE DOCUMENTED. Returns: The value at the given index prior to the atomic operation. r)rrMr]rs r& atomic_orrs  S#Dloo r'c>t||||tjS)zAtomically computes ``x_ref_or_view[idx] ^= val``. Args: x_ref_or_view: The ref to operate on. idx: The indexer to use. mask: TO BE DOCUMENTED. Returns: The value at the given index prior to the atomic operation. r)rrMr^rs r& atomic_xorr'rr' atomic_casc|j|jk7s|j|jk7r td|jr td|jr td|jr tdtj|j|jt j dhfS)Nz1cmp and val must have identical dtypes and shapeszref must be scalar.zcmp must be scalar.zval must be scalar.r)rrjrr1r>r WriteEffect)ref_avalcmp_avalval_avals r&_atomic_cas_abstract_evalr8s ^^x~~%8>>)I H II ^^ * ++ ^^ * ++ ^^ * ++   hnnhnn =@Q@QRS@T?U UUr'c0tj|||S)aPerforms an atomic compare-and-swap of the value in the ref with the given value. Args: ref: The ref to operate on. cmp: The expected value to compare against. val: The value to swap in. Returns: The value at the given index prior to the atomic operation. ) atomic_cas_pr%)rcmprs r&rrEs   3S ))r'cH~~tj||k(||}|ddf|fSrb)r?where)rrrrrnew_vals r&_atomic_cas_discharge_rulerSs/ IIcSj#s +' 4  ##r'max_contiguousc |Srbr*rcr6s r&rere[sr'c |gSrbr*r6rc__s r&rere\sQCr'ctt|ttfs|f}tj |t|S)zTA compiler hint that asserts the ``values`` first values of ``x`` are contiguous. values)rhlistrmax_contiguous_pr%rcrs r&rr^s3 FT5M *YF   qv  77r'c |Srbr*avalr6s r&_max_contiguous_abstract_evalre +r' multiple_ofc |Srbr*rs r&rereksar'c |gSrbr*rs r&rerels!r'clt|tr|fn t|}tj ||S)aTA compiler hint that asserts a value is a static multiple of another. Note that misusing this function, such as asserting ``x`` is a multiple of ``N`` when it is not, can result in undefined behavior. Args: x: The input array. values: A set of static divisors that ``x`` is a multiple of. Returns: A copy of ``x``. r)rhintr multiple_of_pr%rs r&rrns0#63/F9U6]&   Af  --r'c |Srbr*rs r&_multiple_of_abstract_evalr~rr' masked_loadc|j|\}}}}|Jtj||}tj|j |j tjdhfS)Nr) rr,TransformedRefr1r>rjrr ReadEffect)rrr6rrtransformed_refs r&_load_abstract_evalrsp#--j9#z1a   ..sJ?/?00/2G2GH  r'c |j\}tj|jd|j\}}}}t j |g||j}|tjdtj|||g} |U| tjdtjdtjt j||gz } |U| tjdtjdtjt j||gz } tj| S)Nr print_shapes <-  mask=zother=)outvarsrtree_unflattenr7invarsr1pp_varsrpptextrpp_ref_transformspp_varconcat) eqncontextsettingsrdrcrrotherlhsresults r& _load_pp_rulers {{"!(77 jjszz !Zu !gH4I4IJ# ""6"6w:"N O&      g./F       w/0F 6 r'c @|j|\}}}}|j|\}} } } | tj| } tjt j ||||fd|i|tjt j |||| fd|i|fSNr)rr instantiateload_pr%r tree_leaves) primalstangentsrr7 ref_primalrr other_primal ref_tangentr6 other_tangents r& _load_jvpr s/8/B/B7/K,*j$ %.%8%8%B"+q!]'' 6M kk  *j$ !M N   kk  Jm<      r'ctj|tjr%tj|tj|Stj|t j rtj|d|Stj|tjr4tj|tj|j|Stj|tjrtj|d|Stj|t jrtj|d|St|)NrF) r? issubdtypefloatingfullnanr,SEMAPHORE_INTERPRET_DTYPEintegeriinforQrsemaphore_dtyperrjrs r&uninitialized_valuer+s^^E3<<( 88E377E ** ~~e[BBC 88E1e $$ ~~eS[[) 88E399U+// 77 ~~eSXX& 88E5% (( ~~e[889 88E1e $$E""r'Fctd|D}|rtd|D}td|j}tj|||}|S)a DynamicSlice and DynamicUpdateSlice adjust the start index in cases where the requested slice overruns the bounds of the array. This pads the array with uninitialised values such that the requested slice will never overrun. For example, if arr is [1.,2.,3.,4.] and a slice of size 4, start index 2 is requested then the result will be [3.,4.,NaN,NaN] after padding, rather than [1.,2.,3.,4.] from the unpadded array unpad=True performs the inverse operation c3&K|] }d|df ywrr*)rl slice_sizes r&rnz?_pad_values_to_avoid_dynamic_slice_oob_shift..sJ !Z+Jsc34K|]\}}}| | | fywrbr*)rllowhighinteriors r&rnz?_pad_values_to_avoid_dynamic_slice_oob_shift..s-H4T8!D4%(3Hsr*r*)padding_config padding_value)rr+rrpad)rrxunpadr3r4s r&,_pad_values_to_avoid_dynamic_slice_oob_shiftr7s^JkJJ. H8FHHN%BekkB- ''%!/ - /% ,r'T)r6c ~|j|\}}}}t|}|rt|dtjsQt j ||dj} |jtjj| |^} } t| tsJtj|| }td| jDr$| jD]-} t| ts| jdkDs$t!d| j} | D cgc]"} t| t xr | j $}} | D cgc] } t| tr | j"n| "}} t%d| D}t'||}t)j*}t-j.||D cgc]} t1j2| |c} |}t%d|D}||}n2td | jDr|| j}nt ||t1j4|||}d t7|z|fScc} wcc} wcc} w) Nrc3ZK|]#}t|txs |j %ywrbrgrks r&rnz'_load_discharge_rule..rorprsUnimplemented stride support.c3XK|]"}t|tr |jnd$ywrrrtrks r&rnz'_load_discharge_rule..rurvrwc3:K|]}|rdn tdywrrzr|s r&rnz'_load_discharge_rule.. rr~c3>K|]}t|t ywrbrrks r&rnz'_load_discharge_rule.. rrrb)rrrhr NDIndexerrget_transforms_shaperjappendmake_trivial_indexerrrrrristriderrrr7rdefault_int_dtyperrr?astyperr)rrrrr6rrrr ref_shaperrrmrrrrx idx_dtyperrrs r&_load_discharge_rulerHs#!*!4!4Y!?#z4J* :jnh6H6HI**:x{7H7HIIh((==iHI$?C C ## #''_=#D DD [[C Au !((Q,!"ABBCkkGELMz!U++;AGG ;MKMDKLqz!U3AGG:LLLOwOOK 7sK HC((*I   +78aAy !8H O;OOK ; C 9S[[ 99 ckk C  %+ ))D#u %C 3x= # %%+NL 9s)'I%I#I( masked_swapc |j|\}}}}|Jtj||}|j}|j} ||jk7r)t d|jd|jd|d| |jk7rt d| d|jdt j|| tjdhfS)Nz%Invalid shape for `swap`. Ref shape: z. Value shape: z. Transforms: z. z%Invalid dtype for `swap`. Ref dtype: z. Value dtype: r) rr,rrjrrr1r>rr) rrr6rrrrrexpected_output_shapeexpected_output_dtypes r&rrs(22:>#z3   ..sJ?/)//)//cii'  / {; {. B @ cii'  /0E/FG {" &  02GH  r'c n|j\}|jdj|j\}}}}t j |||}t |tjrRtj|tjdtjtj||gStj|g||j}|tjd|tjd|tjdtjtj||g} |U| tjdtjdtjtj||gz } tj| S)Nrrrz, rr)rr7rrrr rhr1DropVarr r r r rr) rrrrdrcrrrx_irs r& _swap_pp_rulerP0sR {{"! ZZ 4>>szzJ!Zd Wa4#8##$ 99  g!>?A BBsG(2G2GH!ggdm ggfo ggdmgghooc7+, &      g./F 6 r'c <|j|\}}}}|j|\}} } } tj| } tjt j ||||fd|i|tjt j ||| |fd|i|fSr)rr rswap_pr%rr) rrrr7rr val_primalrrr6 val_tangents r& _swap_jvprUOs-6-@-@-I**j*d#,#6#6x#@ +q+q##K0+ kk  *j*d!K L   kk  +z;!M N   r'c~|j|\}}}}t|}|rt|dtjsQt j ||dj} |jtjj| |^} } t| tsJtj|| }td| jDre| jD]-} t| ts| jdkDs$t!d| j} t#| D cgc]#\}} t| ts| js|%}}} | D cgc] } t| tr | j$n| "}} t'd| D}t)||}t+j,|||}t/j0||}|0|}t/j2|||}t/j2|||}t/j4||}t+j6|||}t9||}ntd | jDrj|| j}|0|}t/j2|||}t/j2|||}|j:| jj=|}nt |fd t?|dz zz|fScc} }wcc} w) Nr9rc3ZK|]#}t|txs |j %ywrbrgrks r&rnz'_swap_discharge_rule..ororprsr;c3XK|]"}t|tr |jnd$ywrrrtrks r&rnz'_swap_discharge_rule..{rurvrwrc3>K|]}t|t ywrbrrks r&rnz'_swap_discharge_rule..rrrb) rrrhrr?rr@rjrArBrrrrrirCr enumeraterrr7rrr?squeezer expand_dimsr._unpad_values_to_avoid_dynamic_slice_oob_shiftrrr)rrrrr6rrrrrFrrrmrirrrxrout_rs r&_swap_discharge_ruler`ds(229=#z3J* :jnh6H6HI**:x{7H7HIIh((==iHI$?C C ## #''_=#D DD [[C Au !((Q,!"ABBCkkGg& Aq!U#AGG K ELLqz!U3AGG:LLLOwOOK 7sK HC  C; GC ++c; 'C  d IIdC %c IIdC &c //#{ +C  $ $S#\ JE :5+ NE 9S[[ 99 ckk C  d IIdC %c IIdC &c FF3;;  # #C (E  Gs8}q01 13 66; Ms 3(K%"%K+)rrcache_modifiereviction_policyvolatilectj||d\}}tj||||f\} } t j | | |||dS)avReturns an array loaded from the given index. If neither ``mask`` nor ``other`` is specified, this function has the same semantics as ``x_ref_or_view[idx]`` in JAX. Args: x_ref_or_view: The ref to load from. idx: The indexer to use. mask: An optional boolean mask specifying which indices to load. If mask is ``False`` and ``other`` is not given, no assumptions can be made about the value in the resulting array. other: An optional value to use for indices where mask is ``False``. cache_modifier: TO BE DOCUMENTED. eviction_policy: TO BE DOCUMENTED. volatile: TO BE DOCUMENTED. load)rrarb is_volatile)rrrrrr%) rrrrrarbrcrrrrs r&reres`$// sFK%"// j$&)Y #%  r'swaprrb_function_namectj|||\}}tj||||f\}} t j || |dS)zSwaps the value at the given index and returns the old value. See :func:`~jax.experimental.pallas.load` for the meaning of the arguments. Returns: The value stored in the ref prior to the swap. )rrb)rrrrrRr%) rrrrrbrirrrrs r&rgrgsX//S.%#// C0NO)Y I r')rrbc&t|||||d}y)ztStores a value at the given index. See :func:`~jax.experimental.pallas.load` for the meaning of the arguments. storerhN)rg)rrrrrbr6s r&rlrls  =#s!#!r'c|tjk(rtjS|tjk(rtjS|S)zFUgly workaround to support types that don't allow automatic promotion.)r?int4rfloat8_e4m3b11fnuzrrs r& _handle_smallrqs4 chh 88O c$$$ << ,r'c6|jdk7s|jdk7r td|rdnd}|sdnd}|C| td|rtjjntjj }t jt|jt|j}t j|t jrt jnt j} tj|||f|ffdf|| S)aComputes the dot product of two arrays. The inputs can optionally be transposed before computing the product. Depending on the hardware, this can be cheaper than computing the transpose beforehand. Args: a: The left-hand size of the dot product, of shape ``(..., N)``. b: The right-hand size of the dot product, of shape ``(...N, M)``. trans_a: Whether to transpose ``a`` before the product. trans_b: Whether to transpose ``b`` before the product. allow_tf32: Whether to use tf32 precision. Mutually exclusive with ``precision``. precision: Specifies the precision of the dot product. See Also: :func:`jax.numpy.dot` z`a` and `b` must be 2D arrays.rrsz5Only one of allow_tf32 and precision can be specified)r*r*)dimension_numbers precisionpreferred_element_type)ndimrr PrecisionHIGHHIGHESTr? promote_typesrqrr"r'r@float32 dot_general) abtrans_atrans_b allow_tf32rulhs_contract_dimrhs_contract_dimr out_dtypes r&dotrs(ffkqvv{ 5 66!Qq%Q1 N OO&0 ""cmm6K6KI  M!''2M!''4J K%>>%=cii3;;) +-0@/BCXN&  r' reciprocalapproxc0tj||S)Nr) reciprocal_pr%rcrs r&rrs   1V  ,,r'c~|Srbr*rs r&_reciprocal_abstract_evalrs  (r'cNddd}tj|d|||S)NFrc|rOtj|jtjjtjStj|Srb)r?rrErr|rs r& _reciprocalz._reciprocal_lowering_rule.._reciprocal s? ^^AHHS\\2 3 : :3;; GG >>! r'multiple_results)r lower_fun)ctxrcrrs r&_reciprocal_lowering_ruler s/ % = e < 1V r'c4tj|g|ddiS)aPrints values from inside a Pallas kernel. Args: fmt: A format string to be included in the output. The restrictions on the format string depend on the backend: * On GPU, when using Triton, ``fmt`` must not contain any placeholders (``{...}``), since it is always printed before any of the values. * On GPU, when using the experimental Mosaic GPU backend, ``fmt`` must contain a placeholder for each value to be printed. Format specs and conversions are not supported. If a single value is provided, the value may be an array. Otherwise, all values must be scalars. * On TPU, if all inputs are scalars: If ``fmt`` contains placeholders, all values must be 32-bit integers. If there are no placeholders, the values are printed after the format string. * On TPU, if the input is a single vector, the vector is printed after the format string. The format string must end with a single placeholder ``{}``. *args: The values to print. skip_format_checkT)r debug_print)fmtargss r&rrs*   s BT BT BBr'c  d}tjj|D]*\}}}}||dz }|s|r td|s!tdt ||k7r$t d|d|dk(rdndd t |y) NrrszDThe format string should not contain any format specs or conversionszDThe format string should not reference arguments by position or namezThe format string expects z argumentrmz , but got )string Formatterparserr TypeError)rrn_placeholdersr6fieldspec conversions r&check_debug_print_formatr2s.$*$4$4$6$<$.Os0 !Q+,k  A&2s'))rzip)frrnew_argss r&wrap_with_transformsrMs, dJ'( Hr' run_scopedc~~|jSrb)is_high)jaxpravalsr7s r&_run_scoped_is_highrYs V r'ctj||}tj|}|j}t j |d|ji|S)Nr)r1 ClosedJaxprpe lower_jaxprconsts run_scoped_pr%r)rrr7closed_hi_jaxprclosed_lo_jaxprrs r&_run_scoped_to_lojaxr^sN((5/NN?3/  ! !&   F J/*?*? J6 JJr'r*)collective_axesc t|ts|f}tj||f\}}t j t j|t jd||||\}}|Dcgc]}|j} }| Dcgc]*}t|tjr |jn|,} }td| D} t|| }tjd5t!j"|| \} } }dddt%j& |d}tj(||Scc}wcc}w#1swYDxYw)aCalls the function with allocated references and returns the result. The positional and keyword arguments describe which reference types to allocate for each argument. Each backend has its own set of reference types in addition to :class:`jax.experimental.pallas.MemoryRef`. When ``collective_axes`` is specified, the same allocation will be returned for all programs that only differ in their program ids along the collective axes. It is an error not to call the same ``run_scoped`` in all programs along that axis. zpallas run_scoped) debug_infoc3lK|],}t|tjr |jnd.yw)r*N)rhrrr)rlrs r&rnzrun_scoped..s0 !K$>$>?allRGs24FNrr)rhrrrr flatten_funlu wrap_initr get_ref_avalrrrrr mutable_array_checksrtrace_to_jaxpr_dynamicrr%r)rrrkw_types flat_typesin_treeflat_funout_tree_thunkr ref_avalsrref_transformsrr6rrs r&rresb" OU +&(O!..x/@A*g%11ll1&112E23UHFG (N *44Aq~~4)4  ![778aeea? % ."(N ;( ""5)B005AE1fB6P#  ! !."2C 88%5 BBs6E /E7EEc~~|jDchc]@}t|tjr"|jt |j k\s|B}}|j Dcgc]}|jc}|fScc}wcc}wrb)rrhJaxprInputEffect input_indexr constvarsrr)rrreffnonlocal_effectsvs r&_run_scoped_abstract_evalrs{ O  S'22 3ooU__!55   -- (Q!&& (*: :: )s AA>&Bc ~|r tdt|}tj|}t|j}t j |g|dgt|jzz\} } | r tdtj| |} tj|| |d} | d|} | |d} t||Dcgc]4\}}|r+t|tjr| jdnd6}}}t|t|k(sJt|dt||| fScc}}w)Nz:run_scoped discharge does not support collective_axes yet.Fshould_discharge4Cannot handle new consts created by state discharge.rrz != )rrrconvert_constvars_jaxprrrdischarge_staterconvert_invars_to_constvarsrr%rrhr AbstractRefr+)rrrrrr num_consts jaxpr_noconstnum_return_valuesdischarged_body new_constsr return_values ref_outputsshouldrupdatess r&_run_scoped_discharge_rulersn D 9~*,,U3--//0 / ? ?'5'C 4E*EE!/:  > @@22?JO/  # (()-%&'+ %((8($C E FD#z$8I8I'Jkooa  E' E WX &L3w<.S]O(LL & -   Es 9Ec |r tdtj|}t|jt j |gd\}|r tdfd}tj|d|g|S)NzGrun_scoped lowering outside of Pallas does not support collective_axes.Trrct|}j|dDcgc]}|j}}|Dcgc]"}t|j|j $}}t jgg||}|dScc}wcc}wrb)rrrr+rjrr1 eval_jaxpr) lower_fun_argsrr body_avalsr init_valsrrrs r& _lower_funz-_run_scoped_lowering_rule.._lower_funs^$J"1"8"8"EFQ!&&FJF,68$(% DJJ 8I8   or ON OY OC !! "" G8s A>'Br) rrrrrrrrrr) rrrrrrrrrs @@r&_run_scoped_lowering_rulers   ,,U3--//0 / ? ?R$!0/:->@@# ; T :3 F FFr' get_globalcN|j}tj|S)aReturns a global reference that persists across all kernel invocations. Each call to ``get_global`` returns a different and unique reference, but one that is stable across invocations of the kernel body. Args: what: The reference type to allocate. Each backend has its own set of reference types (e.g., :class:`jax.experimental.pallas.mosaic_gpu.SemaphoreType` for GPU). Example:: sem_ref = pl.get_global(plgpu.SemaphoreType.REGULAR) pl.semaphore_signal(sem_ref) pl.semaphore_wait(sem_ref) what)r get_global_pr%)rrs r&rrs&    (     ))r'c|Srbr*rs r&_get_global_abstract_evalr rr'c~~~td)Nz8get_global discharge is not supported in interpret mode.)r)rrrs r&_get_global_discharge_rulers 4@ r'cnt|tjr|j|jfS|dfSr<)rhrrrr)rs r&_get_ref_and_transformsrs.U))* 77CNN "" b.r'ceZdZdZdZy) DeviceIdTypemeshlogicalN)rUrVrWMESHLOGICALr*r'r&rr"s $ 'r'rc |/tjtjtjh}t |t j std|d||j}|r|dj}|rtd|d||jtfd|Dstd|d|dd y) NzCannot z on a non-semaphore Ref: r9z on a non-()-shaped semaphore: c3JK|]}tj|ywrb)r?r")rlsem_type sem_dtypes r&rnz"check_sem_avals..9s$   nnY) s #zMust z$ semaphores of the following types: z. Got r) r, semaphorebarrier_semaphorer&rhrrrrjget_indexer_shaperany)sem_avalsem_transforms_avalsnameallowed_semaphore_types sem_shapers @r&check_sem_avalsr's$%%--  He// 0 wtf$=hZH IInn)$R(::r?r)rrs r&_semaphore_read_abstract_evalrbs%  Y   b#))G"4 55r'c~|j|\}}t|||d}|jtj}dt |z|fS)Nrrb)rrrEr?r@r)rrrrrr sem_values r&_semaphore_read_discharge_rulerjsV)))43 "3 HQK@)syy)) 3x= ) ++r'.semaphore_signal) device_iddevice_id_type core_indexct|\}}tj|tj}|||||g}t j |\}} t j|| |dy)aIncrements the value of a semaphore. This operation can also be performed remotely if ``device_id`` is specified, in which ``sem_or_view`` refers to a Ref located on another device. Note that it is assumed that ``sem_or_view`` is already allocated (e.g. through the proper use of barriers), or else this operation could result in undefined behavior. Args: sem_or_view: A Ref (or view) representing a semaphore. inc: The value to increment by. device_id (optional): Specifies which device to signal. If not specified, ``sem_or_view`` is assumed to be local. device_id_type (optional): The format in which ``device_id`` should be specified. core_index (optional): If on a multi-core device, specifies which core to signal. rp)rr N)rr?asarrayr@rrsemaphore_signal_pr%) rincrr r!rrrrrs r&rrsb4,K8/#z Csyy)# z3 : 6$"//5)Y#r'ctj||\}}}}}t||d|jt jdk7rt d|jt }|tj|} | D]<} | jt jdk7s&t d| j|tjurYt|trI|D]@} t| ts| f} | D]&} |jtj| (Bg|fS|jt j"g|fS)Nsignalr@z$Must signal an int32 value, but got z-`device_id`s must be an int32 value, but got )rrrrr?rrrrrrhr4rrOr1NamedAxisEffectr, comms_effect) rr rr r  value_avaldevice_id_avalcore_index_avaleffsdevice_id_flat_avalsrkk_s r&_semaphore_signal_abstract_evalr1sJy%0 (0(;7++ ;J ?? F GG"3 HQK@) ;889#$99 :y3!]  Gs8}q'89 92 ==r'semaphore_wait) decrementct|\}}tj|tj}||||g}t j |\}}t j|d|iy)a=Blocks execution of the current thread until a semaphore reaches a value. Args: sem_or_view: A Ref (or view) representing a semaphore. value: The target value that the semaphore should reach before unblocking. decrement: Whether to decrement the value of the semaphore after a successful wait. rprN)rr?r#r@rrsemaphore_wait_pr%)rrr?rrrrrs r&r>r>s\,K8/#z ++e399 -% z5) ,$"//5)Y8i8r'ctj||\}}}}t||d|jt jdk7r t dgS)Nwaitr@zMust wait an int32 value.)rrrrr?r)rrr r r*r6s r&_semaphore_wait_abstract_evalrDsX2;2J2J3/( *a(0&97++ 0 11 )r'c ~|j}|jd}tj||\}}}}t j dg} |r$| j t j d| t j dtj|||t j dt j tj||gz } t j| S)Nrr>z[dec]r) rr7rrr r rArr r1r r ) rrrrr3r4r5rr?partss r&_semaphore_wait_pp_eqnrGs ::& K $ tV,    gg % LL!"ggcl7C8ggclgghooeW-.  % 5 r'c~|j|\}}}}t|||d}|jtj}|rt j ||||z \} } n|} | fdt|dz zzdfS)Nrrbrsr*)rrrEr,r&rr;r) rrrrrrrr?rr6r<s r&_semaphore_wait_discharge_rulerI6s(1(;(;I(F%3 E9"3 HQK@) ,,{<< =%&;; ZU*A}M  Gs8}q'89 92 ==r'c tjjd}|in)tt |j |j i}|jD]\}}t|trdtfd|DrOtfd|Dst|d|Dcgc]}| }}t|D] \} }|} tj|| dzd} t!j"|| } | | dz zdk(rB| | zj%dz } t!j&|t!j"|| }nt!j(|| }| | dz zdk(r1t!j*|t!j"||dz }n*t!j,|t!j"|| }|||< ||||<g}D]2}||vr|j/|||j/||4|jDcic] \}}|vr||}}}t||fScc}wcc}}w)N c3&K|]}|v ywrbr*rlr~mesh_axis_sizess r&rnz*_device_id_dict_to_mesh..Us,!"_,c3&K|]}|v ywrbr*rMs r&rnz*_device_id_dict_to_mesh..Xs9!o%9rOz) mixes JAX mesh and Pallas mesh grid axesrsr)r IntegerType get_signlessr4r axis_names mesh_shapeitemsrhrr rrrZmathprodrconstant bit_lengthshruidivsiandiremsirA) mesh_contextdevice_id_dictget_axis_indexi32physical_axis_dict axis_namerr axes_dimensions axis_index axis_sizeinner_mesh_size minor_divisor shift_lenpartial_device_idx device_idxrr/r non_mesh_axesrNs @r&_device_id_dict_to_meshrmJs ##B'#O L # #\%<%<=O&,,. *ni)U#,&/,)9y9 9!kB C     A &! +zz nnS/)"s 4 I1Ici}t|tr1|tjurt d|dt |||\}}|tjurt j|}|d}n |j}t|t|k7r$t dt|dt|dtjjdt|dk(rtjd|fStj tj"fd t%||D|fS|tj&ur||fSt)d |) zNormalizes a device id into a logical device id and axes that don't correspond to JAX mesh axes. The indexing implied by the returned axis dict should be handled by the caller. zN`device_id_type` must be MESH if `device_id` is a dict, got: device_id_type = rr*zANumber of device ids must match the number of mesh axes, but got z ids for a zD mesh.rKrc3xK|]1\}}tj|tj|3ywrb)rmulirX)rlr~rras r&rnz'device_id_to_logical..s3 1 JJq%..a0 1 s7:zUnsupported device id type: )rhr4rrrrmrr mesh_stridesrrrQrRrrX functoolsreduceaddirrr)r^rr r`rlr6rqras @r&device_id_to_logicalrusk- 4 \...  %!%Q (  7|YP^_I}|(((&&y1Jl!..l :#l++  *o k#l*;):JKLLr'delayc~gSrbr*nanoss r&_delay_abstract_evalrzs  )r'c.tj|y)z+Sleeps for the given number of nanoseconds.N)delay_pr%rxs r&rvrvs ,,ur')r#rreturnjax_typing.Array)r#rr}int | jax_typing.Array)rrM)r Any | NonerrM)rr)rcr~rzSequence[int] | intr}r~)F)r}r~)r}None)rzjax_typing.DTypeLike)FFNN)rrrrrz bool | None)rzmlir.LoweringRuleContext)rstrrzjax_typing.ArrayLike) rzCallable[..., Any]rrrzHashable | tuple[Hashable, ...]rrr}r)rzpallas_core.ScratchShaper}r~rb)rs)r%rrDeviceIdr rr!zint | jax_typing.Array | None)r r)rzjax_core.JaxprEqnrzjax_core.JaxprPpContextrzjax_core.JaxprPpSettings)rrr?r)r^pallas_utils.MeshInfo | None)r^rrz5ir.Value | tuple[ir.Value, ...] | dict[Any, ir.Value]r rr}z$tuple[ir.Value, dict[Any, ir.Value]])ryrr}r)__doc__ __future__rcollections.abcrrrenumrrrVrrr jax._src.lax_srcrjax._srcrr r r r1r r r jax_typingrrrrr rrjax._src.interpretersrrrjax._src.lib.mlirrjax._src.lib.mlir.dialectsrjax._src.pallasr,r pallas_utilsjax._src.staterrrrrrrjax.interpretersrrr?rir?safe_mapmap unsafe_mapsafe_zipr unsafe_zipr2r$r r:def_bind_with_tracedef_abstract_evalrBrFrDrIrKEnumrMrrregister_discharge_ruledef_effectful_abstract_evalrrrrrrrrrrrrrrdef_implregister_loweringrrrrrrrr pp_eqn_rulesr primitive_jvpsr+r7partialr]rHrRrrPrUr`rergrlrqrrrrrrrtransformation2rrrrrrto_lojaxrrrregister_partial_discharge_rulerr ref_primitive_ref_allocating_primitivesrOrrrrrrrrrrrrArrayrr4rr$rrr1r9r=rAr>rDrGrIrmrur|rzrvr*r'r&rs) &".$  %)&)$4 ,/17#+/!!    --Z--Z!x!!,/  & V !;< - -$##N3( ""#@A!!-"-499"x!!,/ %7=I%7P6''' 56PQ)) ?* ? @D)@D  ?C  ?C  ?C  ?C  >B  ?C "x!!,/  )) V* V *)((6$7$ &8%%&67*+')?@8##$#""=1 '(}&<=.   !   M *##$,!.f*&&#"7<02C1B1B02. )((0%&1%&P   M *##$,8!.f$&&)((0-71-7`&*d2B<+/#3",0#6;26%%N"x!!,/ "-   16  ! |%>?C0 )6"x!!,/ $ + K -  8:,9,9 ,95,9 ,9  ,9^)) ;* ; , ^>/// = 4))<8G9G2"x!!,/ % ! ##'' 5*( 6''' 5  499 CG8 &8%%&67$)! @##6$6,:'''(89" !! !D (5z7G7G1G1L+M MPTUXZ]`j`p`pZpUpPq q(X''(:;&*# #$"#/#4#404 " " " ! " . "J//!0B &= '? <-E()>"<'''(:;$ &8%%&67$(!239IM9.9BF9"## $ &='?0+A&'>:'''(89" 6)r,M.,MD,M!,M * ,M^ (  W %   r'