ó
    Eñi£ ã                  óR  • S SK Jr  S SKrS SKrS SKrS SKrS SKrS SKJrJ	r	J
r
Jr  S SKrS SKrS SKJr  S SKJr  SSKJr  SSKJr  SS	KJr  SS
KJrJr  SSKJr  SSKJr  SSKJ r J!r!J"r"J#r#J$r$  SSK%J&r&J'r'   " S S\$5      r(\(" 5       RR                  r*\
(       a  S SK+J,r,J-r-  SSKJ.r.  SSK/J0r0  SSK1J2r2  Sr3Sr4S&S jr5\Rl                  Ro                  \8S5      r9 " S S5      r: " S S\;5      r< " S S\#5      r=\R|                   " S  S!5      5       r? " S" S#\&5      r@ " S$ S%\'5      rAg)'é    )ÚannotationsN)ÚAnyÚOptionalÚTYPE_CHECKINGÚUnion)Ú
OrderedSet)ÚModularIndexingé   )Úconfig)ÚComputedBuffer©Útorch_dtype_to_jax)Úget_fused_kernel_nameÚget_kernel_metadata)ÚVé   )ÚBlockPatternMatcher)ÚBackendFeatureÚCSEVariableÚIndentedBufferÚOpOverridesÚPythonPrinter)Ú
SIMDKernelÚSIMDSchedulingc                  ó6   • \ rS rSrSrSS jrSS jrSS jrSrg)	ÚPallasPrinteré    zG
Custom sympy printer for Pallas that handles JAX-specific constructs.
c                óÎ   • U R                  UR                  S   5      nU R                  UR                  S   5      nU R                  UR                  S   5      nSU SU SU S3$ )z!Convert sympy Where to jnp.where.r   r   r
   ú
jnp.where(ú, Ú))ÚdoprintÚargs)ÚselfÚexprÚcÚpÚqs        Ú[/home/mande/repo/quber/.venv/lib/python3.13/site-packages/torch/_inductor/codegen/pallas.pyÚ_print_WhereÚPallasPrinter._print_Where%   s_   € à�L‰L˜Ÿ™ 1™Ó&ˆØ�L‰L˜Ÿ™ 1™Ó&ˆØ�L‰L˜Ÿ™ 1™Ó&ˆØ˜A˜3˜b   2 a S¨Ð*Ð*ó    c                ó˜   • UR                    Vs/ s H  o R                  U5      PM     nnUS   nUSS  H  nSU SU S3nM     U$ s  snf )z7Convert sympy Min to jnp.minimum for JAX compatibility.r   r   Nújnp.minimum(r    r!   ©r#   r"   ©r$   r%   Úargr#   Úresults        r)   Ú
_print_MinÚPallasPrinter._print_Min,   óX   € à-1¯YªYÓ7ªY c—‘˜SÖ!©YˆÐ7Ø�a‘ˆØ˜˜“8ˆCØ# F 8¨2¨c¨U°!Ð4ŠFñ àˆùò	 8ó   �Ac                ó˜   • UR                    Vs/ s H  o R                  U5      PM     nnUS   nUSS  H  nSU SU S3nM     U$ s  snf )z7Convert sympy Max to jnp.maximum for JAX compatibility.r   r   Nújnp.maximum(r    r!   r/   r0   s        r)   Ú
_print_MaxÚPallasPrinter._print_Max4   r5   r6   © N)r%   ú
sympy.ExprÚreturnÚstr)	Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__r*   r3   r9   Ú__static_attributes__r;   r,   r)   r   r       s   † ñô+ô÷r,   r   )ÚCallableÚSequence)ÚIRNode)ÚReductionType)ÚBaseSchedulerNodeÚmainé€   c                ó6   • U [         -   S-
  [         -  [         -  $ )z@Align size to WARPGROUP_SIZE (128) for Mosaic GPU compatibility.r   )ÚWARPGROUP_SIZE)Úsizes    r)   Ú_align_to_warpgrouprO   R   s   € à”NÑ" QÑ&¬>Ñ9¼^ÑKÐKr,   Úkernel_codec                  ó<   • \ rS rSrSr S   S	S jjrSS.S jrSrg)
ÚPallasKernelWrapperé[   z6Wrapper to provide .run() interface for Pallas kernelsNc                óH   • Xl         X l        [        R                  SU5        g )NzPallas kernel path: %s)Ú	kernel_fnÚkernel_pathÚkernel_code_logÚinfo)r$   rU   rV   s      r)   Ú__init__ÚPallasKernelWrapper.__init__^   s    € ð #ŒØ&ÔÜ×ÑÐ5°{ÕCr,   )Ústreamc               ó*   • U R                   " USU0UD6$ )zõ
Execute the Pallas kernel.

Args:
    *args: Arguments to pass to the kernel function
    stream: CUDA stream to pass to the kernel function
    **kwargs: Additional keyword arguments for the kernel

Returns:
    Result of the kernel execution
r[   )rU   )r$   r[   r#   Úkwargss       r)   ÚrunÚPallasKernelWrapper.rune   s   € ð �~Š~˜tÐ=¨FÐ=°fÑ=Ð=r,   )rU   rV   ©N)rU   zCallable[..., Any]rV   úOptional[str])r?   r@   rA   rB   rC   rY   r^   rD   r;   r,   r)   rR   rR   [   s3   † Ù@ð KOðDØ+ðDØ:GõDð !%÷ >ð >r,   rR   c                  ó   • \ rS rSrSrSrg)ÚUnsupportedét   zJException raised when an operation is not supported by the Pallas backend.r;   N)r?   r@   rA   rB   rC   rD   r;   r,   r)   rc   rc   t   s   † ÜTr,   rc   c                  ó	  • \ rS rSrSr\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r	\SuS j5       r
\SuS j5       r\SuS	 j5       r\SuS
 j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SuS j5       r\SvS j5       r \SvS j5       r!\SvS j5       r"\SwS  j5       r#\SxS! j5       r$\  Sy         SzS# jj5       r%\S{S$ j5       r&\S|S% j5       r'\S}S& j5       r(\SuS' j5       r)\SuS( j5       r*\SuS) j5       r+\SuS* j5       r,\SuS+ j5       r-\SuS, j5       r.\SvS- j5       r/\SvS. j5       r0\SvS/ j5       r1\SvS0 j5       r2\SvS1 j5       r3\SuS2 j5       r4\SuS3 j5       r5\SuS4 j5       r6\SvS5 j5       r7\SvS6 j5       r8\SvS7 j5       r9\SuS8 j5       r:\SvS9 j5       r;\SvS: j5       r<\SvS; j5       r=\SvS< j5       r>\SvS= j5       r?\SvS> j5       r@\SvS? j5       rA\S~S@ j5       rB\BrC\SuSA j5       rD\SuSB j5       rE\SuSC j5       rF\SuSD j5       rG\SuSE j5       rH\SuSF j5       rI\SuSG j5       rJ\SuSH j5       rK\SuSI j5       rL\SuSJ j5       rM\SuSK j5       rN\SuSL j5       rO\MrP\SuSM j5       rQ\NrR\SuSN j5       rS\SSO j5       rT\SSP j5       rU\TrV\UrW\SSQ j5       rX\SuSR j5       rY\SSS j5       rZ\SST j5       r[\SSU j5       r\\S€SV j5       r]\S€SW j5       r^\S€SX j5       r_\S€SY j5       r`\S€SZ j5       ra\S€S[ j5       rb\S€S\ j5       rc\S€S] j5       rd\S€S^ j5       re\S€S_ j5       rf\S€S` j5       rg\S€Sa j5       rh\SuSb j5       ri\SuSc j5       rj\S�Sd j5       rk\SvSe j5       rl\SvSf j5       rm\SvSg j5       rn\SuSh j5       ro\SuSi j5       rp\SvSj j5       rq\SvSk j5       rr\SvSl j5       rs\SuSm j5       rt\SvSn j5       ru\SvSo j5       rv\S‚Sp j5       rw\SƒSq j5       rx\SƒSr j5       ry\S„Ss j5       rzStr{g")…ÚPallasKernelOverrideséx   zš
Map element-wise ops to JAX/Pallas operations.

For now, we use the default Python operators which are compatible
with JAX numpy broadcasting semantics.
c                ó   • SU  S3$ )Nzjnp.sin(r!   r;   ©Úxs    r)   ÚsinÚPallasKernelOverrides.sin€   ó   € à˜!˜˜AˆÐr,   c                ó   • SU  S3$ )Nzjnp.cos(r!   r;   ri   s    r)   ÚcosÚPallasKernelOverrides.cos„   rm   r,   c                ó   • SU  S3$ )Nzjnp.tan(r!   r;   ri   s    r)   ÚtanÚPallasKernelOverrides.tanˆ   rm   r,   c                ó   • SU  S3$ )Nz	jnp.sinh(r!   r;   ri   s    r)   ÚsinhÚPallasKernelOverrides.sinhŒ   ó   € à˜1˜#˜QÐÐr,   c                ó   • SU  S3$ )Nz	jnp.cosh(r!   r;   ri   s    r)   ÚcoshÚPallasKernelOverrides.cosh�   rw   r,   c                ó   • SU  S3$ )Nz	jnp.tanh(r!   r;   ri   s    r)   ÚtanhÚPallasKernelOverrides.tanh”   rw   r,   c                ó   • SU  S3$ )Nzjnp.arcsin(r!   r;   ri   s    r)   ÚasinÚPallasKernelOverrides.asin˜   ó   € à˜Q˜C˜qÐ!Ð!r,   c                ó   • SU  S3$ )Nzjnp.arccos(r!   r;   ri   s    r)   ÚacosÚPallasKernelOverrides.acosœ   r�   r,   c                ó   • SU  S3$ )Nzjnp.arctan(r!   r;   ri   s    r)   ÚatanÚPallasKernelOverrides.atan    r�   r,   c                ó   • SU  S3$ )Nzjnp.exp(r!   r;   ri   s    r)   ÚexpÚPallasKernelOverrides.exp¤   rm   r,   c                ó   • SU  S3$ )Nz	jnp.exp2(r!   r;   ri   s    r)   Úexp2ÚPallasKernelOverrides.exp2¨   rw   r,   c                ó   • SU  S3$ )Nz
jnp.expm1(r!   r;   ri   s    r)   Úexpm1ÚPallasKernelOverrides.expm1¬   ó   € à˜A˜3˜aÐ Ð r,   c                ó   • SU  S3$ )Nzjnp.log(r!   r;   ri   s    r)   ÚlogÚPallasKernelOverrides.log°   rm   r,   c                ó   • SU  S3$ )Nz
jnp.log10(r!   r;   ri   s    r)   Úlog10ÚPallasKernelOverrides.log10´   r‘   r,   c                ó   • SU  S3$ )Nz	jnp.log2(r!   r;   ri   s    r)   Úlog2ÚPallasKernelOverrides.log2¸   rw   r,   c                ó   • SU  S3$ )Nz
jnp.log1p(r!   r;   ri   s    r)   Úlog1pÚPallasKernelOverrides.log1p¼   r‘   r,   c                ó   • SU  S3$ )Nz	jnp.sqrt(r!   r;   ri   s    r)   ÚsqrtÚPallasKernelOverrides.sqrtÀ   rw   r,   c                ó   • SU  S3$ )Nzjax.lax.rsqrt(r!   r;   ri   s    r)   ÚrsqrtÚPallasKernelOverrides.rsqrtÄ   s   € à ˜s !Ð$Ð$r,   c                ó   • SU  S3$ )Nzjnp.abs(r!   r;   ri   s    r)   ÚabsÚPallasKernelOverrides.absÈ   rm   r,   c                ó   • SU  S3$ )Nz(-r!   r;   ri   s    r)   ÚnegÚPallasKernelOverrides.negÌ   s   € à�A�3�aˆyÐr,   c                ó   • SU  S3$ )Nz
jnp.floor(r!   r;   ri   s    r)   ÚfloorÚPallasKernelOverrides.floorÐ   r‘   r,   c                ó   • SU  S3$ )Nz	jnp.ceil(r!   r;   ri   s    r)   ÚceilÚPallasKernelOverrides.ceilÔ   rw   r,   c                ó   • SU  S3$ )Nz
jnp.trunc(r!   r;   ri   s    r)   ÚtruncÚPallasKernelOverrides.truncØ   r‘   r,   c                ó   • SU  S3$ )Nz
jnp.round(r!   r;   ri   s    r)   ÚroundÚPallasKernelOverrides.roundÜ   r‘   r,   c                ó   • SU  S3$ )Nzjax.nn.sigmoid(r!   r;   ri   s    r)   ÚsigmoidÚPallasKernelOverrides.sigmoidà   ó   € à    1Ð%Ð%r,   c                ó   • SU  S3$ )Nr8   z, 0)r;   ri   s    r)   ÚreluÚPallasKernelOverrides.reluä   s   € à˜a˜S Ð%Ð%r,   c                ó   • SU  SU S3$ )Nz
jnp.power(r    r!   r;   ©ÚaÚbs     r)   ÚpowÚPallasKernelOverrides.powè   ó   € à˜A˜3˜b   1Ð%Ð%r,   c                ó   • SU  SU S3$ )Nr8   r    r!   r;   r¾   s     r)   ÚmaximumÚPallasKernelOverrides.maximumì   ó   € à˜a˜S  1 # QÐ'Ð'r,   c                ó   • SU  SU S3$ )Nr.   r    r!   r;   r¾   s     r)   ÚminimumÚPallasKernelOverrides.minimumð   rÇ   r,   c                ó   • SU  SU SU S3$ )Nr   r    r!   r;   )Úcondr¿   rÀ   s      r)   ÚwhereÚPallasKernelOverrides.whereô   s   € à˜D˜6  A 3 b¨¨¨1Ð-Ð-r,   c                ó  • U" 5       n[        U[        5      (       aP  [        R                  " U5      (       a  SnO=[        R                  " U5      (       a  US:”  a  SOSnO[        U5      nO[        U5      nSU  SU SU S3$ )zr
Computes body, but only uses the result where mask is true.
Where mask is false, uses the 'other' value instead.
újnp.nanr   újnp.infú-jnp.infr   r    r!   )Ú
isinstanceÚfloatÚmathÚisnanÚisinfÚrepr)ÚmaskÚbodyÚotherr2   Ú	other_strs        r)   ÚmaskedÚPallasKernelOverrides.maskedø   su   € ñ “ˆä�eœU×#Ñ#Ü�zŠz˜%× Ñ Ø%‘	Ü—’˜E×"Ñ"Ø).°«™I¸
‘	ä  ›K‘	ä˜U›ˆIà˜D˜6  F 8¨2¨i¨[¸Ð:Ð:r,   Nc                ó*   • [        U5      nSU  SU S3$ )Nzjnp.asarray(ú	).astype(r!   r   )rj   ÚdtypeÚ	src_dtypeÚuse_compute_typesÚ	jax_dtypes        r)   Úto_dtypeÚPallasKernelOverrides.to_dtype  s#   € ô ' uÓ-ˆ	à˜a˜S 	¨)¨°AÐ6Ð6r,   c                óF   • [        U5      n[        U5      nSU  SU SU S3$ )z=Bitcast a value from one dtype to another with the same size.z)jax.lax.bitcast_convert_type(jnp.asarray(rà   z), r!   r   )rj   rá   râ   rä   Újax_src_dtypes        r)   Úto_dtype_bitcastÚ&PallasKernelOverrides.to_dtype_bitcast  s7   € ô ' uÓ-ˆ	Ü*¨9Ó5ˆà:¸1¸#¸YÀ}ÀoÐUXÐYbÐXcÐcdÐeÐer,   c                ó
  • SSK Jn  [        R                  R                  R                  [        R                  R                  U 5      5        [        R                  R                  U 5      n[        R                  R                  U5      n[        R                  R                  U5      n[        R                  R                  R                  [        R                  R                  XR" U 5      S9n[        R                  Xa5      $ )z>Convert a sympy expression to a JAX array indexing expression.r
   )Úget_bounds_index_expr)Úbounds)Úutilsrì   r   ÚkernelÚused_iter_varsÚupdateÚ_get_used_iter_varsÚprepare_indexingÚrename_indexingÚkexprÚcseÚgenerateÚcomputerf   rå   )r%   rá   rì   ÚpreparedÚrenamedÚidx_strÚvars          r)   Ú
index_exprÚ PallasKernelOverrides.index_expr  s³   € õ 	2ô 	
�‰×Ñ×&Ñ&¤q§x¡x×'CÑ'CÀDÓ'IÔJô —8‘8×,Ñ,¨TÓ2ˆÜ—(‘(×*Ñ*¨8Ó4ˆÜ—(‘(—.‘. Ó)ˆÜ�h‰h�l‰l×#Ñ#Ü�H‰H×Ñ˜gÐ.CÀDÓ.Ið $ð 
ˆô %×-Ñ-¨cÓ9Ð9r,   c                ó  • [        U5      nU[        R                  :X  a  U (       a  S$ S$ [        U [        5      (       aA  [
        R                  " U 5      (       a  g[
        R                  " U 5      (       a
  U S:”  a  S$ S$ SU  SU S	3$ )
z/Convert a constant value to JAX representation.ÚTrueÚFalserÐ   r   rÑ   rÒ   z
jnp.array(z, dtype=r!   )r   ÚtorchÚboolrÓ   rÔ   rÕ   rÖ   r×   )Úvalrá   rä   s      r)   ÚconstantÚPallasKernelOverrides.constant0  sw   € ô ' uÓ-ˆ	Ø”E—J‘JÓÞ �6Ð- gÐ-ä�cœ5×!Ñ!Ü�zŠz˜#�‰Ø Ü�zŠz˜#�‰Ø$'¨!£G�yÐ;°Ð;Ø˜C˜5 ¨¨°1Ð5Ð5r,   c                ó   • SU  S3$ )Nz	jnp.real(r!   r;   ri   s    r)   ÚrealÚPallasKernelOverrides.real>  rw   r,   c                ó   • SU  S3$ )Nz	jnp.imag(r!   r;   ri   s    r)   ÚimagÚPallasKernelOverrides.imagB  rw   r,   c                ó   • SU  S3$ )Nz	jnp.conj(r!   r;   ri   s    r)   ÚconjÚPallasKernelOverrides.conjF  rw   r,   c                ó   • SU  S3$ )Nz
jnp.angle(r!   r;   ri   s    r)   ÚangleÚPallasKernelOverrides.angleJ  r‘   r,   c                ó   • SU  SU  S3$ )z8View complex tensor as real tensor with extra dimension.zjnp.stack([jnp.real(z), jnp.imag(z)], axis=-1)r;   ri   s    r)   Úview_as_realÚ"PallasKernelOverrides.view_as_realN  s   € ð & a S¨°Q°C°|ÐDÐDr,   c                ó   • SU  SU  S3$ )z#View real tensor as complex tensor.Ú(z[..., 0] + 1j * z	[..., 1])r;   ri   s    r)   Úview_as_complexÚ%PallasKernelOverrides.view_as_complexS  s   € ð �1�#Ð% a S¨	Ð2Ð2r,   c                ó   • SU  SU S3$ )Nr  z == r!   r;   r¾   s     r)   ÚeqÚPallasKernelOverrides.eqY  ó   € à�1�#�T˜!˜˜AˆÐr,   c                ó   • SU  SU S3$ )Nr  z != r!   r;   r¾   s     r)   ÚneÚPallasKernelOverrides.ne]  r  r,   c                ó   • SU  SU S3$ )Nr  z < r!   r;   r¾   s     r)   ÚltÚPallasKernelOverrides.lta  ó   € à�1�#�S˜˜˜1ˆ~Ðr,   c                ó   • SU  SU S3$ )Nr  z <= r!   r;   r¾   s     r)   ÚleÚPallasKernelOverrides.lee  r  r,   c                ó   • SU  SU S3$ )Nr  z > r!   r;   r¾   s     r)   ÚgtÚPallasKernelOverrides.gti  r$  r,   c                ó   • SU  S3$ )Nz
jnp.isnan(r!   r;   ri   s    r)   rÖ   ÚPallasKernelOverrides.isnanm  r‘   r,   c                ó   • SU  S3$ )Nz
jnp.isinf(r!   r;   ri   s    r)   r×   ÚPallasKernelOverrides.isinfq  r‘   r,   c                ó   • SU  S3$ )Nzjnp.isfinite(r!   r;   ri   s    r)   ÚisfiniteÚPallasKernelOverrides.isfiniteu  s   € à˜q˜c Ð#Ð#r,   c                ó   • SU  SU S3$ )Nr  z >= r!   r;   r¾   s     r)   ÚgeÚPallasKernelOverrides.gey  r  r,   c                ó   • SU  SU S3$ )Nzjnp.logical_and(r    r!   r;   r¾   s     r)   Úlogical_andÚ!PallasKernelOverrides.logical_and~  ó   € à! !  B q c¨Ð+Ð+r,   c                ó   • SU  SU S3$ )Nzjnp.logical_or(r    r!   r;   r¾   s     r)   Ú
logical_orÚ PallasKernelOverrides.logical_or‚  ó   € à    2 a S¨Ð*Ð*r,   c                ó   • SU  S3$ )Nzjnp.logical_not(r!   r;   ri   s    r)   Úlogical_notÚ!PallasKernelOverrides.logical_not†  ó   € à! !  AÐ&Ð&r,   c                ó   • SU  SU S3$ )Nzjnp.logical_xor(r    r!   r;   r¾   s     r)   Úlogical_xorÚ!PallasKernelOverrides.logical_xorŠ  r8  r,   c                ó   • SU  SU S3$ )Nzjnp.arctan2(r    r!   r;   r¾   s     r)   Úatan2ÚPallasKernelOverrides.atan2�  rÇ   r,   c                ó   • SU  SU S3$ )Nz
jnp.hypot(r    r!   r;   r¾   s     r)   ÚhypotÚPallasKernelOverrides.hypot“  rÃ   r,   c                ó   • SU  SU S3$ )Nz	jnp.fmod(r    r!   r;   r¾   s     r)   ÚfmodÚPallasKernelOverrides.fmod—  s   € à˜1˜#˜R ˜s !Ð$Ð$r,   c                ó   • SU  SU S3$ )Nzjnp.remainder(r    r!   r;   r¾   s     r)   Ú	remainderÚPallasKernelOverrides.remainder›  ó   € à ˜s " Q C qÐ)Ð)r,   c                ó&   • SU  SU SU  SU SU  S3$ )Nz
(jnp.sign(z) * jnp.sign(z) * (jnp.abs(z) // jnp.abs(z))).astype(ú.dtype)r;   r¾   s     r)   ÚtruncdivÚPallasKernelOverrides.truncdivŸ  s.   € ð ˜A˜3˜m¨A¨3¨m¸A¸3¸mÈAÈ3ÈkÐZ[ÐY\Ð\cÐdÐdr,   c                ó   • SU  SU S3$ )Nr  z // r!   r;   r¾   s     r)   ÚfloordivÚPallasKernelOverrides.floordiv¥  r  r,   c                ó   • SU  SU SU S3$ )Nz	jnp.clip(r    r!   r;   )rj   Úmin_valÚmax_vals      r)   ÚclampÚPallasKernelOverrides.clamp©  s   € à˜1˜#˜R ˜y¨¨7¨)°1Ð5Ð5r,   c                ó   • SU  S3$ )Nz	jnp.sign(r!   r;   ri   s    r)   ÚsignÚPallasKernelOverrides.sign°  rw   r,   c                ó   • SU  S3$ )Nzjnp.signbit(r!   r;   ri   s    r)   ÚsignbitÚPallasKernelOverrides.signbit´  s   € à˜a˜S Ð"Ð"r,   c                ó   • SU  S3$ )Nzjax.scipy.special.erf(r!   r;   ri   s    r)   ÚerfÚPallasKernelOverrides.erf¹  s   € à'¨ s¨!Ð,Ð,r,   c                ó   • SU  S3$ )Nzjax.scipy.special.erfc(r!   r;   ri   s    r)   ÚerfcÚPallasKernelOverrides.erfc½  s   € à(¨¨¨1Ð-Ð-r,   c                ó   • SU  S3$ )Nzjax.scipy.special.erfinv(r!   r;   ri   s    r)   ÚerfinvÚPallasKernelOverrides.erfinvÁ  s   € à*¨1¨#¨QÐ/Ð/r,   c                ó   • SU  S3$ )Nzjax.scipy.special.gammaln(r!   r;   ri   s    r)   ÚlgammaÚPallasKernelOverrides.lgammaÅ  ó   € à+¨A¨3¨aÐ0Ð0r,   c                ó   • SU  S3$ )Nzjax.scipy.special.digamma(r!   r;   ri   s    r)   ÚdigammaÚPallasKernelOverrides.digammaÉ  ro  r,   c                ó   • SU  SU  SU  S3$ )Nr   z>.astype(jnp.float64) == 0.0, 1.0, jax.scipy.special.bessel_jn(z&.astype(jnp.float64), v=0)[0]).astype(rR  r;   ri   s    r)   Ú	bessel_j0ÚPallasKernelOverrides.bessel_j0Í  ó+   € ð ˜˜ð +Ø+,¨#ð .Ø�c˜ð"ð	
r,   c                ó   • SU  SU  SU  S3$ )Nr   z>.astype(jnp.float64) == 0.0, 0.0, jax.scipy.special.bessel_jn(z&.astype(jnp.float64), v=1)[1]).astype(rR  r;   ri   s    r)   Ú	bessel_j1ÚPallasKernelOverrides.bessel_j1Ù  rv  r,   c                ó   • SU  SU  S3$ )Nújax.lax.bessel_i0e(ú) * jnp.exp(jnp.abs(ú))r;   ri   s    r)   Úmodified_bessel_i0Ú(PallasKernelOverrides.modified_bessel_i0å  ó   € ð % Q CÐ';¸A¸3¸bÐAÐAr,   c                ó   • SU  SU  S3$ )Nújax.lax.bessel_i1e(r|  r}  r;   ri   s    r)   Úmodified_bessel_i1Ú(PallasKernelOverrides.modified_bessel_i1ë  r€  r,   c                ó   • SU  SU  SU  S3$ )Nr   z == 0.0, 1.0, jnp.sin(z) / r!   r;   ri   s    r)   Úspherical_bessel_j0Ú)PallasKernelOverrides.spherical_bessel_j0ñ  s    € ð ˜A˜3Ð4°Q°C°t¸A¸3¸aÐ@Ð@r,   c                ó   • SU  S3$ )Nr{  r!   r;   ri   s    r)   Úi0eÚPallasKernelOverrides.i0eù  ó   € ð % Q C qÐ)Ð)r,   c                ó   • SU  S3$ )Nr‚  r!   r;   ri   s    r)   Úi1eÚPallasKernelOverrides.i1e   r‹  r,   c                ó   • SU  SU S3$ )Nzjax.scipy.special.gammainc(r    r!   r;   ©rj   Úys     r)   ÚgammaincÚPallasKernelOverrides.gammainc  s   € ð -¨Q¨C¨r°!°°AÐ6Ð6r,   c                ó   • SU  SU S3$ )Nzjax.scipy.special.gammaincc(r    r!   r;   r�  s     r)   Ú	gammainccÚPallasKernelOverrides.gammaincc  s   € ð .¨a¨S°°1°#°QÐ7Ð7r,   c                ó   • SU  SU S3$ )Nzjax.scipy.special.polygamma(z.astype(jnp.int32), r!   r;   r�  s     r)   Ú	polygammaÚPallasKernelOverrides.polygamma  s   € ð .¨a¨SÐ0DÀQÀCÀqÐIÐIr,   c                ó   • SU  S3$ )Nzjax.scipy.special.ndtri(r!   r;   ri   s    r)   ÚndtriÚPallasKernelOverrides.ndtri  s   € ð *¨!¨¨AÐ.Ð.r,   c                ó   • SU  SU S3$ )Nzjax.scipy.special.zeta(r    r!   r;   r�  s     r)   ÚzetaÚPallasKernelOverrides.zeta  s   € ð )¨¨¨2¨a¨S°Ð2Ð2r,   c                ó   • SU  SU S3$ )Nzjax.scipy.special.xlogy(r    r!   r;   r�  s     r)   ÚxlogyÚPallasKernelOverrides.xlogy$  s   € ð *¨!¨¨B¨q¨c°Ð3Ð3r,   c                ó   • SU  SU S3$ )Nzjax.scipy.special.xlog1py(r    r!   r;   r�  s     r)   Úxlog1pyÚPallasKernelOverrides.xlog1py)  s   € ð ,¨A¨3¨b°°°1Ð5Ð5r,   c                ó>   • SU  SU SU  SU  SU SU  SU SU S	U  S
3$ )Nújnp.where(jnp.abs(z) <= 1, jnp.cos(z * jnp.arccos(jnp.clip(z, -1, 1))), jnp.where(z > 1, jnp.cosh(z * jnp.arccosh(jnp.maximum(z, 1.0))), ((-1.0) ** z) * jnp.cosh(z * jnp.arccosh(jnp.maximum(-z
, 1.0)))))r;   ©rj   Úns     r)   Úchebyshev_polynomial_tÚ,PallasKernelOverrides.chebyshev_polynomial_t.  s[   € ð !  ð $Ø�cÐ0°°ð 4Ø˜ð Ø�sÐ5°a°Sð 9Ø˜˜M¨!¨Ð,HÈÈÈ:ð	Wð	
r,   c                óÒ   • SR                  / SPU  PSPU PSPU  PSPU  PSPU  PSPU  PSPU PS	PU PS
PU  PSPU  PSPU  PSPU PSPU PSPU PSPU PSPU  PSPU  PSP5      $ )NÚ r§  z) < 1, jnp.sin((z + 1) * jnp.arccos(jnp.clip(z&, -1, 1))) / jnp.sqrt(jnp.maximum(1 - z**2, 1e-10)), jnp.where(z >= 1, jnp.where(ú == 1, z + 1.0, jnp.sinh((z  + 1) * jnp.arccosh(jnp.maximum(z , 1.0))) / jnp.sqrt(jnp.maximum(z**2 - 1, 1e-10))), jnp.where(z == -1, ((-1.0) ** ú) * (z + 1.0), ((-1.0) ** z) * jnp.sinh((z! + 1) * jnp.arccosh(jnp.maximum(-z**2 - 1, 1e-10)))))©Újoinr¨  s     r)   Úchebyshev_polynomial_uÚ,PallasKernelOverrides.chebyshev_polynomial_u<  sˆ  € ÷	;ó 	;Ð ð 	;  ð 	;ð $ð 	;Ø�sð	;Ø6ð	;Ø78°cð	;ð:(ð	;à() sð	;ð+ð	;ð ˜ð	;ðð	;ð ˜ð		;ð #ð		;ð $% #ð		;ð&ð		;ð
 ˜ð	;ð
 <ð	;ð
 =>¸3ð	;ð
?$ð	;ð %& 3ð	;ð'ð	;ð ˜ð	;ð /ð	;ð 01¨cð	;ð 27ð	;ð 89°cð	;ð:ð	;ð ˜ð	;ð +ð	;ð ,-¨#ð	;ð .Oð	;ð PQÈcð	;ðR$ð	;ð %& 3ð	;ð ':ô	;ð	
r,   c                ó  • SR                  / SPU PSPU  PSPU PSPU  PSPU PSPU  PSPU  PSPU PS	PU  PS
PU  PSPU  PSPU PSPU  PSPU  PSPU  PSPU  PSPU PSPU  PSPU  PSPU  PSPU  PSPU  PSPU  PSP5      $ )Nr­  r   ú == 0, jnp.ones_like(ú), jnp.where(ú	 == 1, 2*ú - 1, jnp.where(ú	 == 2, 4*z**2 - 2*ú	 == 3, 8*z**3 - 4*ú**2 - 4*ú + 1, jnp.where(ú
 == 4, 16*z**4 - 8*ú	**3 - 12*z**2 + 4*ú
 == 5, 32*z	**5 - 16*ú	**4 - 32*z	**3 + 12*ú**2 + 6*z - 1, jnp.zeros_like(ú)))))))r°  r¨  s     r)   Úchebyshev_polynomial_vÚ,PallasKernelOverrides.chebyshev_polynomial_vQ  ó
  € ÷)ó )ˆjð )˜˜ð )Ð0ð )°°ð )ð 4ð )Ø˜ð)Ø$ð)Ø%& Cð)ð(ð)à˜ð)à$ð)à%& Cð)à'/ð)à01¨sð)ð3ð)ð ˜ð)ð %ð)ð &' Cð)ð (0ð)ð 12¨sð)ð 3;ð)ð <=¸#ð)ð>ð)ð ˜ð	)ð &ð	)ð '( Sð	)ð )1ð	)ð 23°ð	)ð 4=ð	)ð >?¸Cð	)ð @Hð	)ð IJÀsð	)ðKð	)ð
 ˜ð)ð
 &ð)ð
 '( Sð)ð
 )2ð)ð
 34°ð)ð
 5>ð)ð
 ?@¸Sð)ð
 AJð)ð
 KLÈð)ð
 MUð)ð
 VWÐTWð)ð
Xð)ð  ˜Sð)ð !(ô)ð	
r,   c                ó  • SR                  / SPU PSPU  PSPU PSPU  PSPU PSPU  PSPU  PS	PU PS
PU  PSPU  PSPU  PS	PU PSPU  PSPU  PSPU  PSPU  PSPU PSPU  PSPU  PSPU  PSPU  PSPU  PSPU  PSP5      $ )Nr­  r   rµ  r¶  r·  r¼  r¹  z**2 + 2*r¸  rº  z**3 + 4*r»  r½  z**4 + 8*r¾  r¿  z	**5 + 16*rÀ  rÁ  z + 1, jnp.zeros_like(rÂ  r°  r¨  s     r)   Úchebyshev_polynomial_wÚ,PallasKernelOverrides.chebyshev_polynomial_wa  rÅ  r,   c                ó6   • [         R                  SU  S3U5      $ ©Nz(2 * z - 1))rf   rª  r¨  s     r)   Úshifted_chebyshev_polynomial_tÚ4PallasKernelOverrides.shifted_chebyshev_polynomial_tq  ó   € ä$×;Ñ;¸eÀAÀ3ÀeÐ<LÈaÓPÐPr,   c                ó6   • [         R                  SU  S3U5      $ rÊ  )rf   r²  r¨  s     r)   Úshifted_chebyshev_polynomial_uÚ4PallasKernelOverrides.shifted_chebyshev_polynomial_uu  rÍ  r,   c                ó6   • [         R                  SU  S3U5      $ rÊ  )rf   rÃ  r¨  s     r)   Úshifted_chebyshev_polynomial_vÚ4PallasKernelOverrides.shifted_chebyshev_polynomial_vy  rÍ  r,   c                ó6   • [         R                  SU  S3U5      $ rÊ  )rf   rÇ  r¨  s     r)   Úshifted_chebyshev_polynomial_wÚ4PallasKernelOverrides.shifted_chebyshev_polynomial_w}  rÍ  r,   c                óÒ   • SR                  / SPU PSPU  PSPU PSPU  PSPU PSPU  PSPU PS	PU  PS
PU  PSPU PSPU  PSPU  PSPU PSPU  PSPU  PSPU  PSPU  PSP5      $ )Nr­  r   rµ  r¶  z == 1, 2 * ú, jnp.where(z == 2, 4 * z**2 - 2, jnp.where(z == 3, 8 * z**3 - 12 * z == 4, 16 * z**4 - 48 * z**2 + 12, jnp.where(z == 5, 32 * z**5 - 160 * z**3 + 120 * ú, jnp.zeros_like(rÂ  r°  r¨  s     r)   Úhermite_polynomial_hÚ*PallasKernelOverrides.hermite_polynomial_h�  s~  € ÷)ó )ˆjð )˜˜ð )Ð0ð )°°ð )ð 4ð )Ø˜ð)Ø&ð)Ø'( cð)ð*ð)à˜ð)à&ð)à'( cð)ð*ð)ð ˜ð)ð 'ð)ð () cð)ð *5ð)ð 67°Cð)ð8ð)ð ˜ð	)ð (ð	)ð )* sð	)ð +6ð	)ð 78°Sð	)ð9ð	)ð
 ˜ð)ð
 (ð)ð
 )* sð)ð
 +7ð)ð
 89°cð)ð
 :Fð)ð
 GHÀSð)ð
Ið)ð  ˜Sð)ð !(ô)ð	
r,   c                óÒ   • SR                  / SPU PSPU  PSPU PSPU  PSPU PSPU  PSPU PS	PU  PS
PU  PSPU PSPU  PSPU  PSPU PSPU  PSPU  PSPU  PSPU  PSP5      $ )Nr­  r   rµ  r¶  r®  rØ  z == 2, z**2 - 1, jnp.where(z == 3, ú
**3 - 3 * z == 4, z
**4 - 6 * z**2 + 3, jnp.where(z == 5, z**5 - 10 * ú**3 + 15 * rÙ  rÂ  r°  r¨  s     r)   Úhermite_polynomial_heÚ+PallasKernelOverrides.hermite_polynomial_he“  s}  € ÷
)ó )ˆjð )˜˜ð )Ð0ð )°°ð )ð 4ð )Ø˜ð)Ø"ð)Ø#$ #ð)ð&ð)à˜ð)à"ð)à#$ #ð)ð&ð)ð ˜ð)ð #ð)ð $% #ð)ð &0ð)ð 12¨sð)ð3ð)ð ˜ð	)ð #ð	)ð $% #ð	)ð &0ð	)ð 12¨sð	)ð3ð	)ð
 ˜ð)ð
 #ð)ð
 $% #ð)ð
 &1ð)ð
 23°ð)ð
 4?ð)ð
 @A¸cð)ð
Bð)ð  ˜Sð)ð !(ô)ð	
r,   c                ó  • SR                  / SPU PSPU  PSPU PSPU  PSPU PSPU  PSPU  PS	PU PS
PU  PSPU  PSPU  PSPU PSPU  PSPU  PSPU  PSPU  PSPU PSPU  PSPU  PSPU  PSPU  PSPU  PSPU  PSP5      $ )Nr­  r   rµ  r¶  z == 1, 1 - rØ  z == 2, (r»  z + 2) / 2, jnp.where(z	 == 3, (-z**3 + 9*z	**2 - 18*z + 6) / 6, jnp.where(z == 4, (z	**4 - 16*z	**3 + 72*z	**2 - 96*z + 24) / 24, jnp.where(z	 == 5, (-z	**5 + 25*z
**4 - 200*z
**3 + 600*z
**2 - 600*z + 120) / 120, jnp.zeros_like(rÂ  r°  r¨  s     r)   Úlaguerre_polynomial_lÚ+PallasKernelOverrides.laguerre_polynomial_l¡  s	  € ÷
)ó )ˆjð )˜˜ð )Ð0ð )°°ð )ð 4ð )Ø˜ð)Ø&ð)Ø'( cð)ð*ð)à˜ð)à#ð)à$% 3ð)à&.ð)à/0¨cð)ð2ð)ð ˜ð)ð %ð)ð &' Cð)ð (0ð)ð 12¨sð)ð 3<ð)ð =>¸3ð)ð?ð)ð ˜ð	)ð $ð	)ð %& 3ð	)ð '0ð	)ð 12¨sð	)ð 3<ð	)ð =>¸3ð	)ð ?Hð	)ð IJÀsð	)ðKð	)ð
 ˜ð)ð
 %ð)ð
 &' Cð)ð
 (1ð)ð
 23°ð)ð
 4>ð)ð
 ?@¸Sð)ð
 AKð)ð
 LMÈ#ð)ð
 NXð)ð
 YZÐWZð)ð
[ð)ð  ˜Sð)ð !(ô)ð	
r,   c                óÒ   • SR                  / SPU PSPU  PSPU PSPU  PSPU PSPU  PSPU PS	PU  PS
PU  PSPU PSPU  PSPU  PSPU PSPU  PSPU  PSPU  PSPU  PSP5      $ )Nr­  r   rµ  r¶  r®  rØ  z == 2, (3 * z**2 - 1) / 2, jnp.where(z == 3, (5 * rÝ  z) / 2, jnp.where(z == 4, (35 * z**4 - 30 * z**2 + 3) / 8, jnp.where(z == 5, (63 * z**5 - 70 * rÞ  z) / 8, jnp.zeros_like(rÂ  r°  r¨  s     r)   Úlegendre_polynomial_pÚ+PallasKernelOverrides.legendre_polynomial_p¯  s~  € ÷
)ó )ˆjð )˜˜ð )Ð0ð )°°ð )ð 4ð )Ø˜ð)Ø"ð)Ø#$ #ð)ð&ð)à˜ð)à'ð)à() sð)ð+ð)ð ˜ð)ð (ð)ð )* sð)ð +5ð)ð 67°Cð)ð8ð)ð ˜ð	)ð )ð	)ð *+¨ð	)ð ,7ð	)ð 89°cð	)ð:ð	)ð
 ˜ð)ð
 )ð)ð
 *+¨ð)ð
 ,7ð)ð
 89°cð)ð
 :Eð)ð
 FGÀCð)ð
Hð)ð  ˜Sð)ð !(ô)ð	
r,   c                ó   • SU  S3$ )Nzjnp.reciprocal(r!   r;   ri   s    r)   Ú
reciprocalÚ PallasKernelOverrides.reciprocal¾  r¹   r,   c                ó   • SU  S3$ )Nzjnp.square(r!   r;   ri   s    r)   ÚsquareÚPallasKernelOverrides.squareÂ  r�   r,   c                ó   • SU  SU SU S3$ )z�Fused multiply-add: a * b + c

JAX doesn't have jnp.fma, so we use the unfused version.
The compiler may still fuse this on supported hardware.
z((r¯  z) + (r}  r;   )r¿   rÀ   r&   s      r)   ÚfmaÚPallasKernelOverrides.fmaÇ  s   € ð �A�3�e˜A˜3˜e A 3 bÐ)Ð)r,   c                ó   • SU  SU S3$ )Nzjnp.copysign(r    r!   r;   r¾   s     r)   ÚcopysignÚPallasKernelOverrides.copysignÐ  s   € à˜q˜c  A 3 aÐ(Ð(r,   c                ó   • SU  SU S3$ )Nzjnp.nextafter(r    r!   r;   r¾   s     r)   Ú	nextafterÚPallasKernelOverrides.nextafterÔ  rP  r,   c                ó   • SU  SU S3$ )Nz
jnp.ldexp(r    r!   r;   r¾   s     r)   ÚldexpÚPallasKernelOverrides.ldexpØ  rÃ   r,   c                ó   • SU  S3$ )Nz
jnp.frexp(r!   r;   ri   s    r)   ÚfrexpÚPallasKernelOverrides.frexpÜ  r‘   r,   c                ó   • SU  S3$ )Nz	jnp.modf(r!   r;   ri   s    r)   ÚmodfÚPallasKernelOverrides.modfà  rw   r,   c                ó   • SU  SU S3$ )Nzjnp.bitwise_and(r    r!   r;   r¾   s     r)   Úbitwise_andÚ!PallasKernelOverrides.bitwise_andå  r8  r,   c                ó   • SU  SU S3$ )Nzjnp.bitwise_or(r    r!   r;   r¾   s     r)   Ú
bitwise_orÚ PallasKernelOverrides.bitwise_oré  r<  r,   c                ó   • SU  SU S3$ )Nzjnp.bitwise_xor(r    r!   r;   r¾   s     r)   Úbitwise_xorÚ!PallasKernelOverrides.bitwise_xorí  r8  r,   c                ó   • SU  S3$ )Nzjnp.bitwise_not(r!   r;   ri   s    r)   Úbitwise_notÚ!PallasKernelOverrides.bitwise_notñ  r@  r,   c                ó   • SU  SU S3$ )Nzjnp.left_shift(r    r!   r;   r¾   s     r)   Ú
left_shiftÚ PallasKernelOverrides.left_shiftõ  r<  r,   c                ó   • SU  SU S3$ )Nzjnp.right_shift(r    r!   r;   r¾   s     r)   Úright_shiftÚ!PallasKernelOverrides.right_shiftù  r8  r,   c                ó¶   • [         R                  R                  R                  SU5      nS[         R                  R                  R	                  U 5       SU S3$ )z)Load the random seed value from a buffer.Úload_seed_offsetr  z[0] + r!   )r   rï   r#   Úseed_offsetÚinput)ÚnameÚoffsetr  s      r)   Ú	load_seedÚPallasKernelOverrides.load_seedþ  sH   € ô —h‘h—m‘m×/Ñ/Ð0BÀFÓKˆØ”1—8‘8—=‘=×&Ñ& tÓ,Ð-¨V°K°=ÀÐBÐBr,   c                ó   • SU  SU SU S3$ )z¼Generate uniform random numbers in [0, 1).

Uses JAX's threefry2x32 PRNG directly for vectorized random generation.
The seed provides the base key, offset provides per-element uniqueness.
zWjax.vmap(lambda o: jax.random.uniform(jax.random.fold_in(jax.random.PRNGKey(jnp.uint32(ú8)), jnp.uint32(o)), (), dtype=jnp.float32))(jnp.asarray(ú!).flatten()).reshape(jnp.asarray(ú).shape)r;   ©Úseedr  s     r)   ÚrandÚPallasKernelOverrides.rand  s,   € ð@Ø@D¸vð FØ"˜8Ð#DÀVÀHÈHðVð	
r,   c                ó   • SU  SU SU S3$ )zºGenerate standard normal random numbers.

Uses JAX's threefry2x32 PRNG directly for vectorized random generation.
The seed provides the base key, offset provides per-element uniqueness.
zVjax.vmap(lambda o: jax.random.normal(jax.random.fold_in(jax.random.PRNGKey(jnp.uint32(r  r  r  r;   r  s     r)   ÚrandnÚPallasKernelOverrides.randn  s,   € ð@Ø@D¸vð FØ"˜8Ð#DÀVÀHÈHðVð	
r,   c                ó&   • SU  SU SU SU SU S3$ )z,Generate random int64 values in [low, high).zWjax.vmap(lambda o: jax.random.randint(jax.random.fold_in(jax.random.PRNGKey(jnp.uint32(z)), jnp.uint32(o)), (), r    z , dtype=jnp.int64))(jnp.asarray(r  r  r;   )r  r  ÚlowÚhighs       r)   Ú	randint64ÚPallasKernelOverrides.randint64$  s>   € ð
@Ø@D¸vÐE]Ð^aÐ]bÐbdÐeiÐdjð kØ"˜8Ð#DÀVÀHÈHðVð	
r,   r;   )rj   r>   r=   r>   )r¿   r>   rÀ   r>   r=   r>   )rÌ   r>   r¿   r>   rÀ   r>   r=   r>   )rÙ   r>   rÚ   zCallable[[], str]rÛ   rÔ   r=   r>   )NT)
rj   r>   rá   útorch.dtyperâ   zOptional[torch.dtype]rã   r  r=   r>   )rj   r>   rá   r)  râ   r)  r=   r>   )r%   r<   rá   r)  r=   r>   )rá   r)  r=   r>   )rj   r>   rY  r>   rZ  r>   r=   r>   )rj   r>   r‘  r>   r=   r>   )rj   r>   r©  r>   r=   r>   )r¿   r>   rÀ   r>   r&   r>   r=   r>   )r  r>   r  r>   r=   r>   )r  r>   r  r>   r=   r>   )
r  r>   r  r>   r%  r>   r&  r>   r=   r>   )|r?   r@   rA   rB   rC   Ústaticmethodrk   ro   rr   ru   ry   r|   r   rƒ   r†   r‰   rŒ   r�   r“   r–   r™   rœ   rŸ   r¢   r¥   r¨   r«   r®   r±   r´   r·   r»   rÁ   rÅ   rÉ   rÍ   rÝ   rå   ré   rý   r  r  r  r  r  r  r  r  r  r"  r&  r)  rÖ   r×   r0  r3  r6  r:  r>  rB  rE  rH  rK  rN  rS  rV  r[  Úclipr^  ra  rd  rg  rj  rm  rq  rt  rx  r~  rƒ  r†  Úi0r‰  Úi1r�  r’  r•  ÚigammaÚigammacr˜  r›  rž  r¡  r¤  rª  r²  rÃ  rÇ  rË  rÏ  rÒ  rÕ  rÚ  rß  râ  rå  rè  rë  rî  rñ  rô  r÷  rú  rý  r   r  r  r	  r  r  r  r  r"  r'  rD   r;   r,   r)   rf   rf   x   s0	  † ñð óó ðð óó ðð óó ðð ó ó ð ð ó ó ð ð ó ó ð ð ó"ó ð"ð ó"ó ð"ð ó"ó ð"ð óó ðð ó ó ð ð ó!ó ð!ð óó ðð ó!ó ð!ð ó ó ð ð ó!ó ð!ð ó ó ð ð ó%ó ð%ð óó ðð óó ðð ó!ó ð!ð ó ó ð ð ó!ó ð!ð ó!ó ð!ð ó&ó ð&ð ó&ó ð&ð ó&ó ð&ð ó(ó ð(ð ó(ó ð(ð ó.ó ð.ð ó;ó ð;ð& ð ,0Ø"&ð	7Øð7àð7ð )ð7ð  ð	7ð
 
ô7ó ð7ð ófó ðfð ó:ó ð:ð  ó6ó ð6ð ó ó ð ð ó ó ð ð ó ó ð ð ó!ó ð!ð óEó ðEð ó3ó ð3ð
 óó ðð óó ðð óó ðð óó ðð óó ðð ó!ó ð!ð ó!ó ð!ð ó$ó ð$ð óó ðð ó,ó ð,ð ó+ó ð+ð ó'ó ð'ð ó,ó ð,ð ó(ó ð(ð ó&ó ð&ð ó%ó ð%ð ó*ó ð*ð óeó ðeð
 óó ðð ó6ó ð6ð €Dð ó ó ð ð ó#ó ð#ð ó-ó ð-ð ó.ó ð.ð ó0ó ð0ð ó1ó ð1ð ó1ó ð1ð ó	
ó ð	
ð ó	
ó ð	
ð óBó ðBð
 óBó ðBð
 óAó ðAð
 
€Bàó*ó ð*ð 
€Bàó*ó ð*ð ó7ó ð7ð
 ó8ó ð8ð €Fà€GàóJó ðJð
 ó/ó ð/ð ó3ó ð3ð ó4ó ð4ð ó6ó ð6ð ó
ó ð
ð ó
ó ð
ð( ó
ó ð
ð ó
ó ð
ð óQó ðQð óQó ðQð óQó ðQð óQó ðQð ó
ó ð
ð" ó
ó ð
ð ó
ó ð
ð ó
ó ð
ð ó&ó ð&ð ó"ó ð"ð ó*ó ð*ð ó)ó ð)ð ó*ó ð*ð ó&ó ð&ð ó!ó ð!ð ó ó ð ð ó,ó ð,ð ó+ó ð+ð ó,ó ð,ð ó'ó ð'ð ó+ó ð+ð ó,ó ð,ð óCó ðCð ó
ó ð
ð ó
ó ð
ð ó
ó ó
r,   rf   c                  óÄ   • \ rS rSr% SrS\S'   S\S'   S\S'   S\S	'   S\S
'   S\S'   S\S'   S\S'   S\S'   S\S'   S\S'   S\S'   S\S'   S\S'   S\S'   S\S'   S\S'   Srg)Ú_CodegenContexti/  z@Bundles local state shared across codegen_kernel helper methods.r   Úcoder>   Úkernel_namer  Úis_tpuÚinterpret_is_cpuÚinterpret_literalú	list[str]Úkernel_paramsÚpure_out_paramsÚoutput_paramsÚsize_var_paramszdict[str, str]Úoutput_buffer_lookupzdict[str, bool]Úaliasable_flagsÚalias_paramsÚpointer_tailÚkernel_input_paramsÚfull_kernel_paramszOrderedSet[str]Únon_alias_out_setz	list[int]Úcopy_output_indicesr;   N)r?   r@   rA   rB   rC   Ú__annotations__rD   r;   r,   r)   r1  r1  /  sj   ‡ áJà
ÓØÓØƒLØÓØÓØÓØÓØÓØÓØ(Ó(Ø$Ó$ØÓØÓØ"Ó"Ø!Ó!Ø&Ó&Ø"Ö"r,   r1  c                  ó¶  ^ • \ rS rSr% Sr\r\rS\	S'   U 4S jr
          S=S jrS>S jrS>S jrS>S	 jrS?S
 jrS@S jrSAS jrSBS jrSAS jrSCS jrSDS jrSES jrSFS jr\SGS j5       rSHS jrSIS jrSES jrSJS jr    SKS jr      SLS jr          SMS jr           SNS jr!          SNS jr"          SNS jr#            SOS jr$SPS jr%        SQS  jr&        SRS! jr'        SSS" jr(STS# jr)        SUS$ jr* SV               SWS% jjr+            SXS& jr,\-R\                  SYS' j5       r/S>S( jr0\-R\                   SV         SZS) jj5       r1\ S[       S\S* jj5       r2 S]     S^S+ jjr3        S_S, jr4        S`S- jr5          SaS. jr6\SbS/ j5       r7SVScS0 jjr8SdS1 jr9      SeS2 jr:\        SfS3 j5       r;SgS4 jr<      SgS5 jr=          ShS6 jr>SiS7 jr?      SiS8 jr@      SiS9 jrA\          SjS: j5       rBSVSkS; jjrCS<rDU =rE$ )lÚPallasKerneliF  aƒ  
Pallas kernel for elementwise operations with support for strided/scatter access.

Strategy:
- Convert index expressions to JAX-compatible array slicing
- Load/store using indexed access: "in_ptrX[slice]" or full-array "in_ptrX[...]"
- Compute expression with Python operators (compatible with jax.numpy broadcasting)
- Generate Python code that defines a Pallas kernel and a host entrypoint.
- Use async_compile.pallas path to compile and load Python code.

For GPU (Mosaic backend):
- Use TMA (Tensor Memory Accelerator) for automatic OOB masking
- Falls back to legacy padding approach for reductions, broadcasting, non-contiguous tensors
zCallable[[sympy.Expr], str]rõ   c                ód  >• [         TU ]  " U0 UD6  [        R                  R	                  5       nUR
                  S:H  U l        U R                  U l        U R                  =(       a    U R                  (       + U l        / U l	        0 U l
        [        5       U l        SU l        [        5       U l        g )NÚcudaF)ÚsuperrY   r   ÚgraphÚget_current_device_or_throwÚtypeÚis_gpuÚuse_emit_pipelineÚuse_warpgroup_paddingÚstore_with_outputÚload_index_exprsr   Úoutputs_need_readÚhas_transposed_loadrð   )r$   r#   r]   ÚdeviceÚ	__class__s       €r)   rY   ÚPallasKernel.__init__Y  s�   ø€ Ü‰Ò˜$Ð) &Ò)ä—‘×4Ñ4Ó6ˆØ—k‘k VÑ+ˆŒð "&§¡ˆÔà%)§[¡[×%O¸×9OÑ9OÔ5OˆÔ"à8:ˆÔà79ˆÔä2<³,ˆÔð $)ˆÔ ä8B»ˆÕr,   c                ó   • g)z)Check array bounds for indirect indexing.Nr;   )r$   r%   rN   ÚlowerÚuppers        r)   Úcheck_boundsÚPallasKernel.check_boundsp  s   � r,   c                ó¾   • U R                  U5      n[        U[        R                  5      (       a  gUR                  (       a  [        U5      $ U R                  U5      $ )a¾  
Convert an index expression to a string suitable for Pallas indexing.

Pallas operates on full arrays, so we need to convert index expressions
to JAX array slicing. For example:
- x0 -> "..." (contiguous access, full array)
- 2*x0 -> "::2" (strided access with stride 2)
- 2*x0 + 1 -> "1::2" (strided access with offset 1, stride 2)

Args:
    index: The indexing expression to convert

Returns:
    The indexing string to use in generated code
ú...)ró   rÓ   ÚsympyÚSymbolÚ
is_Integerr>   Ú_convert_to_jax_slice)r$   ÚindexÚprepared_indexs      r)   Ú_get_index_strÚPallasKernel._get_index_strw  sQ   € ð" ×.Ñ.¨uÓ5ˆô �n¤e§l¡l×3Ñ3ØØ×&×&ä�~Ó&Ð&ð ×-Ñ-¨nÓ=Ð=r,   c                ó   • U R                   (       d  gU R                  U5      nUR                  [        5      (       a;  U R                  R                  U R                  U5      5        U R                  U5      $ [        R                  R                  R                  U5      nU R                  U5      nU R                  R                  U5        [        U5      S:X  a  [        U5      $ [        U5      S:X  Ga  [        [        U5      5      n[         R"                  " X5      n[         R$                  " XC5      nUbš  X-
  n[        R                  R                  R                  U5      nUS:  a  U R                  U5      $ US:X  a  gUS:w  a  U R                  U5      $  ['        U5      nUS:  a  U R                  U5      $  U R                  U5       S3$ X-
  n[        R                  R                  R                  U5      nUS:X  a  XC:X  a  gg[        U5      S:”  aJ  SnU H9  n[         R"                  " X5      n[         R$                  " XC5      nUS:w  d  M7  Sn  O   U(       a  ggg! [(        [*        4 a    U R                  U5      s $ f = f)a  
Convert a sympy index expression to JAX slice notation.

Handles common patterns like:
- stride*var -> ::stride
- stride*var + offset -> offset::stride

For more complex patterns, falls back to explicit indexing.
Uses BlockPatternMatcher for robust pattern matching.
r]  r   r   z::1TF)Úrange_treesrô   Úhasr	   rð   rñ   rò   rõ   r   rJ  ÚsizevarsÚsimplifyÚlenr>   ÚnextÚiterr   Úget_subexpr_involving_symbolÚmatch_affine_block_exprÚintÚ	TypeErrorÚ
ValueError)	r$   rb  Ú	used_varsrü   Úvar_exprÚstrider  Ú
offset_valÚall_unit_strides	            r)   ra  Ú"PallasKernel._convert_to_jax_sliceš  sX  € ð ××Øð ×$Ñ$ UÓ+ˆð �9‰9”_×%Ñ%à×Ñ×&Ñ& t×'?Ñ'?ÀÓ'FÔGð —:‘:˜eÓ$Ð$ô —‘× Ñ ×)Ñ)¨%Ó0ˆà×,Ñ,¨UÓ3ˆ	ð 	×Ñ×"Ñ" 9Ô-äˆy‹>˜QÓä�u“:ÐÜ�‹^˜qÔ ä”t˜I“Ó'ˆCô +×GÒGÈÓSˆHô )×@Ò@ÀÓOˆFàÑ!ØÑ)�ÜŸ™×)Ñ)×2Ñ2°6Ó:�à˜A“:ØŸ:™: eÓ,Ð,à˜Q“;Ø ð ˜Q“;ØŸ:™: eÓ,Ð,ð-Ü!$ V£�JØ! A“~Ø#Ÿz™z¨%Ó0Ð0ð &ð
 Ÿ*™* VÓ,Ð-¨SÐ1Ð1ð Ñ)�ÜŸ™×)Ñ)×2Ñ2°6Ó:�Ø˜Q“; 8£?à ð, ô+ �‹^˜aÓð #ˆOÛ �Ü.×KÒKÈEÓW�Ü,×DÒDÀXÓS�Ø˜Q•;Ø&+�OÙñ !ö àð ð øôA "¤:Ð.ó -ØŸ:™: eÓ,Ò,ð-ús   Æ!I) É)!JÊJc                óÄ   • UR                   nU R                  5       nX#-  nXB:w  a  [        SU 35      eU R                  R	                  U5        U R                  U5      nU$ )a~  
Generate JAX code to compute an index array for strided/complex indexing patterns.

For expressions like `2 * x3 + 32 * x2 + 256 * x1 + 1024 * x0`, we generate
code that computes the flattened index array using broadcasting.

The iteration variables (x0, x1, x2, x3) are already defined as jnp.arange arrays
in the kernel. We just need to convert the sympy expression to JAX code.
z9Pallas backend does not yet support mixed index pattern: )Úfree_symbolsÚ_get_iter_varsrc   rð   rñ   rõ   )r$   rb  rz  Ú	iter_varsrs  Ú	index_strs         r)   Ú_generate_strided_indexÚ$PallasKernel._generate_strided_indexþ  sr   € ð ×)Ñ)ˆØ×'Ñ'Ó)ˆ	ð !Ñ,ˆ	ØÓ$ÜØKÈEÈ7ÐSóð ð
 	×Ñ×"Ñ" 9Ô-ð —J‘J˜uÓ%ˆ	ð Ðr,   c                óH   • [        U R                  R                  5       5      $ )z*Get the set of iteration variable symbols.)r   Úrange_tree_nodesÚkeys)r$   s    r)   r{  ÚPallasKernel._get_iter_vars  s   € ä˜$×/Ñ/×4Ñ4Ó6Ó7Ð7r,   c                ó<   • UR                   U R                  5       -  $ )z4Get iteration variables used in an index expression.)rz  r{  ©r$   rb  s     r)   rò   Ú PallasKernel._get_used_iter_vars   s   € à×!Ñ! D×$7Ñ$7Ó$9Ñ9Ð9r,   c                ó6   • [        U R                  U5      5      $ )z7Check if index expression contains iteration variables.)r  rò   r…  s     r)   Ú_has_iteration_varsÚ PallasKernel._has_iteration_vars$  s   € ä�D×,Ñ,¨UÓ3Ó4Ð4r,   c                óˆ   • UR                    Vs/ s H&  n[        U5      R                  S5      (       d  M$  UPM(     sn$ s  snf )zDGet list of indirect variable symbols (tmp*) in an index expression.Útmp)rz  r>   Ú
startswith)r$   rb  Úss      r)   Ú_get_indirect_varsÚPallasKernel._get_indirect_vars(  s3   € à ×-Ò-ÓJÒ-�a´°Q³×1BÑ1BÀ5×1I—Ñ-ÑJÐJùÒJs   �#?¶?c                ó<   • [        U R                  U5      5      S:„  $ )z6Check if index expression contains indirect variables.r   )rk  rŽ  r…  s     r)   Ú_has_indirect_varsÚPallasKernel._has_indirect_vars,  s   € ä�4×*Ñ*¨5Ó1Ó2°QÑ6Ð6r,   c                ó  • [        U R                  R                  5       5      n/ nU H6  u  p4U R                  UR                  5      nUc  M%  UR                  U5        M8     [        U5      S::  a  U$ [        [        U5      5      $ )a  Get the expected output shape from iteration variables.

Iteration variables are shaped for broadcasting. For 2D outputs:
- First var (e.g., y0) gets shape (1, N) - innermost dimension
- Second var (e.g., x1) gets shape (M, 1) - outermost dimension
The broadcast result is (M, N).
r   )Úlistr�  ÚitemsÚ	_safe_intÚlengthÚappendrk  Úreversed)r$   Ú	var_itemsÚbroadcast_varsÚvar_symÚentryr—  s         r)   Ú_get_expected_output_shapeÚ'PallasKernel._get_expected_output_shape0  s{   € ô ˜×.Ñ.×4Ñ4Ó6Ó7ˆ	ØˆÛ'‰NˆGØ—^‘^ E§L¡LÓ1ˆFØÓ!Ø×%Ñ% fÖ-ñ (ô
 ˆ~Ó !Ó#Ø!Ð!ô
 ”H˜^Ó,Ó-Ð-r,   c                óV  • U R                  U5      nUc  gUu  pEpFn[        U5      S:w  d  [        U5      S:w  a  gU R                  US   5      nU R                  US   5      nUb  Ub  US::  d  US::  a  gUS   n	US   n
U	b  U
c  g[        U R                  R                  5       5      n[        U5      S:  a  g[        S U 5       5      (       a  gUS   S   nUS   S   n[        R                  R                  R                  U5      nU R                  X,5      nU R                  X-5      nUS:w  aw  US:w  aq  X˜:H  =(       a    U
S:H  nU(       d  gU R                  5       nU(       d  g[        Xé-
  5      [        Xê-
  5      :  n[        Xú-
  5      [        Xù-
  5      :  nU=(       a    U$ g)a  Check if buffer access needs transpose.

Transpose on load is needed when:
1. Non-square buffers: dimensions are swapped relative to iteration vars
2. Square buffers: index coefficient pattern indicates transposed access
   (first iteration var has larger coefficient than second)
Fr
   r   r   c              3  ó>   #   • U  H  u  pUR                   v •  M     g 7fr`   ©Úis_reduction)Ú.0Ú_r�  s      r)   Ú	<genexpr>Ú5PallasKernel._is_transposed_access.<locals>.<genexpr>k  s   é € Ð<²)¡h aˆu×!Ö!²)ùs   ‚)Ú_get_buffer_infork  r–  r”  r�  r•  Úanyr   rJ  ri  rj  Ú_get_index_coefficientÚ_has_column_major_outputr¥   )r$   r  rb  rX   r¥  Úbuf_sizeÚactual_stridesÚsize0Úsize1Ús0Ús1rš  Ú	inner_varÚ	outer_varÚinner_coeffÚouter_coeffÚis_standard_row_majorÚoutput_is_column_majorÚinner_matches_s0Úouter_matches_s1s                       r)   Ú_is_transposed_accessÚ"PallasKernel._is_transposed_accessH  s»  € ð ×$Ñ$ TÓ*ˆØ‰<Øà,0Ñ)ˆ�Q¨ô ˆx‹=˜AÓ¤ ^Ó!4¸Ó!9Øà—‘˜x¨™{Ó+ˆØ—‘˜x¨™{Ó+ˆØ‰=˜E™M¨U°a«Z¸5ÀA»:Øà˜AÑˆØ˜AÑˆØ‰:˜™Øô ˜×.Ñ.×4Ñ4Ó6Ó7ˆ	Üˆy‹>˜AÓØô Ñ<±)Ó<×<Ñ<Øð ˜a‘L ‘Oˆ	Ø˜a‘L ‘Oˆ	Ü—‘× Ñ ×)Ñ)¨%Ó0ˆà×1Ñ1°%ÓCˆØ×1Ñ1°%ÓCˆà˜!Ó ¨qÓ 0à$&¡K×$;°B¸!±GÐ!Þ(Øð &*×%BÑ%BÓ%DÐ"Þ)Øô  # ;Ñ#3Ó4´s¸;Ñ;KÓ7LÑLÐÜ" ;Ñ#3Ó4´s¸;Ñ;KÓ7LÑLÐØ#×8Ð(8Ð8àr,   c                ó  • [        U R                  S0 5      n[        R                  " U[        R
                  R                  5      nU Hµ  n[        R
                  R                  U5      nUc  M'  X1;  a  [        U[        5      (       d  MC  [        USS 5      " 5       nUc  M[  [        USS5      nUb  [        U5      S:  a  M|  U R                  US   5      nU R                  US   5      nUc  M©  Uc  M®  Xx:  d  Mµ    g	   g
)z:Check if any output buffer has column-major stride layout.Úoutput_buffersNÚ
get_layoutc                 ó   • g r`   r;   r;   r,   r)   Ú<lambda>Ú7PallasKernel._has_column_major_output.<locals>.<lambda>–  s   € ¸Dr,   ru  r
   r   r   TF)Úgetattrr#   Ú	itertoolsÚchainr   rJ  Úname_to_bufferÚ
get_bufferrÓ   r   rk  r–  )	r$   r½  Ú	buf_namesÚbuf_nameÚout_bufÚlayoutÚ
out_strideÚout_s0Úout_s1s	            r)   r«  Ú%PallasKernel._has_column_major_outputˆ  sé   € ä  §¡Ð,<¸bÓAˆô —O’O N´A·G±G×4JÑ4JÓKˆ	Û!ˆHÜ—g‘g×(Ñ(¨Ó2ˆGØ‰ÙØÓ-´jØœ÷7ñ 7ñ Ü˜W l±LÔAÓCˆFØ‰~ÙÜ  ¨°4Ó8ˆJØÑ!¤S¨£_°qÓ%8ÙØ—^‘^ J¨q¡MÓ2ˆFØ—^‘^ J¨q¡MÓ2ˆFØÓ! fÓ&8¸V½_Ùñ# "ð& r,   c                ó˜  • U R                  U5      nU R                  U5      nU(       a  U(       a  U R                  U5      S4$ U(       a  U R                  U5      S4$ U R	                  U5      nUR                  [        5      =(       a    US:g  nU(       d2  US:w  a,  SU;   d&  UR                  S5      R                  5       (       d  SnXE4$ )z@Get the index expression string and whether it needs flattening.TFr]  ú::Ú-)	r‘  rˆ  Ú_handle_mixed_indexingrõ   rd  rh  r	   ÚlstripÚisdigit)r$   rb  Úhas_indirectÚhas_iter_varsr}  Úneeds_flattens         r)   Ú_get_index_exprÚPallasKernel._get_index_expr£  s¹   € à×.Ñ.¨uÓ5ˆØ×0Ñ0°Ó7ˆæžMØ×.Ñ.¨uÓ5°tÐ;Ð;ÞØ—:‘:˜eÓ$ eÐ+Ð+à×+Ñ+¨EÓ2ˆIð "ŸI™I¤oÓ6×M¸9ÈÑ;MˆMö ! Y°%Ó%7à 	Ó)¨Y×-=Ñ-=¸cÓ-B×-JÑ-J×-LÑ-LØ$(�MØÐ+Ð+r,   c                óF   •  [        U 5      $ ! [        [        4 a     gf = f)z0Convert value to int, returning None on failure.N)rp  rq  rr  )r  s    r)   r–  ÚPallasKernel._safe_int¹  s'   € ð	Ü�s“8ˆOøÜœ:Ð&ó 	Ùð	ús   ‚
 � Ÿ c                óŒ   • SnU H;  nX0R                   ;   d  M  U R                  U R                   U   5      nUc    gX$-  nM=     U$ )zBCompute total numel for given prefixes (e.g., pointwise prefixes).r   N)Únumelsr–  )r$   Úprefixesr2   r'   Únumels        r)   Ú_compute_prefix_numelÚ"PallasKernel._compute_prefix_numelÁ  sG   € àˆÛˆAØ—K‘KÕØŸ™ t§{¡{°1¡~Ó6�Ø‘=ÙØ‘’ñ ð ˆr,   c                óž   • SnU R                    H:  nUR                  (       d  M  U R                  UR                  5      nUc    gX-  nM<     U$ )zCompute total reduction numel.r   N)rg  r£  r–  rß  )r$   r2   Útreerß  s       r)   Ú_compute_reduction_numelÚ%PallasKernel._compute_reduction_numelÌ  sK   € àˆØ×$Ô$ˆDØ× × Ñ ØŸ™ t§z¡zÓ2�Ø‘=ÙØ‘’ñ %ð ˆr,   c                óŒ  ^ • T R                  5       nUb  US:”  a  g/ nT R                  R                   H�  nT R                  U5      nUc    gUu  pVpxn	U	(       d    g[	        USS 5      " 5       n
U
b  SSKnX«R                  :X  a    g[        U 4S jU 5       5      nSU;   a    gUR                  U5        Mƒ     U(       a  [        [        U5      5      S:”  a  gU(       ae  SnUS    H  nXÞ-  nM	     SnT R                   H;  nUR                  (       a  M  T R                  UR                  5      nUc    gUU-  nM=     Xý:w  a  gg)	aV  
Check if TMA (Tensor Memory Accelerator) approach can be used.
TMA works for simple element-wise ops but not for:
- Reductions (need different accumulation patterns)
  TODO: TMA supports float64 for loading but not for reductions
- Broadcasting (inputs have different shapes or output differs)
- Non-contiguous tensors (strided, transposed)
Nr   FÚ	get_dtypec                 ó   • g r`   r;   r;   r,   r)   rÀ  Ú4PallasKernel._can_use_tma_approach.<locals>.<lambda>ò  s   € ¸dr,   r   c              3  óF   >#   • U  H  nTR                  U5      v •  M     g 7fr`   ©r–  ©r¤  r�  r$   s     €r)   r¦  Ú5PallasKernel._can_use_tma_approach.<locals>.<genexpr>ú  s   øé € ÐDº8°a §¡¨q× 1Ð 1º8ùs   ƒ!T)rä  r#   Úinput_buffersr¨  rÂ  r  Úfloat64Útupler˜  rk  r   rg  r£  r–  rß  )r$   Úreduction_numelÚinput_shapesr  rX   Úbuf_objr¬  Ú	buf_numelr­  Úis_contiguousÚ	buf_dtyper  Úshape_tupleÚinput_numelr�  Úoutput_numelrã  rß  s   `                 r)   Ú_can_use_tma_approachÚ"PallasKernel._can_use_tma_approach×  sJ  ø€ ð ×7Ñ7Ó9ˆØÑ&¨?¸QÓ+>Øð %'ˆØ—I‘I×+Ô+ˆDØ×(Ñ(¨Ó.ˆDØ‰|ÙØJNÑGˆG˜y¸-Þ Ùô
   ¨±lÔCÓEˆIØÑ$Ûà§¡Ó-Ù ô  ÔD¹8ÓDÓDˆKØ�{Ó"ÙØ×Ñ Ö,ñ- ,ö2 œC¤
¨<Ó 8Ó9¸AÓ=Øö ØˆKØ! !”_�ØÑ ’ñ %ð ˆLØ×(Ô(�Ø×(×(Ñ(Ø ŸN™N¨4¯:©:Ó6�EØ‘}Ù$Ø  EÑ)’Lñ )ð Ó*Øàr,   c                ó’  • [         R                  R                  U5      nUc  gUR                  5       nSnU H  nU R	                  U5      nXFb  UOU-  nM     / nSn[        USS 5      " 5       n	U	(       a  [        U	SS5      OSn
U
b»  [        [        U5      5       H'  nU R	                  X«   5      nUR                  U5        M)     [        U5      S:X  a  US   b  US   S:w  a  SnO[[        U5      S:”  aL  Sn[        [        U5      S-
  S	S	5       H-  nX{   nUb  XÍ:w  a  SnU R	                  X;   5      nUc  M)  XÞ-  nM/     X#XGU4$ )
z~Get buffer metadata (buf_obj, buf_size, buf_numel, actual_strides, is_contiguous).

Returns None if the buffer doesn't exist.
Nr   Tr¾  c                 ó   • g r`   r;   r;   r,   r)   rÀ  Ú/PallasKernel._get_buffer_info.<locals>.<lambda>)  s   € ¸r,   ru  r   Féÿÿÿÿ)	r   rJ  rÆ  Úget_sizer–  rÂ  Úrangerk  r˜  )r$   r  ró  r¬  rô  r�  Úsvalr­  rõ  rÊ  Ú
buf_strideÚiÚactual_strideÚexpected_strideÚdim_sizes                  r)   r¨  ÚPallasKernel._get_buffer_info  s]  € ô
 —'‘'×$Ñ$ TÓ*ˆØ‰?ØØ×#Ñ#Ó%ˆØˆ	ÛˆAØ—>‘> !Ó$ˆDØÑ!1™°qÑ8ŠIñ ð
  "ˆØˆä˜ ,±Ô=Ó?ˆÞ8>”W˜V X¨tÔ4ÀDˆ
àÑ!Üœ3˜x›=Ö)�Ø $§¡¨z©}Ó =�Ø×%Ñ% mÖ4ñ *ô
 �8‹} Ó!Ø! !Ñ$Ñ0°^ÀAÑ5FÈ!Ó5KØ$)�MøÜ�X“ Ó"Ø"#�Üœs 8›}¨qÑ0°"°bÖ9�AØ$2Ñ$5�MØ$Ñ,°Ó0PØ(-˜Ø#Ÿ~™~¨h©kÓ:�HØÓ+Ø'Ñ3šñ :ð  )¸]ÐJÐJr,   c                ó   • U R                  U5      n/ nU HT  nX@R                  ;   d  M  U R                  U   nU R                  UR                  5      nUc  MC  UR	                  U5        MV     SnU H  nXx-  nM	     Xr4$ )zNCompute expected output numel and used vars from iteration variables in index.r   )rò   r�  r–  r—  r˜  )	r$   rb  rs  Úused_range_lengthsrü   r�  Ú
length_valrù  Úls	            r)   Ú _compute_output_numel_from_indexÚ-PallasKernel._compute_output_numel_from_indexA  sŠ   € ð ×,Ñ,¨UÓ3ˆ	àÐÛˆCØ×+Ñ+Õ+Ø×-Ñ-¨cÑ2�Ø!Ÿ^™^¨E¯L©LÓ9�
ØÓ)Ø&×-Ñ-¨jÖ9ñ ð ˆÛ#ˆAØÑŠLñ $ð Ð&Ð&r,   c                óÜ   • [        5       nU H[  n[        R                  " X5      n[        R                  " XT5      nUc  SnU R	                  U5      nUR                  Ub  UOU5        M]     U$ )zD
Extract coefficients of iteration variables from index expression.
r   )r   r   rn  ro  r–  Úadd)r$   rb  rs  Úcoefficientsrü   rt  ru  Úcoefs           r)   Ú_get_index_coefficientsÚ$PallasKernel._get_index_coefficientsU  sj   € ô $.£<ˆÛˆCÜ*×GÒGÈÓSˆHÜ(×@Ò@ÀÓOˆFØ‰~Ø�Ø—>‘> &Ó)ˆDØ×Ñ TÑ%5™T¸6ÖBñ ð Ðr,   c                óV  • S/n[        U5      S:”  aR  Sn/ n[        [        U5      S-
  SS5       H1  nUR                  SU5        U R                  X   5      nUc  M-  Xh-  nM3     U(       a  [	        U5      n	U H
  n
X©;  d  M
    g   g[	        S U 5       5      nU H
  n
X«;  d  M
    g   g)zB
Check if access pattern requires gather (non-standard striding).
r   rÿ  r   Tc              3  ó.   #   • U  H  oc  M  Uv •  M     g 7fr`   r;   ©r¤  r�  s     r)   r¦  Ú5PallasKernel._check_gather_pattern.<locals>.<genexpr>ƒ  s   é € Ð*V²n°¯1©1²nùs   ‚Œ	F)rk  r  Úinsertr–  r   )r$   r¬  r­  rõ  r  Úexpected_stridesr  r  r  Úexpected_stride_setr  Úactual_stride_sets               r)   Ú_check_gather_patternÚ"PallasKernel._check_gather_patterne  sÁ   € ð ˜3Ðäˆx‹=˜1ÓØˆOØ!ÐÜœ3˜x›=¨1Ñ,¨b°"Ö5�Ø ×'Ñ'¨¨?Ô;ØŸ>™>¨(©+Ó6�ØÓ'Ø#Ñ/’Oñ	 6ö ä",Ð-=Ó">ÐÛ$�ØÕ2Ùñ %ð ô !+Ñ*V±nÓ*VÓ VÐÛ$�ØÕ0Ùñ %ð r,   c                óH  ^ • US:w  d  U(       a  X44$ T R                  U5      nUc  X44$ Uu  pgp‰n
T R                  U5      u  p¼T R                  5       nT R                  X,5      nT R	                  XyX®5      n[        U 4S jU 5       5      n[        U5      [        U5      :  =(       a8    [        U5      S:„  =(       a#    US:„  =(       a    [        U5      [        U5      :„  n[        R                  R                  5       R                  S:H  nU
(       + =(       a    [        S U	 5       5      n[        S U 5       5      nU=(       a    U(       + =(       a    X‹:H  nUS:”  aC  X‹:w  d  U(       d  U(       a0  [        U5      S:”  a!  U(       d  U(       d  T R                  U5      S4$ X44$ )	zó
Check if buffer access needs strided indexing due to size mismatch or gather patterns.

This handles cases like:
- Pooling operations where input/output have different sizes
- im2col-like gather patterns
- Transposed or strided buffer access
r]  c              3  óV   >#   • U  H  nTR                  U5      S :w  d  M  S v •  M      g7f©r   Nrë  rì  s     €r)   r¦  Ú7PallasKernel._needs_strided_indexing.<locals>.<genexpr>¬  s#   øé € Ð OªH q¸¿¹ÀqÓ8IÈQÑ8N§¡ªHùs   ƒ) 	)r   r   Útpuc              3  ó(   #   • U  H  oS Lv •  M
     g 7fr`   r;   r  s     r)   r¦  r"  ¶  s   é € ð <
Ú#1˜a�T�M¢>ùó   ‚c              3  óZ   #   • U  H!  n[        U[        [        -  5      (       + v •  M#     g 7fr`   )rÓ   rp  rÔ   )r¤  r&   s     r)   r¦  r"  ¹  s!   é € ÐUÊÀ1¤J¨q´#¼±+Ó$>× >Ñ >Êùs   ‚)+T)r¨  r  r{  r  r  Úsumrk  r   rJ  rK  rL  Úallr©  r~  )r$   r  rb  r}  r×  rX   ró  r¬  rô  r­  rõ  rù  rs  Úall_iter_varsr  Úhas_non_unit_stridesÚbuf_effective_dimsÚnot_all_vars_usedr4  Úis_known_non_contiguousÚhas_symbolic_coefÚskip_for_non_contiguouss   `                     r)   Ú_needs_strided_indexingÚ$PallasKernel._needs_strided_indexingŠ  s›  ø€ ð  ˜Ó¦ØÐ+Ð+à×$Ñ$ TÓ*ˆØ‰<ØÐ+Ð+àFJÑCˆ˜9°mØ"&×"GÑ"GÈÓ"NÑˆØ×+Ñ+Ó-ˆØ×3Ñ3°EÓEˆð  $×9Ñ9Ø mó 
Ðô
 !Ô O©HÓ OÓOÐä�	‹NœS Ó/Ñ/÷ /Ü�I“ Ñ"÷/à" QÑ&÷/ô �I“¤ X£Ñ.ð	 	ô —‘×4Ñ4Ó6×;Ñ;¸uÑDˆØ&3Ô"3÷ #
¼ñ <
Ù#1ó<
ó 9
Ðô  ÑUÉÓUÓUÐà#×P¨F¬
×P°yÑ7Pð 	 ð ˜1ÓØÓ*Ö.?ÖCWÜ�I“ Ó"Þ+Þ%à×/Ñ/°Ó6¸Ð<Ð<àÐ'Ð'r,   c                ó”  • U(       d  US:X  a  X44$ [         R                  R                  U5      nUc  X44$ UR                  5       n[	        U5      S:X  a  SU4$ [	        U5      S:”  a5  U R                  U5      nU(       d  US4$ SU;   a  U R                  U5      S4$ U R                  (       a  SU;   a  U R                  U5      S4$ X44$ )zP
Adjust index expression based on buffer shape (0-dim scalar, multi-dim, etc.).
r]  r   r   TrÐ  )r   rJ  rÆ  r   rk  rˆ  r~  rM  )r$   r  rb  r}  r×  ró  r¬  rÖ  s           r)   Ú_adjust_index_for_buffer_shapeÚ+PallasKernel._adjust_index_for_buffer_shapeÊ  sØ   € ö ˜I¨Ó.ØÐ+Ð+ä—'‘'×$Ñ$ TÓ*ˆØ‰?ØÐ+Ð+à×#Ñ#Ó%ˆô ˆx‹=˜AÓØ˜-Ð'Ð'ô ˆx‹=˜1ÓØ ×4Ñ4°UÓ;ˆMÞ Ø  $�Ð&Ø˜Ó"à×3Ñ3°EÓ:¸DÐ@Ð@ð �;�;˜4 9Ó,Ø×/Ñ/°Ó6¸Ð<Ð<àÐ'Ð'r,   c                óØ  • U(       d  X44$ [         R                  R                  U5      nUc  X44$ UR                  5       n[	        U5      nUS:  a  X44$ U R                  U5      n[	        U5      S:w  a  X44$ [        [        U5      5      n	[        R                  " X)5      n
U R                  [        R                  " X©5      5      nUb  US::  a  X44$ [         R                  R                  R                  X*-
  5      n [        U5      nUS:  d  XÛ:¼  a  X44$ U R                  US   5      nUb  Xë-  S:w  a  X44$ U R"                  R%                  U	5      nUc  X44$ U R                  UR&                  5      nSnU H!  nU R                  U5      nUc  X44s  $ UU-  nM#     Ub	  UU-  U:w  a  X44$ SUS-
  -  nUS:X  a  U SU 3nUS4$ U U SU 3nUS4$ ! [        [         4 a    X44s $ f = f)zË
Try to emit multi-dim slice notation instead of flatten + gather.

For a buffer with shape (d0, ..., dk) and index `stride * var + offset`,
emit `buf[:, ..., :, offset::stride]` when stride divides dk.
r
   r   r   rÿ  z:, rÐ  F)r   rJ  rÆ  r   rk  rò   rl  rm  r   rn  r–  ro  ri  rj  rp  rq  rr  r�  Úgetr—  )r$   r  rb  r}  r×  ró  r¬  Úndimrs  rü   rt  ru  r  rv  Úlast_dimr�  Ú
var_lengthrô  r�  ÚdÚprefixÚ	slice_strs                         r)   Ú_try_multidim_sliceÚ PallasKernel._try_multidim_sliceð  s9  € ö ØÐ+Ð+ä—'‘'×$Ñ$ TÓ*ˆØ‰?ØÐ+Ð+à×#Ñ#Ó%ˆÜ�8‹}ˆØ�!‹8ØÐ+Ð+ð ×,Ñ,¨UÓ3ˆ	Üˆy‹>˜QÓØÐ+Ð+ä”4˜	“?Ó#ˆÜ&×CÒCÀEÓOˆØ—‘Ü×7Ò7¸ÓFó
ˆð ‰>˜V q›[ØÐ+Ð+ä—‘×!Ñ!×*Ñ*¨5Ñ+;Ó<ˆð	,Ü˜V›ˆJð ˜‹>˜ZÓ1ØÐ+Ð+à—>‘> (¨2¡,Ó/ˆØÑ˜xÑ0°AÓ5ØÐ+Ð+ð ×%Ñ%×)Ñ)¨#Ó.ˆØ‰=ØÐ+Ð+Ø—^‘^ E§L¡LÓ1ˆ
Øˆ	ÛˆAØ—‘˜qÓ!ˆAØ‰yØ Ð/Ò/Ø˜‰NŠIñ	 ð
 Ñ ¨fÑ!4¸	Ó!AØÐ+Ð+à˜$ ™(Ñ#ˆØ˜‹?Ø!˜( " V HÐ-ˆIð ˜%ÐÐð "˜( : ,¨b°°Ð9ˆIØ˜%ÐÐøôA œ:Ð&ó 	,ØÐ+Ò+ð	,ús   Ã>G ÇG)Ç(G)c                ó.  • U(       a\  UR                  [        R                  5      =(       d    UR                  [        R                  5      nU(       a  SU S3OUnU SU S3$ U SU S3nUS:X  a#  U R	                  X#5      (       a  SU S3nS	U l        U$ )
z3
Build the load expression based on indexing mode.
r  z).astype(jnp.int64)z[...].flatten()[Ú]Ú[r]  zjnp.transpose(r!   T)rh  r^  ÚMinÚMaxrº  rS  )	r$   Úbufr  rb  r}  r×  Ú
has_minmaxÚidxÚ	load_exprs	            r)   Ú_build_load_exprÚPallasKernel._build_load_expr;  sœ   € ö àŸ™¤5§9¡9Ó-×E°·±¼5¿9¹9Ó1EˆJÞ8B�A�i�[Ð 3Ñ4È	ˆCØ�UÐ*¨3¨%¨qÐ1Ð1ð ˜%˜q  ¨1Ð-ˆIð ˜EÓ! d×&@Ñ&@À×&MÑ&MØ,¨Y¨K°qÐ9�	Ø+/�Ô(àÐr,   c                ó8  ^• UR                  S5      (       d  U$ [        U4S jU R                  R                   5       5      nU(       aP  [        R
                  R                  U5      mTb.  TR                  5       n[        U5      S:X  a  US   S:X  a  SU S3$ U$ )z�
Squeeze (N,1) intermediate buffers when kernel has 1D graph inputs.

This avoids wrong broadcasting: (N,) op (N,1) -> (N,N) instead of (N,)
rD  c              3  óê   >#   • U  Hh  nUR                  S 5      (       + =(       aF    [        R                  R                  U5      =mSL=(       a    [	        TR                  5       5      S:H  v •  Mj     g7f)rD  Nr   )rŒ  r   rJ  rÆ  rk  r   )r¤  rÈ  ró  s     €r)   r¦  ÚBPallasKernel._maybe_squeeze_intermediate_buffer.<locals>.<genexpr>`  sk   øé € ð 
ò 4�ð ×#Ñ# EÓ*Ô*÷ -ÜŸG™G×.Ñ.¨xÓ8Ð8�ÀÐE÷-ä�G×$Ñ$Ó&Ó'¨1Ñ,ô-ò 4ùs   ƒA0A3r
   rÿ  r   zjnp.squeeze(z
, axis=-1))	rŒ  r©  r#   rî  r   rJ  rÆ  r   rk  )r$   r  rG  Úhas_1d_inputr¬  ró  s        @r)   Ú"_maybe_squeeze_intermediate_bufferÚ/PallasKernel._maybe_squeeze_intermediate_bufferV  s˜   ø€ ð �‰˜u×%Ñ%ØÐô ô 
ð !ŸI™I×3Ò3ó	
ó 
ˆö Ü—g‘g×(Ñ(¨Ó.ˆGØÑ"Ø"×+Ñ+Ó-�Ü�x“= AÓ%¨(°2©,¸!Ó*;Ø)¨)¨°JÐ?Ð?àÐr,   c                óô  • [         R                  R                  U5      nUb  [        UR	                  5       5      S:w  a  U$ U R                  UR	                  5       S   5      nUc  U$ UR                  S5      (       a  U$ [         R                  R                  U5      nUb  UR                  (       d  U$ SnU R                  R                   H�  n[         R                  R                  U5      n	U	c  M'  [        U	R	                  5       5      S:”  d  MF  U	R	                  5        V
s/ s H  o R                  U
5      PM     nn
[        S U 5       5      (       a    OSnM‘     Ub  [        U5      S::  a  U$ U R                  U5      n[        U5      S:w  a  U$ [        [        U5      5      nXÀR                  ;  a  U$ U R                  U   nU R                  UR                   5      U:w  a  U$ U R                  R#                  5        VVs/ s H;  u  pïU R                  UR                   5      U:X  d  M&  UR$                  (       a  M9  UPM=     nnn[        U5      S:w  a  U$ ['        U5       VV
s/ s H  u  noªU:X  d  M  UPM     nnn
[        U5      S:w  a  U$ US   nU[        U5      S-
  :X  a  U$ S/[        U5      -  nSUU'   U SSR)                  [+        [,        U5      5       S	3$ s  sn
f s  snnf s  sn
nf )
zGReshape 1D buffers (e.g., batch norm mean) for higher-dim broadcasting.Nr   r   rD  c              3  ó(   #   • U  H  oS Lv •  M
     g 7fr`   r;   r  s     r)   r¦  Ú:PallasKernel._maybe_broadcast_1d_buffer.<locals>.<genexpr>‰  s   é € Ð;ªl¨ •}ªlùr%  rÿ  ú	.reshape(r    r!   )r   rJ  rÆ  rk  r   r–  rŒ  rç  Úis_floating_pointr#   rî  r(  rò   rl  rm  r�  r—  r•  r£  Ú	enumerater±  Úmapr>   )r$   r  rb  rG  ró  Ú
buf_lengthrá   Úref_buf_sizerÈ  Ú	other_bufr�  rs  Úused_varr�  ÚvÚeÚmatching_varsr  Úmatching_dimsÚaxis_posÚreshape_dimss                        r)   Ú_maybe_broadcast_1d_bufferÚ'PallasKernel._maybe_broadcast_1d_bufferp  sÇ  € ô —'‘'×$Ñ$ TÓ*ˆØ‰?œc '×"2Ñ"2Ó"4Ó5¸Ó:ØÐà—^‘^ G×$4Ñ$4Ó$6°qÑ$9Ó:ˆ
ØÑØÐð �?‰?˜5×!Ñ!ØÐÜ—‘×!Ñ! $Ó'ˆØÑ U×%<×%<ØÐð ˆØŸ	™	×/Ô/ˆHÜŸ™×*Ñ*¨8Ó4ˆIØÓ$¬¨Y×-?Ñ-?Ó-AÓ)BÀQÕ)FØ;D×;MÑ;MÔ;OÓPÒ;O°a§¡¨qÖ 1Ñ;O�ÐPÜÑ;©lÓ;×;Ñ;ÙØ#’ñ 0ð Ñ¤3 |Ó#4¸Ó#9ØÐð ×,Ñ,¨UÓ3ˆ	Üˆy‹>˜QÓØÐÜœ˜Y›Ó(ˆØ×0Ñ0Ó0ØÐð ×%Ñ% hÑ/ˆØ�>‰>˜%Ÿ,™,Ó'¨:Ó5ØÐð
 ×-Ñ-×3Ñ3Ô5ô
â5‘�Ø�~‰~˜aŸh™hÓ'¨:Ñ5ó à>?¿n½n÷ Ù5ð 	ñ 
ô
 ˆ}Ó Ó"ØÐô (1°Ô'>ÔRÒ'>™t˜q !ÀzÁ/ŸÑ'>ˆÑRÜˆ}Ó Ó"ØÐà  Ñ#ˆØ”s˜<Ó(¨1Ñ,Ó,ØÐà�sœS Ó.Ñ.ˆØ!#ˆ�XÑØ�˜I d§i¡i´´C¸Ó0FÓ&GÐ%HÈÐJÐJùòQ  Qùó*
ùó Ss$   Ä&K)È	%K.È2K.ÉK.É-K4É=K4c                óž  • US:w  d  U(       a  X#4$ U R                  U5      nU R                  5       nU R                  U5      n[        US5      (       a  UR                  O	[        5       U-  nXv-
  nU(       a  [        U5      S::  a  X#4$ Sn	U R                  R                  5        Hˆ  u  p«U R                  U5      nU(       d  M  XÆ:w  a  M&  U R                  U5      n[        US5      (       a  UR                  O	[        5       U-  nXÎ:w  d  Xv:X  a  Mm  U R                  X«U5      (       a  M†  Sn	  O   U	(       a  U R                  U5      S4$ X#4$ )zê
Check for im2col-like patterns where store uses block variables but load doesn't.

For cat/expand patterns, both load and store prepared indices share block vars.
For im2col patterns, store compresses to block vars but load doesn't.
r]  rz  r   FT)ró   r{  rò   Úhasattrrz  r   rk  rQ  r•  Ú_check_load_is_strided_inputr~  )r$   rb  r}  r×  rc  r|  Ústore_orig_varsÚstore_prep_varsÚnew_varsÚhas_im2col_patternrÈ  Ú
load_indexÚload_orig_varsÚ	prep_loadÚload_prep_varss                  r)   Ú_check_im2col_patternÚ"PallasKernel._check_im2col_pattern²  sl  € ð ˜Ó¦ØÐ+Ð+à×.Ñ.¨uÓ5ˆØ×'Ñ'Ó)ˆ	Ø×2Ñ2°5Ó9ˆô �~ ~×6Ñ6ð ×'Ò'ä“Øñ	ˆð
 #Ñ4ˆö œ3˜Ó/°1Ó4ØÐ+Ð+ð #ÐØ$(×$9Ñ$9×$?Ñ$?Ö$AÑ ˆHØ!×5Ñ5°jÓAˆNÞ!Ùð Ó0Ùð ×-Ñ-¨jÓ9ˆIô ˜9 n×5Ñ5ð ×&Ò&ä“\Øñ	ˆNð Ó/°?Ó3UÙð ×4Ñ4Ø n÷ó ð &*Ð"Ùñ5 %Bö8 Ø×/Ñ/°Ó?ÀÐEÐEàÐ'Ð'r,   c                ó`  • [         R                  R                  U5      nUc  g[        USS 5      " 5       nUc  g[        USS5      nUc  gUR	                  5       n/ nU H[  n	[
        R                  " X)5      n
[
        R                  " X©5      nUc  M4  U R                  U5      nUR                  Ub  UOU5        M]     [        5       n[        U5       HJ  u  pïU R                  X~   5      nUb  US:”  d  M#  U R                  U5      nUR                  Ub  UOU5        ML     [        U5      U:H  $ )zL
Check if load coefficients match buffer strides (strided input vs im2col).
NFr¾  c                 ó   • g r`   r;   r;   r,   r)   rÀ  Ú;PallasKernel._check_load_is_strided_input.<locals>.<lambda>ù  s   € °Dr,   ru  r   )r   rJ  rÆ  rÂ  r   r   rn  ro  r–  r˜  r   rU  r  )r$   rÈ  rj  rk  rD  rÊ  Úbuf_stridesÚ	buf_sizesÚload_coeffsrü   rt  r  Úint_coefÚbuf_stride_setr  r�  r  Úint_ss                     r)   re  Ú)PallasKernel._check_load_is_strided_inputï  s   € ô �g‰g× Ñ  Ó*ˆØ‰;Øä˜˜l©LÔ9Ó;ˆØ‰>Øä˜f h°Ó5ˆØÑØà—L‘L“Nˆ	ð ˆÛ!ˆCÜ*×GÒGÈ
ÓXˆHÜ&×>Ò>¸xÓMˆDØÓØŸ>™>¨$Ó/�Ø×"Ñ"¨xÑ/C¡8ÈÖNñ "ô $›ˆÜ˜kÖ*‰DˆAØ—~‘~ i¡lÓ3ˆHØÑ 8¨a¥<ØŸ™ qÓ)�Ø×"Ñ"¨EÑ,=¡5À1ÖEñ	 +ô ˜+Ó&¨.Ñ8Ð8r,   c                ó  • U R                   (       a  gU R                  U5      nUc  gUu  p4p5n[        U5      S:w  d  [        U5      S:w  a  gU R                  US   5      nU R                  US   5      nUS   nUS   n	Ub  U	b  X‰:  a  Ub  Ub  US:”  a  US:”  d  gU R                  R
                   HL  n
U R                  U
5      nUc  M  Uu      p<n[        U5      S:w  a  M1  US   nUS   nUc  M@  Uc  ME  XÞ:  d  ML    g   g)zÛ
Check if output needs transpose for column-major storage.

Transpose on store is needed when:
- Output has column-major stride (s0 < s1)
- But input(s) have row-major stride
- And we haven't already transposed on load
Fr
   r   r   T)rS  r¨  rk  r–  r#   rî  )r$   r  rX   r¥  r¬  r­  r®  r¯  r°  r±  Úinp_nameÚinp_infoÚinp_stridesÚinp_s0Úinp_s1s                  r)   Ú_check_store_needs_transposeÚ)PallasKernel._check_store_needs_transpose  s+  € ð ×#×#Øà×$Ñ$ TÓ*ˆØ‰<Øà,0Ñ)ˆ�Q¨Üˆ~Ó !Ó#¤s¨8£}¸Ó'9Øà—‘˜x¨™{Ó+ˆØ—‘˜x¨™{Ó+ˆØ˜AÑˆØ˜AÑˆð ‰NØ‘Ø“ØÑ!ØÑ!Ø˜“	Ø˜“	àð Ÿ	™	×/Ô/ˆHØ×,Ñ,¨XÓ6ˆHØÑÙØ&.Ñ#ˆAˆq�! !Ü�;Ó 1Ó$ÙØ  ‘^ˆFØ  ‘^ˆFØÓ! fÓ&8¸V½_Ùñ 0ð r,   c                ó”   • SU S3/nU(       a  UR                  U SU S35        U$ UR                  U SU SU SU SU S3
5        U$ )	zµ
Build store expression for full array assignment.

Handles scalar broadcast, shape matching, and optional transpose.
Returns a list of lines to emit (variable assignment + store).
ú_val = jnp.asarray(r!   ú[...] = jnp.full(z8.shape, _val) if _val.ndim == 0 else jnp.transpose(_val)z3.shape, _val) if _val.ndim == 0 else (_val.reshape(z.shape) if _val.size == z".size else jnp.broadcast_to(_val, z.shape)))r˜  )r$   ÚoutÚvalueÚneeds_transposeÚliness        r)   Ú_build_full_array_store_exprÚ)PallasKernel._build_full_array_store_exprK  s‹   € ð ' u g¨QÐ/Ð0ˆÞØ�L‰LØ�%ð Ø˜5ð !+ð,ôð ˆð �L‰LØ�%ð Ø˜5ð !&Ø&) UÐ*BÀ3À%ð H/Ø/2¨e°8ð=ôð ˆr,   c                ó¸  • US:X  a#  U R                  U5      nU R                  XU5      $ U(       a  US:X  a  SOSn	U SU SU SU	 SU S	U S
3/$ U R                  U5      n
[        R                  R                  U5      nUbG  UR                  5       n[        U5      S:”  a(  U R                  U5      (       d  U R                  XS5      $ U
(       a€  US:X  a  SOSn	SU S3/nSU SU S3nUS:X  aE  U R                  R                  U5        U S3nUR                  U SU SU SU	 SU SU S
35        U$ UR                  U SU SU 35        U$ U SU SU 3/$ )zŠ
Build the store expression based on indexing mode.
mode can be None (set) or "atomic_add" (accumulate).
Returns a list of lines to emit.
r]  Ú
atomic_addr  Úsetú[...] = z[...].flatten().at[(z).flatten()].z(jnp.asarray(z).flatten()).reshape(ú.shape)r   Frƒ  r!   z
(jnp.full(z%.shape, _val) if _val.ndim == 0 else Ú_aliasr  z.flatten()).reshape(rA  z] = )r€  r‰  r‘  r   rJ  rÆ  r   rk  rˆ  rR  r  r˜  )r$   r…  r  rb  r†  r}  r×  Úmoder‡  Ú
scatter_oprÕ  rD  r¬  rˆ  Ú
value_exprÚalias_params                   r)   Ú_build_store_exprÚPallasKernel._build_store_exprd  sÉ  € ð ˜Óà"×?Ñ?ÀÓEˆOØ×4Ñ4°SÀÓQÐQæà"&¨,Ó"6™¸EˆJà�%�x ˜uÐ$8¸¸À=ÐQ[ÐP\ð ]Ø$˜gÐ%:¸3¸%¸wðHðð ð ×.Ñ.¨uÓ5ˆÜ�g‰g× Ñ  Ó&ˆà‰?Ø—|‘|“~ˆHÜ�8‹}˜qÓ ¨×)AÑ)AÀ%×)HÑ)Hà×8Ñ8¸ÀUÓKÐKæà"&¨,Ó"6™¸EˆJØ*¨5¨'°Ð3Ð4ˆEà˜Y˜KÐ'LÈUÈGÐSTÐUð ð �|Ó#à×&Ñ&×*Ñ*¨3Ô/Ø!$  V˜n�Ø—‘Ø�e˜8 K =Ð0DÀYÀKÈ}Ð]gÐ\hÐhiØ!�lÐ"6°s°e¸7ðDôð ˆLð —‘ ˜u A i [°°Z°LÐAÔBØˆLà�%�q˜˜ 4¨ wÐ/Ð0Ð0r,   c           
     óf  • UR                  SS5      nU R                  R                  U5        U S3nUS:X  a  SOSnU(       aw  US   n	US   n
US	   n/ n[        [	        U5      5       H,  nXÚ:X  a  UR                  U	5        M  UR                  S
5        M.     SR                  U5      nU SU SU SU SU S3
$ US   n	US   nUS   n[        R                  R                  U5      nUb  [	        UR                  5       5      OSn[	        U5      [	        U5      -   n[	        U R                  5      nUS-
  nUU:H  =(       a    UU:H  nU(       a{  U VVs/ s H  u  nnUPM
     nnn[	        U5      n[	        U5      nUS:”  a  US:”  a  SU-  nSU-  nU	 SU SU S3nOU	nUR                  U5        UR                  S U 5       5        O;U Vs/ s H  nSPM     nnUR                  U	5        UR                  S U 5       5        SR                  U5      nU SU SU SU SU S3
$ s  snnf s  snf )zBBuild store expression for scatter operations (indirect indexing).Úis_point_scatterFr�  rŒ  r  r�  Úindirect_varÚindirect_dimÚoutput_shapeÚ0r    rŽ  z	[...].at[z].r  r!   Údims_beforeÚ
dims_afterr   r   úNone, ú, NonerA  r]  r@  c              3  ó*   #   • U  H	  u  pUv •  M     g 7fr`   r;   )r¤  Úvar_namerN   s      r)   r¦  Ú9PallasKernel._build_scatter_store_expr.<locals>.<genexpr>ã  s   é € ÐIºj©N¨H�xºjùs   ‚Ú:c              3  ó&   #   • U  H  nS v •  M	     g7f)r¤  Nr;   )r¤  r¥  s     r)   r¦  r£  è  s   é € Ð7ªJ q�sªJùs   ‚)r6  rR  r  r  rk  r˜  r±  r   rJ  rÆ  r   r�  Úextend)r$   r…  r†  Úscatter_infor  r‘  r˜  r”  r’  r™  rš  r›  Úindex_partsÚdimÚindex_tupler�  rž  rD  Úoutput_ndimÚnum_iter_vars_in_storeÚtotal_kernel_iter_varsÚremaining_dimsÚis_element_wiser¢  rN   Ú	n_leadingÚ
n_trailingÚleading_onesÚtrailing_nonesÚindirect_reshapedr¥  s                                  r)   Ú_build_scatter_store_exprÚ&PallasKernel._build_scatter_store_exprŸ  sˆ  € ð (×+Ñ+Ð,>ÀÓFÐð 	×Ñ×"Ñ" 3Ô'Ø˜˜V�nˆð # lÓ2‘U¸ˆ
æà'¨Ñ7ˆLØ'¨Ñ7ˆLØ'¨Ñ7ˆLð ˆKÜœS Ó.Ö/�ØÓ&Ø×&Ñ& |Ö4à×&Ñ& sÖ+ñ	 0ð Ÿ)™) KÓ0ˆKØ�U˜( ; -¨y¸¸ÀRÈ
À|ÐSTÐUZÐT[Ð[\Ð]Ð]ð $ NÑ3ˆØ" =Ñ1ˆØ! ,Ñ/ˆ
ô �g‰g× Ñ  Ó&ˆØ-0©_”c˜#Ÿ,™,›.Ô)À!ˆä!$ [Ó!1´C¸
³OÑ!CÐÜ!$ T×%:Ñ%:Ó!;ÐØ$ q™ˆð # nÑ4÷ AØ&Ð*@Ñ@ð 	ö
 á:EÔFº+©¨°$›8¹+ˆKÑFô ˜KÓ(ˆIÜ˜Z›ˆJØ˜1‹} ¨a£Ø'¨)Ñ3�Ø!)¨JÑ!6�Ø'3 n°A°l°^À3À~ÐFVÐVWÐ$XÑ!à$0Ð!Ø×ÑÐ0Ô1à×ÑÑI¹jÓIÕIñ )4Ó4ª 1›3©ˆKÐ4Ø×Ñ˜|Ô,Ø×ÑÑ7©JÓ7Ô7à—i‘i Ó,ˆàˆe�8˜K˜=¨	°+°¸bÀÀÈAÈeÈWÐTUÐVð	
ùó+ Gùò  5s   ÅH(ÇH.c                ó   • U R                   R                  U5      n[        R                  R	                  U5      nX R
                  U'   U R                  U5      u  pVU R                  XXV5      u  pVU R                  XXV5      u  pVU R                  XXV5      u  pVU R                  X1X%U5      nU(       d)  US:X  a#  U R                  X5      nU R                  XU5      nU R                  R                  U R                  UUS9$ )Nr]  ©rá   )r#   r  r   rJ  rç  rQ  rØ  r0  r3  r=  rH  rN  ra  rö   r÷   rø   )r$   r  rb  rD  rá   r}  r×  rG  s           r)   ÚloadÚPallasKernel.loadï  s  € à�i‰i�o‰o˜dÓ#ˆÜ—‘×!Ñ! $Ó'ˆð ',×Ñ˜dÑ#ð $(×#7Ñ#7¸Ó#>Ñ ˆ	ð $(×#?Ñ#?Ø˜ó$
Ñ ˆ	ð
 $(×#FÑ#FØ˜ó$
Ñ ˆ	ð
 $(×#;Ñ#;Ø˜ó$
Ñ ˆ	ð
 ×)Ñ)¨#°UÀ}ÓUˆ	ö  ¨eÓ!3Ø×?Ñ?ÀÓPˆIà×7Ñ7¸ÀYÓOˆIà�x‰x× Ñ Ø�L‰LØØð !ð 
ð 	
r,   c                óz	  ^ ^^$^%• T R                  T5      nT R                  R                  U5        [        U5      S:X  a  T R	                  T5      $ UU 4S jn[        X#SS9nU Vs/ s H
  oS" U5      PM     nnT R	                  T R                  T5      5      nT R                  T5      nU V	s/ s H  n	[        U	5      PM     n
n	U Vs0 s H  n[        U5      U" U5      _M     nn[        U5      S:X  aº  [        U
5      S:X  a«  US   n[        U5      nUT R                  ;   =(       a    T R                  U   R                  nU(       ae  UT R                  ;   aS  T R                  U   nUR                  nT R                  U5      nST R	                  U5       S3nUR                  UU5      nU$ Sn[        U
5      S:”  a!  T R                  5       U-
  n[        U5      S:H  nU(       a÷  S[        U5      -   nU
 H+  nS	[        U5      -  nU S
U S3nUR                  UU5      nM-     [        U5       H§  u  nn[        U5      nUT R                  ;   d  M#  T R                  U   nUR                  nT R                  U5      nS/U-  nT R	                  U5      UUS-   '   SR                  U5      nST R	                  U5       SU S3nUR                  UU5      nM©     U$ / nU H  nUR!                  U" U5      SU45        M     U H  n	UR!                  U" U	5      SU	45        M     UR#                  S SS9  [        U5       HÎ  u  nn[        U5      nUT R                  ;   d  M#  T R                  U   nUR                  nT R                  U5      nU" U5      m%ST R	                  U5       S3n[%        U%4S jU 5       5      n[%        U%4S jUR'                  5        5       5      nUU-   nUS:”  a  SU-  n U SU  S3nUR                  UU5      nMÐ     U
 H�  nUU   m$[%        U$4S jU 5       5      n![%        U$4S jU 5       5      nU!S:”  a  US:”  a  SU!-  n"SU-  n#U SU" SU# S3nO*U!S:”  a  SU!-  n"U SU" S3nOUS:”  a  SU-  n#U SU# S3nOUnUR                  UU5      nM’     U$ s  snf s  sn	f s  snf )a9  
Handle indexing with both indirect variables and iteration variables.

For example, x[indices, :] generates index = i0 + stride * tmp0
where tmp0 is loaded from indices and i0 is the iteration variable.

We need to convert this to JAX advanced indexing with proper broadcasting.
When there are multiple iteration variables, they need different shapes
to form an outer product (grid) rather than broadcasting together.

Special case: For gather operations where a single iteration variable
and single indirect variable have the same extent, they should be
element-wise aligned, not broadcast into an outer product.

PyTorch advanced indexing semantics: When multiple indirect indices have
the same shape, they are paired element-wise (not outer product), and
the combined result dimension appears at the FRONT of the output.
r   c                ó8   >• TR                  TU [        S5      S9$ )NÚinf)Údefault)rª  rÔ   )rü   rb  r$   s    €€r)   Ú_coeffÚ3PallasKernel._handle_mixed_indexing.<locals>._coeff6  s   ø€ Ø×.Ñ.¨u°cÄ5ÈÃ<Ð.ÐPÐPr,   T©ÚkeyÚreverser   újnp.arange(r!   Fz, 1z.reshape(-1Ú1r    z
).reshape(rm  Úindirectc                ó   • U S   $ )Nr   r;   ri   s    r)   rÀ  Ú5PallasKernel._handle_mixed_indexing.<locals>.<lambda>›  s   € ¨!¨Aª$r,   c              3  ó6   >#   • U  H  oT:  d  M
  S v •  M     g7fr!  r;   ©r¤  r&   Ú	var_coeffs     €r)   r¦  Ú6PallasKernel._handle_mixed_indexing.<locals>.<genexpr>´  s   øé € Ð%N²¨AÀIÁ§a¡a²ùó   ƒ	�	c              3  ó6   >#   • U  H  oT:  d  M
  S v •  M     g7f)r
   Nr;   rÊ  s     €r)   r¦  rÌ  µ  s   øé € ð *Ú7˜!¸y¹=—A‘AÒ7ùrÍ  r   z[:r@  c              3  ó6   >#   • U  H  oT:”  d  M
  S v •  M     g7fr!  r;   ©r¤  r&   Úindirect_coeffs     €r)   r¦  rÌ  Å  s   øé € ÐI¢{ !¸.Ñ6HŸA™A¢{ùrÍ  c              3  ó6   >#   • U  H  oT:  d  M
  S v •  M     g7fr!  r;   rÐ  s     €r)   r¦  rÌ  Æ  s   øé € ÐJª 1¸>Ñ7IŸQ™QªùrÍ  rŸ  rA  r]  z...]z[...)rò   rð   rñ   rk  rõ   Úsortedrô   rŽ  r>   r�  r£  r—  Úreplacer{  rU  r±  r˜  Úsortr'  Úvalues)&r$   rb  Úused_iter_vars_setr¿  rð   rü   Úiter_coeffsr}  Úindirect_var_symsÚsymÚindirect_varsr�  Úindirect_coeffsr¢  Úis_reduction_varÚrange_entryÚ
range_sizeÚrenamed_sizeÚarange_exprÚpaired_indirectÚunused_iter_varsÚn_output_dimsr™  Útrailing_onesÚreshape_exprr  Úshape_partsÚ	shape_strÚall_componentsÚn_trailing_iterÚn_trailing_indirectr±  Útrailing_dimsr°  Úleading_nonesr³  rÑ  rË  s&   ``                                  @@r)   rÒ  Ú#PallasKernel._handle_mixed_indexing  s8  û€ ð& "×5Ñ5°eÓ<Ðð 	×Ñ×"Ñ"Ð#5Ô6äÐ!Ó" aÓ'Ø—:‘:˜eÓ$Ð$ö
	Qô  Ð 2ÈÑMˆÙ.<Ó=ªn s�v˜c–{©nˆÐ=ð —J‘J˜t×3Ñ3°EÓ:Ó;ˆ	Ø ×3Ñ3°EÓ:ÐÙ->Ó?Ò-> cœ˜SžÑ->ˆÐ?ñ 7HÓHÒ6G°œ3˜q›6¡6¨!£9Ò,Ñ6GˆÐHô ˆ~Ó !Ó#¬¨MÓ(:¸aÓ(?Ø  Ñ#ˆCÜ˜3“xˆHà�t×,Ñ,Ñ,×X°×1FÑ1FÀsÑ1K×1XÑ1Xð ö  à˜$×/Ñ/Ó/Ø"&×"7Ñ"7¸Ñ"<�KØ!,×!3Ñ!3�Jà#'×#7Ñ#7¸
Ó#C�LØ$/°·
±
¸<Ó0HÐ/IÈÐ"K�KØ )× 1Ñ 1°(¸KÓ H�IØ Ð ð  ˆÜˆ}Ó Ó!à#×2Ñ2Ó4Ð7IÑIÐô "Ð"2Ó3°qÑ8ˆOæð ¤ NÓ 3Ñ3ˆMó !.�Ø %¬¨NÓ(;Ñ ;�Ø". ¨{¸=¸/ÈÐK�Ø%×-Ñ-¨l¸LÓI’	ñ !.ô $ NÖ3‘��3Ü˜s›8�Ø˜$×/Ñ/Õ/Ø"&×"7Ñ"7¸Ñ"<�KØ!,×!3Ñ!3�Jà#'×#7Ñ#7¸
Ó#C�Lð $' %¨-Ñ"7�KØ)-¯©°LÓ)A�K  A¡Ñ&Ø $§	¡	¨+Ó 6�Ià% d§j¡j°Ó&>Ð%?¸zÈ)ÈÐTUÐVð  ð !*× 1Ñ 1°(¸KÓ H’Iñ# 4ð& Ðð
 ˆÛ!ˆCØ×!Ñ!¡6¨#£;°¸Ð"<Ö=ñ "ã$ˆCØ×!Ñ!¡6¨#£;°
¸CÐ"@ÖAñ %à×Ñ¡¸ÐÑ=ô   Ö/‰FˆAˆsÜ˜3“xˆHØ�d×+Ñ+Õ+Ø"×3Ñ3°CÑ8�Ø(×/Ñ/�
à#×3Ñ3°JÓ?�Ù" 3›K�	à +¨D¯J©J°|Ó,DÐ+EÀQÐG�ô #&Ô%N±Ó%NÓ"N�Ü&)ô *Ø.×5Ñ5Ô7ó*ó 'Ð#ð -Ð/BÑB�
à “>Ø$,¨zÑ$9�MØ%0 M°°M°?À!Ð"D�Kà%×-Ñ-¨h¸ÓD’	ñ5 0ó: *ˆLØ,¨\Ñ:ˆNô ÔI¡{ÓIÓIˆIÜÔJ©ÓJÓJˆJð ˜1‹} ¨a£Ø (¨9Ñ 4�Ø!)¨JÑ!6�Ø". ¨q°°¸sÀ>ÐBRÐRSÐT‘Ø˜Q“Ø (¨9Ñ 4�Ø". ¨q°°¸tÐD‘Ø˜a“Ø!)¨JÑ!6�Ø". ¨t°NÐ3CÀ1ÐE‘à+�à!×)Ñ)¨,¸ÓEŠIñ+ *ð. Ðùò} >ùò
 @ùò Is   Á&R.Â.R3Ã	R8c           	     óú  • Ub  US:w  a  [        SU S35      eU R                  R                  U5      nU R                  R	                  U5        [
        R                  R                  U5      nUS L=(       a    [        UR                  5       5      S:H  nU(       a  SU S3U SU SU S	3/nOŽU R                  X!5      n	U	b?  U R                  R                  U R                  U5      5        U R                  XSX‘U5      /nO;U R                  U5      u  p«U R!                  X*U5      u  p«U R#                  XQX#X«U5      nU H:  nU R$                  R'                  U5        U R(                  R+                  X\45        M<     g )
NrŒ  zpallas store mode 'z' not supportedr   rƒ  r!   r„  z2.shape, _val) if _val.ndim == 0 else _val.reshape(r�  )rc   r#   ÚoutputÚstore_buffer_namesr  r   rJ  rÆ  rk  r   Ú_detect_scatter_patternrð   rñ   rò   rµ  rØ  rn  r•  ÚstoresÚ	writelinerP  r˜  )r$   r  rb  r†  r‘  r…  rD  Ú	is_scalarÚstore_linesr§  r}  r×  Úlines                r)   ÚstoreÚPallasKernel.storeÚ  s‚  € ð
 Ñ ¨Ó 4ÜÐ 3°D°6¸ÐIÓJÐJØ�i‰i×Ñ˜tÓ$ˆØ×Ñ×#Ñ# DÔ)ô �g‰g× Ñ  Ó&ˆØ˜t�O×@¬¨C¯L©L«NÓ(;¸qÑ(@ˆ	æà% e W¨AÐ.Ø�%Ð(¨¨Ð-_Ð`cÐ_dÐdkÐlð‰Kð  ×7Ñ7¸ÓDˆLàÑ'à×#Ñ#×*Ñ*¨4×+CÑ+CÀEÓ+JÔKà×2Ñ2°3¸|ÐSWÓXð‘ð
 ,0×+?Ñ+?ÀÓ+FÑ(�	ð ,0×+EÑ+EØ mó,Ñ(�	ð
 #×4Ñ4Ø˜u¨YÀtó�ó  ˆDØ�K‰K×!Ñ! $Ô'à×"Ñ"×)Ñ)¨3¨+Ö6ò  r,   c                ó¤   • U R                  U5      nUS:X  a  [        R                  " X5      n [        U5      $ ! [        [
        4 a    Us $ f = f)z=Get integer coefficient of a variable in an index expression.r   )Úcoeffr^  Údiffrp  rq  rr  )rb  rü   r¾  rû  s       r)   rª  Ú#PallasKernel._get_index_coefficient
	  sO   € ð
 —‘˜CÓ ˆØ�A‹:Ü—J’J˜uÓ*ˆEð	Ü�u“:ÐøÜœ:Ð&ó 	ØŠNð	ús   ¯
: ºAÁAc                ó  • U R                  U5      n[        U5      S:w  a  gUS   n[        U5      n[        U R	                  X5      5      nUS:X  a  gU R                  U5      (       d  U R                  X%U5      $ U R                  XU5      $ )zDDetect scatter operation pattern. Returns scatter info dict or None.r   Nr   )rŽ  rk  r>   rp  rª  rˆ  Ú_detect_point_scatterÚ_detect_iter_scatter)r$   rb  Úoutput_nameÚindirect_symsÚindirect_symr™  rÑ  s          r)   rò  Ú$PallasKernel._detect_scatter_pattern	  s“   € ð ×/Ñ/°Ó6ˆÜˆ}Ó Ó"Øà$ QÑ'ˆÜ˜<Ó(ˆÜ! $×"=Ñ"=¸eÓ"RÓSˆØ˜QÓØð ×'Ñ'¨×.Ñ.Ø×-Ñ-¨kÈÓXÐXð ×(Ñ(¨¸nÓMÐMr,   c                ó€  • U(       d  g [         R                  R                  U5      nUR                  5        Vs/ s H  n[	        U5      PM     nn[        U5      S:  a  gSn[        U5      S-
  n[        [        U5      S-
  SS5       H  n	X7:X  a  U	n  OXvU	   -  nM     UU/ / SUS.$ s  snf ! [
         a     gf = f)z&Detect single-element scatter pattern.Nr
   r   rÿ  T©r™  rš  r�  rž  r˜  r›  )r   rJ  rÆ  r   rp  Ú	Exceptionrk  r  )
r$   r  r™  rÑ  rD  r�  r›  Ú
cumulativerš  r©  s
             r)   rÿ  Ú"PallasKernel._detect_point_scatter,	  sÖ   € ö Øð	Ü—'‘'×$Ñ$ [Ó1ˆCØ,/¯L©L¬NÓ;ªN qœC žF©NˆLÐ;ô ˆ|Ó˜qÓ Øð ˆ
Ü˜<Ó(¨1Ñ,ˆÜœ˜\Ó*¨QÑ.°°BÖ7ˆCØÓ+Ø"�ÙØ sÑ+Ñ+ŠJñ	 8ð )Ø(ØØØ $Ø(ñ
ð 	
ùò! <øÜó 	Ùð	ús"   Š1B0 »B+ÁB0 Â+B0 Â0
B=Â<B=c                óÀ  ^• U R                  U5      n/ nU H  n[        U R                  X5      5      nUS:”  d  M%  X`R                  ;   d  M6  U R	                  U R                  U   R
                  5      nUc    gUR                  [        U5      Xx45        M�     UR                  TUS45        UR                  S SS9  [        U4S j[        U5       5       S5      n	U	c  gSn
[        XYS-   S 5       H  u  p·nXz:w  a    gX¨-  n
M     X::w  a  gTU	USU	  VVVs/ s H	  u  pËoÜU4PM     snnnXYS-   S  VVVs/ s H	  u  pËoÜU4PM     snnnS	SS
.$ s  snnnf s  snnnf )z0Detect scatter pattern with iteration variables.r   Nrÿ  c                ó   • U S   $ )Nr   r;   ri   s    r)   rÀ  Ú3PallasKernel._detect_iter_scatter.<locals>.<lambda>^	  s   €  A a¢Dr,   TrÁ  c              3  óD   >#   • U  H  u  nu  n  o2T:X  d  M  Uv •  M     g 7fr`   r;   )r¤  r  r  r¥  r™  s       €r)   r¦  Ú4PallasKernel._detect_iter_scatter.<locals>.<genexpr>b	  s$   øé € ÐRÒ&9‘?�1‘l�t˜Q À\Ñ=Q�Q‰QÒ&9ùs   ƒ —	 r   Fr  )rò   rp  rª  r�  r–  r—  r˜  r>   rÕ  rl  rU  r™  )r$   rb  r™  rÑ  rð   Úall_varsrü   rû  r—  Úindirect_posÚexpectedr¥  r©  r  s     `           r)   r   Ú!PallasKernel._detect_iter_scatterM	  s  ø€ ð ×1Ñ1°%Ó8ˆð 02ˆÛ!ˆCÜ˜×3Ñ3°EÓ?Ó@ˆEØ�q�y˜S×$9Ñ$9Õ9ØŸ™¨×(=Ñ(=¸cÑ(B×(IÑ(IÓJ�Ø‘>ÙØ—‘¤ S£¨5Ð 9Ö:ñ "ð 	�‰˜ ~°rÐ:Ô;Ø�‰™.°$ˆÑ7ô ÜR¤i°Ô&9ÓRØó
ˆð ÑØð ˆÜ (¨ÀÑ2BÐ2DÐ)EÖ FÑˆA�fØÓ ÙØÑŠHñ !Gð Ó%Øð )Ø(Ø2:¸=¸LÑ2IÕJÒ2I¡w q¨Q ›FÑ2IÓJØ19ÈÑ:JÐ:LÑ1MÕNÒ1M¡g a¨A˜q›6Ñ1MÓNØ %Ø ñ
ð 	
ùô KùÜNs   ÄEÄ9Ec           	     óP  ^ • T R                   (       d   eUS:X  a  T R                  X5      $ [        U[        5      (       a  [	        S5      eX#U4nUT R
                  R                  ;   a  T R
                  R                  U   $ SSSSSSS	S
.n[        / SQ5      n[        U 4S jU 5       5      nT R                  U5      n	T R                  5       n
[        S T R                  R                  5        5       5      nUS:X  a(  U(       a  U	(       a  U
(       a  SU SU	 S3nGO<SU S3nGO4US;   GaŒ  Xc   nU=(       a    U	=(       a    U	S:„  =(       a    U
nU(       GaW  US:”  GaP  SnT R                  (       Ga0  [        [!        T R                  R#                  5       5      5      nT R                  R                  5        VVs/ s H  u  nnUR$                  (       d  M  UPM     nnnU(       a»  US   nUR'                  U5      nUS:w  a  T R)                  U5      OSnUc  SnT R                  R                  5        VVs/ s H  u  nnUR$                  (       a  M  UPM     nnnU(       a>  US   nUR'                  U5      nUS:w  a  T R)                  U5      OSnUc  SnUU:”  a  SOSnU SU SU S3nOªU SU S3nO¡X6;   as  U=(       a    U	SL=(       a    U	S:„  =(       a    U
nU=(       a    US:„  =(       a    U	SL nU(       a  Xc   nSU SU SU	 SU
 S3	nOFU(       a  Xc    SU S3nO4Xc    SU S3nO)[	        SU S[+        UR-                  5       5       S35      eT R
                  R/                  T R0                  UUS 9nUT R
                  R                  U'   U$ s  snnf s  snnf )!a  
Generate code for reduction operations in JAX/Pallas.

Reductions in Pallas work by:
1. Loading the input data into the kernel
2. Applying JAX reduction operations (jnp.sum, jnp.max, etc.)
3. Storing the reduced result

The reduction happens over the loaded block of data.
Úwelford_reducezHTuple reductions (e.g., welford_combine) not supported in Pallas backendzjnp.sumzjnp.prodzjnp.maxzjnp.minzjnp.anyz
jnp.argmaxz
jnp.argmin)r'  ÚprodÚmaxÚminr©  ÚargmaxÚargmin)rj   r‘  Úzc              3  ó@   >#   • U  H  oTR                   ;   v •  M     g 7fr`   )rÝ  )r¤  r'   r$   s     €r)   r¦  Ú)PallasKernel.reduction.<locals>.<genexpr>©	  s   øé € ÐIÒ6H° §¡Ö,Ò6Hùs   ƒc              3  óP   #   • U  H  u  pUR                   (       d  M  S v •  M     g7fr!  r¢  )r¤  rü   r�  s      r)   r¦  r  °	  s    é € ð 
Ú =‘*�#À×ASÕAS�A‰AÒ =ùs   ‚&�	&Úxor_sumzjnp.bitwise_xor.reduce(rS  z, -1), axis=-1)r!   )r  r  r   r   rÿ  Nr  z, axis=zpallas_partial_reduce(r    z	, axis=0)zReduction type 'z8' not yet supported in Pallas backend. Supported types: z	, xor_sumr¸  )Úinside_reductionÚwelford_reduce_fallbackrÓ   rð  rc   rö   Úreduction_cacher   r©  rà  rä  r'  r�  r•  rQ  rl  rm  rÖ  r£  rû  r–  r”  r‚  r÷   rø   )r$   rá   râ   Úreduction_typer†  Ú	cache_keyÚreduction_opsÚpointwise_prefixesÚhas_pointwiseÚpointwise_numelrñ  Ún_reduction_dimsÚreduction_exprÚreduction_opÚis_partial_reductionÚreduction_axisrj  rü   r�  Úreduction_varsÚr_varÚr_coeffÚr_strideÚpw_varsÚpw_varÚpw_coeffÚ	pw_strideÚis_symbolic_partialr2   s   `                            r)   Ú	reductionÚPallasKernel.reductionz	  s/  ø€ ð" ×$×$Ð$Ð$ð Ð-Ó-Ø×/Ñ/°Ó=Ð=ä�eœU×#Ñ#ÜØZóð ð
 °Ð6ˆ	Ø˜Ÿ™×0Ñ0Ó0Ø—8‘8×+Ñ+¨IÑ6Ð6ð ØØØØØ"Ø"ñ
ˆô (ªÓ8ÐÜÔIÑ6HÓIÓIˆð *.×)CÑ)CÐDVÓ)WˆØ)-×)FÑ)FÓ)Hˆô ñ 
Ø $× 5Ñ 5× ;Ñ ;Ô =ó
ó 
Ðð ˜YÓ&Þ¦¶_Ø#:¸5¸'ÀÈ?ÐJ[Ð[jÐ!k’à#:¸5¸'ÀÐ!C’ØÐ3Ô3ð )Ñ8ˆLð ÷ $Ø#÷$à# aÑ'÷$ð $ð	 !÷ $Ð(8¸1Ô(<ð "$�Ø×(×(Ð(ä!%¤d¨4×+@Ñ+@×+GÑ+GÓ+IÓ&JÓ!K�Jð +/×*?Ñ*?×*EÑ*EÔ*Gô&â*G™J˜C Ø ×-Õ-÷ Ù*Gð #ñ &ö
 &Ø .¨qÑ 1˜à",×"2Ñ"2°5Ó"9˜Ø>EÈ»l 4§>¡>°'Ô#:ÐPQ˜Ø#Ñ+Ø'(˜Hð /3×.CÑ.C×.IÑ.IÔ.Kô#â.K¡
  UØ#(×#5Õ#5÷  Ù.Kð  ñ #ö
 #Ø%,¨Q¡Z˜FØ'1×'7Ñ'7¸Ó'?˜HØDLÐPQÃM¨¯©°xÔ(@ÐWX˜IØ(Ñ0Ø,- 	ð 3;¸YÓ2F©QÈB˜NØ$0 >°°5°'¸ÀÐ@PÐPQÐ!R‘ð %1 >°°5°'¸Ð!;‘ØÓ,ð ÷ $Ø#¨4Ð/÷$à# aÑ'÷$ð $ð	 !ð ×RÐ"2°QÑ"6×R¸?ÈdÐ;Rð  ö $ð  -Ñ<�à#9¸,¸ÀrÈ%ÈÐPRÐSbÐRcÐceÐfuÐevÐvwÐ!x‘Þ$à$1Ñ$AÐ#BÀ!ÀEÀ7È)Ð!T‘ð %2Ñ$AÐ#BÀ!ÀEÀ7È!Ð!L‘äØ" >Ð"2ð 3$Ü$(¨×);Ñ);Ó)=Ó$>Ð#?¸yðJóð ð —‘×"Ñ"Ø�L‰LØØð #ð 
ˆð /5ˆ�‰× Ñ  Ñ+ØˆùóU&ùó#s   Æ?NÇNÈ<N"ÉN"c                ó€   • [         R                  R                  U 5      nUR                  5       nUR	                  5       $ r`   )r   rJ  rÆ  r¾  rõ  )Úbuffer_namerD  rÊ  s      r)   Ú_buffer_is_contiguousÚ"PallasKernel._buffer_is_contiguous
  s1   € ä�g‰g× Ñ  Ó-ˆØ—‘Ó!ˆØ×#Ñ#Ó%Ð%r,   c                ó`  • [        5       nU R                  R                  5       u  p4  nU Vs/ s H  ofR                  PM     nnU Vs/ s H  oˆR	                  S5      (       d  M  UPM     n	nU Vs/ s H  oˆR	                  S5      (       d  M  UPM     n
n[        U R                  R                  R                  5       5      nU Vs/ s H  oˆU;   d  M
  UPM     nnU
(       d  [        S5      eU R                  R                  R                  5        VVs0 s H  u  pÞ[        U[        5      (       d  M  Xí_M      nnnU=(       d    Sn[        R                  R                  5       R                   S:H  n[        R                  R                  5       R                   S:H  nU(       a  SOSn0 nU	 H  nS	UU'   M
     U	 Vs/ s H  nUU   (       d  M  U S
3PM     nnU Vs/ s H  oˆR	                  S5      (       d  M  UPM     nnUU-   nUU-   n[        UR                  5        VVs/ s H  u  nnU(       a  M  UPM     snn5      nU(       d  U(       a  [#        [%        ['        U
5      5      5      nO([)        U
5       VVs/ s H  u  nnUU;   d  M  UPM     nnn[+        SA0 SU_SU_SU_SU_SU_SU_SU	_SU
_SU_SU_SU_SU_SU_SU_SU_SU_SU_6nUU l        U R/                  U5        [        5       nUR1                  5          U R3                  UU5        U R4                  R6                   H  n UR9                  [        U 5      5        M     SSS5        U R                  R                  5       u  p4  nU Vs/ s H  ofR                  PM     nn[        U R                  R                  R                  5       5      nU Vs/ s H  oˆU;   d  M
  UPM     snUl        U Vs/ s H  oˆR	                  S5      (       d  M  UPM     snUl        UUR<                  -   Ul        UU-   Ul         SU SS RC                  UR@                  5       S!3n!UR9                  U!5        UR1                  5          UR6                   HT  n [        U [        5      (       a!  UR9                  U RE                  5       5        M9  UR6                  RG                  U 5        MV     U RH                   H)  u  n"n#U"UR@                  ;   d  M  UR9                  U#5        M+     SSS5        UR9                  S"5        U S#3n$/ n%S$['        UR:                  5      -   n&[)        UR>                  5       H8  u  nnUU;   d  UR	                  S%5      (       d  M$  U%RG                  UU&-   5        M:     U%(       a  S&S RC                  S' U% 5       5      -   S(-   n'OS)n'[#        [%        S$['        UR:                  5      -   5      5      n(S&S RC                  S* U( 5       5      -   S(-   n)UR9                  S+U) S,U' S-35        S.S//UR:                  -   UR>                  -   n*UR9                  SU$ S&S RC                  U*5       S!35        UR1                  5          UR9                  S05        UR9                  S15        UR9                  S25        UR9                  S-5        UR9                  S35        UR9                  S45        UR9                  S55        UR9                  S25        UR9                  S-5        UR9                  S65        UR9                  S75        UR9                  S8S RC                  UR>                  5      -   S9-   5        UR9                  S-5        / n+[)        URJ                  5       H˜  u  n,nUR	                  S5      (       aN  URM                  US:5      (       a5  U S
3n-UR>                  RO                  U-5      n.U+RG                  U.U,45        Mh  Mj  UR>                  RO                  U5      n.U+RG                  U.U,45        Mš     S RC                  S; U+ 5       5      n// n0UR:                   H  n1U0RG                  U1 S<U1 35        M     U0(       a  S=U S>S RC                  U05       S?3n2OU S@3n2U RP                  =(       a#    U RR                  =(       a    U RU                  5       n3U3(       a  U RW                  UU25        O8U RP                  (       a  U RY                  UU25        OU R[                  UU2U+U/5        SSS5        U R]                  UU$5        UR_                  5       $ s  snf s  snf s  snf s  snf s  snnf s  snf s  snf s  snnf s  snnf ! , (       d  f       GN"= fs  snf s  snf s  snf ! , (       d  f       GNˆ= f! , (       d  f       N“= f)Ba}  
Generate the complete Pallas kernel code as a Python string.

This includes:
- Import statements for JAX/Pallas
- The kernel function that operates on refs
- The main wrapper function that handles PyTorch<->JAX conversions via DLPack

Args:
    name: Optional kernel name (will use placeholder if not provided)

Returns:
    str: Complete Python source code for the Pallas kernel
Úout_ptr)r=  Ú
in_out_ptrz2Pallas backend requires at least one output bufferú<KERNEL_NAME>r#  Úcpur   r  Tr�  )r>  Úin_ptrr2  r3  r4  r5  r6  r8  r9  r:  r;  r<  r=  r>  r?  r@  rA  rB  rC  Núdef z_kernel(r    ú):r­  Ú_jit_wrapperr
   r>  r  c              3  ó8   #   • U  H  n[        U5      v •  M     g 7fr`   ©r>   ©r¤  rj   s     r)   r¦  Ú.PallasKernel.codegen_kernel.<locals>.<genexpr>®
  s   é € Ð,Lº^¸¬S°¯V¨Vº^ùó   ‚ú,)z()c              3  ó8   #   • U  H  n[        U5      v •  M     g 7fr`   rF  rG  s     r)   r¦  rH  ²
  s   é € Ð0PÂ¸A´°Q·°ÂùrI  z+@functools.partial(jax.jit, static_argnums=z, donate_argnums=r!   Ú
out_shapesÚ
out_dtypeszout_shapes_pallas = tuple(z&    jax.ShapeDtypeStruct(shape, dtype)z3    for shape, dtype in zip(out_shapes, out_dtypes)z%indexer = lambda n: lambda i: [i] * nzout_specs_pallas = tuple(z,    pl.BlockSpec(shape, indexer(len(shape)))zin_specs_pallas = tuple(z0    pl.BlockSpec(i.shape, indexer(len(i.shape)))z    for i in [r@  Fc              3  ó4   #   • U  H  u  pU S U 3v •  M     g7f)z: Nr;   )r¤  r  Úos      r)   r¦  rH  Õ
  s   é € Ð)PÂK¹&¸1¨Q¨C¨r°!°­+ÂKùs   ‚Ú=zfunctools.partial(z	_kernel, z),z_kernel,r;   )0r   r#   Úpython_argdefsr  rŒ  r   ri  rÖ  ÚRuntimeErrorr½  r•  rÓ   r>   r   rJ  rK  rL  r”  r  rk  rU  r1  Úaliasable_out_ptrsÚ_codegen_importsÚindentÚ_codegen_iteration_varsrø   Ú_linesrô  r;  r?  r@  rA  r±  rÓ  r˜  rP  r:  r6  rb  rM  rN  rú  Ú_codegen_jit_wrapper_tmaÚ_codegen_jit_wrapper_legacy_gpuÚ_codegen_jit_wrapper_cpu_tpuÚ_codegen_main_entryÚgetvalue)4r$   r  r2  Úarg_defsÚ	call_argsr¥  r¿   r8  r'   r9  r:  Úsize_var_namesr;  ÚouterÚinnerr<  r3  r4  r5  r6  r=  Úparamr>  r?  r@  rA  ÚflagrB  rC  rF  ÚctxÚkernel_bodyr÷  Úkernel_signaturer=  Ú
store_lineÚjit_wrapper_nameÚdonate_indicesÚbase_offsetÚdonate_literalÚstatic_argnumsÚstatic_argnums_literalÚwrapper_paramsÚalias_pairsÚout_idxÚ
alias_nameÚ	input_idxÚalias_map_literalÚpartial_argsÚsv_paramÚ
kernel_argÚuse_tmas4                                                       r)   Úcodegen_kernelÚPallasKernel.codegen_kernel!
  s	  € ô Óˆð %)§I¡I×$<Ñ$<Ó$>Ñ!ˆ˜Q Ù)1Ó2ª AŸœ©ˆÐ2Ù&3ÓO¢m ·|±|ÀI×7NŸ1¡mˆÐOá$ó
Ú$�!¯©Ð5N×(O�A‘}ð 	ð 
ô $ D§I¡I×$6Ñ$6×$=Ñ$=Ó$?Ó@ˆÙ&3ÓK¢m ¸NÑ7JŸ1¡mˆÐKÞÜÐSÓTÐTð !%§	¡	× 8Ñ 8× >Ñ >Ô @ô 
â @‘�Ü˜%¤×%ó ˆEŠLÙ @ð 	ñ  
ð ×-˜oˆÜ—‘×4Ñ4Ó6×;Ñ;¸uÑDˆÜŸ7™7×>Ñ>Ó@×EÑEÈÑNÐÞ&6™F¸GÐà+-ˆÛ$ˆEØ%)ˆO˜EÓ"ñ %ñ +:ó
Ú*9 ¸_ÈUÕ=SÓˆuˆg�VÓ©/ð 	ð 
ñ %ó
Ú$�!¯©Ð5M×(N�A‘}ð 	ð 
ð +¨\Ñ9ÐØ)¨MÑ9ÐÜ&Ø$3×$9Ñ$9Ô$;ÔHÒ$;‘j�d˜DÄ4�TÑ$;ÒHó
Ðö žvÜ"&¤u¬S°Ó-?Ó'@Ó"AÑô "+¨=Ô!9ô#â!9‘I�C˜ØÐ,Ñ,÷ Ù!9ð  ñ #ô ò 
Ùð
á#ð
ñ ð
ñ .ð	
ñ
 0ð
ñ (ð
ñ ,ð
ñ (ð
ñ ,ð
ñ "6ð
ñ ,ð
ñ &ð
ñ &ð
ñ !4ð
ñ  2ð
ñ  0ð!
ñ" !4ð#
ˆð& #2ˆÔà×Ñ˜cÔ"ô
 %Ó&ˆØ×ÑÕ!Ø×(Ñ(¨°cÔ:àŸ™×+Ô+�Ø×%Ñ%¤c¨$£iÖ0ñ ,÷ "ð %)§I¡I×$<Ñ$<Ó$>Ñ!ˆ˜Q Ù)1Ó2ª AŸœ©ˆÐ2Ü# D§I¡I×$6Ñ$6×$=Ñ$=Ó$?Ó@ˆÙ*7ÓOª- QÀÑ;NŸq©-ÑOˆÔá$ó
Ú$�!¯©Ð5M×(N�A‘}ñ
ˆÔð #/°×1AÑ1AÑ"AˆÔØ!-°Ñ!=ˆÔð �;�-˜x¨¯	©	°#×2HÑ2HÓ(IÐ'JÈ"ÐMð 	ð 	�‰Ð'Ô(à�[‰[�]Ø#×*Ô*�Ü˜d¤C×(Ñ(Ø—N‘N 4§;¡;£=Ö1à—K‘K×&Ñ& tÖ,ñ	 +ð (,×'=Ô'=Ñ#�˜Ø˜c×4Ñ4Õ4Ø—N‘N :Ö.ñ (>÷ ð 	�‰�rÔØ)˜]¨,Ð7ÐØˆØœ#˜c×1Ñ1Ó2Ñ2ˆÜ" 3×#:Ñ#:Ö;‰IˆC�Ø˜Ó$¨¯©¸×)FÓ)FØ×%Ñ% c¨KÑ&7Ö8ñ <ö Ø  4§9¡9Ñ,L¹^Ó,LÓ#LÑLÈtÑS‰Nà!ˆNÜœe A¬¨C×,?Ñ,?Ó(@Ñ$@ÓAÓBˆØ!$ t§y¡yÑ0PÁÓ0PÓ'PÑ!PÐSWÑ!WÐØ�‰ð'Ø'=Ð&>Ð>OØÐ˜að!ô	
ð ˜<Ð(¨3×+>Ñ+>Ñ>À×AXÑAXÑXð 	ð 	�‰˜Ð.Ð/¨q°·±¸>Ó1JÐ0KÈ2ÐNÔOØ�[‰[�]Ø�N‰NÐ7Ô8Ø�N‰NÐCÔDØ�N‰NÐPÔQØ�N‰N˜3ÔØ�N‰NÐBÔCØ�N‰NÐ6Ô7Ø�N‰NÐIÔJØ�N‰NÐPÔQØ�N‰N˜3ÔØ�N‰NÐ5Ô6Ø�N‰NÐMÔNØ�N‰NÐ+¨d¯i©i¸×8OÑ8OÓ.PÑPÐSVÑVÔWØ�N‰N˜3Ôà13ˆKÜ!*¨3×+<Ñ+<Ö!=‘�˜Ø—?‘? 9×-Ñ-Ø&×*Ñ*¨4°×7Ñ7Ø(, v¨V _˜
Ø$'×$;Ñ$;×$AÑ$AÀ*Ó$M˜	Ø#×*Ñ*¨I°wÐ+?Ö@ñ 8ð
 !$× 7Ñ 7× =Ñ =¸dÓ C�IØ×&Ñ&¨	°7Ð';Ö<ñ ">ð !%§	¡	Ñ)PÁKÓ)PÓ PÐð ˆLØ×/Ô/�Ø×#Ñ# x j°°(°Ð$<Ö=ñ 0ö Ø1°+°¸iÈÏ	É	ÐR^ÓH_ÐG`Ð`bÐc‘
à +˜}¨HÐ5�
ð —‘×W × 6Ñ 6×W¸4×;UÑ;UÓ;Wð ö Ø×-Ñ-¨c°:Õ>Ø——Ø×4Ñ4°S¸*ÕEà×1Ñ1Ø˜ [Ð2Cô÷[ ðb 	× Ñ  Ð&6Ô7Ø�}‰}‹Ðùòu 3ùÚOùò
ùò
 Lùó 
ùò
ùò
ùó Iùó#÷B "Ö!üò 3ùâOùò
÷ Ž]ú÷D �]ús°   ­d<ÁeÁ%eÁ1eÂeÃ	eÃeÄeÄ4eÆ=eÇ	eÇeÇ;eÈ%e 
È7e 
É?e&Êe&ÌA
e,Í;e>Ï	fÏfÏ"fÏ?fÑ.BfÓ7fÙ2J få,
e;æ
fæ
f-c                ó    • SnUR                   (       a  US-  nUS-  nOUR                  (       d  US-  nUR                  R                  USS9  g )Na=  
import functools
import math
import torch
import jax
import jax.numpy as jnp
from jax.experimental import pallas as pl
from torch._inductor.runtime.runtime_utils import (
    pallas_gpu_align_output_specs, pallas_gpu_pad_inputs,
    pallas_gpu_unpad_results, pallas_partial_reduce,
    torch_dtype_to_jax_runtime,
)
z
import jax.exportz8
from torch_tpu._internal.pallas import tpu_torch_pallasz8
from jax.experimental.pallas import mosaic_gpu as plgpuT©Ústrip)r4  r5  r2  Úsplice)r$   rd  Úimportss      r)   rT  ÚPallasKernel._codegen_importsð
  sP   € ðˆð �:�:ØÐ,Ñ,ˆGØÐRÑR‰GØ×%×%ØÐRÑRˆGØ�‰�‰˜ tˆÒ,r,   c           
     óÚ  ^^• U R                   (       a"  U R                  (       d  U R                  (       d  g UR                  S5        [	        U R                   R                  5        Vs/ s HI  n[        UR                  [        [        R                  45      (       d  M4  [        UR                  5      PMK     sn5      mU4S jn/ nUR                  (       a@  UR                  R                  UR                  S   5      nU(       a  UR                  U5        UR                  U R                   R"                  5        Su  pxU H  nU" U5      n	U	S   (       d  M  U	u  px  O   [%        U R                   R'                  5       5      n
/ nS n[)        U
5       HC  u  mu  pÞU R+                  UR                  5      nUb	  Xø:X  a  TnM/  UR                  TXÞU45        ME     [-        U5      n[)        U
5       GHð  u  mu  pÞXÐR                  ;  a  M  [/        U5      nUR                  nU R1                  U5      nU R3                  U5      nU R+                  U5      nUc   U(       a  US:”  ay  TU:w  as  [5        U4S j[)        U5       5       S 5      nUbQ  U R7                  UUU5      nS/U-  nUUU'   SR9                  U5      nS	U S
3nUR                  U SU SU S
35        Mí  UR                  U SU S
35        GM  U(       aO  [-        U5      S:”  a@  Xø:X  a;  SR9                  S U 5       5      nS	U S
3nUR                  U SU SU S
35        GM]  US:”  av  TU:w  ap  [5        U4S j[)        U5       5       5      nU R7                  UUU5      nS/U-  nUUU'   SR9                  U5      nS	U S
3nUR                  U SU SU S
35        GMÙ  UR                  U SU S
35        GMó     g s  snf )Nz*# Define iteration variables as JAX arraysc                ó  >• [         R                  R                  U 5      nUb  [        UR	                  5       5      S::  a  g[        S UR	                  5        5       5      n[        R                  " U5      nUT;   a  X#4$ S$ )Nr   ©NNc              3  ó†   #   • U  H7  n[        U[        [        R                  45      (       a  [        U5      OUv •  M9     g 7fr`   )rÓ   rp  r^  ÚIntegerr  s     r)   r¦  ÚYPallasKernel._codegen_iteration_vars.<locals>._get_nd_shape_if_matches.<locals>.<genexpr>  s5   é € ð â'�Aô % Q¬¬e¯m©mÐ(<×=Ñ=”�A”À1ÔDÚ'ùs   ‚?A)r   rJ  Útry_get_bufferrk  r   rð  rÕ   r  )rÈ  rD  Úshaperß  Úiter_lengthss       €r)   Ú_get_nd_shape_if_matchesÚFPallasKernel._codegen_iteration_vars.<locals>._get_nd_shape_if_matches  sr   ø€ Ü—'‘'×(Ñ(¨Ó2ˆCØ‰{œc #§,¡,£.Ó1°QÓ6Ø!Üñ àŸ™œóó ˆEô —I’I˜eÓ$ˆEØ%*¨lÓ%:�E�>ÐLÀÐLr,   r   r‚  r   c              3  óH   >#   • U  H  u  nu  n    nUT:X  d  M  Uv •  M     g 7fr`   r;   ©r¤  r  Úvidxr¥  rF  s       €r)   r¦  Ú7PallasKernel._codegen_iteration_vars.<locals>.<genexpr>O  s.   øé € ð â6OÑ 2 ¡? D¨!¨Q°Ø# s™{÷ ™AÚ6Oùs   ƒ"™	"rÅ  r    rÄ  r!   z = rS  z = jnp.arange(c              3  ó8   #   • U  H  n[        U5      v •  M     g 7fr`   rF  r  s     r)   r¦  rŽ  j  s   é € Ð%KÒ6J°¤c¨!§f fÒ6JùrI  c              3  óF   >#   • U  H  u  nu  n    o2T:X  d  M  Uv •  M     g 7fr`   r;   rŒ  s       €r)   r¦  rŽ  n  s)   øé € ð %Ú0IÑ,˜!™_˜d A q¨!ÐUXÉ[—A‘AÒ0Iùs   ƒ!˜	!)r�  rM  rð   rô  r   rÖ  rÓ   r—  rp  r^  r„  r:  r<  r6  r˜  r¦  r#   rî  r”  r•  rU  r–  rk  r>   rô   rõ   rl  Ú_broadcast_axis_idxr±  )r$   re  rd  r\  r‰  Úcandidate_buf_namesrÈ  Úreshape_target_shapeÚreshape_target_numelr2   rš  r›  Útotal_var_idxrœ  r�  r  Únum_broadcast_dimsr¢  r—  Úrenamed_lengthÚ
length_strÚbroadcast_idxÚaxis_idxrç  rè  ÚarangerF  rˆ  s                             @@r)   rV  Ú$PallasKernel._codegen_iteration_vars  sð  ù€ ð
 ×%×%¨d¯k¯k¸d×>Q×>QØà×ÑÐJÔKô "ð ×.Ñ.×5Ñ5Ô7óâ7�AÜ˜aŸh™h¬¬e¯m©mÐ(<×=ó ”�A—H‘H–Ù7ñó
ˆõ		Mð !ÐØ××Ø×/Ñ/×3Ñ3°C×4EÑ4EÀaÑ4HÓIˆHÞØ#×*Ñ*¨8Ô4Ø×"Ñ" 4§9¡9×#:Ñ#:Ô;à5?Ñ2ÐÛ+ˆHÙ-¨hÓ7ˆFØ�a�y‰yØ=CÑ:Ð$Ùñ	 ,ô ˜×.Ñ.×4Ñ4Ó6Ó7ˆ	àˆØˆÜ%.¨yÖ%9Ñ!ˆCÑ!�'ØŸ™¨¯©Ó5ˆJØÑ%¨*Ó*LØ #’à×%Ñ% s¨G¸JÐ&GÖHñ &:ô ! Ó0Ðä%.¨y×%9Ñ!ˆCÑ!�'Ø×1Ñ1Ó1ÙÜ˜7“|ˆHØ—\‘\ˆFØ!×1Ñ1°&Ó9ˆNØŸ™ NÓ3ˆJØŸ™¨Ó/ˆJàÑ!æ(Ø*¨QÓ.Ø˜}Ó,ä$(ôä6?ÀÔ6Oóð
 ó%�Mð %Ñ0Ø#'×#;Ñ#;Ø*¨MÐ;Mó$˜ð (+ eÐ.@Ñ&@˜Ø0:˜ HÑ-Ø$(§I¡I¨kÓ$:˜	Ø#.¨z¨l¸!Ð!<˜Ø#×-Ñ-Ø'˜j¨¨F¨8°9¸Y¸KÀqÐIôñ !Ø×%Ñ%¨¨
°.ÀÀÈAÐ&NÔOÚö %ÜÐ,Ó-°Ó1ØÓ6à ŸI™IÑ%KÑ6JÓ%KÓK�	Ø& z l°!Ð4�Ø×%Ñ%¨¨
°#°f°X¸YÀyÀkÐQRÐ&S×TØ# aÓ'¨C°=Ó,@Ü $ô %Ü09¸.Ô0Ió%ó !�ð  ×3Ñ3Ø" MÐ3Eó�ð  #˜eÐ&8Ñ8�Ø(2�˜HÑ%Ø ŸI™I kÓ2�	Ø& z l°!Ð4�Ø×%Ñ%¨¨
°#°f°X¸YÀyÀkÐQRÐ&S×Tà×%Ñ%¨¨
°.ÀÀÈAÐ&N×Oòw &:ùò[s   Á)3O(Â O(c                ó‚   • [        S U  5       5      n[        S U  5       5      nU=(       a    UnU(       a  U$ US-
  U-
  $ )Nc              3  ó^   #   • U  H#  u  p  n[        U5      R                  S 5      v •  M%     g7f©ÚrN©r>   rŒ  ©r¤  r¥  r[  s      r)   r¦  Ú3PallasKernel._broadcast_axis_idx.<locals>.<genexpr>…  s+   é € ð !
Ú5C¡z q¨Q°ŒC�‹F×Ñ˜c×"Ð"²^ùs   ‚+-c              3  óh   #   • U  H(  u  p  n[        U5      R                  S 5      (       + v •  M*     g7frŸ  r¡  r¢  s      r)   r¦  r£  ˆ  s.   é € ð !
Ú9G©:¨1°°A”�A“×!Ñ! #Ó&×&Ñ&ºùs   ‚02r   )r©  )r›  r™  r–  Úhas_reduction_varsÚhas_pointwise_varsÚis_mixeds         r)   r‘  Ú PallasKernel._broadcast_axis_idx|  s\   € ô !ñ !
Ù5Có!
ó 
Ðô !ñ !
Ù9Gó!
ó 
Ðð &×<Ð*<ˆÞØ Ð Ø! AÑ%¨Ñ5Ð5r,   c                ó  • UR                   nUR                  nUR                  nUR                  S5        UR                  S5        UR                  S5        UR                  S5        UR                  S5        U H  nUR                  SU S35        M     UR                  S5        UR                  S	5        UR                  S
5        U Vs/ s H  ow S3PM	     nnU Vs/ s H  ow S3PM	     n	nU Vs/ s H  ow S3PM	     n
nU Vs/ s H  ow S3PM	     nnUR                  S5        UR                  S5        X‰-   nX«-   n[	        [        U5      5       Vs/ s H  nSU 3PM
     nnSR                  Xß-   5      nUR                  SSR                  U5       SU S35        UR                  5          UR                  S5        UR                  S5        UR                  5          UR                  S5        UR                  S5        UR                  S5        [        [        XŠ5      5       H$  u  nu  nnUR                  SU SU SU S35        M&     UR                  S5        UR                  S5        [        U5       H  u  nnUR                  SU S35        M     UR                  S5        UR                  S5        X«-   nUR                  S5      R                  5       nUR                  U SSR                  U5       S35        UR                  S5        UR                  S 5        UR                  S!5        [        X›5       H  u  nnUR                  S"U SU S#35        M      UR                  S$5        UR                  S5        UR                  S%5        S S S 5        UR                  S5        UR                  S&5        UR                  S'5        S S S 5        UR                  S5        UR                  S(5        UR                  S)5        [        U
5       H"  u  nnXN   nUR                  S*U S+U S,35        M$     [        U5       H  u  nnUR                  S*U S-U S.35        M      U H  nUR                  S*U S/35        M     UR                  S5        UR                  S05        UR                  S15        UR                  S5        UR                  S25        UR                  S35        UR                  5          UR                  S45        UR                  S55        UR                  S65        S S S 5        UR                  S75        U H  nUR                  S8U S935        M     UR                  S5        UR                  S5        UR                  S:5        UR                  S;5        g s  snf s  snf s  snf s  snf s  snf ! , (       d  f       GN8= f! , (       d  f       GN= f! , (       d  f       N¿= f)<Nz6# Use lax.fori_loop with TMA for automatic OOB maskingzfrom jax import laxz"_tile_size = 128  # Warpgroup sizez_orig_out_shapes = out_shapesz_max_numel = 0z_max_numel = max(_max_numel, ú.size)úfor shape in out_shapes:z2    _max_numel = max(_max_numel, math.prod(shape))z8_num_tiles = (_max_numel + _tile_size - 1) // _tile_sizeÚ_gmemÚ_smemr­  z4# Wrapper kernel using lax.fori_loop with direct TMAÚ	_barrier_r    zdef _tma_kernel(z, *, rC  zdef _tile_body(_tile_idx, _):z$_tile_start = _tile_idx * _tile_sizez5# TMA load inputs from GMEM to SMEM (OOB auto-masked)zplgpu.copy_gmem_to_smem(z%.at[pl.ds(_tile_start, _tile_size)], z, _barrier_r!   z # Wait for TMA loads to completezplgpu.barrier_wait(_barrier_z# Compute on SMEM tilesÚ,r  z7# TMA store outputs from SMEM to GMEM (OOB auto-masked)zplgpu.commit_smem()zplgpu.copy_smem_to_gmem(z$.at[pl.ds(_tile_start, _tile_size)])zplgpu.wait_smem_to_gmem(0)zreturn Nonez# Iterate over all tilesz.lax.fori_loop(0, _num_tiles, _tile_body, None)zA# Build SMEM scratch shapes for inputs, outputs, and TMA barriersz_scratch_shapes = {}z_scratch_shapes['z'] = plgpu.SMEM((_tile_size,), rR  z*'] = plgpu.SMEM((_tile_size,), out_dtypes[ú])z"'] = plgpu.Barrier(num_arrivals=1)z4# Create flattened output specs aligned to tile sizezV_flat_out_specs, _ = pallas_gpu_align_output_specs(out_shapes, out_dtypes, _tile_size)z## Call plgpu.kernel with TMA kernelz_result = plgpu.kernel(z_tma_kernel,zout_shape=_flat_out_specs,zscratch_shapes=_scratch_shapes,ú)(ú    z.flatten(),z$# Reshape results to original shapesz:return pallas_gpu_unpad_results(_result, _orig_out_shapes))r2  r@  r:  rô  r  rk  r±  rU  rU  ÚzipÚrstripr|  )r$   rd  rv  r2  r@  r:  rb  r'   Úgmem_input_paramsÚgmem_output_paramsÚsmem_input_paramsÚsmem_output_paramsÚwrapper_kernel_paramsÚall_smem_paramsr  Úbarrier_paramsÚscratch_paramsÚgmem_inÚsmem_inr¥  Úkernel_call_argsrU   Úgmem_outÚsmem_outÚ
smem_paramÚ
orig_paramÚbarrier_params                              r)   rX  Ú%PallasKernel._codegen_jit_wrapper_tma�  s~  € Ø�x‰xˆØ!×5Ñ5ÐØ×)Ñ)ˆð 	�‰ÐOÔPØ�‰Ð,Ô-Ø�‰Ð;Ô<Ø�‰Ð6Ô7à�‰Ð'Ô(Û(ˆEØ�N‰NÐ:¸5¸'ÀÐHÖIñ )à�‰Ð1Ô2Ø�‰ÐKÔLà�‰ÐQÔRá2EÓFÒ2E¨Q˜s %›[Ñ2EÐÐFÙ3@ÓA²=¨a  5›k±=ÐÐAÙ2EÓFÒ2E¨Q˜s %›[Ñ2EÐÐFÙ3@ÓA²=¨a  5›k±=ÐÐAà�‰�rÔØ�‰ÐMÔNà 1Ñ FÐØ+Ñ@ˆÜ38¼Ð=PÓ9QÔ3RÓSÒ3R¨a˜I a S›/Ñ3RˆÐSØŸ™ ?Ñ#CÓDˆà�‰Ø˜tŸy™yÐ)>Ó?Ð@ÀÀnÐEUÐUWÐXô	
ð �[‰[�]Ø�N‰N˜2ÔØ�N‰NÐ:Ô;Ø—‘•Ø—‘ÐEÔFØ—‘˜rÔ"à—‘ÐVÔWÜ-6ÜÐ)Ó=ö.Ñ)�AÑ)˜ ð —N‘NØ2°7°)Ð;`ÐahÐ`iÐitÐuvÐtwÐwxÐyöñ.ð —‘˜rÔ"Ø—‘ÐAÔBÜ%Ð&7Ö8‘D�A�qØ—N‘NÐ%AÀ!ÀÀAÐ#FÖGñ 9ð —‘˜rÔ"Ø—‘Ð8Ô9Ø#4Ñ#IÐ Ø&×-Ñ-¨cÓ2×8Ñ8Ó:�	Ø—‘ ) ¨A¨d¯i©iÐ8HÓ.IÐ-JÈ!ÐLÔMà—‘˜rÔ"Ø—‘ØMôð —‘Ð4Ô5Ü*-Ð.@Ö*UÑ&�H˜hØ—N‘NØ2°8°*¸B¸x¸jÐHlÐmöñ +Vð —‘Ð;Ô<Ø—‘˜rÔ"Ø—‘˜}Ô-÷E ðH �N‰N˜2ÔØ�N‰NÐ5Ô6Ø�N‰NÐKÔL÷S ðX 	�‰�rÔØ�‰ØOô	
ð 	�‰Ð-Ô.Ü&Ð'8Ö9‰MˆAˆzØ,Ñ/ˆJØ�N‰NØ# J <Ð/NÈzÈlÐZaÐböñ :ô
 'Ð'9Ö:‰MˆAˆzØ�N‰NØ# J <Ð/YÐZ[ÐY\Ð\^Ð_öñ ;ó ,ˆMØ�N‰NØ# M ?Ð2TÐUöñ ,ð
 	�‰�rÔØ�‰ÐMÔNØ�‰Ødô	
ð 	�‰�rÔØ�‰Ð<Ô=Ø�‰Ð0Ô1Ø�[‰[�]Ø�N‰N˜>Ô*Ø�N‰NÐ7Ô8Ø�N‰NÐ<Ô=÷ ð 	�‰�tÔÛ(ˆEØ�N‰N˜T % ¨Ð4Ö5ñ )à�‰�sÔà�‰�rÔØ�‰Ð=Ô>Ø�‰ÐSÕTùòG GùÚAùÚFùÚAùò T÷ –ú÷ Ž]ú÷P �]úsO   ÃV;Ã#W Ã7WÄW
ÅWÆ63W&Ç)F:WÎ#;W&Ô4W8×
W#	×W&×&
W5×8
Xc                ó¼  • UR                   nUR                  nSSR                  U5       S3nUR                  S5        UR                  S5        U H  nUR                  SU S35        M     UR                  S5        UR                  S	5        UR                  S
5        UR                  S5        UR                  S5        UR                  S5        UR                  S5        UR                  SU S35        UR                  S5        UR                  S5        UR                  SU-   5        UR                  S5        UR                  S5        UR                  S5        UR                  S5        UR                  S5        UR                  S5        UR                  S5        UR                  S5        UR                  S5        UR                  SU S35        UR                  S5        UR                  S5        UR                  S 5        UR                  S!5        UR                  S"5        UR                  S#U-   5        UR                  S$5        UR                  S%5        UR                  S&5        UR                  S'5        UR                  S(5        UR                  S)5        UR                  S*5        UR                  S+U 35        UR                  S,5        UR                  S-5        UR                  S.5        UR                  S"5        UR                  S#U-   5        UR                  S$5        UR                  S%5        UR                  S/5        g )0NrA  r    r@  z7# Check if all tensors have same size (no broadcasting)z_all_sizes = []z_all_sizes.append(rª  r«  z'    _all_sizes.append(math.prod(shape))z_unique_sizes = set(_all_sizes)zH_can_pad = len(_unique_sizes) == 1 and all(s > 1 for s in _unique_sizes)r­  zif _can_pad:z5    # All tensors same size - safe to flatten and padz+    _padded_inputs = pallas_gpu_pad_inputs(r!   zZ    _aligned_out_specs, _is_scalar = pallas_gpu_align_output_specs(out_shapes, out_dtypes)z    _result = plgpu.kernel(z        z%        out_shape=_aligned_out_specs,z    )(*_padded_inputs)zD    return pallas_gpu_unpad_results(_result, out_shapes, _is_scalar)zelse:zA    # Different sizes - check if it's a reduction (scalar output)z)    _out_numel = math.prod(out_shapes[0])r²  z    if _out_numel <= 1:zG        # Scalar output (reduction) - pad inputs but keep scalar outputz/        _padded_inputs = pallas_gpu_pad_inputs(z#        _aligned_out_specs = tuple(z.            jax.ShapeDtypeStruct(shape, dtype)z;            for shape, dtype in zip(out_shapes, out_dtypes)z	        )z        _result = plgpu.kernel(z            z)            out_shape=_aligned_out_specs,z        )(*_padded_inputs)z        return _resultz	    else:zP        # Non-scalar output with broadcasting - broadcast inputs to output shapez%        _target_shape = out_shapes[0]z        _broadcasted = [z>            jnp.broadcast_to(_inp, _target_shape) for _inp in z	        ]z<        _padded_inputs = pallas_gpu_pad_inputs(_broadcasted)z^        _aligned_out_specs, _is_scalar = pallas_gpu_align_output_specs(out_shapes, out_dtypes)zH        return pallas_gpu_unpad_results(_result, out_shapes, _is_scalar))r2  r@  r±  rô  )r$   rd  rv  r2  r@  Ú
input_listrb  s          r)   rY  Ú,PallasKernel._codegen_jit_wrapper_legacy_gpu  sß  € ð �x‰xˆØ!×5Ñ5ÐØ˜Ÿ™Ð#6Ó7Ð8¸Ð:ˆ
ð
 	�‰ÐPÔQØ�‰Ð(Ô)Û(ˆEØ�N‰NÐ/°¨w°fÐ=Ö>ñ )à�‰Ð1Ô2Ø�‰Ð@ÔAØ�‰Ð8Ô9Ø�‰ØVô	
ð 	�‰�rÔØ�‰�~Ô&Ø�‰ÐNÔOØ�‰ÐDÀZÀLÐPQÐRÔSØ�‰Øhô	
ð 	�‰Ð4Ô5Ø�‰�z JÑ.Ô/Ø�‰Ð>Ô?Ø�‰Ð/Ô0Ø�‰ØRô	
ð 	�‰�wÔØ�‰ØOô	
ð 	�‰ÐBÔCØ�‰�vÔØ�‰Ð0Ô1Ø�‰ØUô	
ð 	�‰ÐHÈÈÐTUÐVÔWØ�‰Ð<Ô=Ø�‰ÐGÔHØ�‰ÐTÔUØ�‰�{Ô#Ø�‰Ð8Ô9Ø�‰�~¨
Ñ2Ô3Ø�‰ÐBÔCØ�‰Ð3Ô4Ø�‰Ð/Ô0Ø�‰�{Ô#Ø�‰Ø^ô	
ð 	�‰Ð>Ô?Ø�‰Ð1Ô2Ø�‰ØLÈZÈLÐYô	
ð 	�‰�{Ô#Ø�‰ÐUÔVØ�‰Ølô	
ð 	�‰Ð8Ô9Ø�‰�~¨
Ñ2Ô3Ø�‰ÐBÔCØ�‰Ð3Ô4Ø�‰ØVõ	
r,   c                ó*  • UR                   nUR                  S5        UR                  SU-   5        UR                  S5        UR                  S5        UR                  S5        UR                  SUR                   S35        UR                  S5        UR                  U(       a  S	U S
3OS5        UR                  S5        UR                  (       a.  UR                  SSR	                  UR                  5       S35        UR                  S5        g )Nzreturn pl.pallas_call(r²  z     out_shape=out_shapes_pallas,z    out_specs=out_specs_pallas,z    in_specs=in_specs_pallas,z    interpret=r¯  z    grid=(1,),z    input_output_aliases={ z },z    input_output_aliases={},r±  r    r!   )r2  rô  r6  r@  r±  )r$   rd  rv  ro  rs  r2  s         r)   rZ  Ú)PallasKernel._codegen_jit_wrapper_cpu_tpuV  så   € ð �x‰xˆØ�‰Ð/Ô0Ø�‰�v 
Ñ*Ô+Ø�‰Ð9Ô:Ø�‰Ð8Ô9Ø�‰Ð6Ô7Ø�‰˜¨×(=Ñ(=Ð'>¸aÐ@ÔAØ�‰Ð'Ô(Ø�‰æð +Ð+<Ð*=¸TÑBà/ô	
ð
 	�‰�tÔØ×"×"Ø�N‰N˜T $§)¡)¨C×,CÑ,CÓ"DÐ!EÀQÐGÔHØ�‰�sÕr,   c                ól   • UR                   (       a  U R                  X5        g U R                  X5        g r`   )r4  Ú_codegen_main_entry_tpuÚ_codegen_main_entry_default)r$   rd  rh  s      r)   r[  Ú PallasKernel._codegen_main_entryo  s%   € Ø�:�:Ø×(Ñ(¨Õ?à×,Ñ,¨SÕCr,   c                ó	  • UR                   nUR                  S5        UR                   S3nUR                  nUR                  SU SSR                  UR                  5       S35        UR                  5          UR                  S5        UR                  S5        UR                  S	5        / nUR                   H1  nUR                  U S
U SU S35        UR                  U S35        M3     UR                   H1  nUR                  U S
U SU S35        UR                  U S35        M3     UR                  SSR                  UR                   V	s/ s H	  n	SU	 S3PM     sn	5      -   S-   5        / n
UR                   Ht  n	UR                  R                  U	5      nUb>  [        R                  R                  U5      nUb  U
R                  [        U5      5        M_  U
R                  SU	 S35        Mv     UR                  SSR                  U
5      -   S-   5        SS/nUR!                  UR"                  5        UR!                  U5        UR                  SU SSR                  U5       S35        UR                  SU SUR                  S    S35        UR                  SU S35        UR                  5          UR                  SU S 35        S S S 5        [%        UR                  5      [%        UR                  5      -   nUR                  S!SR                  U5       S"35        UR                  S#5        UR                  5          UR                   Hv  n	UR                  R                  U	5      nUb=  [        R                  R                  U5      nUb  UR                  S$U	 S%U< S&35        M^  UR                  S$U	 S%U	 S'35        Mx     S S S 5        UR                  S"5        UR                  S(U S)35        UR&                   H)  nUR                  U   nUR                  U S*U S+35        M+     S S S 5        g s  sn	f ! , (       d  f       GN{= f! , (       d  f       N�= f! , (       d  f       g = f),Nr­  Ú_mainrB  r  r    ú, stream=None):ú)jax.config.update('jax_enable_x64', True)újax.clear_caches()z+# Build JAX placeholders for export tracingz$_placeholder = jax.ShapeDtypeStruct(z#.shape, torch_dtype_to_jax_runtime(z.dtype))Ú_placeholderúout_shapes = (útuple(r�  rJ  útorch_dtype_to_jax_runtime(rR  úout_dtypes = (rL  rM  zexported = jax.export.export(z, platforms=['tpu'])(r!   zkernel_key = 'z_' + '_'.join(str(s) for s in r   z.if not tpu_torch_pallas.lookup_custom_kernel('z', kernel_key):z)tpu_torch_pallas.register_custom_kernel('z/', kernel_key, exported.mlir_module_serialized)zinput_tensors = [r@  zoutput_shape_tensors = [ztorch.empty(z.shape, dtype=z, device='tpu'),z.dtype, device='tpu'),zTresults = tpu_torch_pallas.call_custom_kernel(input_tensors, output_shape_tensors, 'z', kernel_key)z.copy_(results[r°  )r2  rô  r3  r±  rA  rU  r>  r˜  r?  r:  r<  r6  r   rJ  rç  r   r¦  r;  r”  rC  )r$   rd  rh  r2  Ú	main_nameÚkernel_name_strÚall_jax_input_namesrq  Úptrr  Údtype_exprsrÈ  rá   Úwrapper_placeholder_argsÚinput_tensor_namesrF  Úout_names                    r)   rÌ  Ú$PallasKernel._codegen_main_entry_tpuu  s  € ð �x‰xˆØ�‰�rÔØ—‘Ð' uÐ-ˆ	ØŸ/™/ˆØ�‰Ø�9�+˜Q˜tŸy™y¨×)?Ñ)?Ó@ÐAÀÐQô	
ð �[‰[�]Ø�N‰NÐFÔGØ�N‰NÐ/Ô0ð �N‰NÐHÔIØ"$ÐØ!×.Ô.�
Ø—‘Ø!�lÐ"FØ!�lÐ"EÀjÀ\ÐQYð[ôð $×*Ñ*¨j¨\¸Ð+FÖGñ /ð ×'Ô'�Ø—‘Ø�eÐ?Ø�eÐ>¸s¸eÀ8ðMôð $×*Ñ*¨c¨U°,Ð+?Ö@ñ (ð �N‰NØ Ø—)‘)À×@QÒ@QÓRÒ@Q¸˜v d V¨7Ó3Ñ@QÑRÓSñTàñôð
 &(ˆKØ×)Ô)�Ø×3Ñ3×7Ñ7¸Ó=�ØÑ'ÜŸG™G×-Ñ-¨hÓ7�EØÑ(Ø#×*Ñ*Ô+=¸eÓ+DÔEÙ Ø×"Ñ"Ð%@ÀÀÀgÐ#NÖOñ *ð �N‰NÐ+¨d¯i©i¸Ó.DÑDÀtÑKÔLð )5°lÐ'CÐ$Ø$×+Ñ+¨C×,?Ñ,?Ô@Ø$×+Ñ+Ð,?Ô@Ø�N‰NØ/Ø#Ð$ð %Ø—I‘IÐ6Ó7Ð8¸ð;ôð �N‰NØ  Ð 1ð 2,Ø,/×,=Ñ,=¸aÑ,@Ð+AÀðJôð �N‰NØ@ÀÐ@QÐQ`Ðaôð —‘•Ø—‘ðØ'Ð(Ð(WðYô÷ ô "& c×&6Ñ&6Ó!7¼$¸s×?OÑ?OÓ:PÑ!PÐØ�N‰NÐ.¨t¯y©yÐ9KÓ/LÐ.MÈQÐOÔPð �N‰NÐ5Ô6Ø—‘•Ø×-Ô-�DØ"×7Ñ7×;Ñ;¸DÓA�HØÑ+Ü !§¡× 1Ñ 1°(Ó ;˜Ø Ñ,Ø ŸN™NØ".¨t¨f°NÀ5Á)ÐK[Ð \ôñ %Ø—N‘NØ& t f¨N¸4¸&Ð@VÐWöñ .÷ ð �N‰N˜3Ôà�N‰Nðà#Ð$ Nð4ôð ×.Ô.�Ø×,Ñ,¨SÑ1�Ø—‘ ( ¨?¸3¸%¸rÐBÖCñ /÷s ˆ]ùò. S÷@ –ú÷ •ú÷G �]úsS   Á:CQ1ÅQ	Å'EQ1ËQËA8Q1ÍBQ ÏA'Q1Ñ	Q1Ñ
Q	ÑQ1Ñ 
Q.	Ñ*Q1Ñ1
Q?c                ó¸  ^• UR                   nUR                  S5        UR                   S3nUR                  SU SSR                  UR                  5       S35        UR                  5          UR                  S5        UR                  S5        UR                  S	5        UR                  (       aJ  UR                  S
5        UR                   H)  nU R                  UUUR                  UR                  S9  M+     UR                  S5        UR                   H6  nUR                  S5      (       d  M  U R                  X6UR                  SS9  M8     UR                  S5        UR                   H6  nUR                  S5      (       d  M  U R                  X6UR                  SS9  M8     UR                  S5        UR                  SSR                  UR                   Vs/ s H	  nSU S3PM     sn5      -   S-   5        / nUR                   Ht  nUR                  R                  U5      n	U	b>  [        R                   R#                  U	5      n
U
b  UR%                  ['        U
5      5        M_  UR%                  SU S35        Mv     UR                  SSR                  U5      -   S-   5        0 mUR                   H  nU S3TU'   M     UR                   H  nU S3TU'   M     SS/nUR)                  UR*                  5        UR)                  U4S jUR,                   5       5        UR                  SU SSR                  U5       S35        UR                  S 5        UR.                  (       aJ  UR                  S!5        UR.                   H)  nUR                  U   nUR                  U S"U S#35        M+     S S S 5        g s  snf ! , (       d  f       g = f)$Nr­  rÐ  rB  r  r    rÑ  z/# Enable JAX x64 mode for float64/int64 supportrÒ  rÓ  z*# Convert Torch -> JAX for donated outputs)Ú
contiguousz+# Convert Torch -> JAX for in-place tensorsr>  Fz!# Convert Torch -> JAX for inputsrA  Tz-# Prepare output metadata from PyTorch tensorrÕ  rÖ  r�  rJ  r×  rR  rØ  Ú_jaxrL  rM  c              3  ó.   >#   • U  H
  nTU   v •  M     g 7fr`   r;   )r¤  r  Úarg_name_maps     €r)   r¦  Ú;PallasKernel._codegen_main_entry_default.<locals>.<genexpr>  s   øé € ð %Ú/F t�˜TÖ"Ò/Fùs   ƒzres = r!   zjax.block_until_ready(res)z9result_values = res if isinstance(res, tuple) else (res,)z'.copy_(torch.from_dlpack(result_values[z])))r2  rô  r3  r±  rA  rU  r>  Ú_emit_torch_to_jaxr4  r5  r?  rŒ  r:  r<  r6  r   rJ  rç  r˜  r   r¦  r;  r@  rC  )r$   rd  rh  r2  rÙ  rq  rÜ  r  rÝ  rÈ  rá   Úwrapper_call_argsrF  rà  ræ  s                 @r)   rÍ  Ú(PallasKernel._codegen_main_entry_defaultÜ  s�  ø€ ð �x‰xˆØ�‰�rÔØ—‘Ð' uÐ-ˆ	Ø�‰Ø�9�+˜Q˜tŸy™y¨×)?Ñ)?Ó@ÐAÀÐQô	
ð �[‰[�]Ø�N‰NÐLÔMØ�N‰NÐFÔGØ�N‰NÐ/Ô0Ø××Ø—‘ÐKÔLØ"%×"2Ô"2�Jð ×+Ñ+ØØ"ØŸ
™
Ø#&×#7Ñ#7ð	 ,ó ñ #3ð �N‰NÐHÔIØ×'Ô'�Ø—>‘> ,×/Ó/Ø×+Ñ+¨D°s·z±zÈeÐ+ÓTñ (ð �N‰NÐ>Ô?Ø×'Ô'�Ø—>‘> (×+Ó+Ø×+Ñ+¨D°s·z±zÈdÐ+ÓSñ (ð �N‰NÐJÔKØ�N‰NØ Ø—)‘)À×@QÒ@QÓRÒ@Q¸˜v d V¨7Ó3Ñ@QÑRÓSñTàñôð
 &(ˆKØ×)Ô)�Ø×3Ñ3×7Ñ7¸Ó=�ØÑ'ÜŸG™G×-Ñ-¨hÓ7�EØÑ(Ø#×*Ñ*Ô+=¸eÓ+DÔEÙ Ø×"Ñ"Ð%@ÀÀÀgÐ#NÖOñ *ð �N‰NÐ+¨d¯i©i¸Ó.DÑDÀtÑKÔLØ+-ˆLØ!×.Ô.�
Ø.8¨\¸Ð+>�˜ZÓ(ñ /à×'Ô'�Ø'* e¨4 L�˜SÓ!ñ (ð ".¨|Ð <ÐØ×$Ñ$ S×%8Ñ%8Ô9Ø×$Ñ$ô %Ø/2×/FÒ/Fó%ô ð �N‰N˜VÐ$4Ð#5°Q°t·y±yÐARÓ7SÐ6TÐTUÐVÔWØ�N‰NÐ7Ô8Ø×&×&Ø—‘ØOôð ×2Ô2�CØ"×0Ñ0°Ñ5�HØ—N‘NØ#˜*Ð$KÈCÈ5ÐPSÐTöñ 3÷} ˆ]ùò@ S÷A �]ús.   Á/COÄ6AOÆAOÇ"OÇ2GOÏOÏ
Oc               óL   • U(       a  SOSnU R                  U SU U S35        g )Nz.detach().contiguous()z	.detach()z_jax = jax.dlpack.from_dlpack(r!   )rô  )r2  r¢  r4  rã  Úsuffixs        r)   rè  ÚPallasKernel._emit_torch_to_jax)  s-   € ö .8Ñ)¸[ˆØ�‰˜(˜Ð#AÀ(ÀÈFÈ8ÐSTÐUÕVr,   c                ó2  • [         R                  R                  nU R                  R	                  5       u  pE  nU Vs/ s H  owR
                  PM     nnU V	s/ s H  o™R                  S5      (       d  M  U	PM     n
n	[        [        [        U5      5      n[        U S0 5      nU
 V	s/ s H/  n	UR                  U	S5      (       d  M  X¸R                  U	5         PM1     nn	U SSR                  XÛ-   5       S3nUR                  U5        gs  snf s  sn	f s  sn	f )z7Generate the Python code that calls this Pallas kernel.r=  rS  Fz.run(r    r!   N)r   rJ  Úwrapper_coder#   rQ  r  rŒ  r”  rV  r>   rÂ  r6  rb  r±  rô  )r$   r  ÚnodeÚwrapperr]  r^  r¥  r¿   Úkernel_param_namesr'   r9  Úcall_arg_strsÚ	aliasableÚalias_call_argsÚkernel_calls                  r)   Úcall_kernelÚPallasKernel.call_kernel0  sü   € ä—'‘'×&Ñ&ˆØ$(§I¡I×$<Ñ$<Ó$>Ñ!ˆ˜Q Ù.6Ó7ªh¨Ÿfœf©hÐÐ7Ù&8ÓTÒ&8 ¿L¹LÈ×<SŸ1Ñ&8ˆÐTÜœS¤ iÓ0Ó1ˆÜ˜DÐ"6¸Ó;ˆ	ñ %ó
â$�Ø�}‰}˜Q ×&ó 7ˆM×2Ñ2°1Ó5Ô6Ù$ð 	ð 
ð ˜˜e D§I¡I¨oÑ.MÓ$NÐ#OÈqÐQˆØ×Ñ˜+Õ&ùò 8ùÚTùò
s   ½D
ÁDÁ5DÂ'DÃD)	rS  rS  rM  rQ  rR  rP  rN  rO  rð   )
r%   r<   rN   r<   rX  r  rY  r  r=   ÚNone)rb  r<   r=   r>   )r=   r   )rb  r<   r=   r   )rb  r<   r=   r  )rb  r<   r=   zlist[sympy.Symbol])r=   r”  )r  r>   rb  r<   r=   r  )r=   r  )rb  r<   r=   útuple[str, bool])r  r   r=   úOptional[int])rÞ  r   r=   rû  )r=   rû  )r  r>   r=   z*Optional[tuple[Any, Any, Any, list, bool]])rb  r<   r=   ztuple[int, OrderedSet])rb  r<   rs  r   r=   r   )
r¬  r”  r­  r”  rõ  r  r  r   r=   r  )
r  r>   rb  r<   r}  r>   r×  r  r=   rú  )rD  r>   r  r>   rb  r<   r}  r>   r×  r  r=   r>   )r  r>   rG  r>   r=   r>   )r  r>   rb  r<   rG  r>   r=   r>   )rb  r<   r}  r>   r×  r  r=   rú  )rÈ  r>   rj  r<   rk  r   r=   r  )r  r>   r=   r  )r…  r>   r†  r   r‡  r  r=   r7  r`   )r…  r>   r  r>   rb  r<   r†  r   r}  r>   r×  r  r‘  r   r=   r7  )r…  r>   r†  r   r§  zdict[str, Any]r  r>   r‘  r   r=   r>   )r  r>   rb  r<   r=   r   )
r  r>   rb  r<   r†  r   r‘  r   r=   rù  )r   )rb  r<   rü   zsympy.Symbolr¾  úint | floatr=   rü  )r­  )rb  r<   r  r>   r=   úOptional[dict[str, Any]])r  r>   r™  r>   rÑ  rp  r=   rý  )rb  r<   r™  r>   rÑ  rp  r=   rý  )
rá   r)  râ   r)  r"  rH   r†  ú+Union[CSEVariable, tuple[CSEVariable, ...]]r=   rþ  )r9  r>   r=   r  )r  ra   r=   r>   )rd  r1  r=   rù  )re  r   rd  r1  r=   rù  )r›  zlist[tuple[int, Any, Any, Any]]r™  rp  r–  rp  r=   rp  )rd  r1  rv  r>   r=   rù  )
rd  r1  rv  r>   ro  zlist[tuple[int, int]]rs  r>   r=   rù  )rd  r1  rh  r>   r=   rù  )
r2  r   r¢  r>   r4  r  rã  r  r=   rù  )r  r>   rð  zOptional[IRNode]r=   rù  )Fr?   r@   rA   rB   rC   rf   Ú	overridesÚpallas_pexprrõ   rD  rY   rZ  rd  ra  r~  r{  rò   rˆ  rŽ  r‘  rž  rº  r«  rØ  r*  r–  rà  rä  rú  r¨  r  r  r  r0  r3  r=  rH  rN  ra  rn  re  r€  r‰  r•  rµ  Útyping_extensionsÚoverrider¹  rÒ  rø  rª  rò  rÿ  r   r6  r:  rx  rT  rV  r‘  rX  rY  rZ  r[  rÌ  rÍ  rè  r÷  rD   Ú__classcell__)rU  s   @r)   rF  rF  F  s  ø‡ ñð &€IØ)5€EÐ&Ó5õEð.8Øð8Ø&0ð8Ø9=ð8ØFJð8à	ô8ô!>ôFbôHô<8ô:ô5ôKô7ô.ô0>ô@ô6,ð, óó ðô	ô	ô>ô@(KðT'Øð'à	ô'ð(ØðØ,6ðà	ôð #àð#ð ð#ð ð	#ð
 !ð#ð 
ô#ðJ>(àð>(ð ð>(ð ð	>(ð
 ð>(ð 
ô>(ð@$(àð$(ð ð$(ð ð	$(ð
 ð$(ð 
ô$(ðLI àðI ð ðI ð ð	I ð
 ðI ð 
ôI ðVàðð ðð ð	ð
 ðð ðð 
ôô6ð4@KØð@KØ *ð@KØ7:ð@Kà	ô@KðD;(Øð;(Ø,/ð;(Ø@Dð;(à	ô;(ðz&9Øð&9Ø)3ð&9ØEOð&9à	ô&9ôP2ðhØðØ*ðØ=Aðà	ôðB ð91àð91ð ð91ð ð	91ð
 ð91ð ð91ð ð91ð ð91ð 
õ91ðvN
àðN
ð ðN
ð %ð	N
ð
 ðN
ð ðN
ð 
ôN
ð` ×Ñó&
ó  ð&
ôP@ðD ×ÑàLPð-7Øð-7Ø *ð-7Ø3>ð-7ØFIð-7à	ô-7ó  ð-7ð^ àEFð
Øð
Ø ,ð
Ø7Bð
à	ô
ó ð
ð 57ðNØðNØ.1ðNà	!õNð*
Øð
Ø.1ð
ØCFð
à	!ô
ðB+
Øð+
Ø/2ð+
ØDGð+
à	!ô+
ðZ_àð_ð ð_ð &ð	_ð
 ;ð_ð 
5ô_ðB ó&ó ð&ö
Mô^-ð*uPØ)ðuPØ0?ðuPà	ôuPðn ð6Ø7ð6àð6ð  ð6ð 
ó	6ó ð6ô&vUðpL
Ø"ðL
Ø03ðL
à	ôL
ð\àðð ðð +ð	ð
 ðð 
ôô2DðeDØ"ðeDØ69ðeDà	ôeDðNKØ"ðKØ69ðKà	ôKðZ ðWØðWØ(+ðWØ59ðWØJNðWà	óWó ðW÷'ô 'r,   rF  c                  óF   • \ rS rSr\r\SS j5       r        SS jrSr	g)ÚPallasSchedulingiE  c                ó6   • [        [        R                  /5      $ r`   )r   r   ÚREDUCE_TO_SINGLE_ELEMENT)ÚclsrT  s     r)   Úget_backend_featuresÚ%PallasScheduling.get_backend_featuresH  s   € ô œ>×BÑBÐCÓDÐDr,   c                ó¾  • [         R                  R                  nXR                  ;   a  UR                  U   $ [        R
                  R                  (       a$  [        U[        R
                  R                  5      OSn[        R                  " UR                  S5      5      R                  5       S S nUS:X  a  SU 3nOSU SU 3nXtR                  U'   UR                  SU5      n[        5       nUR                  SU< S	35        UR                  US
S9  UR                  S5        [!        X$5      u  pšU	 SU
 3nUR#                  XxR%                  5       U5        U$ )Nr­  zutf-8é   ÚfusedÚpallas_r¥  r?  zasync_compile.pallas(z, r'''Tr{  z''')Ú
)r   rJ  rï  Úsrc_to_kernelr   ÚtritonÚdescriptive_namesr   ÚhashlibÚsha256ÚencodeÚ	hexdigestrÔ  r   rô  r}  r   Údefine_kernelr\  )r$   Úsrc_codeÚnode_schedulerï   rñ  Ú
fused_nameÚkernel_hashr3  Úcompile_wrapperÚoriginsÚdetailed_originsÚmetadata_comments               r)   r  ÚPallasScheduling.define_kernelN  sO  € ô —'‘'×&Ñ&ˆØ×,Ñ,Ó,Ø×(Ñ(¨Ñ2Ð2ô �}‰}×.×.ô " -´·±×1PÑ1PÔQàð 	ô
 —n’n X§_¡_°WÓ%=Ó>×HÑHÓJÈ2ÈAÐNˆØ˜Ó Ø# K =Ð1‰Kà# J <¨q°°Ð>ˆKØ*5×Ñ˜hÑ'ð ×#Ñ# O°[ÓAˆä(Ó*ˆØ×!Ñ!Ð$9¸+¹ÈÐ"OÔPØ×Ñ˜x¨tÐÑ4Ø×!Ñ! &Ô)ä$7¸Ó$OÑ!ˆØ%˜Y bÐ)9Ð(:Ð;ÐØ×Ñ˜k×+CÑ+CÓ+EÐGWÔXàÐr,   r;   N)rT  ztorch.devicer=   zOrderedSet[BackendFeature])r  r>   r  zSequence[BaseSchedulerNode]rï   rF  r=   r>   )
r?   r@   rA   rB   rF  Úkernel_typeÚclassmethodr	  r  rD   r;   r,   r)   r  r  E  sF   † Ø€KàóEó ðEð
"àð"ð 3ð"ð ð	"ð
 
÷"r,   r  )rN   rp  r=   rp  )BÚ
__future__r   Údataclassesr  rÃ  rÕ   r  Útypingr   r   r   r   r^  r  Útorch.utils._ordered_setr   Útorch.utils._sympy.functionsr	   r­  r   Úirr   Úruntime.runtime_utilsr   rî   r   r   Úvirtualizedr   Úblock_analysisr   Úcommonr   r   r   r   r   Úsimdr   r   r   r"   r   Úcollections.abcrE   rF   rG   Úops_handlerrH   Ú	schedulerrI   ÚMAIN_SUFFIXrM   rO   Ú_loggingÚgetArtifactLoggerr?   rW   rR   rR  rc   rf   Ú	dataclassr1  rF  r  r;   r,   r)   Ú<module>r5     s  ðÝ "ã Û Û Û Û ß 6Ó 6ã ã Ý /Ý 8å Ý Ý 6ß >Ý Ý /÷õ ÷ -ô�Mô ñ< ‹×&Ñ&€ö ß2åÝ+Ý-ð €ð
 €ôLð —.‘.×2Ñ2°8¸]ÓK€÷>ñ >ô2U�,ô Uôt

˜Kô t

ðn ×Ñ÷#ð #ó ð#ô,|''�:ô |''ô~O+�~õ +r,   